一个block里面分配多少的线程效果比较好

加速计算 CUDA
1

我想问一下，有没有一些经验告诉大家，比如说规模是500，一个block里面分配32个线程比较好；规模是1000，一个block里面分配128个线程比较好，。。。。。。我只是打个比方而已，大家有什么经验或者建议吗？

2

这是一个相关很多因素的问题，一般的典型值是256或者512这样。

祝您编码愉快~

3

好的。我发现用我自己写的程序测试，当矩阵是10050的时候，一个block分配32个线程比分配512个线程要快一些；当矩阵是600121的时候，同样是一个block分配32个线程比分配512个线程要快一些。我想是不是因为我的数据太小了，因为仅仅是行并行嘛，即使是600的时候，用2个block就够了，而32的时候，可以使用更多的block，这样并行效果是不是好一些。

4

一个block选取多少threads，一般是为了occupancy最大化选定的。

一个block选取32个threads，在绝大多数情况下都是偏少的，因为一个SM上同时加载的block数量有限，按照programming guide的附录中Table-F2的数据，在fermi上，一个SM上最大resident blocks数量是8，而最大的resident warps是48（也就是最大resident threads为1536）。那么要尽量多地加载threads，那么每个block最少需要48/8=6个warps，也就是32*6=192threads。

当一个block线程数量大于等于192的时候，能保证occupancy不因为同时resident block数量的限制而下降。

同时，occupancy还受到kernel复杂程度的影响，比如复杂的kernel占用了过多的寄存器，那么可能一个SM上只加载了少量的线程，就已经受到寄存器数量限制，而无法继续加载线程。shared memory等限制因素类似。

虽然确实有通过降低occupancy，来优化性能的例子，但是一般的习惯和建议是通过提高occupancy来提高性能。（提高occupancy可以更方便地掩盖各种延迟，以及尽可能饱满地使用计算单元。）

至于您的测试结果，我相信是真实的，但是影响因素可能在其他方面。

欢迎莅临cudazone，祝您编码愉快~

5

好的。又学习了一些技巧~

6

CUDA_Occupancy_Calculator.xls
这个EXCEL里的东西可以用来根据你的GPU计算能力计算出占有率，寄存器分配，shared memory分配情况，可以在你执行KERNEL以前做一个简单的理论分析，也可以用visual profile来分析你的内核，可以根据这些结果来设置你的block大小

7

好的。等我看看。我在上一个帖子“如何打破memory bandwidth limited”的最后一个问题，能提些建议吗？或者给我一些关于合并访问的资料，我认为解决合并访问才是我那几个KERNEL加速的关键。

8 
__global__ void KernelV2(int *dev_ZZ,double *dev_V,double *dev_VKK)
{
   int k=blockIdx.x*512+threadIdx.x;
   int i=blockIdx.y;
   if((i<XROW)&&(k<XROW))
   {
   for(int j=0;j<(XCOL*(XCOL-1)/2);j++)
   {
   dev_V[i*XROW+k]=dev_V[i*XROW+k]+dev_ZZ[i+j*XROW]*dev_ZZ[k+j*XROW]*dev_VKK[j];
   }
   }
}

拿你以前的代码来看下吧，dev_v[i*XROW+k]使得warp内线程连续读取，这个没问题
dev_zz[i+j*XROW]是一个广播的过程，就是说warp内所有的线程都会得到这个值，
dev_zz[k+j*XROW]使得warp内线程连续读取，这个也没问题
dev_vkk[j]也是一个广播的过程
现在有个问题就是global memory无法广播，所以dev_zz[i+j*XROW]和dev_vkk[j]的访问可能是性能瓶颈

解决方法：尝试用texture memory，把dev_zz和dev_vkk绑定到纹理内存（就用线性纹理一维的就可以），然后用tex1Dfetch来获取它的值（因为dev_zz和dev_vkk不需要写入，因此可以用纹理）

9

好的，因为没有用过Texture memory ,所以暂时没有往那方面想，网上比较多的是矩阵相乘以及其优化的例子，所以我首先想到的shared memory，但是fermi构架下，shared memory是48KB或16KB，如果数据大的话，不能全部加载进去，肯定得用其他办法，想想就头疼。纹理内存的大小是多大呀？

10

你可以先用texture尝试一下，对于texture的广播方式我也没搞明白，从programming guide上看只说了打包的数据可以在1次操作中广播到多个独立变量中，我也没搞清楚这个广播过程是什么意思…不过这种非合并访问的情况,texture的性能要强于global memory

11

纹理内存大小没有限制，只是限制了纹理三维的大小。推荐你可以用cudaArray来绑定二维纹理内存试试看，因为纹理内存的地址计算时在kernel外完成的，用一维的话你还得加上i*XROW+k这种计算

12

好的，我去GUIDE上看看有没有详细点儿的说明