GPU 内存访问和使用 (CUDA）

GPU memory access and usage (CUDA)

本文关键字：CUDA 内存 访问 GPU      更新时间：2023-10-16

我是 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. 我使用这些邻居进行一些计算：
    data[i][j] = data[i][j] + a * data[i+1][j] + b * data[i+1][j+1] + ...

  1. 然后我需要等到所有其他N * N - 1线程执行这些计算并同步data数组。

  2. 我迭代了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数组保留在全局内存中，并在每个线程中从内核访问？据我了解，它相对较慢，我需要它尽可能快地计算。

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

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

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

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

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

1) 填满共享内存数组。

2)__syncthreads();

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

选择block_size，以便共享阵列可以容纳在每个块的共享内存中。网格大小将是原始数组和block_size的比率

使用共享内存的这种所谓的 3D 模板图案有一个官方示例：5_Domain_Specific/FDTD3d。由于第三维是按顺序计算的，因此将其简化为 2D 应该很容易。所以是的，即使使用外部k循环，将瓦片+"光环"层放入共享内存也是一个好主意，因为即使在单次k迭代期间，多个线程也将访问相同的坐标i, j。尽管必须仔细进行基准测试，但缓存在 GPU 上变得越来越强大。我已经看到 3D 模板代码由于硬件的良好 L1 缓存而没有从共享内存中受益。

通常，需要在k迭代之间进行网格范围的同步。这通常是通过每次迭代启动一次内核来实现的(请参阅示例)。

使用协作组 API 的当今网格同步功能，可以在内核内进行迭代，但由于用于协作启动的块数量有限，这将更难实现。这意味着每个块可能需要按顺序计算多个图块(或对于共享内存来说太大的图块)，类似于 3D 循环，这将破坏在单次启动中进行计算的目的，因为数据不会保存在共享内存中。此外，不能保证这实际上比使用多次启动更快。

有一些算法策略可以最小化块之间的同步/通信量。它们通常应用于使用 MPI 的多节点 HPC 代码(此处：MPI 节点<-> CUDA 块)。例如，可以通过重叠瓷砖(创建与外部"重影"或"光晕"层不同的重复层)来进行冗余计算，然后在每次"内部迭代"(阻止k迭代)时缩小有效图块大小。这样，可以在没有合作启动的情况下在单个内核启动中实现多个但不是全部k迭代。在 GPU 上进行此优化的一个问题是，块内的并行度随着内部迭代而下降。

相关文章：