在多 GPU 系统中使用 CUDA 迭代一维数组

Question

过去几个月我一直在研究并行编程，现在我正在尝试使我的应用程序适应多 GPU 平台。问题是我仍然不太明白如何使用多个 GPU 遍历数组。

我是否需要将我的主数组分成更小的子数组并将每个子数组发送到每个 GPU，或者有一种方法可以让每个 GPU 在数组的一个片段中迭代？我有这个应用程序的串行和单 GPU 版本，我一直在尝试使用不同的方法来解决这个问题并使其适应多 GPU，但 none return 相同结果与前两个版本相同。我不知道我还能做什么，所以我的结论是我不了解如何在多 GPU 系统中遍历数组。有谁可以帮助我吗？

我的代码运行 N 次迭代，并且在每次迭代中它遍历我的数组（代表一个网格）中的每个值并为其计算一个新值。

这是我的代码现在的样子的草图：

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>

#define DIM     24
#define BLOCK_SIZE 16
#define SRAND_VALUE 585

__global__ void random(int* t, int* newT){

    int iy = blockDim.y * blockIdx.y + threadIdx.y + 1;
    int ix = blockDim.x * blockIdx.x + threadIdx.x + 1;
    int id = iy * (dim+2) + ix;

    if (iy <= DIM && ix <= DIM) {
        if (t[id] % 2 == 0)
            newT[id] = t[id]*3;
        else
            newT[id] = t[id]*5;
    }
}

int main(int argc, char* argv[]){
    int i,j, devCount;
    int *h_test, *d_test, *d_tempTest, *d_newTest;
    size_t gridBytes;

    cudaGetDeviceCount(&devCount);

    gridBytes = sizeof(int)*(DIM)*(DIM);
    h_test = (int*)malloc(gridBytes);

    srand(SRAND_VALUE);
    #pragma omp parallel for private(i,j)
        for(i = 1; i<=DIM;i++) {
            for(j = 1; j<=DIM; j++) {
                h_test[i*(DIM)+j] = rand() % 2;
            }
        }

    if (devCount == 0){
        printf("There are no devices in this machine!");
        return 1; // if there is no GPU, then break the code
    }

    dim3 blockSize(BLOCK_SIZE, BLOCK_SIZE,1);
    int  linGrid = (int)ceil(DIM/(float)BLOCK_SIZE);
    dim3 gridSize(linGrid,linGrid,1);

    dim3 cpyBlockSize(BLOCK_SIZE,1,1);
    dim3 cpyGridRowsGridSize((int)ceil(DIM/(float)cpyBlockSize.x),1,1);
    dim3 cpyGridColsGridSize((int)ceil((DIM+2)/(float)cpyBlockSize.x),1,1);

    else if (devCount == 1){

        cudaMalloc(&d_test, gridBytes);
        cudaMalloc(&d_tempTest, gridBytes);
        cudaMalloc(&d_newTest, gridBytes);

        cudaMemcpy(d_test, h_test, gridBytes, cudaMemcpyHostToDevice);

        for (iter = 0; iter < DIM; iter ++){
            random<<<gridSize, blockSize>>>(d_test, d_newTest);

            d_tempTest = d_test;
            d_test = d_newTest;
            d_newTest = d_tempTest;
        }

        cudaMemcpy(h_test, d_test, gridBytes, cudaMemcpyDeviceToHost);

        return 0;
    }

    else{
        int nThreads, tId, current;
        omp_set_num_threads(devCount);

        for (iter = 0; iter < DIM; iter ++){

            #pragma omp parallel private(tId, h_subGrid, ) shared(h_grid, gridBytes)
            {
                tId = omp_get_thread_num();
                cudaSetDevice(tId);

                cudaMalloc(&d_test, gridBytes);
                cudaMalloc(&d_tempTest, gridBytes);
                cudaMalloc(&d_newTest, gridBytes);

                cudaMemcpy(d_grid, h_grid, gridBytes, cudaMemcpyHostToDevice);

                ******// What do I do here//******

            } 
        }
        return 0;
    }
}

提前致谢。

Answer 1

简短的回答：是的，您应该将数组分成每个 GPU 的子数组。

详情：每个GPU都有自己的显存。在您的代码中，您为每个 GPU 上的整个数组分配内存，并将整个数组复制到每个 GPU。现在您可以对数组的一个子集进行操作。但是当你想复制回来时，你需要确保只复制每个数组的更新部分。从一开始就更好的方法是只复制要在特定 GPU 上更新的数组部分。

解决方法：将multiGPU部分修改成类似下面的样子（如果gridBytes%devCount != 0需要保证不遗漏元素，我的代码片段没有检查这个）

int gridBytesPerGPU = gridBytes/devCount;
cudaMalloc(&d_test, gridBytesPerGPU);
cudaMalloc(&d_newTest, gridBytesPerGPU );

cudaMemcpy(d_test, &h_test[tId*gridBytesPerGPU], gridBytesPerGPU, cudaMemcpyHostToDevice); // copy only the part of the array that you want to use on that GPU
// do the calculation
cudaMemcpy(&h_test[tId*gridBytesPerGPU], d_newTest, gridBytesPerGPU, cudaMemcpyDeviceToHost);

现在您只需要计算出合适的块和网格大小即可。见下文 (c)。如果您对该部分有疑问，请在评论中提问，我将扩展此答案。

除此之外，您的代码中还有一些我不理解的部分：

a) 为什么需要交换指针？

b) 您运行内核部分多次，但 for 循环中的代码不依赖于计数器。为什么？我想念什么？

for (iter = 0; iter < DIM; iter ++){
    random<<<gridSize, blockSize>>>(d_test, d_newTest);

    d_tempTest = d_test;
    d_test = d_newTest;
    d_newTest = d_tempTest;
}

c) 这个简单内核的网格和块大小的计算看起来有点复杂（我在阅读你的问题时跳过了它）。我会将问题视为一维问题，然后一切看起来都会简单得多，包括您的内核。

在多 GPU 系统中使用 CUDA 迭代一维数组

Iterate on 1D array using CUDA in multi-GPU system

c

arrays

cuda

nvidia

multi-gpu