感谢奈奈同学工作之余, 百忙之中的回复.
楼主的代码(double)实际上在profiler下, 并不会报告bank conflict了. 该行为从计算能力5.0+就可以观察到(包括楼主的pascal的显卡), 但没有在计算能力7.5(图灵)+上测试过.
你可以和楼主用一张6.1的卡在profiler下观察到这点.
我们(你, 我, W)当初的笔记和讨论实际上并不非常适用于楼主今日的话题. 楼主主要是用于连续顺序访问的. 这种情况的确不会, 但是我不知道原因, 手册也没有做出解释.
我个人的猜测是shared memory和相关的单元在处理的时候, 使用了一个很小的缓冲区, 将32个8B里面的前面的16个8B组合了一下, 然后后面的16B个8B组合了一下, 有点像按照half warp(本例)拆分了访问一样. 或者对于一次性的非4B的访问(指真正的访存指令时候的超过8B, 例如用double和float2都可以激活这种效果), 例如8B的访问的时候, 按照half warp进行, 而对于标准的4B访问, 按照warp进行.
我建议奈奈同学分别使用4B, 8B, 16B访问来测试profiler的bank conflict报告情况. (即: float, float2, float4, 或者float, double, double2), 这3种类型(以及他们对应的整数类型)均会导致编译器生成LDS, LDS.64, LDS.128这种指令. 丹丹可以猜测和验证是否4B是按照warp来的, 8B按照half warp来的, 而16B按照1/4 warp来的. 无论验证或者否定, 都将会是一个有趣的发现, 不是吗?
看到你给出的链接里面的手抄笔记的熟悉的字体, 感慨万千, 虽然当年参与讨论的人已经有一个不在了, 但是该网站保存了足够美好的回忆.
S写于祖国的边疆
感谢奈奈与屠戮人神 两位前辈的不吝指导。
我做了上机测试,确实如同屠戮人神所说,没有出现bank conflict。且在profiler里记录了每个warp的lstores transaction 为2次。
我在nvida的一份关于volta架构的文档里看到相关的说明(
https://on-demand.gputechconf.com/gtc/2018/presentation/s81006-volta-architecture-and-performance-optimization.pdf)(见66页),大致意思是:
对于8B的数据,会分两个阶段处理数据,第一个阶段处理前半个warp,后一个阶段处理后半个warp。
估计pascal架构也采用这种方式。我大致猜测一下其中的细节,希望指正(我对编译形成的汇编指令不熟悉),第一个阶段处理前半个warp,16个8B,然后某种指令拆解成32个4B,对应32个bank,每个bank存储4B的数据,这里就只需进行一次transaction,同理后半个warp。这就是我所理解profiler为什么得出每个warp进行2次transaction的原因。那么推测在load过程中,也会前半个warp读取32*4B的数据,然后指令组合成16个8B数据。
在这样的基础上,我做了个试验,试验问题如下:
我要处理256组由4个double类型组成的数据(例如 struct{double x,y,z,w}我简化成了double数组),每个线程处理一组数据,我希望使用shared memory达到合并访问的目的。
第一次尝试定义如下:
__shared__ double a[32][32+1];//每行存32/4=8组数据,+1是为了填充
使用数据方式为
double rowid = threadIdx.x/8;
double colid = (threadIdx.x - rowid*

*4;
double x=a[rowid][colid+0];
double y=a[rowid][colid+1];
double z=a[rowid][colid+2];
double w=a[rowid][colid+3];
此实profiler检测到了bank conflict
我做了第二次尝试:
__shared__ double a[64][16+1];//每行存16/4=4组数据,+1是为了填充
使用数据方式为
double rowid = threadIdx.x/4;
double colid = (threadIdx.x - rowid*4)*4;
double x=a[rowid][colid+0];
double y=a[rowid][colid+1];
double z=a[rowid][colid+2];
double w=a[rowid][colid+3];
此实profiler没有检测到了bank conflict
此结果验证了前面的理解。