GPU计算基础：GPU架构、CUDA核心、显存层次与并行计算模型

核心原理：从“多核”到“众核”的架构分野

要理解GPU为何能加速AI与高性能计算，首先需要看清CPU与GPU在硬件设计哲学上的根本差异。CPU由少数几个（通常8~64个）功能强大的核心构成，每个核心拥有大容量缓存、复杂的分支预测与乱序执行单元，擅长处理延迟敏感、逻辑复杂的串行任务。而GPU则采用“众核”架构，以NVIDIA A100为例，它集成了6912个CUDA核心，每个核心都非常精简，没有复杂的控制单元，依赖大规模线程并行来隐藏访存延迟。

这种差异源于两种芯片的目标不同：CPU追求单线程性能最大化，GPU追求总吞吐量最大化。GPU将更多晶体管用于浮点运算单元（FP32/FP64/Tensor Core），而非控制逻辑和缓存。下图展示了GPU流式多处理器（SM）内部的典型结构，这是理解CUDA编程模型的关键硬件视图。

CUDA核心与线程层次模型

CUDA（Compute Unified Device Architecture）是NVIDIA推出的并行计算平台和编程模型。在CUDA中，线程（Thread）是最小的执行单位，多个线程组成线程块（Block），多个线程块组成网格（Grid）。硬件执行时，GPU以Warp（32个线程）为基本调度单位，一个Warp中的线程在同一个SM上执行相同的指令（SIMT，单指令多线程）。

这种层次化设计对性能有决定性影响：

• 线程块内的线程可以通过共享内存（Shared Memory）通信，并通过__syncthreads()同步。
• 不同线程块之间无法直接通信，只能通过全局内存（Global Memory）间接交互。
• Warp divergence（分支发散）：若Warp内线程执行不同分支，则分支会被串行执行，导致性能下降。

在编写内核函数（Kernel）时，开发者需要显式指定网格和线程块的维度。例如，一个典型的向量加法内核启动配置如下：

// 内核定义
__global__ void vecAdd(float *A, float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    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);

上述代码中，blockIdx.x、blockDim.x、threadIdx.x是CUDA内置变量，分别表示当前线程块的索引、线程块维度大小和线程在线程块内的索引。这种“网格-线程块-线程”的三层结构，使得GPU能够灵活地将海量线程映射到有限的SM上执行。

显存层次结构与数据搬运策略

GPU的显存层次是性能优化的核心战场。从寄存器到全局内存，每一级的容量、延迟和带宽差异巨大。下表总结了A100各存储层次的典型参数：

在实际优化中，最核心的原则是减少对全局内存的访问。全局内存虽然带宽很高，但延迟是共享内存的20倍以上。常用的优化手段包括：

1. 数据复用：将需要反复使用的数据块先加载到共享内存，例如矩阵分块乘法的tile策略。
2. 合并访问（Coalesced Access）：让同一Warp的32个线程访问连续的内存地址，使硬件将多次访问合并为一次或少数几次内存事务。
3. 使用向量化类型：如float4，一次读取16字节，减少指令数。

并行计算模型：从数据并行到任务并行

GPU最擅长的是数据并行（Data Parallelism）——同一操作应用于大量数据元素。典型的例子包括图像像素处理、矩阵元素运算、神经网络中的张量操作等。CUDA的编程模型天然适合数据并行：每个线程处理一个或多个数据元素。

除了数据并行，GPU也支持任务并行（Task Parallelism），即不同线程块执行不同的内核或不同的代码路径。但任务并行在GPU上通常效率较低，因为GPU的调度器倾向于将资源分配给同质的Warp。因此，实际工程中应优先考虑数据并行，将任务分解为可并行处理的数据块。

一个经典的并行计算模式是归约（Reduction），例如计算一个大型数组的总和。朴素实现会让每个线程读取一个元素并累加，但更好的做法是采用分块归约：

__global__ void reduce(float *input, float *output, 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 stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            sdata[tid] += sdata[tid + stride];
        }
        __syncthreads();
    }
    
    // 每个线程块输出一个部分和
    if (tid == 0) output[blockIdx.x] = sdata[0];
}

上述归约算法将时间复杂度从O(n)降低到O(log n)的并行步数，同时通过共享内存减少了全局内存访问次数。在实际应用中，还需要考虑Warp级归约（使用__shfl_down_sync指令）来进一步减少同步开销。

案例：矩阵乘法优化中的显存与并行策略

矩阵乘法（GEMM）是深度学习与科学计算中最核心的运算之一。以C = A × B为例，其中A为M×K矩阵，B为K×N矩阵。朴素实现中，每个线程计算C的一个元素，需要读取A的一行和B的一列，导致大量重复的全局内存访问。

优化的关键思路是分块（Tiling）：将矩阵划分为大小为BLOCK_SIZE × BLOCK_SIZE的子块，每个线程块负责计算一个输出子块。计算过程中，将A和B对应的子块加载到共享内存，实现数据复用。以下是一个简化版本的核心逻辑：

#define BLOCK_SIZE 16

__global__ void matMulTiled(float *A, float *B, float *C, int M, int N, int K) {
    __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 < (K + BLOCK_SIZE - 1) / BLOCK_SIZE; tile++) {
        // 协作加载A和B的子块到共享内存
        if (row < M && tile * BLOCK_SIZE + threadIdx.x < K)
            As[threadIdx.y][threadIdx.x] = A[row * K + tile * BLOCK_SIZE + threadIdx.x];
        else
            As[threadIdx.y][threadIdx.x] = 0.0f;
            
        if (col < N && tile * BLOCK_SIZE + threadIdx.y < K)
            Bs[threadIdx.y][threadIdx.x] = B[(tile * BLOCK_SIZE + threadIdx.y) * N + col];
        else
            Bs[threadIdx.y][threadIdx.x] = 0.0f;
            
        __syncthreads();
        
        // 计算当前子块的乘积
        for (int k = 0; k < BLOCK_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }
    
    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}

这个分块矩阵乘法将全局内存访问次数从O(M×N×K)降低到O(M×N×K/BLOCK_SIZE)，显著提升了计算访存比。在A100上，配合Tensor Core使用wmma或mma指令，可以达到接近理论峰值的性能。

常见误区与性能陷阱