CUDA中简单的理解线程管理

原创已于 2026-03-05 11:20:52 修改 · 529 阅读

12 ·

本内容遵循CC 4.0 BY-SA版权协议

GEO检测

标签

#开发语言 #c++ #人工智能

于 2026-03-04 17:36:21 首次发布

CUDA分享专栏收录该内容

3 篇文章

订阅专栏

CUDA中简单的理解线程管理

一、核函数

如果你先我写的太罗嗦，可以直接直接往下拉，到第五章中，第2节的例子。

通过上一篇（CUDA简单的基础概念），我们知道：

一个核函数只能有一个grid，一个grid可以有很多个块，每个块可以有很多的线程：

所以，我们看到核函数的定义的时候：

语法：核函数名<<<网格配置, 线程块配置>>>(参数列表);

即：

kernel_name<<<grid,block>>>(argument list);

二、核函数中的参数

举个例子，如下面这段代码：

const int numElements = 1024;
const int threadsPerBlock = 256;
const int numBlocks = (numElements + threadsPerBlock - 1) / threadsPerBlock;
addKernel<<<numBlocks, threadsPerBlock>>>(d_a, d_b, d_c, numElements);

这段代码里根据数据元素总数 numElements 和每个线程块的线程数 threadsPerBlock 计算出所需的线程块数 numBlocks，这样就能确保所有数据都能被处理。即：

每个数据元素都能在一个独立线程中被处理或计算。

注：我们知道每个线程都执行同样的一段串行代码，那么怎么让这段相同的代码对应不同的数据呢？首先第一步就是让这些线程彼此区分开，才能对应到相应从线程，使得这些线程也能区分自己的数据。如果线程本身没有任何标记，那么没办法确认其行为。
依靠下面两个内置结构体确定线程标号：

blockIdx（线程块在线程网格内的位置索引）
threadIdx（线程在线程块内的位置索引）

这样每个线程一个独立的编号了。

三、核函数中参数的理解

以上，就可以明白，核函数中的参数：

网格配置：网格中的线程块的数量

线程块配置：线程块中线程的数量

详解

如果你有点困扰，别急，我们再举个例子：

对于核函数

kernel_name<<<grid,block>>>(argument list);

我们可以使用，dim3类型的grid维度和block维度配置内核，也可以使用int类型的变量，或者常量直接初始化。

假设我们使用常量4，8来初始化：

kernel_name<<<4,8>>>(argument list);

那么线程布局就是下面这样的：

四、代码中查看

我们来写个代码，看看打印输出：

代码内容：

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void checkIndex(void)
{
    printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d) blockDim:(%d,%d,%d)gridDim(%d,%d,%d)\n",
           threadIdx.x,threadIdx.y,threadIdx.z,
           blockIdx.x,blockIdx.y,blockIdx.z,blockDim.x,blockDim.y,blockDim.z,
           gridDim.x,gridDim.y,gridDim.z);

}

void printCheckIndex()
{
    int nElem=6;
    dim3 block(3);
    dim3 grid((nElem+block.x-1)/block.x);
    printf("grid.x %d grid.y %d grid.z %d\n",grid.x,grid.y,grid.z);
    printf("block.x %d block.y %d block.z %d\n",block.x,block.y,block.z);
    checkIndex<<<grid,block>>>();
    cudaDeviceReset();

}

我们执行printCheckIndex()函数，输出内容如下：

threadIdx:(0,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)gridDim(2,1,1)
threadIdx:(1,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)gridDim(2,1,1)
threadIdx:(2,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)gridDim(2,1,1)
threadIdx:(0,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)gridDim(2,1,1)
threadIdx:(1,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)gridDim(2,1,1)
threadIdx:(2,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)gridDim(2,1,1)

grid.x 2 grid.y 1 grid.z 1
block.x 3 block.y 1 block.z 1

chcekIndex函数中的变量解释：

threadIdx：
含义：threadIdx 表示线程在其所在线程块内的索引。它是一个三维向量，通过 x、y 和 z 三个分量来唯一标识线程块内的每个线程。
解释：
threadIdx.x：线程在当前线程块 x 维度上的索引。在一维线程块布局中，x 索引用于区分不同的线程。例如，如果一个线程块中有 256 个线程按一维排列，threadIdx.x 的值将从 0 到 255。
threadIdx.y：线程在当前线程块 y 维度上的索引。当使用二维线程块布局时，y 索引与 x 索引一起确定线程的位置。例如，对于一个 16x16 的二维线程块，threadIdx.y 的值范围是 0 到 15。
threadIdx.z：线程在当前线程块 z 维度上的索引。三维线程块布局在一些复杂的并行算法中会用到，z 索引与 x 和 y 索引共同确定线程在三维空间中的位置。例如，在一个 4x4x4 的三维线程块中，threadIdx.z 的值范围是 0 到 3。

blockIdx：
含义：blockIdx 表示线程块在网格（grid）中的索引。同样是一个三维向量，用于在网格中唯一标识每个线程块。
解释：
blockIdx.x：线程块在网格 x 维度上的索引。如果网格是一维的，blockIdx.x 用于区分不同的线程块。例如，如果有 10 个线程块组成一维网格，blockIdx.x 的值将从 0 到 9。
blockIdx.y：线程块在网格 y 维度上的索引。在二维网格布局中，y 索引与 x 索引一起确定线程块的位置。例如，对于一个 5x5 的二维网格，blockIdx.y 的值范围是 0 到 4。
blockIdx.z：线程块在网格 z 维度上的索引。三维网格布局可用于处理更复杂的数据结构或并行任务，z 索引与 x 和 y 索引共同确定线程块在三维网格中的位置。例如，在一个 2x2x2 的三维网格中，blockIdx.z 的值范围是 0 到 1。

blockDim：
含义：blockDim 定义了线程块在每个维度上的大小，即每个线程块包含的线程数量。
输出解释：
blockDim.x：线程块在 x 维度上的大小，即 x 方向上的线程数量。例如，如果 blockDim.x = 256，表示线程块在 x 方向上有 256 个线程。
blockDim.y：线程块在 y 维度上的大小，即 y 方向上的线程数量。例如，对于一个 32x32 的二维线程块，blockDim.y = 32。
blockDim.z：线程块在 z 维度上的大小，即 z 方向上的线程数量。例如，在一个 8x8x8 的三维线程块中，blockDim.z = 8。

gridDim：
含义：gridDim 定义了网格在每个维度上的大小，即网格中包含的线程块数量。
输出解释：
gridDim.x：网格在 x 维度上的大小，即 x 方向上的线程块数量。例如，如果 gridDim.x = 10，表示在一维网格中有 10 个线程块。
gridDim.y：网格在 y 维度上的大小，即 y 方向上的线程块数量。例如，对于一个 4x4 的二维网格，gridDim.y = 4。
gridDim.z：网格在 z 维度上的大小，即 z 方向上的线程块数量。例如，在一个 3x3x3 的三维网格中，gridDim.z = 3。

五、CUDA中线性位置的计算

1、线性位置

在 CUDA 中，线性位置（linear position）通常指将三维的线程索引（threadIdx 和 blockIdx）映射到一维的索引位置，以便访问线性存储的数据结构（如数组）。

我们回到第一节的代码：

const int numElements = 1024;
const int threadsPerBlock = 256;
const int numBlocks = (numElements + threadsPerBlock - 1) / threadsPerBlock;
addKernel<<<numBlocks, threadsPerBlock>>>(d_a, d_b, d_c, numElements);

这种关于线性位置的计算，就是一维网格和一维线程块的简单情况，这种也是实际应用中比较常见的。

假设 gridDim.x = numBlocks，blockDim.x = numThreadsPerBlock，线性位置计算为：

linearIndex=blockIdx.x∗blockDim.x+threadIdx.x

这种简单形式更直观，易于理解，常用于处理一维数据数组。例如，在对长度为 N 的数组进行并行操作时，如果每个线程处理一个数组元素，可以设置 numThreadsPerBlock 为一个合适的值（如 256），numBlocks 为 (N + numThreadsPerBlock - 1) / numThreadsPerBlock，每个线程通过上述公式计算出自己要处理的数组元素索引。

2、举个例子：

你有一个图像处理的核函数，你想对一个宽度和高度为2048x2048像素的图像进行处理，你想让每个线程块处理16x16像素的区域

__global__ void processImage(float* input, float* output, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < width && y < height) {
        int index = x + y * width;
        output[index] = input[index] * 2.0f; // 示例操作：将输入值乘以2
    }
}

int main() {
    // 省略CUDA内存分配和错误检查等初始化代码...

    dim3 threadsPerBlock(16, 16, 1); // 每个线程块处理16x16像素
    dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x,
                   (height + threadsPerBlock.y - 1) / threadsPerBlock.y,
                   1); // 计算需要的线程块数量

    processImage<<<numBlocks, threadsPerBlock>>>(d_input, d_output, width, height);
    // 省略CUDA错误检查和内存回收代码...
}

注：dim3类型是用来表示三维网格和线程块的大小。这对于并行计算非常重要，因为它决定了你的GPU上运行的线程的数量和布局。dim3有三个成员：x、y和z，分别代表网格或线程块在X、Y、Z方向上的大小。

我们逐行解释下，你就明白了：

这个图象像素的元素个数是2048*2048=4194304个。

每个线程块中分配16*16=256个元素，即dim3 threadsPerBlock(16, 16, 1);

而需要多少个线程块呢？应该是Width宽度上线程块数量*Height高度上线程块数量，这个总的线程块数量包含每一个元素，就能保证每个元素线程id的唯一性了。那么，width方向上就应该是width/16，即width/threadsPerBlock.x,那么如果采用向上取整就能避免漏处理，于是，宽度维度计算就是(width + threadsPerBlock.x - 1) / threadsPerBlock.x。同样的，Height方向就是(width + threadsPerBlock.x - 1) / threadsPerBlock.x。因为深度方向为1，所以块的数量就是128*128=16384，每个块的元素是256，那么总的元素数量就是16384*256=4194304个。

那么在核函数中，（int x = blockIdx.x * blockDim.x + threadIdx.x）和（int y = blockIdx.y * blockDim.y + threadIdx.y;）该如何理解呢？

x 方向线性位置计算：
blockIdx.x * blockDim.x：这部分计算出当前线程所在线程块在 x 方向上，距离网格起始位置的线程偏移量。例如，如果 blockIdx.x = 2，blockDim.x = 16，那么当前线程块在 x 方向上距离起始位置已经偏移了 2 * 16 = 32 个线程。
+ threadIdx.x：这部分加上线程在当前线程块内 x 方向上的索引。所以，blockIdx.x * blockDim.x + threadIdx.x 就得到了线程在整个网格 x 方向上的线性位置。例如，threadIdx.x = 5，那么这个线程在 x 方向上的线性位置就是 32 + 5 = 37。
y 方向线性位置计算：
blockIdx.y * blockDim.y：与 x 方向类似，这部分计算出当前线程所在线程块在 y 方向上，距离网格起始位置的线程偏移量。例如，如果 blockIdx.y = 3，blockDim.y = 16，那么当前线程块在 y 方向上距离起始位置已经偏移了 3 * 16 = 48 个线程。
+ threadIdx.y：再加上线程在当前线程块内 y 方向上的索引。所以，blockIdx.y * blockDim.y + threadIdx.y 就得到了线程在整个网格 y 方向上的线性位置。例如，threadIdx.y = 7，那么这个线程在 y 方向上的线性位置就是 48 + 7 = 55。

核函数中，下面这段代码，该如何理解呢？

  if (x < width && y < height) {
        int index = x + y * width;
        output[index] = input[index] * 2.0f; // 示例操作：将输入值乘以2
    }

因为我们参数float* input一般是一个一维数组，图像一般是二维数组 data[width * height]，这就需要一个映射，我们，通过上述计算得到的 x 和 y 线性位置，可以直接映射到数组的索引。例如，如果数组是按行优先存储，那么 y * width + x 就可以得到对应元素在数组中的线性索引。

到这里，简单的线性位置，你应该理解一点了，多练习一下，应该就能理解了。

3、多维的计算公式

对于，多维的线性位置计算，大家可以自行研究下：

注：假设我们有以下配置：
gridDim：网格维度，gridDim.x、gridDim.y、gridDim.z 分别表示网格在 x、y、z 维度上的线程块数量。
blockDim：线程块维度，blockDim.x、blockDim.y、blockDim.z 分别表示线程块在 x、y、z 维度上的线程数量。
threadIdx：线程在其所在线程块内的索引，threadIdx.x、threadIdx.y、threadIdx.z 分别表示线程在 x、y、z 维度上的索引。

计算线性位置 linearIndex 的公式为：
linearIndex = threadIdx.z+blockDim.z∗(threadIdx.y+blockDim.y∗threadIdx.x)+
blockDim.z∗blockDim.y∗blockDim.x∗

(blockIdx.z+gridDim.z∗(blockIdx.y+gridDim.y∗blockIdx.x))

本文原创作者：冯一川（csdn：ifeng12358），未经作者授权同意，请勿转载。