问Cuda gridDim和blockDim

EN

• blockDim.x,y,z给出了块中的线程数，在特定的direction
• gridDim.x,y,z中给出了网格中的块数，在特定的direction
• 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网络研讨会。例如，这两个：

使用CUDA的

• GPU计算-简介(2010)使用CUDA的GPU计算基础知识介绍将通过代码示例演练来说明概念。使用CUDA Advanced 1 (2010)第一级优化技术的
• GPU计算无需任何先前的图形处理器计算经验，例如全局内存优化和处理器利用率。我们将使用真实的代码示例

来说明概念

如果你想更好地理解这些概念，2个小时是值得的。

网格跨度循环的一般主题在here中有一些详细的介绍。

转译自CUDA Programming Guide

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

blockIdx:此变量包含网格中的块索引。

blockDim:此变量，包含块的尺寸。

threadIdx:此变量包含块中的线程索引。

您似乎对CUDA的线程层次结构有点困惑；简而言之，对于内核，将有1个网格(我总是将其可视化为3维立方体)。它的每个元素都是一个块，因此声明为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 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个块)

在这个源代码中，我们甚至有4个threds，内核函数可以访问所有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个数组。

有关threadIdx、blockIdx、blockDim、gridDim的含义，请参阅cuda详细信息页面。

在这段源代码中，

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

10+0*2=0

21+0*2=1

30+1*2=2

41+1*2=3

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

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

**内核第一次调用**

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

-2 \f25 4+2*2=8 -2\f6循环

-3 \f25 loop-3\f6：-3\f25 8+2*2=12 -3\f6(但此值为-3\f25 false -3\f6，-3\f25 while out!)

**内核二次调用**

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

-2 \f25 5+2*2=9 -2\f6循环

-3 \f25 loop-3\f6：-3\f25 9+2*2=13 -3\f6(但此值为-3\f25 false -3\f6，-3\f25 while out!)

**内核第三次调用**

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

-2 \f25 loop-2\f6：-2\f25 6+2*2=10 -2\f6(但此值为-2\f25 false -2\f6，输出时！)

**内核第四次调用**

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

-2 \f25 loop-2\f6：-2\f25 7+2*2=11 -2\f6(但此值为-2\f25 false -2\f6，输出时！)

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

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