system
1
最近改写了之前的一个程序，用shared mem来代替global mem，想得到一些加速。结果一直有问题，debug时检测到global mem读写冲突，但我自己看不出来相关语句的问题。怀疑是shared mem的变量作为参数传递时需要特殊处理？示意代码如下：
const int ThreadPerBlock=32;
global void kernal(double* A， double* B）
{
shared a[ThreadPerBlock], b[ThreadPerBlock];
a[threadIdx.x] = A[threadIdx.x];
b[threadIdx.x] = B[threadIdx.x];
handle(a,b,threadIdx.x);
…
}
device void handle(double* A, double* B, int id)
{
…
double C = A[id]*B[id];
…
}
代码大概这个样子，然后在 handle函数里会检测到内存冲突。向大家请教一下这样处理正不正确。
system
2
楼主请问：
（1）您的blocks中的线程数目是32个嘛？
（2）您在compute_20,sm_20或者更高上编译，再运行有无问题？
谢谢。
system
3
以及，补充一下2#横扫斑竹，一般而言对shared memory初始化内容之后需要__syncthreads(),如果LZ这里只有32个线程每block，的确可以不加。
特此简要说明。
system
4
就您发的代码来说，如果32个线程的话，是无问题的。
但就怕您不是32个线程，越界了，或者您还有没发的代码，问题在那些代码里面。
建议您检查是否是>32个线程，以及，尽量将您发的代码补全。
感谢深夜来访。
system
5
不好意思我自己弄错了。这个函数被多次调用，有时候传递的是shared变量，有时候传递的是global变量。我没注意到。
我自己出了问题，却麻烦多位版主回复，不好意思。
注意到了别的帖子里版主们发的代码就是加同步的。如果我可以保证每个线程只访问共享内存中唯一的位置，就不需要同步了吧。
借此问一下另外一个共享内存的问题。我的程序里线程之间独立性很好，但每个线程的任务比较复杂，需要很多的局部变量。程序受制于寄存器数目，occupy很低，我能否把一部分局部变量放到共享内存中，以缓解寄存器压力，提高occupy？没见过类似的说法，不知道这想法可不可行。
用编译器指定寄存器数目或者减少程序中的局部变量试过，都不够好。
楼主您好，
在2.0+的硬件上，Shared memory和寄存器的读写峰值带宽差距要比1.x时代更大。所以在2.0+的硬件上，编译器尝试对shared memory进行缓冲，使用寄存器做它的手工管理的cache。
而__syncthreads()可以保证编译器正常的flush这个Cache. 所以没有同步可能会导致您的代码发生错误。
而楼主您说的，每个线程都访问Shared memory中的不同位置（可能是您准备用来做您的手工寄存器的缓冲）。这个行为可能没有任何好处。
一是我之前说的，编译器会尝试你相反方向的缓冲，如果编译器能确定每个线程的访问都是独立的，而且访问的元素下标都能确定，在你没有__syncthreads的情况下，你的手工的这种“缓冲”将被会完全优化掉。编译器拒绝这样做的。
二是，local memory有专业用途用来干这个，而您不应该缓冲到shared memory.
您的想法不可行。