如何正确合并从全局内存写入全局内存？

Question

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

我的计算环境是

• CPU：英特尔至强X5690 3.46GHz的* 2EA
• OS：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％的内核而言]”。

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

用于详细信息代码的更多信息，The result of an experiment different from CUDA Occupancy Calculator

Answer 1

聚结存储器加载和存储在CUDA是一个非常简单的概念 - 在相同的经纱需要从/加载或存储线程成合适地对准，在连续的字记忆。

CUDA中的变形大小为32，并且由相同块内的线程形成变形，从而使得threadIdx.{xyz}的x尺寸变化最快，y次最快，并且z最慢（从功能上来说，这是与数组中的列主要排序相同）。

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

一个简单的黑客攻击，提高凝聚将是解决在列优先的顺序记忆，所以：

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

... 

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

二维块和网格上较普遍的办法，以实现在你表现出什么样的精神凝聚在的问题是这样的：

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] = ...; 

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