CUDA编程模型：线程层次、内存模型与同步机制

核心原理：从线程到网格的三级抽象

CUDA将GPU的执行资源抽象为三层结构：线程（Thread）、线程块（Block）与网格（Grid）。这种设计并非随意为之，而是为了同时兼顾两件事：一是让程序员能够以细粒度控制并行逻辑，二是让硬件能够在大规模并行下高效调度。一个内核（Kernel）启动时，实际上是以网格为单位提交任务，网格由若干线程块组成，每个线程块又包含若干线程。硬件调度器以线程块为基本分配单位，将块分发到不同的流式多处理器（SM）上执行，而块内的线程则进一步划分为更小的线程束（Warp，通常32线程）进行SIMT执行。

理解这一层次的关键在于：块内线程可以通过共享内存协作，块间线程无法直接通信。因此，线程块既是逻辑分组，也是资源隔离的边界。程序员在设计并行算法时，必须预先决定块的大小与网格的维度，这直接决定了问题的分解方式与硬件利用率。

// 典型的CUDA内核启动语法
dim3 blockDim(256);          // 每个线程块256个线程
dim3 gridDim((N + 255) / 256); // 网格中线程块数量
kernel<<<gridDim, blockDim>>>(data, N);

内存模型：五层存储的可见性与性能差异

CUDA的内存模型与线程层次紧密耦合。每个线程拥有私有寄存器和本地内存；每个线程块拥有共享内存，块内所有线程可见；整个网格拥有全局内存、常量内存和纹理内存，所有线程可见。这种分层设计直接反映了物理硬件的缓存结构：寄存器位于SM内部，访问延迟约为1个周期；共享内存是SM上的片上SRAM，延迟约20～30周期；全局内存则位于显存（DRAM）中，延迟高达400～800周期。

编程时最需要警惕的是数据局部性。如果算法反复从全局内存读取同一份数据，性能将受制于显存带宽。正确的做法是：先将数据从全局内存加载到共享内存，由块内线程协作处理，最后再写回全局内存。这种“分块（Tiling）”技术是矩阵乘法、卷积等算子优化的核心手段。

同步机制：协作的规则与边界

CUDA提供两级同步原语：块内同步使用__syncthreads()，它确保块内所有线程到达该屏障后才能继续执行；网格级同步没有直接的原语，通常通过内核结束（kernel boundary）隐式实现，或借助原子操作与全局内存标志位手动构造。跨块同步代价极高，应尽量避免。

除了显式屏障，原子操作（如atomicAdd）用于对全局内存或共享内存中的同一地址进行不可分割的读-改-写操作。原子操作保证了并发安全，但过度使用会造成严重的竞争冲突，导致串行化。例如，在归约（reduction）操作中，每个线程块先内部归约到共享内存，再由块0的线程原子累加到全局结果，这种两级策略能有效减少原子冲突。

// 块内归约 + 原子操作示例
__global__ void reduce_kernel(float *input, float *result, int n) {
    __shared__ float sdata[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + tid;
    
    sdata[tid] = (i < n) ? input[i] : 0.0f;
    __syncthreads();  // 确保所有线程完成共享内存写入
    
    // 树形归约
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) sdata[tid] += sdata[tid + s];
        __syncthreads();  // 防止读取未更新的数据
    }
    
    if (tid == 0) atomicAdd(result, sdata[0]);  // 跨块原子累加
}

案例：向量加法的完整实现与剖析

以最基础的向量加法y = x + y为例，展示线程层次与内存协作的完整流程。假设向量长度N=10000，选择块大小256，网格大小ceil(10000/256)=40。每个线程处理一个元素，直接使用全局内存读写。这是最简单的实现，但并非最优——它没有利用共享内存。

__global__ void vector_add(float *x, float *y, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        y[idx] = x[idx] + y[idx];
    }
}

int main() {
    int n = 10000;
    size_t bytes = n * sizeof(float);
    float *h_x, *h_y;
    float *d_x, *d_y;
    
    // 分配主机内存并初始化
    h_x = (float*)malloc(bytes);
    h_y = (float*)malloc(bytes);
    // ... 初始化数据 ...
    
    // 分配设备内存
    cudaMalloc(&d_x, bytes);
    cudaMalloc(&d_y, bytes);
    
    // 拷贝数据到设备
    cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, bytes, cudaMemcpyHostToDevice);
    
    // 启动内核
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;
    vector_add<<<grid_size, block_size>>>(d_x, d_y, n);
    
    // 拷贝结果回主机
    cudaMemcpy(h_y, d_y, bytes, cudaMemcpyDeviceToHost);
    
    // 释放内存
    cudaFree(d_x); cudaFree(d_y);
    free(h_x); free(h_y);
    return 0;
}

这个案例揭示了几个关键点：边界检查（if (idx < n)）是必须的，因为网格大小是向上取整的，可能超出实际数据范围；内存拷贝是显式的，主机与设备内存物理隔离；内核启动是异步的，cudaMemcpy会隐式同步，但多个内核之间的同步需要显式使用cudaDeviceSynchronize()。

误区与复盘：常见设计错误

• 误区一：块大小越大越好。实际上，块大小受限于SM的最大线程数（通常1024），且过大的块会降低占用率，因为每个SM能同时驻留的线程总数有限。块大小256～512通常能平衡资源利用与调度开销。
• 误区二：忽略共享内存的bank冲突。共享内存被划分为32个bank，当同一线程束内多个线程访问同一bank的不同地址时，会发生冲突并串行化。例如，访问sdata[tid * 2]会导致2路冲突。应尽量让相邻线程访问相邻地址。
• 误区三：过度使用原子操作。原子操作在低竞争下开销尚可，但在高竞争场景下性能急剧下降。应优先使用块内归约减少全局原子操作的次数。
• 误区四：忘记同步导致数据竞争。在共享内存中写入后立即读取，如果没有__syncthreads()，部分线程可能读到旧值。这是最难调试的并发错误之一，建议使用compute-sanitizer工具检测。