CUDA编程基础：GPU架构、线程层次、内存层次、核函数编写与编译

CUDA（Compute Unified Device Architecture）是NVIDIA推出的通用并行计算平台。要编写高效的CUDA程序，首先需要理解GPU的硬件架构、线程组织方式、内存体系，以及核函数（Kernel）的编写与编译流程。本节将围绕这四个核心主题展开，帮助你建立CUDA编程的底层认知框架。

核心原理：GPU架构与线程层次

GPU采用SIMT（单指令多线程）执行模型，即同一时刻，大量线程执行同一条指令，但操作各自不同的数据。现代NVIDIA GPU（如Ampere、Hopper架构）由多个流多处理器（SM）组成，每个SM内部包含若干CUDA核心、共享内存、寄存器文件和调度器。

CUDA的线程层次结构分为三层：线程（Thread）→ 线程块（Block）→ 网格（Grid）。线程是执行的最小单位；线程块由多个线程组成，块内线程可通过共享内存协作，并能通过__syncthreads()同步；网格由多个线程块组成，块之间无法直接同步，只能通过全局内存或原子操作间接通信。

关键硬件映射关系：

线程块被调度到SM上执行，一个SM可同时驻留多个线程块。
线程块内的线程被进一步划分为Warp（通常32个线程），Warp是GPU调度和执行的基本单位。
同一Warp内的线程必须执行相同指令，若出现分支分歧（divergence），则串行执行各分支路径，导致性能下降。

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

CUDA内存模型是性能优化的核心。按访问速度和物理位置，内存层次从慢到快依次为：全局内存（Global Memory）→ 本地内存（Local Memory）→ 共享内存（Shared Memory）→ L1/L2缓存 → 寄存器（Register）。全局内存容量最大（可达数十GB），但延迟高达数百周期；共享内存容量小（每SM通常几十KB到上百KB），但延迟极低（约20-30周期）；寄存器速度最快，但数量有限（每线程最多255个）。

内存类型	作用域	生命周期	访问延迟	典型用途
全局内存	所有线程	程序运行期间	400-600周期	主数据存储、跨块通信
共享内存	块内线程	块执行期间	20-30周期	数据复用、块内协作
寄存器	单个线程	线程执行期间	1周期	局部变量、中间计算
常量内存	所有线程	程序运行期间	缓存命中时≈寄存器	只读参数、系数表

优化要点：当多个线程访问全局内存中相邻地址时，硬件会合并为一次事务（Memory Coalescing），大幅提升带宽。因此，设计数据布局时应尽量让相邻线程访问相邻内存地址。

核函数编写与编译流程

核函数是运行在GPU上的函数，使用__global__修饰符声明。编写核函数需遵循以下规则：

返回类型必须为void，不能使用可变参数。
通过内置变量threadIdx、blockIdx、blockDim、gridDim获取线程索引。
核函数不能调用主机端函数（如printf除外，但需谨慎使用）。
核函数是异步的，主机端调用后立即返回，需用cudaDeviceSynchronize()等待完成。

编译流程如下：nvcc编译器将.cu文件中的主机代码与设备代码分离，设备代码编译为PTX（并行线程执行）中间表示，再通过ptxas生成SASS（流处理器汇编）代码。最终链接生成可执行文件。

// 向量加法核函数示例
__global__ void vectorAdd(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 threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
    cudaDeviceSynchronize();
    // ... 回收资源 ...
    return 0;
}

案例：矩阵乘法优化

以矩阵乘法C = A × B为例，展示线程层次与内存层次的实际应用。朴素实现中，每个线程计算C的一个元素，需要读取A的一行和B的一列，导致大量重复的全局内存访问。通过分块（Tiling）技术，将矩阵划分为子块，利用共享内存缓存子块数据，可显著减少全局内存访问次数。

// 使用共享内存的矩阵乘法核函数（分块大小为16×16）
__global__ void matMulShared(const float* A, const float* B, float* C, int N) {
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = by * 16 + ty;
    int col = bx * 16 + tx;
    float sum = 0.0f;
    
    for (int k = 0; k < (N + 15) / 16; ++k) {
        if (row < N && k * 16 + tx < N)
            As[ty][tx] = A[row * N + k * 16 + tx];
        else
            As[ty][tx] = 0.0f;
        if (col < N && k * 16 + ty < N)
            Bs[ty][tx] = B[(k * 16 + ty) * N + col];
        else
            Bs[ty][tx] = 0.0f;
        __syncthreads();
        
        for (int i = 0; i < 16; ++i)
            sum += As[ty][i] * Bs[i][tx];
        __syncthreads();
    }
    
    if (row < N && col < N)
        C[row * N + col] = sum;
}

该实现将全局内存访问量从O(N³)降低到O(N³/16)，在N=1024时性能提升可达10倍以上。实际项目中，还需结合向量化加载（float4）、无Bank冲突的共享内存布局、循环展开等技巧进一步优化。

误区与复盘

常见误区一：认为线程块越大越好。实际上，线程块大小受SM资源限制（寄存器、共享内存），过大的块会导致占用率下降，反而降低性能。通常256-512线程/块是较优区间。

常见误区二：忽略Warp分歧。在核函数中写if (threadIdx.x % 2 == 0)这类分支，会导致Warp内一半线程执行，另一半等待，性能减半。应尽量让分支条件基于Warp边界（如threadIdx.x / 32）。

常见误区三：忘记同步。使用共享内存时，若未在写入后调用__syncthreads()，会产生数据竞争，导致结果不确定。调试此类问题极为困难，建议在开发初期就养成同步习惯。

复盘建议：在编写CUDA程序时，先用cudaOccupancyMaxActiveBlocksPerMultiprocessor查询理论占用率，再用nvprof或Nsight Compute分析实际瓶颈。优先优化全局内存访问模式，其次考虑共享内存复用，最后才调整线程块大小。记住：GPU性能优化的本质是隐藏延迟——通过足够的并行线程掩盖内存访问延迟，而不是减少计算量。