各位大，弱问：
global void kernel(i32 *src)
{
shared int a[16];
…
for(int x=0; x<128; x+=16) {
a[threadIdx.x] = src[x+threadIdx.x];
__syncthreads();
…
}
…
}
请问shared mem和sync在这里，是说“对于所有线程，执行完毕x为某个值的取数”or“对于所有线程，执行完循环”？谢谢！
楼主您好，关于您的问题中的诸多疑惑，我为您解答：
（1）关于shared memory对于“所有线程”的问题:
“所有线程”，这里实际上上“一个block里的所有线程”而不是您的grid的所有线程。
同样的，shared memory的内容和有效性，只是对当前block来说的。
在您同意了（1）后，我继续说一下第二点，
（2）关于__syncthreads()同步：
同理，__syncthreads()也只对当前block里的线程们有效。它的作用是，这些线程[*1]会在__syncthreads()处互相等待，等大家都执行完了__syncthreads()，才会继续往下执行。对于你的问题，显然是大家都完成了对src的1次读取和a的1次写入后，才继续执行[*2]。
注释：
[*1] 实际上并不是要求所有线程都参与，__syncthreads()才能继续的，只要warp中有1-31个线程参与了, 那么就如同32个线程都参与了。
[*2] 这里实际上还有__syncthreads()的另外一个作用，除了保证同步，它还保证之前对shared memory的写入实际发生了。
感谢您莅临CUDAZone China,
祝您愉快！
谢谢！
对于第2点，“src的1次读取和a的1次写入”是针对X等于某个值来说的吧？而不是0-127（step=16）
LZ您好：
1：一般认为，__syncthreads()在什么位置，所有的线程就停止在这个地方等待。因此，根据您提供的代码，__syncthreads()下面还有省略号表示的其他内容，所以此时会停在这里，等待同步之后，再继续执行完循环。
2：一般地，使用shared memory的时候，都是先用一部分线程读入shared memory，然后其他线程再用的。那么，因为warp间的相对执行，是乱序的，所以为了保证读入的数据线程真正执行过了，需要在读取数据之后，加上线程同步，然后再使用。您的代码中，循环内部的部分便是如此。
大致如上，祝您编码愉快！
我不认为在序列
{
a[…] = src[…]; //A处
__syncthreads();
…
}
中的A处会有127次循环写入，显然这是一次，一次大家每人都写一次，下次循环才能大家都写第二次。
谢谢二位，本帖结案啦~:)
您客气了，您的莅临使得本论坛增光不少。
盼望您的下次光临！
x==0的时候读写一轮，x==16的时候读写下一轮。每一轮都有更新一次a，考虑到a[16]，应该有16个线程更新a，然后block内所有/部分线程再使用（因为您“…”了这里，所以无法得知具体怎么使用的，不过不影响这里讨论。）
其实这就是5# 横扫千军 斑竹说的意思，我重新叙述了一下。
善哉，已经结案了，祝LZ好运！