GPU架构基础：从硬件到编程模型的映射

核心原理：GPU为何“宽”而不“深”

CPU与GPU的设计哲学差异，决定了它们各自适用的计算场景。CPU拥有少量但强大的核心（通常4~16个），每个核心配备大容量缓存、分支预测器和复杂的乱序执行逻辑，目标是最大化单线程延迟的优化。GPU则相反，它集成了数千个轻量级核心，将大量晶体管用于计算单元而非控制逻辑，目标是最大化总吞吐量。这种“宽而浅”的架构，使得GPU在面对海量可并行数据时，能够以远高于CPU的速率完成运算。

理解GPU架构的关键，在于把握其层次化的组织方式。从物理硬件到编程模型，CUDA将GPU抽象为三个层级：线程（Thread）→ 线程块（Block）→ 网格（Grid）。这种抽象并非凭空设计，而是精确映射了GPU硬件的执行单元结构：线程映射到CUDA核心，线程块映射到流多处理器（SM），网格则对应整个GPU设备。

下图展示了从CUDA编程模型到GPU物理硬件的映射关系，以及各级存储器的位置与角色：

线程层次结构：从Grid到Thread的调度逻辑

CUDA的线程层次结构是编程模型的核心抽象。一个内核启动时，会创建一个Grid，Grid由多个Block组成，每个Block包含若干Thread。这种两级并行结构的设计初衷，是为了兼顾全局并行度与局部通信效率：Block内的线程可以高效协作（通过共享内存和同步屏障），而不同Block之间则相互独立，由硬件自动调度到不同的SM上执行。

在实际编程中，我们通过内核配置语法来指定Grid和Block的维度：

// 内核定义
__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

// 内核启动配置
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, N);

上述代码中，blockIdx.x、blockDim.x和threadIdx.x是CUDA提供的内置变量，分别表示当前Block在Grid中的索引、Block的维度大小以及当前Thread在Block中的索引。通过公式 idx = blockIdx.x * blockDim.x + threadIdx.x，我们可以将一维的线程索引映射到数据数组的对应位置。

硬件层面，Block是调度到SM的最小单位。当一个Block被分配到某个SM后，其中的线程会被进一步划分为Warp（通常32个线程为一组）。Warp是GPU实际执行指令的单位——同一个Warp中的所有线程在同一时刻执行同一条指令（SIMD风格），如果出现分支分歧，则不同路径会被串行执行，造成性能损失。

内存层次结构：各级存储器的角色与权衡

GPU的内存层次结构直接决定了数据访问的性能。理解各级存储器的特点，是编写高效CUDA程序的前提。下表总结了主要内存类型的关键属性：

在性能优化中，最核心的原则是最大化数据复用，减少全局内存访问。共享内存是程序员最常使用的优化手段——它相当于一个由程序员管理的缓存，允许Block内的线程以极低的延迟交换数据。例如，在矩阵乘法中，如果不使用共享内存，每个输出元素需要读取A的一行和B的一列，导致大量重复的全局内存访问；而使用共享内存分块（Tiling）后，每个数据块只从全局内存加载一次，后续计算全部在共享内存中完成，可将全局内存访问量减少一个数量级。

案例：向量求和中的内存访问优化

我们以一个简单的向量求和为例，展示如何利用共享内存减少全局内存访问。假设需要计算长度为N的数组所有元素之和，采用两阶段归约策略：

__global__ void reduceKernel(float *input, float *partialSums, int N) {
    __shared__ float sdata[256];  // 共享内存缓冲区
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tid;
    
    // 阶段1：每个线程加载一个元素到共享内存
    sdata[tid] = (idx < N) ? input[idx] : 0.0f;
    __syncthreads();  // 确保所有线程完成加载
    
    // 阶段2：块内归约（树形归约）
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            sdata[tid] += sdata[tid + stride];
        }
        __syncthreads();
    }
    
    // 每个Block输出一个部分和
    if (tid == 0) {
        partialSums[blockIdx.x] = sdata[0];
    }
}

该内核中，每个Block处理256个元素，通过__syncthreads()同步屏障确保共享内存中的数据在归约前完整就绪。最终，每个Block输出一个部分和，再由CPU或第二个内核将这些部分和累加得到最终结果。这种两阶段归约将全局内存访问次数从N次降低到N/256次（仅初始加载），大幅提升了内存带宽利用率。

误区与复盘：常见性能陷阱

if (idx % 2 == 0) {
    // 偶数路径
} else {
    // 奇数路径
}