56

I get what blockDim is, but I have a problem with gridDim. Blockdim gives the size of the block, but what is gridDim? On the Internet it says gridDim.x gives the number of blocks in the x coordinate.

How can I know what blockDim.x * gridDim.x gives?

How can I know that how many gridDim.x values are there in the x line?

For example, consider the code below:

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;
}

I know that tid starts with 0. The code then has tid+=blockDim.x * gridDim.x. What is tid now after this operation?

4

4 回答 4

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

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

在 CUDA 文档中，这些变量在此处定义

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

您可能需要考虑参加NVIDIA 网络研讨会页面上提供的几个介绍性 CUDA 网络研讨会。例如，这两个：

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

如果您想更好地理解这些概念，这将花费 2 个小时。

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

于 2013-05-18T00:43:06.413 回答
55

摘自CUDA 编程指南：

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 个块）

于 2013-05-18T00:30:37.737 回答
1

在这个源代码中，我们甚至有 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 个线程 ex)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 loop: 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 我无法上传图片，因为声誉低。

于 2015-03-16T11:46:37.657 回答
0

首先，从CUDA官方文档看这张图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) 中变化

因此，当我们有和时，让我们尝试计算x 方向的索引。根据图中，判断你是哪个block，给定block位置的时候判断你是哪个线程。因此，我们有：threadIdx.xblockIdx.xblockIdx.xthreadIdx.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。我在上面一步一步展示的原因。

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

于 2020-03-25T04:53:03.777 回答