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

CUDA 编程模型的核心思想，是将 GPU 视为一个高度并行的协处理器。CPU（主机，Host）负责逻辑控制和串行任务，GPU（设备，Device）负责大规模数据并行计算。要真正驾驭 GPU 的算力，必须理解三个相互交织的层次：线程层次（如何组织并行单元）、内存层次（数据如何在不同速度的存储间流动），以及核函数（如何编写并启动设备端代码）。这三者共同决定了程序的正确性、性能上限与可扩展性。

3.1 线程层次：从 Grid 到 Thread 的组织逻辑

GPU 执行并行任务的最小单位是 线程（Thread）。但单个线程无法管理成千上万个并发执行单元，因此 CUDA 引入了三层线程组织结构：

Grid（网格）：一次核函数启动所创建的所有线程的集合。Grid 是线程组织的最高层。
Block（线程块）：Grid 内的一个子集，包含若干线程。Block 内的线程可以通过共享内存协作，并通过栅栏（__syncthreads()）同步。
Thread（线程）：实际执行计算的最小单元。

每个线程在全局范围内拥有唯一的 全局线程 ID，在 Block 内拥有唯一的 局部线程 ID。这种层级设计并非随意为之，它直接映射到 GPU 的物理硬件架构：一个 Block 会被调度到单个流多处理器（SM）上执行，Block 内的线程共享 SM 上的计算资源和共享内存；而一个 Grid 则横跨多个 SM。这种映射关系使得程序能够透明地扩展到不同规模的 GPU 上——只要 Block 数量足够多，小 GPU 和大 GPU 都能高效运行。

线程 ID 的计算是编写 CUDA 程序的基本功。对于一维结构，全局线程 ID 通常为 blockIdx.x * blockDim.x + threadIdx.x。对于二维或三维结构，则需要同时考虑各维度的索引。下图展示了 Grid、Block 与 Thread 的层级关系及索引计算方式：

在实际编程中，Block 的维度和大小是性能调优的关键参数。Block 内线程数量通常取 128、256 或 512 等 32 的倍数（一个 Warp 包含 32 个线程），以保证 Warp 完整填充且无浪费。

3.2 内存层次：数据流动的“高速公路”系统

CUDA 的内存模型是分层的，不同层次的内存具有不同的容量、速度和可见性。理解这一层次结构，是优化访存性能的核心。下表总结了各内存类型的关键属性：

内存类型	作用域	生命周期	速度	典型用途
寄存器（Register）	单个线程	线程内	最快	局部变量、中间计算结果
共享内存（Shared Memory）	单个 Block	Block 内	极快（~1.5TB/s）	Block 内线程间数据共享、数据复用
全局内存（Global Memory）	所有线程 + 主机	程序运行期间	较慢（~200GB/s）	主数据存储，CPU 与 GPU 交换数据
常量内存（Constant Memory）	所有线程（只读）	程序运行期间	有缓存加速	只读参数、查找表
纹理内存（Texture Memory）	所有线程（只读）	程序运行期间	有缓存加速	图像处理、空间局部性强的访问
局部内存（Local Memory）	单个线程	线程内	慢（实际在全局内存中）	寄存器溢出时使用

数据在 CPU 与 GPU 之间的传输通过 PCIe 总线（或 NVLink）进行，这是整个系统中速度最慢的环节。因此，一个重要的编程原则是：尽量减少 CPU 与 GPU 之间的数据拷贝，尽量在设备端完成所有计算。在 GPU 内部，数据从全局内存加载到共享内存，再由线程从共享内存中高速读取，是典型的优化模式（即“分块（Tiling）”技术）。

3.3 核函数编写与调用：从声明到启动

核函数（Kernel）是在设备端执行的函数，使用 __global__ 关键字声明。核函数的编写和调用遵循一套严格的语法规则，下面通过一个向量加法的完整示例来展示标准流程：

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

// 主机端调用核函数
int main() {
    int n = 1 << 20;  // 1,048,576 个元素
    size_t bytes = n * sizeof(float);
    
    // 1. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);
    
    // 2. 将输入数据从主机拷贝到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
    
    // 3. 配置核函数执行参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    // 4. 启动核函数
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
    
    // 5. 同步并检查错误
    cudaDeviceSynchronize();
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("CUDA Error: %s\n", cudaGetErrorString(err));
    }
    
    // 6. 将结果拷贝回主机
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
    
    // 7. 释放设备内存
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    return 0;
}

核函数调用使用三括号语法 <<<gridDim, blockDim, sharedMemSize, stream>>>，其中 gridDim 和 blockDim 可以是 dim3 类型（支持三维），sharedMemSize 指定动态共享内存字节数（默认为 0），stream 指定执行流（默认为默认流）。核函数是异步的，CPU 不会等待 GPU 执行完毕，因此需要 cudaDeviceSynchronize() 或 cudaMemcpy 来保证同步。

最佳实践：核函数内部应避免使用 printf 进行调试（性能极低），推荐使用 cuda-gdb 或 compute-sanitizer 进行调试和内存检查。同时，核函数应设计为“无状态”的——不依赖任何全局变量，所有数据通过参数传入，以保证可重入性和可扩展性。

3.4 案例：矩阵乘法的分块优化

矩阵乘法是深度学习中最基础也最核心的运算。朴素实现中，每个线程计算输出矩阵的一个元素，需要从全局内存读取 A 的一行和 B 的一列，访存效率极低。通过分块（Tiling）技术，将数据先加载到共享内存，可以大幅减少全局内存访问次数。

假设输出矩阵大小为 M×N，分块大小为 BLOCK_SIZE=16。每个 Block 负责计算一个 16×16 的输出子块，需要从全局内存加载 A 的一个 16×K 子块和 B 的一个 K×16 子块到共享内存。这样，每个全局内存数据被 BLOCK_SIZE 个线程复用，全局内存访问量减少为原来的 1/BLOCK_SIZE。

__global__ void matMulTiled(const float *A, const float *B, float *C, int M, int N, int K) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    // 遍历 K 方向的子块
    for (int tile = 0; tile < (K + BLOCK_SIZE - 1) / BLOCK_SIZE; ++tile) {
        // 协作加载 A 和 B 的子块到共享内存
        if (row < M && tile * BLOCK_SIZE + threadIdx.x < K)
            As[threadIdx.y][threadIdx.x] = A[row * K + tile * BLOCK_SIZE + threadIdx.x];
        else
            As[threadIdx.y][threadIdx.x] = 0.0f;
            
        if (col < N && tile * BLOCK_SIZE + threadIdx.y < K)
            Bs[threadIdx.y][threadIdx.x] = B[(tile * BLOCK_SIZE + threadIdx.y) * N + col];
        else
            Bs[threadIdx.y][threadIdx.x] = 0.0f;
        
        __syncthreads();  // 确保所有线程完成加载
        
        // 计算局部乘积
        #pragma unroll
        for (int k = 0; k < BLOCK_SIZE; ++k) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        
        __syncthreads();  // 防止下一轮覆盖共享内存
    }
    
    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}

这个案例展示了 CUDA 编程的核心思维：将大问题分解为可并行的小块，利用共享内存实现数据复用，通过同步保证正确性。实际工程中，cuBLAS 库对矩阵乘法做了更深度的优化（如寄存器分块、向量化加载、Warp 级调度等），性能远超手写实现，但理解分块原理是阅读和调优高性能库的基础。

3.5 常见误区与调试复盘

初学者在编写 CUDA 程序时，常会遇到以下几类问题，这里逐一剖析：

误区一：忽略边界检查。当数据总量不是 Block 大小的整数倍时，最后一个 Block 的部分线程会访问越界内存。这可能导致程序崩溃或产生“幽灵数据”。对策：始终在核函数入口处添加 if (idx < n) 之类的边界判断。

误区二：过度使用 __syncthreads()。该函数是 Block 级同步屏障，但若放在条件分支中（如 if (threadIdx.x < 10) __syncthreads();），会导致死锁，因为部分线程永远无法到达屏障。对策：确保所有线程都执行到 __syncthreads()，或使用 if (threadIdx.x < 10) { ... } __syncthreads(); 的结构。

误区三：主机端与设备端内存混用。CPU 代码不能直接解引用设备指针（如 d_A），否则会报段错误。对策：严格区分 h_（主机）和 d_（设备）前缀，并通过 cudaMemcpy 显式传输数据。

调试 CUDA 程序时，推荐使用 compute-sanitizer（原 cuda-memcheck）检测内存越界和未初始化访问。该工具能精确定位出错线程的 ID 和访问地址，是排查疑难杂症的利器。此外，cuda-gdb 支持断点调试核函数，但需注意 GPU 调试的并发特性——断点会暂停整个 Grid 的执行。

最后，一个实用的性能分析流程是：先用 ncu（Nsight Compute）获取核函数的性能指标（如占用率、内存吞吐量、Warp 停滞原因），再针对瓶颈进行优化。通常，优化优先级为：先减少全局内存访问（分块、合并访问），再优化计算指令（使用快速数学函数、减少分支发散），最后调整 Block 大小和 Grid 维度以提升占用率。