3. GPU编程模型

核心原理：异构计算与线程层次模型

GPU编程模型的核心思想是将CPU作为主机（Host），负责逻辑控制与串行任务；GPU作为设备（Device），承担大规模并行计算。CUDA与OpenCL均采用单指令多线程（SIMT）执行模型，但真正决定性能上限的是其线程层次结构与内存层次结构的配合方式。

CUDA将线程组织为三层结构：网格（Grid）→ 线程块（Block）→ 线程（Thread）。一个Grid包含多个Block，每个Block包含多个Thread。硬件调度以Block为单位分配到SM（流式多处理器），而实际执行以Warp（32线程）为最小调度单位。OpenCL则对应为NDRange → Work-Group → Work-Item，概念完全对应，只是命名不同。

理解线程层次的关键在于：Block内的线程可以通过共享内存协作，并支持同步操作；而Block之间无法直接通信，只能通过全局内存间接交互。这一约束决定了并行算法的分解方式——数据分块必须与Block边界对齐。

方法流程：CUDA编程五步法与内存管理

CUDA/OpenCL程序的开发遵循固定流程，核心是数据搬运与内核执行的分离。以CUDA为例，标准流程如下：

1. 分配设备内存：调用 cudaMalloc 在显存中分配输入/输出缓冲区。
2. 主机→设备拷贝：使用 cudaMemcpy 将输入数据从CPU内存复制到GPU显存。
3. 配置内核启动参数：指定Grid维度和Block维度，如 kernel<<<gridDim, blockDim>>>(args)。
4. 执行内核：GPU按Block调度到SM，线程并行执行核函数。
5. 设备→主机回拷：将计算结果复制回CPU内存，并释放设备内存。

内存管理是性能调优的核心战场。CUDA内存体系按速度从快到慢排列为：寄存器 → 共享内存 → L1/L2缓存 → 全局内存 → 主机内存。其中共享内存（__shared__）是用户可控的片上存储，延迟约为全局内存的1/20，但容量有限（通常48KB/SM）。OpenCL对应为 local 内存，需在核函数参数中显式声明大小。

// CUDA 向量加法示例：展示线程索引计算与内存分配
__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 main() {
    float *d_a, *d_b, *d_c;
    int n = 1 << 20;  // 1M 元素
    size_t bytes = n * sizeof(float);
    
    cudaMalloc(&d_a, bytes);   // 步骤1
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);  // 步骤2
    
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;
    vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);  // 步骤3-4
    
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);  // 步骤5
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

案例实战：矩阵乘法优化中的内存复用

以经典的分块矩阵乘法为例，展示如何利用共享内存减少全局内存访问。朴素实现中，每个输出元素需要读取A的一行和B的一列，全局内存访问量为 2×N³ 次。采用分块（Tiling）策略后，每个Block负责计算一个 BLOCK_SIZE × BLOCK_SIZE 的输出子矩阵，将A和B的对应分块加载到共享内存，计算复用次数提升 BLOCK_SIZE 倍。

#define BLOCK_SIZE 16
__global__ void matMulTiled(float *A, float *B, float *C, int N) {
    __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 k = 0; k < N; k += BLOCK_SIZE) {
        // 协作加载分块到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + k + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(k + threadIdx.y) * N + col];
        __syncthreads();  // 确保分块加载完成
        
        for (int i = 0; i < BLOCK_SIZE; i++)
            sum += As[threadIdx.y][i] * Bs[i][threadIdx.x];
        __syncthreads();  // 防止覆盖共享内存
    }
    C[row * N + col] = sum;
}

该案例中，__syncthreads() 是Block内同步的关键，防止线程间数据竞争。当 BLOCK_SIZE=16 时，每个Block有256个线程，共享内存占用 2×16×16×4B=2KB，远低于48KB限制。实际测试中，相比朴素版本，分块优化可获得5~10倍加速比，主要收益来自全局内存访问量减少至 2×N³/BLOCK_SIZE。

误区与复盘：线程配置与内存访问的常见陷阱