对齐访问

加速计算 CUDA
1

这个问题是一个常见的问题，但是自己之前重来没关注过。程序中要对一个二维 数组(MN)操作，以一维的形式读取.
每个bock内处理一行的数据。我现在把一行的数据读入到shared memory中。
int tid=threadIdx.x;
int bid=blockIdx.x;
shared int pre[M];
for(int i=tid;i<M;i+=blockDim.x)
{
int index=bidM+i;
pre=d_Pre[index];//d_Pre来自显存，是一个M*N的数组
。。。。。。
}

对于上述操作，profile给出的提示是：low global memory load efficieny。效率值好像9.8%左右吧，具体不记得了。
然后我处理完之后写回到d_Pre[index]中，提示是：low global memory restore efficieny。
对于这种情况的话，怎么去实现对齐访问呢？

ps：今天优化程序，确实发现了合理使用存储器的好处，可能一个小改动，性能提升就比较明显的。但是对于全局内存的对齐访问，我就只好请教各位大大了。

2

楼主您好，根据您的写法看，读取d_pre[index]实际上是：
读取d_pre[blockIdx.x * M + threadIdx.x]

按理说只要block形状是(M,1,1)这种，而且M是32的倍数。那么这个是合并的。从您的信息中我未能找到导致“不合并”的原因。

可能的原因（1）：您在做profiling的exe和您的源代码不一致。请问您是在测试最新的编译出的代码吗？

可能的原因（2）：您的block形状是否比较奇葩？例如只有1个线程的block? 建议发一下形状。
:slight_smile:

3

block中的线程设置为32的倍数的。不会有您说的奇葩的设置的。我再看看我的源码。多谢您的解答。