block个数与处理数据量有关吗？

[sdtougao]

各位好，
我现在在做基于GPU的优化算法，遇到一些基本的问题，请高手指教。
（1）如何选择thread、block个数？
感觉很多时候都是凭感觉在试，不知道有没有什么通用的规则。
例如，要处理4MB=4*1024*1024数据，选择thread个数为256，那么是否意味着block的个数一定要选择为(4*1024*1024)/256=16384 吗？如果不是应该怎么选？
（2）block、grid维数的具体含义？
我的GPU是GTX260+，通过Device Query可知，block的最大尺寸是：512*512*64，不知道这到底是指什么？是说该block的X方向最多包含512个thread，Y方向最多包含512个thread，Z方向最多包含64个thread吗？

都是一些最基本的概念问题，一直比较困扰，请牛人们多指点，在线等，谢谢！

[悟之思语]

1、warp该是block中的单位.不同block中的线程不会组成一个warp的.
2、不同核的分配是按block分配的.同一block的线程在一个sp上执行.不同的block才可能在不同的sp上运行.

不同线程的不一致较大?有根据threadid的if判断?

[sdtougao]

例如，threadperblock = 256;
datalen = 1<<22; //4M
blockpergrid = 64;

在kernel中
tid = threadId.x;
i = blockId.x * blockDim.x + threadId.x;
if( i<datalen ) s[tid] = g_idata[i];

问题是，tid的最大值是256，当i大于256时，是否程序会把g_idata[i]的值自动付给下一个block的tid呢？

[ukyolei]

[Quote=引用 2 楼 sdtougao 的回复:]
例如，threadperblock = 256;
datalen = 1<<22; //4M
blockpergrid = 64;

在kernel中
tid = threadId.x;
i = blockId.x * blockDim.x + threadId.x;
if( i<datalen ) s[tid] = g_idata[i];

问题是，tid的最大值是256，当i……
[/Quote]
会的

[sdtougao]

[Quote=引用 3 楼 ukyolei 的回复:]
引用 2 楼 sdtougao 的回复:
例如，threadperblock = 256;
datalen = 1<<22; //4M
blockpergrid = 64;

在kernel中
tid = threadId.x;
i = blockId.x * blockDim.x + threadId.x;
if( i<datalen ) s[tid] = g_idata[i];……
[/Quote]
能否具体说明一下，在给定数据量、线程个数时如何确定block个数吗

[无心人_过过小日子]

[Quote=引用 4 楼 sdtougao 的回复:]
能否具体说明一下，在给定数据量、线程个数时如何确定block个数吗
[/Quote]

block尺寸(每block中thread数)一般取128/256等32的倍数,
而block数(grid尺寸)取比要处理数据总数(并行处理数)/block尺寸,并向上取整以保证block数*block尺寸>=要处理数据总数.

[sdtougao]

[Quote=引用 5 楼 l7331014 的回复:]
而block数(grid尺寸)取比要处理数据总数(并行处理数)/block尺寸,并向上取整以保证block数*block尺寸>=要处理数据总数.
[/Quote]
请问，您所指的block数（即grid尺寸）是指三个维度上的吗，
还是上面的例子，当处理数据总数是4M时，假设block尺寸（即每block中thread数）=256，那么block数至少为16384，那么应该怎么具体选block数呢？是dim3 grid（256,64,1）这种，只要使block三个维度乘积大于16384就行吗

[无心人_过过小日子]

[Quote=引用 6 楼 sdtougao 的回复:]
引用 5 楼 l7331014 的回复:
而block数(grid尺寸)取比要处理数据总数(并行处理数)/block尺寸,并向上取整以保证block数*block尺寸>=要处理数据总数.

请问，您所指的block数（即grid尺寸）是指三个维度上的吗，
还是上面的例子，当处理数据总数是4M时，假设block尺寸（即每block中thread数）=256，那么block数至少为16384，……
[/Quote]

总数够就行.

[shiyongshuai]

可以分几步：
step1. 确定block的维数，即线程数（threadsPerBlock）。一般这样的问题要根据你做怎么样的计算进行处理，对显存访问较多的线程，block中线程数应该尽量多，这样可以掩盖访存的延迟；如果单线程的计算量很大，建议充分利用共享缓存，线程数一般存在一个最佳的中间值，需要测试才能确定，例如做MD5时，我在GTX9800上线程数128时是最快的（只使用共享缓存）。注意：建议尽量使用一维的block，这个比二维快些。

step2. 确定总的计算量(neededThreads)，也就是你的工程总共需要多少个线程才能计算完？
step3. 确定block个数（blocksNum）。
blocksNum = neededThreads / threadsPerBlock
如果你的计算量很大，blocksNum超过了硬件允许的单维数量限制（比如GTX480为65536），拆分成二维
（grid.x, grid.y）,其中grid.x = 16 * x。(grid.x * grid.y) < blocksNum < (grid.x * grid.y)

[wobuzhidaotashishui]

[Quote=引用 1 楼 jbjwpzyl3611421 的回复:]

1、warp该是block中的单位.不同block中的线程不会组成一个warp的.
2、不同核的分配是按block分配的.同一block的线程在一个sp上执行.不同的block才可能在不同的sp上运行.

不同线程的不一致较大?有根据threadid的if判断?
[/Quote]同一个block的线程在一个SM上，所以可以在8个SP上运行。而且SM安排每个warp中的线程进入8个sp已达到并行效果。

[wobuzhidaotashishui]

[Quote=引用 2 楼 sdtougao 的回复:]

例如，threadperblock = 256;
datalen = 1<<22; //4M
blockpergrid = 64;

在kernel中
tid = threadId.x;
i = blockId.x * blockDim.x + threadId.x;
if( i<datalen ) s[tid] = g_idata[i];

问题是，tid的最大值是256，当……
[/Quote]if( i<datalen )没有用处吧，i最大也就是16384，绝对小于4M。

[wobuzhidaotashishui]

[Quote=引用 6 楼 sdtougao 的回复:]

引用 5 楼 l7331014 的回复:
而block数(grid尺寸)取比要处理数据总数(并行处理数)/block尺寸,并向上取整以保证block数*block尺寸>=要处理数据总数.

请问，您所指的block数（即grid尺寸）是指三个维度上的吗，
还是上面的例子，当处理数据总数是4M时，假设block尺寸（即每block中thread数）=256，那么block数至少为16384……
[/Quote]随便，总线程比总数据多，加上边界条件，可以让一个线程处理一个数据。比总数据少，可以线程里加循环处理多个数据，循环步长最好32的倍数，因为一个warp是32。

[ddweidong]

[Quote=引用 8 楼 shiyongshuai 的回复:]

可以分几步：
step1. 确定block的维数，即线程数（threadsPerBlock）。一般这样的问题要根据你做怎么样的计算进行处理，对显存访问较多的线程，block中线程数应该尽量多，这样可以掩盖访存的延迟；如果单线程的计算量很大，建议充分利用共享缓存，线程数一般存在一个最佳的中间值，需要测试才能确定，例如做MD5时，我在GTX9800上线程数128时是最快的（只使用共享缓存）。注意：……
[/Quote]

将的很透彻，前几天在这个问题上摸索了好久...