如何正确地将全局内存中的写入合并到全局内存中？

Question

请理解我，但我不懂英语。

我的计算环境是

• CPU：英特尔至强 x5690 3.46Ghz * 2EA
• 操作系统：CentOS 5.8
• VGA：Nvidia Geforce GTX580（CC 为 2.0）

我已经阅读了 CUDA C 编程指南上有关“合并内存访问”的文档。 但我不能在我的情况下应用它们。

我有 32x32 块/网格和 16x16 线程/块。 这意味着如下代码。

dim3 grid(32, 32);
dim3 block(16,16);

kernel<<<grid, block>>>(...);

那么，我该如何使用合并后的内存访问呢？

我在下面的内核中使用了代码。

int i = blockIdx.x*16 + threadIdx.x;
int j = blockIdx.y*16 + threadIdx.y;

...

global_memory[i*512+j] = ...;

我使用常量 512 因为线程总数是 512x512 个线程：它是 grid_size x block_size。

但是，我从 Visual Profiler 看到“全局内存存储效率低[平均 9.7%，内核占计算的 100%]”。

Helper 说使用合并的内存访问。 但是，我不知道我应该使用内存的索引上下文。

详细代码，The result of an experiment different from CUDA Occupancy Calculator

Answer 1

在 CUDA 中合并内存加载和存储是一个非常简单的概念 - 同一个 warp 中的线程需要从内存中适当对齐的连续单词加载或存储。

在 CUDA 中，warp 大小为 32，warp 由同一块内的线程形成，按顺序排列，threadIdx.{xyz} 的 x 维度变化最快，y 次之快，z 最慢（从功能上讲，这与数组中的列主要排序相同）。

您发布的代码未实现合并内存存储，因为同一经线中的线程以 512 个字的间距存储，而不是在所需的 32 个连续字内。

改善合并的一个简单技巧是按列主要顺序寻址内存，因此：

int i = blockIdx.x*16 + threadIdx.x;
int j = blockIdx.y*16 + threadIdx.y;

...

global_memory[i+512*j] = ...;

按照您在问题中展示的精神，在 2D 块和网格上实现合并的更通用方法如下：

   tid_in_block = threadIdx.x + threadIdx.y * blockDim.x;
   bid_in_grid = blockIdx.x + blockIdx.y * gridDim.x;
   threads_per_block = blockDim.x * blockDim.y;

   tid_in_grid = tid_in_block + thread_per_block * bid_in_grid;

   global_memory[tid_in_grid] = ...;

最合适的解决方案将取决于您未描述的代码和数据的其他细节。