CUDA中的gridDim和blockDim

Question

我理解 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 是多少？

- ehah

Answer 1

• blockDim.x,y,z 分别表示块（block）在特定方向上的线程数。
• gridDim.x,y,z 分别表示网格（grid）在特定方向上的块数。
• blockDim.x * gridDim.x 表示网格中线程总数（这里是在 x 方向上）。

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

在 CUDA 文档中，这些变量在这里进行了定义。

特别地，在 x 维度上的总线程数 (gridDim.x*blockDim.x) 小于我要处理的数组大小时，通常会创建一个循环，并让线程网格遍历整个数组。在这种情况下，处理完一次循环迭代后，每个线程都必须移动到下一个未处理的位置，即 tid+=blockDim.x*gridDim.x;。实际上，整个线程网格以网格宽度为单位跳跃穿过 1-D 数组数据。这个主题有时被称为“网格步进循环”，在这个博客文章中有进一步讨论。

如果您想更好地了解这些概念，可以考虑参加介绍性的 CUDA 网络研讨会，例如前四个单元。如果您想更好地理解这些概念，那么这将是值得投入 4 小时的时间。

- Robert Crovella

---

也许应该提一下 threadIdx.x 和 blockIdx.x 的 x 方向是完全独立的，可以用于不同的目的。例如，x 和 y 可以组合使用。 - Sebastian

Answer 2

从CUDA编程指南改述：

gridDim：此变量包含网格的维度。

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

blockDim：此变量包含块的维度。

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

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

说到这里，通常只使用块和网格的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 的范围为 [0 ~ 128)

blockIdx.x 的范围为 [0 ~ 10)

blockDim.x 等于 128

gridDim.x 等于 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个块）。

- alrikai

---

具体的例子非常有帮助，谢谢。很多人只是重复解释gridDim、blockIdx等的定义，但例子是至关重要的。 - cmo

先生，不好意思打扰您，但在最后一句话中，您说“每个块最多可以有2^10个线程，每个线程块最多可以有2^16-1个块”，但是应该是：“每个块最多可以有2^10个线程，每个网格最多可以有2^16-1个块”，对吗？ - meJustAndrew

Answer 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>>>或add<5,2>>>？ 因为我们必须创建合理数量的线程，如果N大于10，例如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（块数*线程数）。

在添加内核函数时，我们可以访问线程的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 由于声望不足，我无法上传图片。

- jh Kim

Answer 4

首先，看一下CUDA官方文档中的线程块网格图

通常，我们使用内核函数如下：

__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 确定您所在的块，而 threadIdx.x 在给定块的位置时确定您所在的线程。因此，我们有：

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的范围执行了。我已经逐步解释了上述原因。

我希望这个答案能够帮助解决您的困惑。

- Peter Hou