我正在尝试将我的网格从一维网格扩展到二维网格。有没有办法做到这一点?
这是我当前的代码:
int idx = threadIdx.x + blockDim.x * blockIdx.x;
在#include
列表中,我有以下定义:
#define BLOCKS_PER_GRID 102
#define THREADS_PER_BLOCK 1024
我正在尝试将我的网格从一维网格扩展到二维网格。有没有办法做到这一点?
这是我当前的代码:
int idx = threadIdx.x + blockDim.x * blockIdx.x;
在#include
列表中,我有以下定义:
#define BLOCKS_PER_GRID 102
#define THREADS_PER_BLOCK 1024
假设您希望每个块有 1024 个线程,则可以轻松地将块重塑为 2D。
32 x 32 = 1024;
所以你的块看起来像这样:
dim3 Block(32,32); //1024 threads per block. Will only work for devices of at least 2.0 Compute Capability.
我不知道您的确切要求是什么,但通常块数不是固定的(正如您在宏中定义的那样)。块的数量取决于输入数据的大小,因此网格可以动态缩放。
与你的情况一起,你有很多选择,但你的网格最接近的最佳尺寸是17 x 6
or 6 x 17
。
dim3 Grid(17,6);
现在您可以使用以下参数调用内核:
kernel<<<Grid,Block>>>();
在内核内部,线程的二维索引计算如下:
int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
或者,如果您遵循行/列约定而不是 x/y,则:
int row = blockIdx.y * blockDim.y + threadIdx.y;
int column = blockIdx.x * blockDim.x + threadIdx.x;
您还可以拥有一维线程块的 2D 网格,以绕过每个网格维度 65535 个块的限制(对于 cc3.0 之前的设备)。这可能是一种更简单的方法,可以将根本上的一维问题扩展到极限之外,而无需为数据引入二维数组表示。
假设我们有一个DATA_ELEMENTS
参数定义为内核将处理的元素数量(每个线程一个元素)。如果DATA_ELEMENTS
大于 65535*1024,那么如果每个线程只处理 1 个元素,则无法使用一维网格处理它们。
你可以让你的THREADS_PER_BLOCK
参数保持不变。您在内核中的线程索引计算将更改为:
int idx = threadIdx.x + (blockDim.x * ((gridDim.x * blockIdx.y) + blockIdx.x));
您需要确保使用以下内容调整内核计算:
if (idx < DATA_ELEMENTS){
(kernel code)
}
您的网格尺寸如下:
dim3 grid;
if (DATA_ELEMENTS > (65535*THREADS_PER_BLOCK)){ // create a 2-D grid
int gridx = 65535; // could choose another number here
int gridy = ((DATA_ELEMENTS+(THREADS_PER_BLOCK-1))/THREADS_PER_BLOCK)/gridx;
if ((((DATA_ELEMENTS+(THREADS_PER_BLOCK-1))/THREADS_PER_BLOCK)%gridx) != 0) gridy++;
grid.x=gridx;
grid.y=gridy;
grid.z=1;
}
else{ // create a 1-D grid
int gridx = (DATA_ELEMENTS+(THREADS_PER_BLOCK-1))/THREADS_PER_BLOCK;
grid.x=gridx;
grid.y=1;
grid.z=1;
}
你会启动你的内核:
kernel<<<grid, THREADS_PER_BLOCK>>>(...);
解决此类问题的另一种方法是创建某个维度的一维网格(假设网格中的线程总数为NUM_THREADS_PER_GRID
),并让每个线程处理数据元素数组中的多个元素,使用类似 for 循环或 while 循环的东西:
while (idx < DATA_ELEMENTS) {
(code to process an element)
idx += NUM_THREADS_PER_GRID
}
我喜欢罗伯特上面的解决方案。我对他的第一个解决方案的唯一评论是,gridx
当DATA_ELEMENTS > (65535*THREADS_PER_BLOCK)
. 原因是,如果数据元素的数量是65535*THREADS_PER_BLOCK + 1
,并且gridx
是65535
,则65535*2*THREADS_PER_BLOCK
启动,所以几乎一半的线程将什么都不做。如果gridx
更小,那么什么都不做的线程就会更少。