c++ - GPU 内存访问和使用 (CUDA)

Question

我是 CUDA 的新手，对内存类型以及如何正确使用它们感到很困惑。

我想做的很简单。我有一个包含 NxN 元素的二维数组 data。 N = 4096。

设i,j 为当前元素。我需要访问它的所有邻居: (i + m, j + n) 和 m,n = -1,0,1。 对于 i = 0，i-1 变为 N-1。 j=0 也是如此。

我使用这些邻居做了一些计算:

(1) 数据[i][j] = 数据[i][j] + a * 数据[i+1][j] + b * 数据[i+ 1][j+1] ...

然后我需要等到所有其他 N * N - 1 线程执行这些计算并同步 data 数组。我将它迭代 K 次。 我不确定是否应该在内核内或外部进行迭代，因为我不知道它如何影响内存。

按照我的理解，我的内核应该是这样的:

__global__ void my_kernel(float* data, int rows, int cols)
{
    int i = (blockIdx.y * blockDim.y + threadIdx.y) % rows; 
    int j = (blockIdx.x * blockDim.x + threadIdx.x) % cols; 

    i = (i >= 0) ? i : rows + i;
    j = (j >= 0) ? j : cols + j;

    int t = i * sizeof(float) + j;

    for(int m = -1; m <= 1; m ++) 
    {
        for(int n = -1; n <= 1; n ++) 
        {
            if(m != 0 || n != 0) 
            {
                int s = (i + m) * sizeof(float) + (j + n);
                data[t] += data[s]/2;  // just an example 
            }
        }
    }   
}

...

int k = 1000;
int rows = 4096;
int cols = 4096;

dim3 block(8,8);
dim3 grid ( (cols + block.x -1)/block.x, (rows + block.y -1)/block.y );

for(int i = 0; i < k; i++) {
    my_kernel<<<grid, block>>>(reinterpret_cast<float*>(mat.data), rows, cols );
}

我不明白的是内存在这里是如何工作的。 这是否意味着 data 数组保留在全局内存中并在每个线程中从内核访问？据我了解，它相对较慢，我需要它尽可能快地计算。

同时，共享内存被限制为每个 block 48KB，data 数组不适合，因为它的大小是 64mb(4*4096*4096 字节)。此外，它会将计算限制在单个 block 内。

不过，我实际上并不需要在每个线程中使用整个 data 数组。我只需要 9 个元素。我的直觉说它应该在不访问全局内存的情况下工作。那可能吗？是否有意义？如果是这样，我该怎么做？

Answer 1

是的，可以使用共享内存实现您想要的东西，是的，它应该加速您的代码，因为您的问题是内存限制的。

两年前我成功地实现了这样的东西。为了弄清楚细节，我查看了矩阵乘法示例(其中使用了共享内存)。它可以在 CUDA Samples 文件夹中找到。

请注意，在您的情况下，主要区别在于您的相邻 block 应该重叠(1 行或 1 列)(当您在每个 block 的边界计算结果时，您将需要它)。实现非常简单，可以在 1 个内核调用中完成，您:

1) 填充共享内存数组。

2) __syncthreads();

3) 使用共享内存数组执行必要的计算。

选择block_size 以便共享数组可以适合每个 block 的共享内存。网格大小将是您的原始数组和 block_size 的比率