📚 课程目录第3章 CUDA编程模型

CUDA编程模型：核函数、线程组织与网格-块层次结构

核心原理：从“主机调用”到“设备执行”的编程范式

CUDA 编程模型的核心思想是将 GPU 视为一个高度并行的协处理器（Device），由 CPU（Host）负责逻辑控制与串行任务，而将数据密集型、可并行的计算任务以核函数（Kernel）的形式提交到 GPU 上执行。核函数使用 __global__ 修饰符声明，它不能直接调用，必须通过主机端代码使用特殊的尖括号语法启动。核函数的执行遵循 单指令多线程（SIMT） 模型：同一段代码被大量线程同时执行，但每个线程通过内置变量（如 threadIdx、blockIdx）区分自己的身份，从而处理不同的数据。

理解 CUDA 编程模型的关键在于把握两个层次的抽象：线程组织层次与内存层次。线程组织层次从细到粗依次为：线程（Thread）、线程块（Block）、网格（Grid）。每个线程块内的线程可以通过共享内存（Shared Memory）协作，并支持同步操作（__syncthreads()）；而不同线程块之间则无法直接同步，只能通过全局内存（Global Memory）间接通信。这种层次化设计使得 CUDA 程序既能适应不同规模的 GPU（小 GPU 可串行执行多个块，大 GPU 可并行执行更多块），又能保证块内的高效协作。

核函数声明限定符速查

__global__：在设备端执行，可从主机端调用（也可从计算能力 3.5+ 的设备端调用），必须返回 void。
__device__：在设备端执行，只能从设备端调用。
__host__：在主机端执行，即普通 C++ 函数（默认限定符）。

方法流程：网格与块的配置策略

核函数启动时，主机端通过执行配置 <<<gridDim, blockDim, sharedMemSize, stream>>> 指定网格和块的维度。其中 gridDim 和 blockDim 均为 dim3 类型，最多支持三维。合理的配置需要综合考虑数据规模、硬件资源上限与线程协作需求。

配置线程块时，需遵循以下约束：

每个线程块的最大线程数通常为 1024（部分新架构为 1024），即 blockDim.x * blockDim.y * blockDim.z ≤ 1024。
网格各维度的上限为 2³¹ - 1（x 维）和 65535（y、z 维）。
线程块大小建议取 32 的倍数（一个 Warp 的大小），如 128、256、512，以充分利用硬件调度单元并避免资源浪费。
若块内需要共享内存协作，则块大小不宜过大，否则会限制每个 SM 上可驻留的块数量，降低占用率。

典型的配置流程如下：

确定总线程数：根据数据规模 N 计算需要的总线程数，通常使总线程数 ≥ N。
选择块大小：根据算法特性选择 128 或 256 作为块大小（经验值），兼顾占用率与同步开销。
计算网格大小：gridSize = (N + blockSize - 1) / blockSize，向上取整。
处理边界：在核函数内使用 if (idx < N) 进行边界检查，防止越界访问。

// 一维向量加法示例
__global__ void vectorAdd(const float *A, const float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {          // 边界检查
        C[idx] = A[idx] + B[idx];
    }
}

// 主机端启动配置
int N = 1 << 20;            // 1,048,576 个元素
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

案例实战：矩阵乘法的线程组织设计

以 C = A × B（M×K 乘以 K×N）为例，展示二维网格与二维块的配置方法。每个线程负责计算输出矩阵 C 中的一个元素，因此需要 M × N 个线程。采用二维组织方式：blockIdx.y 对应行方向（M），blockIdx.x 对应列方向（N）。

配置参数	推荐值	说明
blockDim	(16, 16)	256 线程/块，适合共享内存分块
gridDim	(N/16, M/16)	假设 M、N 能被 16 整除
共享内存	2 × 16×16 × 4B	双缓冲分块，减少全局内存访问

__global__ void matMulShared(const float *A, const float *B, float *C,
                             int M, int K, int N) {
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];

    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0.0f;

    for (int tile = 0; tile < (K + 15) / 16; ++tile) {
        // 协作加载 A 和 B 的分块到共享内存
        if (row < M && tile * 16 + threadIdx.x < K)
            As[threadIdx.y][threadIdx.x] = A[row * K + tile * 16 + threadIdx.x];
        else
            As[threadIdx.y][threadIdx.x] = 0.0f;

        if (col < N && tile * 16 + threadIdx.y < K)
            Bs[threadIdx.y][threadIdx.x] = B[(tile * 16 + threadIdx.y) * N + col];
        else
            Bs[threadIdx.y][threadIdx.x] = 0.0f;

        __syncthreads();  // 确保分块加载完成

        for (int k = 0; k < 16; ++k)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];

        __syncthreads();  // 防止共享内存被覆盖
    }

    if (row < M && col < N)
        C[row * N + col] = sum;
}

该案例展示了线程组织的核心技巧：线程索引映射到数据坐标、块内协作加载数据到共享内存、双同步点保护共享内存数据一致性。通过将全局内存访问次数从 O(K) 降低到 O(K/16)，大幅提升了访存效率。

误区与复盘：常见配置错误及优化思路

误区一：忽略边界检查。当 N 不是 blockSize 的整数倍时，直接访问会导致越界。必须使用 if (idx < N) 保护。

误区二：块大小设置过大或过小。块大小超过 1024 直接编译失败；块过小（如 32）则块数过多，调度开销增大。经验值 128~256 通常最优。

误区三：过度使用 __syncthreads()。同步点会阻塞块内所有线程，若分支发散导致部分线程无法到达同步点，将造成死锁。

复盘一个典型调优案例：某图像处理核函数初始配置为 <<<N/64, 64>>>，每个线程处理一个像素。通过性能分析发现占用率仅为 37%，原因是块太小导致每个 SM 上块数过多，资源分配碎片化。调整为 <<<N/256, 256>>> 后，占用率提升至 75%，性能提升 1.8 倍。进一步使用共享内存进行行缓存后，性能再提升 40%。

最佳实践清单：

块大小取 128 或 256，并确保为 32 的倍数。
总线程数略大于数据规模，用边界检查保证安全。
优先使用二维/三维组织以匹配数据空间结构。
需要块内协作时，使用共享内存 + __syncthreads()，但同步次数应尽量少。
使用 cudaOccupancyMaxPotentialBlockSize() 辅助确定最优块大小。

理解网格与块的层次关系是 CUDA 编程的基石。掌握线程索引到数据坐标的映射方法，合理配置块大小与网格维度，并规避边界与同步陷阱，才能写出正确且高效的并行程序。在实际开发中，应结合硬件规格（如 SM 数量、共享内存容量）与性能分析工具（如 Nsight Compute）反复迭代，找到针对特定 GPU 的最优配置。