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

3. CUDA编程模型：从线程到数据并行的核心抽象

CUDA 编程模型是 GPU 计算的基石。它并非要求开发者直接管理数千个硬件线程，而是提供了一套分层抽象，让开发者以“数据并行”的思维编写核函数（Kernel），由硬件和运行时自动将其映射到海量计算核心上。理解这套模型，是后续进行数据驱动迭代调优的前提——因为性能瓶颈往往源于线程组织方式与数据访问模式的不匹配。

3.1 核心抽象层次：线程、线程块与网格

CUDA 将执行模型组织为三层结构，每一层都有明确的语义和硬件映射关系：

线程（Thread）：执行核函数的最小单元。每个线程有独立的程序计数器、寄存器状态和局部变量。
线程块（Block）：一组线程的集合，块内线程可通过共享内存（Shared Memory）协作，并通过栅栏（Barrier）同步。一个块会被调度到单个流式多处理器（SM）上，块内线程数通常限制在 1024 以内。
网格（Grid）：核函数启动时创建的所有线程块集合。网格中的块之间无全局同步，只能通过全局内存或原子操作进行轻量级通信。

这种分层设计的关键在于：块内协作成本低（纳秒级），块间通信成本高（微秒级甚至更高）。因此，算法设计时应尽量让数据在块内完成交换。

  硬件映射关系：网格 → GPU 设备；线程块 → SM（一个 SM 可同时驻留多个块）；线程 → SM 内的 CUDA 核心。块内线程数（blockDim）和网格尺寸（gridDim）共同决定了总线程数。

3.2 核函数定义与线程索引映射

核函数使用 __global__ 修饰符定义，通过 <<<gridDim, blockDim>>> 语法启动。在核函数内部，开发者通过内置变量 threadIdx、blockIdx、blockDim、gridDim 计算当前线程的全局唯一 ID。这是数据并行映射的核心：每个线程负责处理一个或多个数据元素。

// 一维映射：处理长度为 N 的数组
__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;  // 全局线程ID
    if (idx < N) {  // 边界检查
        C[idx] = A[idx] + B[idx];
    }
}

// 启动配置：256 线程/块，共 (N+255)/256 个块
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

对于二维/三维数据（如图像、矩阵），CUDA 支持 dim3 类型的网格和块维度，使得索引表达更自然：

// 二维矩阵：每个线程处理一个元素
__global__ void matrixAdd(const float* A, const float* B, float* C, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int idx = row * width + col;
    C[idx] = A[idx] + B[idx];
}

dim3 blockDim(16, 16);          // 16×16 = 256 线程
dim3 gridDim((width+15)/16, (height+15)/16);
matrixAdd<<<gridDim, blockDim>>>(d_A, d_B, d_C, width);

性能提示：线程块大小通常选择 128～512 的倍数（如 256），以保证每个线程束（Warp，32线程）满额且负载均衡。块大小过小会导致调度开销占比高，过大则可能降低 SM 上并发块的数量。

3.3 共享内存与同步机制

共享内存（Shared Memory）是块内线程共享的高速缓存（延迟约 20～30 周期，远低于全局内存的 400～600 周期）。它常用于：

数据复用：多个线程读取同一数据区域时，先协作加载到共享内存，减少全局内存访问。
块内通信：线程间交换中间结果，避免通过全局内存。
合并访问优化：将非合并的全局访问转换为共享内存中的合并访问。

共享内存在核函数内使用 __shared__ 修饰符声明，其生命周期仅存在于块执行期间。由于块内线程执行进度可能不一致，CUDA 提供了 __syncthreads() 栅栏同步：所有线程必须到达该点后才能继续执行，防止数据竞争。

// 示例：使用共享内存进行矩阵分块乘法（简化版）
__global__ void matMulShared(const float* A, const float* B, float* C, int N) {
    __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;

    for (int tile = 0; tile < N / BLOCK_SIZE; ++tile) {
        // 协作加载数据到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * BLOCK_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * BLOCK_SIZE + threadIdx.y) * N + col];
        __syncthreads();  // 确保所有线程完成加载

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

        __syncthreads();  // 防止下一轮覆盖前仍有线程在读取
    }
    C[row * N + col] = sum;
}

关键陷阱：过度使用 __syncthreads() 会导致线程束分化（Warp Divergence）和性能下降。若块内分支导致部分线程跳过同步点，程序会死锁。务必确保所有线程都能到达同步点，或使用 __syncwarp() 进行更细粒度的同步。

3.4 内存层次与数据局部性

CUDA 内存模型是分层的，不同内存的容量、速度、作用域差异巨大。理解这些差异是数据驱动调优的核心基础：

内存类型	作用域	生命周期	速度（相对）	典型用途
寄存器	单个线程	线程执行期间	最快	局部变量、中间计算
共享内存	线程块	块执行期间	约 20～30 周期	块内数据复用、协作通信
全局内存	所有线程 + 主机	程序运行期间	400～600 周期	主数据存储、跨块通信
常量内存	所有线程（只读）	程序运行期间	有缓存加速	核函数参数、查找表
纹理内存	所有线程（只读）	程序运行期间	有缓存加速	图像处理、空间局部性数据

在数据驱动迭代调优中，访存模式往往比计算量更影响性能。例如，全局内存访问应尽量满足“合并访问”条件——即同一线程束的 32 个线程访问连续的内存地址。若访问不连续，则会产生多次事务，浪费带宽。

3.5 典型误区与调优复盘

  误区 1：线程数越多越好。实际中，过大的网格会导致线程块驻留数量下降，上下文切换开销增加。应通过 Occupancy Calculator 或实验找到每个 SM 的最佳块数。

  误区 2：忽略边界检查。当数据长度不是块大小的整数倍时，未加 if (idx < N) 会导致越界访问，产生难以调试的随机错误。

  误区 3：共享内存使用不当。共享内存容量有限（通常 48KB～164KB/SM），过度使用会降低块驻留数量。应优先优化数据复用率，而非盲目增加共享内存占用。

复盘案例：某图像滤波核函数，初始版本每个线程直接读取全局内存中的邻域像素，导致大量冗余访问。优化后，将图像分块加载到共享内存，每个线程负责加载一个像素并同步，滤波时从共享内存读取邻域。结果显示：全局内存访问量减少 85%，核函数耗时从 3.2ms 降至 1.1ms，加速比约 2.9 倍。关键收益来自数据局部性的提升，而非计算量的减少。