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

Question

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

我是否需要将我的主数组分成更小的子数组并将每个子数组发送到每个 GPU，或者有一种方法可以让每个 GPU 在数组的一个片段中迭代？我有这个应用程序的串行和单 GPU 版本，我一直在尝试使用不同的方法来解决这个问题并使其适应多 GPU，但它们都没有返回与前两个版本相同的结果。我不知道我还能做什么，所以我的结论是我不了解如何在多 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);

现在您只需要计算合适的 block 和网格大小。见下文 (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) 这个简单内核的网格和 block 大小的计算看起来有点复杂(我在阅读你的问题时跳过了它)。我会将问题视为一维问题，然后一切看起来都会简单得多，包括您的内核。