1

在过去的几个月里，我一直在研究并行编程，现在我正在尝试使我的应用程序适应多 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;
    }
}

提前致谢。

4

1 回答 1

1

简短的回答：是的，您应该将您的阵列划分为每个 GPU 的子阵列。

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

解决方法：将multiGPU部分修改为如下（需要确保不会遗漏元素if 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）这个简单内核的网格和块大小的计算看起来有点复杂（我在阅读你的问题时跳过了它）。我会将这个问题视为一维问题，然后一切看起来都会简单得多，包括您的内核。

于 2015-07-16T20:57:44.840 回答