0

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

我的计算环境是

  • CPU：英特尔至强 x5690 3.46Ghz * 2EA
  • 操作系统：CentOS 5.8
  • VGA：英伟达 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%]”。

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

有关详细代码的更多信息，与 CUDA 占用计算器不同的实验结果

4

1 回答 1

2

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

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

您发布的代码没有实现合并内存存储，因为同一经线中的线程以 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] = ...;

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

于 2013-03-17T11:12:18.077 回答