使用cuda 非32位数据的共享内存使用方法

  • 7 replies
  • 1288 views
使用cuda 非32位数据的共享内存使用方法
« 于: 十一月 12, 2020, 03:59:01 pm »
共享内存的数据类型通常为float,因为某种特殊需求我需要读取double的数据
__shared__ double a[256];
软硬件环境:计算力6.1的pascal 架构显卡 cuda 10.1
问题:
1. 对于一个block大小为256的块,每个线程直接访问对应数组的元素(如下),会出现bank conflict吗?
     a[threadIdx.x]=1.0;
2.如果会出现bank conflict,应该怎么处理?
3.如果不会,按照手册里说maxwell架构之后的,共享内存的bank size 都是4字节,他是通过什么方式避免了bank conflict?
4.不同架构下,处理__shared__ double a[256];这种数据访问的处理方式有什么不同?

Re: 使用cuda 非32位数据的共享内存使用方法
« 回复 #1 于: 十一月 12, 2020, 06:38:39 pm »
共享内存的数据类型通常为float,因为某种特殊需求我需要读取double的数据
__shared__ double a[256];
软硬件环境:计算力6.1的pascal 架构显卡 cuda 10.1
问题:
1. 对于一个block大小为256的块,每个线程直接访问对应数组的元素(如下),会出现bank conflict吗?
     a[threadIdx.x]=1.0;
2.如果会出现bank conflict,应该怎么处理?
3.如果不会,按照手册里说maxwell架构之后的,共享内存的bank size 都是4字节,他是通过什么方式避免了bank conflict?
4.不同架构下,处理__shared__ double a[256];这种数据访问的处理方式有什么不同?
个人愚见:来源https://blog.csdn.net/wd1603926823/article/details/78326570 若有不对请指出,相互学习
1、现在的卡一般都是warp访问,32个banks,每个bank内32bit,所以我假设你的卡也是如此。那么会出现conflict,因为数据是64bit占2个banks。也就是a[0]放在bank0和bank1简称b0、b1;a[1]--b2、b3;依次类推,a[15]在b30、b31;a[16]在b0、b1,其实类推到此时就可以看到线程0和线程16即不同的线程访问了同一个bank0,拟继续类推会发现更严重的conflict。你的线程0、16、32、64、128即5个不同线程都访问了bank0。按照conflict的定义:不同线程访问同一个bank会出现conflict,所以此处存在conflict。
2、怎么处理,我刚刚推了一下,可以通过补位的方式,至于补多少,你可以看完这个网址上的介绍自己推。假设每16个数补4个无意义的数extend,那么也就是a[0]放在bank0和bank1简称b0、b1;a[1]--b2、b3;依次类推,a[15]在b30、b31;extend[0]在b0,b1;extend[1]在b2、b3以此类推;a[16]在b8、b9;a[17]在b10、b11...可以看到至少目前没有发生conflict,后面你自己慢慢推吧。其实你自己像这个网址中这样用一张草稿推更简单明了。
3、4、前面的推明白了,这两点也知道了

Re: 使用cuda 非32位数据的共享内存使用方法
« 回复 #2 于: 十一月 12, 2020, 08:29:13 pm »
个人愚见:来源https://blog.csdn.net/wd1603926823/article/details/78326570 若有不对请指出,相互学习
1、现在的卡一般都是warp访问,32个banks,每个bank内32bit,所以我假设你的卡也是如此。那么会出现conflict,因为数据是64bit占2个banks。也就是a[0]放在bank0和bank1简称b0、b1;a[1]--b2、b3;依次类推,a[15]在b30、b31;a[16]在b0、b1,其实类推到此时就可以看到线程0和线程16即不同的线程访问了同一个bank0,拟继续类推会发现更严重的conflict。你的线程0、16、32、64、128即5个不同线程都访问了bank0。按照conflict的定义:不同线程访问同一个bank会出现conflict,所以此处存在conflict。
2、怎么处理,我刚刚推了一下,可以通过补位的方式,至于补多少,你可以看完这个网址上的介绍自己推。假设每16个数补4个无意义的数extend,那么也就是a[0]放在bank0和bank1简称b0、b1;a[1]--b2、b3;依次类推,a[15]在b30、b31;extend[0]在b0,b1;extend[1]在b2、b3以此类推;a[16]在b8、b9;a[17]在b10、b11...可以看到至少目前没有发生conflict,后面你自己慢慢推吧。其实你自己像这个网址中这样用一张草稿推更简单明了。
3、4、前面的推明白了,这两点也知道了


感谢奈奈同学工作之余, 百忙之中的回复.

楼主的代码(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写于祖国的边疆




Re: 使用cuda 非32位数据的共享内存使用方法
« 回复 #3 于: 十一月 13, 2020, 10:08:47 am »

感谢奈奈同学工作之余, 百忙之中的回复.

楼主的代码(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写于祖国的边疆
好久不见。
如果如你所说“4B是按照warp来的, 8B按照half warp来的, 而16B按照1/4 warp来访问”那么楼主的例子的确不会出现conflict,因为16个线程内的确没有不同线程访问同一个bank。是我轻率了,我还以为现在的卡都是warp访问。看来理解得还不够透彻才导致了浅薄的推导。虽然我的卡是7.5的,但我也会抽空验证一下你的猜测。 
另外,再次感叹日月如梭,看下时间都过了三年,向你们请教的日子让我成长很多,包括现在重新学习CUDA,回看当时的记录仍旧有新的感慨和领悟。现在自己也会遇到问题,会查会问,无果时会自己推论、验证、思考反复循环,过程艰难又快乐。依旧感谢你们当初的指导。

Re: 使用cuda 非32位数据的共享内存使用方法
« 回复 #4 于: 十一月 13, 2020, 11:24:11 am »

感谢奈奈同学工作之余, 百忙之中的回复.

楼主的代码(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*8)*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
此结果验证了前面的理解。

Re: 使用cuda 非32位数据的共享内存使用方法
« 回复 #5 于: 十一月 13, 2020, 11:28:22 am »
好久不见。
如果如你所说“4B是按照warp来的, 8B按照half warp来的, 而16B按照1/4 warp来访问”那么楼主的例子的确不会出现conflict,因为16个线程内的确没有不同线程访问同一个bank。是我轻率了,我还以为现在的卡都是warp访问。看来理解得还不够透彻才导致了浅薄的推导。虽然我的卡是7.5的,但我也会抽空验证一下你的猜测。 
另外,再次感叹日月如梭,看下时间都过了三年,向你们请教的日子让我成长很多,包括现在重新学习CUDA,回看当时的记录仍旧有新的感慨和领悟。现在自己也会遇到问题,会查会问,无果时会自己推论、验证、思考反复循环,过程艰难又快乐。依旧感谢你们当初的指导。

很感谢您的指导

我做过一点测试,能验证屠戮人神的说法,结果在4楼。

Re: 使用cuda 非32位数据的共享内存使用方法
« 回复 #6 于: 十一月 13, 2020, 01:24:13 pm »
感谢奈奈与屠戮人神 两位前辈的不吝指导。

我做了上机测试,确实如同屠戮人神所说,没有出现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*8)*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
此结果验证了前面的理解。

感谢帮我们测试实验. 感谢提供了NV的GTC上的资料的说法, 我看到了.

有几个问题, 一个是自带的内置类型(例如float4, double2)都有__align__修饰, 自己定义的看似一样的东西可能没有(例如你定义了4个a,b,c,d的float的struct). 二是最大的硬件指令的后缀有.128的(16B), 超过这个大小就会被拆分, 无论是否有align.

继续测试一下.128的(16B的)是否按1/4的warp?

Re: 使用cuda 非32位数据的共享内存使用方法
« 回复 #7 于: 十一月 13, 2020, 01:36:19 pm »
好久不见。
如果如你所说“4B是按照warp来的, 8B按照half warp来的, 而16B按照1/4 warp来访问”那么楼主的例子的确不会出现conflict,因为16个线程内的确没有不同线程访问同一个bank。是我轻率了,我还以为现在的卡都是warp访问。看来理解得还不够透彻才导致了浅薄的推导。虽然我的卡是7.5的,但我也会抽空验证一下你的猜测。 
另外,再次感叹日月如梭,看下时间都过了三年,向你们请教的日子让我成长很多,包括现在重新学习CUDA,回看当时的记录仍旧有新的感慨和领悟。现在自己也会遇到问题,会查会问,无果时会自己推论、验证、思考反复循环,过程艰难又快乐。依旧感谢你们当初的指导。

对了, 你正好有7.5的卡的话, 测试一下7.5上"诡异"的shared memory行为. 7.x的卡, 很多方面的测试都看起来像是只有16个shared memory的banks每SM, 而不是32个. 这点和NV的文档不符合(NV说从2.0+的计算能力以后, 都是32个banks/SM).

注意NV在计算能力8.6的文档中, 变相的从另外一个角度, 多多少少的承认了这一点. 如果这是真的, 则是最大的虚假宣传(7.5)丑闻.

我最近没有时间, 欢迎你们来
« 最后编辑时间: 十一月 13, 2020, 01:37:32 pm 作者 屠戮人神 »