Cuda gridDim 和 blockDim答案

【问题标题】：Cuda gridDim and blockDimCuda gridDim 和 blockDim
【发布时间】：2013-05-13 05:25:05
【问题描述】：

我知道blockDim 是什么，但我对gridDim. Blockdim 给出块的大小有疑问，但gridDim 是什么？在互联网上，它说gridDim.x 给出了 x 坐标中的块数。

我怎么知道blockDim.x * gridDim.x 提供了什么？

我如何知道 x 行中有多少个 gridDim.x 值？

例如，考虑下面的代码：

int tid = threadIdx.x + blockIdx.x * blockDim.x;
double temp = a[tid];
tid += blockDim.x * gridDim.x;

while (tid < count)
{
    if (a[tid] > temp)
    {
       temp = a[tid];
    }
    tid += blockDim.x * gridDim.x;
}

我知道 tid 以 0 开头。然后代码有 tid+=blockDim.x * gridDim.x。操作后tid 现在是什么？

【问题讨论】：

标签： cuda

【解决方案1】：

blockDim.x,y,z 给出块中的线程数，在特定方向
gridDim.x,y,z 给出网格中的块数，在特定方向
blockDim.x * gridDim.x 给出网格中的线程数（在本例中为 x 方向）

块和网格变量可以是 1、2 或 3 维的。处理一维数据时通常只创建一维块和网格。

在 CUDA 文档中，这些变量被定义为 here

特别是，当 x 维度 (gridDim.x*blockDim.x) 中的总线程小于我希望处理的数组的大小时，通常的做法是创建一个循环并拥有线程网格在整个数组中移动。在这种情况下，在处理一次循环迭代之后，每个线程必须移动到下一个未处理的位置，由tid+=blockDim.x*gridDim.x; 给出。实际上，整个线程网格正在跳过一维数据数组，一个网格-一次宽度。这个主题，有时称为“网格跨步循环”，在blog article 中有进一步讨论。

您可能需要考虑参加NVIDIA webinar page 上提供的一些介绍性 CUDA 网络研讨会。比如这2个：

GPU Computing using CUDA C – An Introduction (2010) 简介使用 CUDA C 进行 GPU 计算的基础知识。通过代码示例的演练进行说明。没有先前的 GPU 计算需要经验
GPU 计算使用 CUDA C – 高级 1 (2010) 第一级优化技术，例如全局内存优化，以及处理器利用率。将使用真实代码说明概念示例

如果您想更好地理解这些概念，花 2 个小时会很合适。

here 详细介绍了网格跨步循环的一般主题。

【讨论】：

@alrikai 当您删除它时，我正在为您的答案添加一些 cmets。（后来你重新发布了，我猜。）你是第一个，你的答案很好。
是的，我写到一半时不小心把它贴出来了（哎呀）

【解决方案2】：

转述自CUDA Programming Guide：

gridDim：此变量包含网格的尺寸。

blockIdx：此变量包含网格内的块索引。

blockDim：这个变量，包含块的尺寸。

threadIdx：此变量包含块内的线程索引。

您似乎对 CUDA 的线程层次结构有些困惑；简而言之，对于内核，将有 1 个网格（我总是将其可视化为 3 维立方体）。它的每个元素都是一个块，因此声明为dim3 grid(10, 10, 2); 的网格将总共有 10*10*2 个块。反过来，每个块都是线程的 3 维立方体。

话虽如此，通常只使用块和网格的 x 维度，这看起来就像您问题中的代码正在做的那样。如果您使用一维数组，这一点尤其重要。在这种情况下，您的 tid+=blockDim.x * gridDim.x 行实际上是网格中每个线程的唯一索引。这是因为您的 blockDim.x 将是每个块的大小，而您的 gridDim.x 将是块的总数。

所以如果你启动一个带参数的内核

dim3 block_dim(128,1,1);
dim3 grid_dim(10,1,1);
kernel<<<grid_dim,block_dim>>>(...);

那么在你的内核中有threadIdx.x + blockIdx.x*blockDim.x，你实际上会拥有：

threadIdx.x range from [0 ~ 128)

blockIdx.x range from [0 ~ 10)

blockDim.x equal to 128

gridDim.x equal to 10

因此，在计算 threadIdx.x + blockIdx.x*blockDim.x 时，您的值将在 [0, 128) + 128 * [1, 10) 定义的范围内，这意味着您的 tid 值的范围为 {0, 1, 2, ..., 1279}。当您想将线程映射到任务时，这很有用，因为这为内核中的所有线程提供了唯一标识符。

但是，如果你有

int tid = threadIdx.x + blockIdx.x * blockDim.x;
tid += blockDim.x * gridDim.x;

那么您基本上将拥有：tid = [0, 128) + 128 * [1, 10) + (128 * 10)，并且您的 tid 值范围为 {1280, 1281, ..., 2559} 我不确定这在哪里相关，但这完全取决于您的应用程序以及如何将线程映射到数据。这种映射对于任何内核启动都是非常重要的，您是决定应该如何完成它的人。当你启动你的内核时，你指定了网格和块的维度，你是必须强制映射到内核中数据的人。只要不超过硬件限制（对于现代卡，每个块最多可以有 2^10 个线程，每个网格最多可以有 2^16 - 1 个块）

【讨论】：

具体的例子很有帮助，谢谢。很多人只是重复gridDim、blockIdx等的定义，但例子很重要。
对不起，先生，但在最后一句你说过每个块最多可以有 2^10 个线程，每个线程最多可以有 2^16 - 1 个块，但应该不是吗：每个块最多可以有 2^10 个线程，每个 grid 最多可以有 2^16 - 1 个块

【解决方案3】：

在这个源代码中，我们甚至有 4 个线程，内核函数可以访问所有 10 个数组。怎么样？

#define N 10 //(33*1024)

__global__ void add(int *c){
    int tid = threadIdx.x + blockIdx.x * gridDim.x;

    if(tid < N)
        c[tid] = 1;

    while( tid < N)
    {
        c[tid] = 1;
        tid += blockDim.x * gridDim.x;
    }
}

int main(void)
{
    int c[N];
    int *dev_c;
    cudaMalloc( (void**)&dev_c, N*sizeof(int) );

    for(int i=0; i<N; ++i)
    {
        c[i] = -1;
    }

    cudaMemcpy(dev_c, c, N*sizeof(int), cudaMemcpyHostToDevice);

    add<<< 2, 2>>>(dev_c);
    cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost );

    for(int i=0; i< N; ++i)
    {
        printf("c[%d] = %d \n" ,i, c[i] );
    }

    cudaFree( dev_c );
}

为什么我们不创建 10 个线程例如）add<<<2,5>>> or add<5,2>>> 因为我们必须创建相当少量的线程，如果 N 大于 10 ex) 33*1024。

此源代码是这种情况的示例。数组是 10，cuda 线程是 4。如何仅通过 4 个线程访问所有 10 个数组。

查看cuda详细信息中threadIdx、blockIdx、blockDim、gridDim的含义页面。

在本源代码中，

gridDim.x : 2    this means number of block of x

gridDim.y : 1    this means number of block of y

blockDim.x : 2   this means number of thread of x in a block

blockDim.y : 1   this means number of thread of y in a block

我们的线程数是 4，因为 2*2(blocks * thread)。

在添加内核函数中，我们可以访问线程的0、1、2、3索引

->tid = threadIdx.x + blockIdx.x * blockDim.x

①0+0*2=0

②1+0*2=1

③0+1*2=2

④1+1*2=3

如何访问索引 4、5、6、7、8、9 的其余部分。 while循环中有计算

tid += blockDim.x + gridDim.x in while

** 第一次调用内核 **

-1 循环：0+2*2=4

-2 循环：4+2*2=8

-3 循环：8+2*2=12（但这个值是假的，虽然出来了！）

** 内核的第二次调用 **

-1 循环：1+2*2=5

-2 循环：5+2*2=9

-3 循环：9+2*2=13（但这个值是假的，虽然出来了！）

** 第三次内核调用 **

-1 循环：2+2*2=6

-2 循环：6+2*2=10（但这个值是假的，虽然出来了！）

**第四次内核调用**

-1 循环：3+2*2=7

-2 循环：7+2*2=11（但这个值是假的，虽然出来了！）

所以，0、1、2、3、4、5、6、7、8、9的所有索引都可以通过tid值访问。

请参阅此页面。 http://study.marearts.com/2015/03/to-process-all-arrays-by-reasonably.html 我无法上传图片，因为声誉低。

【讨论】：

【解决方案4】：

首先，从CUDA official document看这个图Grid of thread blocks

通常，我们这样使用内核：

__global__ void kernelname(...){
    const id_x = blockDim.x * blockIdx.x + threadIdx.x;
    const id_y = blockDim.y * blockIdx.y + threadIdx.y;
    ...
}

// invoke kernel
// assume we have assigned the proper gridsize and blocksize
kernelname<<<gridsize, blocksize>>>(...)

一些变量的含义：

gridsize每个网格的块数，对应gridDim

blocksize每个块的线程数，对应blockDim

threadIdx.x 在 [0, blockDim.x) 中变化

blockIdx.x 在 [0, gridDim.x) 中变化

所以，当我们有threadIdx.x 和blockIdx.x 时，让我们尝试计算x 方向 的索引。根据figure，blockIdx.x决定你是哪个block，threadIdx.x在给定block位置的时候决定你是哪个线程。因此，我们有：

which_blk = blockDim.x * blockIdx.x; // which block you are
final_index_x = which_blk + threadIdx.x; // based on the given block, we can have the final location by adding the threadIdx.x

即：

final_index_x = blockDim.x * blockIdx.x + threadIdx.x;

与上面的示例代码相同。

同样，我们可以分别得到y或z方向的索引。

如我们所见，我们通常不会在我们的代码中使用gridDim，因为这个信息是作为blockIdx的范围来执行的。相反，我们必须使用blockDim，尽管此信息是作为threadIdx 的范围执行的。我在上面一步一步展示的原因。

我希望这个答案可以帮助您解决困惑。

【讨论】：