CUDA编程模型：线程层次、内存层次、核函数编写与调用

CUDA 编程模型的核心价值在于：它把 GPU 的海量并行计算能力，封装为一套对开发者相对友好的抽象。理解这套模型，是编写高效 CUDA 程序的第一步。本章聚焦三个最基础也最重要的概念：线程层次（Thread Hierarchy）、内存层次（Memory Hierarchy）以及核函数（Kernel）的编写与调用方式。

核心原理：从线程到网格的组织方式

GPU 执行并行任务的基本单位是线程（Thread）。但单个线程的能力有限，必须将成千上万个线程组织起来，才能发挥 GPU 的规模优势。CUDA 采用了一种两层结构的线程组织方式：网格（Grid） 和 线程块（Block）。

线程（Thread）：执行核函数的最小单元，拥有独立的寄存器文件和局部内存。
线程块（Block）：由多个线程组成，块内线程可以通过共享内存（Shared Memory）协作，并通过栅栏（Barrier）同步。
网格（Grid）：由多个线程块组成，一个核函数对应一个网格。网格内的块之间无法直接同步，协作成本较高。

这种两层结构并非随意设计，它直接映射到 GPU 的物理硬件架构：线程块被调度到流多处理器（SM）上执行，块内的线程进一步划分为更小的线程束（Warp）（通常为 32 个线程），这是硬件实际调度和执行的最小单位。理解这一点，有助于解释为什么块大小通常设置为 32 的倍数。

关键索引变量：在核函数内部，我们通过内置变量来定位当前线程的全局身份：

threadIdx.x：线程在块内的索引。
blockIdx.x：块在网格内的索引。
blockDim.x：块内线程的维度大小。
gridDim.x：网格内块的维度大小。

全局线程 ID 的计算公式为：int tid = blockIdx.x * blockDim.x + threadIdx.x; 这是所有 CUDA 程序中最常见的代码行之一。

方法流程：核函数的编写与调用

核函数是运行在 GPU 上的函数，使用 __global__ 关键字声明。它的编写和调用遵循一套固定的流程，与普通 CPU 函数有显著区别。

核函数编写要点

返回类型必须是 void，不能直接返回数据到主机端。
通过参数传递数据指针（通常指向设备内存）。
内部通过 threadIdx、blockIdx 等内置变量计算数据索引。
不能调用主机端函数（如 printf 在特定架构下除外），不能使用 C++ 标准库中的大部分功能。

调用流程

核函数的调用语法非常特殊，使用三括号语法 <<<grid, block>>> 指定执行配置：

// 核函数定义
__global__ void vectorAdd(const float *A, const float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

// 主机端调用
int main() {
    // ... 分配和拷贝内存（cudaMalloc, cudaMemcpy）...
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
    // ... 将结果拷回主机（cudaMemcpy）并释放内存 ...
}

上述代码展示了完整的核函数调用模式。其中 blocksPerGrid 的计算使用了向上取整的技巧，确保覆盖所有数据元素。调用后，主机端代码会继续执行，而 GPU 异步执行核函数。如果需要等待 GPU 完成，需要调用 cudaDeviceSynchronize()。

案例/场景：向量加法与矩阵乘法

向量加法是最简单的入门案例，但实际工程中，矩阵乘法更能体现线程组织和内存优化的价值。以矩阵乘法 C = A × B 为例，最直接的实现是每个线程计算输出矩阵中的一个元素：

__global__ void matMulKernel(const float *A, const float *B, float *C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; ++k) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

这个朴素版本虽然正确，但性能不佳，因为每个线程都直接访问全局内存，导致大量重复的内存读取。优化方向包括：

使用共享内存：将 A 和 B 的块加载到共享内存中，减少全局内存访问次数。
线程块大小调整：通常使用 16×16 或 32×32 的块，与共享内存容量匹配。
内存合并访问：确保相邻线程访问相邻内存地址，提高内存带宽利用率。

性能调优经验：对于矩阵乘法，使用共享内存分块（Tiling）技术，可以将性能提升数倍甚至一个数量级。核心思想是：每个线程块负责计算输出矩阵的一个子块，先将所需的 A 和 B 子块加载到共享内存，再进行计算。这样，每个全局内存数据被加载到共享内存后，可以被块内多个线程复用。

误区/复盘：常见错误与调试技巧

初学者在编写 CUDA 程序时，经常会遇到一些典型问题。理解这些误区，可以显著减少调试时间。

常见误区	现象	正确做法
忘记检查 `cudaMalloc` 返回值	程序崩溃或产生随机结果	使用 `cudaError_t` 检查每个 CUDA API 调用的返回值
核函数内数组越界	计算结果错误，且难以定位	在核函数内添加边界检查：`if (i < n)`
主机和设备内存混用	段错误或非法内存访问	明确区分 `cudaMalloc` 分配的设备指针和 `malloc` 分配的主机指针
忘记同步	主机读取到未完成计算的数据	在需要结果时调用 `cudaDeviceSynchronize()` 或使用 `cudaMemcpy`（隐式同步）
块大小设置过大	核函数启动失败（Launch Failure）	块大小不超过 1024，且考虑寄存器数量限制

调试建议：CUDA 程序的调试比 CPU 程序更复杂。推荐使用 cuda-gdb 或 compute-sanitizer（原 cuda-memcheck）工具。其中 compute-sanitizer 可以自动检测越界访问、未初始化数据等问题，是排查内存错误的利器。在开发初期，建议在每次核函数调用后检查 cudaGetLastError()，快速定位启动失败的原因。

复盘一个典型的调试场景：假设向量加法程序输出全零。首先检查 cudaMemcpy 的方向是否正确（主机到设备 vs 设备到主机）；其次检查核函数中的索引计算是否正确；最后检查是否忘记调用 cudaDeviceSynchronize()。按照这个顺序排查，通常能快速找到问题所在。