第三章 CUDA编程模型：核函数与线程层次结构

CUDA编程模型的核心思想是将GPU视为一个高度并行的协处理器，通过编写称为“核函数”（Kernel）的特殊函数，让数千个线程同时执行相同的计算任务。理解核函数的编写与调用方式，以及线程层次结构（Grid、Block、Thread）和内置变量（threadIdx、blockIdx、blockDim、gridDim）的语义，是掌握CUDA编程的第一道门槛。本章将从概念原理、代码实操、案例拆解和常见误区四个维度，系统性地构建你对CUDA执行模型的基础认知。

上图展示了CUDA中Grid、Block、Thread三层结构的完整映射关系。在CUDA编程模型中，Grid是最大的线程组织单位，由多个Block组成；每个Block又包含多个Thread。这种层次化设计不仅让程序员能够灵活地组织并行任务，更重要的是它直接映射到GPU的物理硬件架构——Block被调度到流多处理器（SM）上执行，而Thread则是GPU硬件执行的基本单位。理解这张图，你就掌握了CUDA并行编程的“世界观”。

3.1 核函数（Kernel）的编写与调用

核函数是CUDA编程的入口，它使用__global__修饰符声明，在设备端（GPU）执行，由主机端（CPU）调用。核函数的编写规则与普通C函数有显著差异，理解这些差异是避免程序崩溃或产生错误结果的前提。

3.1.1 核函数的核心原理

核函数本质上是一个并行执行模式的声明：当你调用一个核函数时，CUDA运行时会在GPU上创建指定数量的线程，每个线程都执行相同的函数体代码。这种模式被称为“单指令多线程”（SIMT）。核函数必须满足以下约束：

• 返回类型必须为void：核函数无法直接向主机返回计算结果，所有数据交换必须通过显存（全局内存）进行。
• 必须使用__global__修饰：这是区分核函数与普通函数的关键标志。
• 不支持可变参数：核函数的参数个数在编译时必须确定。
• 不能使用C++标准库（如printf、malloc等），但可以使用CUDA提供的数学函数（如sinf、sqrtf等）。
• 不能调用主机端函数，但可以调用其他__device__函数。

3.1.2 核函数的编写与调用流程

编写一个完整的核函数程序需要遵循以下操作流程：

1. 编写核函数：使用__global__声明，函数体内部通过内置变量确定当前线程的全局ID。
2. 分配设备内存：使用cudaMalloc在GPU上分配显存空间。
3. 数据拷贝：使用cudaMemcpy将输入数据从主机内存拷贝到设备显存。
4. 配置执行参数：通过<<<gridDim, blockDim>>>语法指定Grid和Block的维度。
5. 调用核函数：主机端发起调用，GPU异步执行。
6. 结果回传：使用cudaMemcpy将计算结果从设备拷贝回主机。
7. 释放资源：使用cudaFree释放显存。

// 示例：向量加法核函数
__global__ void vectorAdd(float *A, float *B, float *C, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;  // 计算全局线程ID
    if (idx < n) {  // 边界检查，防止越界
        C[idx] = A[idx] + B[idx];
    }
}

// 主机端调用
int main() {
    int n = 1024;
    size_t bytes = n * sizeof(float);
    
    // 1. 分配主机内存并初始化
    float *h_A = (float*)malloc(bytes);
    float *h_B = (float*)malloc(bytes);
    float *h_C = (float*)malloc(bytes);
    
    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);
    
    // 3. 数据拷贝到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
    
    // 4. 配置执行参数并调用核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
    
    // 5. 结果回传
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
    
    // 6. 释放资源
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);
    return 0;
}

3.2 线程层次结构：Grid、Block、Thread

CUDA的线程层次结构是理解并行粒度控制的关键。它允许程序员以三维方式组织线程，从而自然地映射到数据的多维结构（如矩阵、体素等）。

3.2.1 层次结构的核心原理

CUDA的线程组织分为三个层级，每个层级都有明确的硬件对应关系：

这种层次化设计带来了两个关键优势：可扩展性和局部性。可扩展性体现在：同一份核函数代码，在只有几个SM的低端GPU和拥有上百个SM的高端GPU上都能正确运行，只是并行度不同。局部性体现在：同一个Block内的线程可以通过共享内存高效通信，而不同Block之间的通信则必须通过全局内存，代价高得多。

3.2.2 如何选择合适的线程层次配置

选择合适的Grid和Block维度是CUDA性能优化的核心决策之一。以下是判断和选择的方法：

1. 确定Block大小：通常选择128、256或512个线程。经验法则：Block大小应为32的倍数（一个Warp的大小），以充分利用硬件调度效率。
2. 计算Grid大小：根据总数据量计算：blocksPerGrid = (totalThreads + threadsPerBlock - 1) / threadsPerBlock。
3. 考虑硬件限制：每个SM能同时驻留的Block数量有限（通常8-32个），Block太大可能导致资源浪费，太小则无法隐藏内存延迟。
4. 使用多维组织：对于二维数据（如图像），使用dim3 block(16, 16)；对于三维数据（如体素），使用dim3 block(8, 8, 8)。

3.3 内置变量：threadIdx、blockIdx、blockDim、gridDim

内置变量是CUDA提供给程序员在核函数内部使用的只读变量，它们描述了当前线程在层次结构中的位置。正确理解和使用这些变量，是编写正确核函数的基础。

3.3.1 内置变量的核心原理

这四个内置变量都是dim3类型（实际上在设备端是uint3结构体），包含x、y、z三个分量。它们的语义如下：

• threadIdx：当前线程在所属Block内的索引，范围从0到blockDim-1。
• blockIdx：当前Block在Grid内的索引，范围从0到gridDim-1。
← 上一章 📖 返回目录 下一章 →