📚 课程目录第2章 GPU计算基础

GPU计算基础：GPU架构、CUDA核心与并行计算原理

2.1 从“多而慢”到“快而多”：GPU的设计哲学

CPU与GPU的根本差异，源于两者面向的负载类型不同。CPU擅长处理延迟敏感型任务——单条指令执行极快，分支预测、乱序执行、大容量缓存都是为了降低单线程延迟。而GPU面向的是吞吐量敏感型任务：不在乎某一条线程多快，而在乎每秒能处理多少条线程。因此，GPU将大量晶体管用于算术逻辑单元（ALU）而非控制逻辑与缓存，形成了“多核宽SIMD”的物理结构。

核心认知：CPU是“一个聪明人做复杂事”，GPU是“一万个普通人做简单事”。当任务可拆分为大量相互独立的小任务时，GPU的吞吐量优势可达CPU的数十倍。

2.2 CUDA核心与线程层次结构

NVIDIA GPU的最小计算单元是CUDA核心（也称流处理器）。一个CUDA核心在一个时钟周期内可执行一次浮点或整数运算。但CUDA核心并非独立工作，而是以流多处理器（SM）为单位组织。以NVIDIA A100为例，每个SM包含64个FP32 CUDA核心、32个FP64核心、4个Tensor Core，以及共享内存、寄存器文件和线程调度器。

CUDA的线程层次结构是理解并行编程的关键，从细到粗分为三层：

线程（Thread）： 最基础的执行单元，每个线程执行相同的核函数（Kernel）代码，但通过线程ID处理不同数据。
线程块（Block）： 一组线程（通常128~1024个），块内线程可通过共享内存通信，并可同步（__syncthreads()）。
网格（Grid）： 一个核函数启动时创建的所有线程块集合。网格中的块之间无通信，独立调度。

硬件执行时，SM以线程束（Warp）为基本调度单位——32个线程组成一个warp，SIMD式地执行同一条指令。若warp内出现分支分歧（部分线程走if，部分走else），则两个分支串行执行，这是GPU性能优化中需要避免的典型问题。

2.3 并行计算原理：延迟隐藏与吞吐量

GPU之所以能高效利用硬件资源，核心机制是延迟隐藏。当某个warp因访问全局内存而等待时（延迟约400~800个时钟周期），SM的线程调度器立即切换到另一个已就绪的warp执行计算。只要SM中有足够的活动warp（通常需要至少6~8个），内存延迟就被“掩盖”了，ALU始终处于忙碌状态。

一个典型的CUDA核函数执行流程如下：

// 向量加法示例：C = A + B，N = 1,000,000
__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // 全局线程ID
    if (i < N) {
        C[i] = A[i] + B[i];   // 每个线程处理一个元素
    }
}

// 主机端启动核函数
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

执行流程可拆解为：

步骤1： 主机将输入数据从CPU内存拷贝到GPU显存（cudaMemcpy）。
步骤2： 主机启动核函数，网格被分配到各SM。
步骤3： SM将线程块拆分为warp，调度器持续发射指令。
步骤4： 每个线程计算自己的全局索引，读取A[i]、B[i]，相加后写回C[i]。
步骤5： 核函数结束，结果拷回主机。

性能关键指标：占用率（Occupancy）= 活动warp数 / SM最大warp数。占用率越高，延迟隐藏能力越强。但过高的占用率可能导致寄存器溢出，需在两者间平衡。

2.4 案例：矩阵乘法的并行化

矩阵乘法是深度学习中最基础的计算（全连接层、卷积本质都是矩阵乘）。设C = A × B，其中A为M×K，B为K×N。串行实现需要M×N×K次乘加运算。GPU并行化的核心思路：每个线程计算C中的一个元素。

朴素实现中，每个线程读取A的一行和B的一列，做K次乘加。但这种方式对全局内存带宽消耗极大——每个线程重复读取A和B的数据。优化策略是使用共享内存分块（Tiling）：

// 分块矩阵乘法：每个线程块计算 16×16 的子矩阵
#define TILE_SIZE 16
__global__ void matMulTiled(float *A, float *B, float *C, int M, int K, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    for (int t = 0; t < (K + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // 协作加载A和B的分块到共享内存
        if (row < M && t * TILE_SIZE + threadIdx.x < K)
            As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE_SIZE + threadIdx.x];
        else
            As[threadIdx.y][threadIdx.x] = 0.0f;
        if (col < N && t * TILE_SIZE + threadIdx.y < K)
            Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            Bs[threadIdx.y][threadIdx.x] = 0.0f;
        __syncthreads();  // 确保分块加载完成
        
        for (int k = 0; k < TILE_SIZE; k++)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();  // 防止下一轮覆盖未用完的数据
    }
    if (row < M && col < N)
        C[row * N + col] = sum;
}

分块后，每个线程对全局内存的访问次数从2K次降至2K/TILE_SIZE次，大幅降低带宽压力。实际测试中，16×16分块比朴素实现快5~10倍。

2.5 常见误区与排查要点

误区	后果	正确做法
线程数越多越快	线程过多导致上下文切换开销，占用率过高寄存器溢出	根据SM资源（寄存器、共享内存）计算最优线程块大小
忽略内存合并访问	全局内存带宽利用率低至10%以下	相邻线程访问相邻地址（`threadIdx.x`对应连续内存）
频繁使用`if`分支	warp内分支分歧，串行执行多个路径	尽量让warp内线程走同一分支，或使用预计算掩码
忽略`__syncthreads()`位置	数据竞争，结果不确定	在共享内存读写之间正确放置同步点
主机与设备间频繁拷贝	PCIe带宽成为瓶颈（约16~32 GB/s）	尽量在设备端完成所有计算，减少拷贝次数

重要提醒：核函数启动是异步的，若需调试或计时，必须显式调用cudaDeviceSynchronize()。同时，务必检查每个CUDA API调用的返回值（cudaError_t），使用cudaGetLastError()捕获核函数启动错误，避免“静默失败”导致数据错误。

理解GPU架构与CUDA编程模型，是后续学习Tensor Core、CUDA库（cuBLAS、cuDNN）以及深度学习框架底层优化的基础。掌握本节内容后，你应能解释“为什么GPU适合并行计算”“CUDA线程如何映射到硬件”以及“如何初步评估一个核函数的性能瓶颈”。