4. GPU体系结构与CUDA编程入门

CPU擅长复杂控制与低延迟串行任务，而GPU的设计目标完全不同：它把大量晶体管投入到算术运算单元，用海量线程隐藏访存延迟，从而获得极高的吞吐量。理解GPU体系结构，是写出高效CUDA程序的前提。本节从硬件架构、线程层次、内存层次三个维度切入，最后落到核函数与线程索引的编程实践。

一、GPU硬件架构：从SM到CUDA Core

一块现代GPU由多个流多处理器（SM，Streaming Multiprocessor）组成，SM是GPU的基本调度与执行单元。每个SM内部包含若干CUDA Core（整数/浮点运算单元）、若干特殊功能单元（SFU）、寄存器文件、共享内存/L1缓存，以及负责线程束调度的Warp Scheduler。

GPU采用SIMT（单指令多线程）执行模型：32个线程组成一个线程束（Warp），同一Warp内的线程在同一周期执行同一条指令，但可以操作不同的数据。当Warp内线程因条件分支走向不同路径时，会发生分支发散（Divergence），硬件只能串行执行各分支路径，导致性能下降。因此，让同一Warp内的线程尽量走相同分支，是优化的重要原则。

GPU通过延迟隐藏而非降低延迟来提升性能：当一个Warp等待全局内存返回数据时，调度器立刻切换到其他就绪Warp继续执行。因此，占用率（Occupancy）——即SM上活跃Warp数与最大可支持Warp数之比——直接影响GPU能否充分隐藏延迟。

二、线程层次：Grid、Block、Thread

CUDA用三层结构组织线程：线程（Thread）→ 线程块（Block）→ 网格（Grid）。一个核函数启动时对应一个Grid，Grid被划分为多个Block，每个Block包含若干Thread。Block是独立调度的单位，同一个Block内的线程可以协作（通过共享内存和__syncthreads()同步），而不同Block之间无法直接同步。

线程索引由内置变量提供：threadIdx表示Block内线程编号，blockIdx表示Grid内Block编号，blockDim表示Block维度大小，gridDim表示Grid维度大小。对于一维索引，全局线程ID为：

int idx = blockIdx.x * blockDim.x + threadIdx.x;

对于二维数据（如图像），常用二维Block与二维Grid：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int offset = y * width + x;

三、内存层次：从寄存器到全局内存

CUDA内存层次按访问速度与作用域分为：寄存器（线程私有，最快）、共享内存（Block内共享，片上）、L1/L2缓存、全局内存（Device全局，容量大但延迟高）、常量内存和纹理内存。理解各层级的延迟差异，是性能优化的关键。

四、核函数与线程索引实战

核函数用__global__修饰，在主机端用<<<grid, block>>>语法启动。下面是一个向量加法的完整示例：

__global__ void vecAdd(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 main() {
    int N = 1 << 20;
    size_t bytes = N * sizeof(float);
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);
    // ... 拷贝数据到设备 ...
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
    cudaDeviceSynchronize();
    // ... 拷贝结果回主机 ...
}

这段代码体现了CUDA编程的标准流程：分配设备内存 → 拷贝数据到设备 → 启动核函数 → 同步 → 拷贝结果回主机 → 释放内存。其中gridSize向上取整保证覆盖所有元素，核函数内的if (idx < N)做边界保护。

五、常见误区与复盘

• Block大小随意设置：非32倍数会导致Warp内线程浪费，建议取128或256。
• 忽略边界检查：当N不是blockSize整数倍时，越界线程会读写非法地址，必须加if (idx < N)。
• 频繁主机-设备拷贝：PCIe带宽远低于显存带宽，应尽量让数据驻留设备端，减少拷贝次数。
• 共享内存滥用：共享内存虽快，但过度使用会降低Occupancy，需在复用收益与占用率之间权衡。
• 分支发散：同一Warp内线程走不同分支会串行化，尽量让分支条件与Warp边界对齐。