假设要把 a[64] 存入共享内存 ,于是在共享内存中声明b，如下：
shared floatb;
来读入数组a.
如果在程序中,一个线程能读入两个数组a的元素，也就是线程0读入a[0]和a[1]如下：
b[2i]=a[2i];
b[2i+1]=a[2i+1];其中i是线程的 threadIdx.x。其他项以此类推。
那么就会在线程0读入a[0]和线程8读入a[16]时出现冲突。为了解决这个问题，很多代码中是这样来做的：
b[2i+i/16]=a[2i];
b[2i+1+i/16]=a[2i+1];也就是当读入第a[16]时，不是写在bank0里的b[16]而是写入bank1中的b[17]，也就是集体向右移，把b[16]空了出来。这样就避开了冲突。
我的疑问是：这样的话等于第一个half warp 中的16个线程工作完毕时 ，a[31]应该被写入b[32]，那么当第二个half warp 开始工作的时候 ，线程16是不是把a[32]存入b[33]呢？也就是说是不是只要把a[16]存入b[17]时做一次右移，以后就不会在发生冲突，还是每存入16个元素就要右移一次，空出一个呢？
LZ您好：
1：您说的是16bank的情况，现在流行的GPU架构已经是32bank了，请您注意。
2：按照您的写法，线程8依然会将a[16]写入b[16],这与您的叙述是不相符的，请您检查您的写法是否背离了您的意图。
3：这不过是一个写入64个float数据的shared memory赋值么，干嘛整这么折腾。就一组线程连续赋值过去就完了的事情，既合并访问，又没有bank conflict。非要折腾成一个线程先读偶数的，再读奇数的，还要padding，来回折腾。你就是直接2-way bank conflict，也不过就是写64个float，又能慢到那里去？况且直接上连bank conflict都木有。
请您仔细考虑您的写法是否是值得的。
祝您好运！