最近看编程指南里面的广播机制有些地方不懂，希望版主指点。。编程指南里面讲：
1）：对于计算能力1.x的： 每个线程以线程ID tid为索引，以s为步长从数组中访问一个32位字是一个常见的模式:
extern __shared__ float shared[];
float data=shared[BaseIndex+s*tid];
这种情况下，当s * n是存储体数的倍数时，线程tid和tid+n访问同一存储体，或等价地，当n是16/d的整数倍时，其中d是16和n的最大公约数。因此，只有半束长度小于等于16/d时没有存储体冲突，此时只有d=1，也就是说s是奇数。
2）：对于计算能力为2.x的：
特别地，这意味着不像计算能力1.x的设备，如下的方式访问char数组没有存储体冲突：
extern __shared__ float shared[];
char data=shared[BaseIndex+s*tid];
但是这个在计算能力1.x的设备上是有存储体冲突的。这两个地方到底是怎么回事？我估计对这个理解得不对，麻烦斑竹指点一下。。谢谢
LZ您好：
1：您在文字叙述中提到“n”但是您的示例代码中并无“n”，我表示未能看明白您的叙述。
2：是否有bank conflict这个是容易看出的，1.x的硬件的shared memory有16个bank，2.x和3.x都是32bank的。那么您只需对应half-warp（1.x）和warp（其他）线程观察其访问shared memory的行为即可，如果有多个线程访问同一bank管辖的不同位置的数据，那么会发生bank conflict；如果有多个线程访问同一个bank管辖的相同位置的数据，那么会启动广播机制。至于哪些情况下哪个架构会/不会发生bank conflict，具体分析即可。
3：以及计算能力3.x的硬件（kepler）的shared memory还引入了64-Bit（8B）模式，这使得每个bank的最大可用宽度增加了一倍。您可以在CUDA 5.0的Programming guide的F.5.3章节看到详细的叙述。
最后，友情建议您以英文原版CUDA C Programming Guide为准，以避免各种翻译环节造成的混淆。
祝您好运~
谢谢斑竹。。确实是看的中文。。。这个问题大致明白了。。后面不懂的地方在看看英文描述。。
嗯嗯，不客气的，欢迎您常来论坛参与讨论~
祝您编码顺利！