CUDA:grid和block维度设计

[_梦魇花葬]

按照CUDA的执行模型，grid中的各个block会被分配到GPU的各个SM中执行。下面我们给出一些建议，如何确定合适的Grid和block尺寸。在设计时，应该优先考虑block的尺寸，而grid的尺寸一般来说越大越好。
在Tesla架构GPU的每个SM中，至少有6个active warp才能有效地隐藏流水线延迟。此外，如果所有的active warp都来自同一block，当这个block中的线程进行存储器访问或者同步时，执行单元就会闲置。基于以上原因，最好让每个SM上拥有至少2个active block。
一个SM上的active warp和active block数量计算方法如下：
（1）确定每个SM使用的资源数量
使用nvcc的--keep编译选项，或者在.cu编译规则（cuda build rule）中选择保留中间文件，得到.cubin文件。用写字板打开.cubin文件，在每个内核函数的开始部分，可以看到：

lmem=0
smem=256
reg =8

其中，lmem和reg分别代表内核函数中每个线程使用的local memory数量和寄存器数量，smem代表每个block使用的shared memory数量。也就是说：这个内核函数的每个线程使用了0Byte local memory，8个寄存器文件（每个寄存器文件的大小是32bit）；每个block使用了256Byte的shared memory。
（2）根据硬件确定SM上的可用资源
可以用SDK中的deviceQuery获得每个SM钟的资源。要注意到：在程序编译时，要使目标代码和目标硬件版本与实际使用的硬件一致（使用-arch、-gencode和-code编译选项）。不同的架构，限制也是不一样的。
（3）计算每个block使用的资源，并确定active block和active warp数量
假设每个block中有64个线程，每个block使用256 Byte shared memory，8个寄存器文件，那么有：
1、每个block使用的shared memory是：256Byte
2、每个block使用的寄存器文件数量：8*64=512
3、每个block中的warp数量：64/32=2
然后根据每个block使用的资源，就可以计算出由每个因素限制的最大active block数量。这里，假设在G80/G92 GPU中运行：
1、由shared memory数量限制的active block数量：16384/256=64
2、由寄存器数量限制的active block数量：8192/512=16
3、由warp数量限制的active block数量：24/2=12
4、每个SM中的最大active block数量：8
这样就确定了grid中active block的数量。各位小伙伴，可以去试一试！

[jswanghaodz]

好东西看看fffffff

[longye1986]

收缴啦，非常好

[agan很努力]

牛！我学习学习啊@

[hugh_z]

learning

[cattpon]

好复杂的样子~

[hugh_z]

learning

[「已注销」]

学习了，谢谢分享！！！

[nettman]

[YCMyTot]

学习了，谢谢分享了!!

[okkk]

[u010185989]

想求教楼主个别的问题呀，为什么我写的cuda程序中 __syncthreads();显示未定义呀，其它的都能正常运行，就是这点怎么都解决不了，望楼主能抽空答疑解惑一下，多谢了， 话说，cuda sample的 __syncthreads();是可以正常使用的，但是各种配置，各种头文件添加都无济于事的感觉， 求教楼主呀！

[wyx_sky]

1、由shared memory数量限制的active block数量：16384/256=64 我记得shared memory一般是48kb，比如shared memory我使用7500bytes，这是所有的threads都用的啊，与threads数量无关。 版主可以对比着deviceQuery的回显进行说明吗？ 十分感谢！！

[wenqingkaituozhe]

学习了，谢谢楼主

[aries刘]

楼主 好厉害，现在正在学习，多谢啊

[wyx_sky]

引用 16 楼 wyx_sky 的回复:

1、由shared memory数量限制的active block数量：16384/256=64 我记得shared memory一般是48kb，比如shared memory我使用7500bytes，这是所有的threads都用的啊，与threads数量无关。 版主可以对比着deviceQuery的回显进行说明吗？ 十分感谢！！

没有具体测试过，但应该可行。 小规模数据测试，为了测试最佳计算性能，可以对block作为循环变量进行测试，测试计算用时。

[wyx_sky]

一开始不理解warp，所以看不懂楼主帖子。 在结合下面几篇文章后，总算明白。谢谢 http://www.360doc.com/content/13/0819/19/12278894_308350397.shtml# http://blog.csdn.net/poisonchry/article/details/16332369 http://www.cnblogs.com/dubing/archive/2011/06/21/2085742.html

[wyx_sky]

昨天的实验，发现楼主说的不对。 设备情况： Device 1: "Tesla K20m" CUDA Capability Major/Minor version number: 3.5 Total amount of global memory: 4800 MBytes (5032706048 bytes) (13) Multiprocessors, (192) CUDA Cores/MP: 2496 CUDA Cores Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) 代码情况： 每个block中有64个线程，每个block使用0 Byte shared memory，63个寄存器文件，那么有： 1、每个block使用的shared memory是：0Byte 2、每个block使用的寄存器文件数量：63*64=4032 3、每个block中的warp数量：64/32=2 1、由shared memory数量限制的active block数量：没有使用shared memory 2、由寄存器数量限制的active block数量：65536/63/64=16 3、由warp数量限制的active block数量：2048/32=64/2=32 4、每个SM中的最大active block数量：192 寄存器说明，3.5版本，编译选项加上--ptxas-options=-v，可以看见：ptxas info : Used 63 registers, 4248 bytes cumulative stack size, 1240 bytes cmem[0] 计算情况，单位是毫秒： num_blocks=16 num_threads=32 time=4899.6 num_blocks=16 num_threads=64 time=4285.0 num_blocks=32 num_threads=32 time=2530.7 num_blocks=32 num_threads=64 time=2270.1 num_blocks=64 num_threads=32 time=1349.7 num_blocks=64 num_threads=64 time=1247.1 num_blocks=96 num_threads=32 time=946.5 num_blocks=96 num_threads=64 time=908.3 num_blocks=128 num_threads=32 time=758.5 num_blocks=128 num_threads=64 time=732.2 num_blocks=160 num_threads=32 time=644.4 num_blocks=160 num_threads=64 time=640.0 num_blocks=192 num_threads=32 time=554.8 num_blocks=192 num_threads=64 time=563.5 num_blocks=224 num_threads=32 time=810.3 num_blocks=224 num_threads=64 time=781.5 num_blocks=256 num_threads=32 time=751.4 num_blocks=256 num_threads=64 time=796.1 num_blocks=288 num_threads=32 time=704.2 num_blocks=288 num_threads=64 time=729.2 num_blocks=320 num_threads=32 time=649.5 num_blocks=320 num_threads=64 time=659.6 发现根据楼主说的来计算，不知道是为什么。

[linxxx3]

"如果所有的active warp都来自同一block，当这个block中的线程进行存储器访问或者同步时，执行单元就会闲置" 这句话依据在哪？

[hello_hi_hi]

这是理论上的算法。