GPU架构与CUDA编程模型

要真正理解CUDA，首先需要回答一个根本问题：GPU为什么能比CPU快那么多？答案并不在于“频率更高”或“核心更多”这么简单，而在于两者设计哲学的根本差异。CPU是为低延迟优化的——它拥有庞大的缓存、复杂的乱序执行逻辑和分支预测器，目的是让单条指令尽快完成；而GPU是为高吞吐量优化的——它把大量晶体管用于计算单元（ALU），用极简单的控制逻辑和较小的缓存，同时运行成千上万个线程来掩盖访存延迟。这种“以数量换延迟”的策略，使得GPU在数据并行任务上拥有数量级的优势。

2.1 GPU硬件架构：从SM到CUDA Core

NVIDIA GPU的基本计算单元是流多处理器（Streaming Multiprocessor, SM）。一个GPU由多个SM组成，每个SM内部又包含若干CUDA Core（即浮点/整数运算单元）、共享内存、寄存器文件、调度器和缓存。以Ampere架构的A100为例，它拥有108个SM，每个SM包含64个FP32 CUDA Core，总计6912个核心。但请注意：“核心数”并不等于“线程数”，CUDA Core是物理执行单元，而线程是软件抽象，一个CUDA Core可以按时间片执行多个线程的指令。

GPU执行指令的最小调度单位是Warp——通常由32个线程组成。每个Warp内的线程在同一时刻执行同一条指令（SIMT，单指令多线程），如果出现分支分歧（即线程走不同路径），则Warp会串行执行各分支，导致性能损失。一个SM中的Warp调度器负责在每个时钟周期选择一个可执行的Warp发射指令。下图展示了GPU硬件与CUDA线程层次的对应关系：

上图展示了GPU硬件与CUDA软件线程层次的对应关系。每个Block（线程块）会被调度到某个SM上执行，一个SM可以同时驻留多个Block（取决于资源限制）。Block内的线程通过共享内存和同步屏障（__syncthreads()）协作；而不同Block之间无法直接通信，只能通过全局内存间接交互。这种设计使得GPU可以透明地扩展——无论硬件上有多少个SM，同一个CUDA程序都能运行，只是并行度不同。

2.2 CUDA线程层次：Grid、Block与Thread

CUDA的线程组织是三层结构：Grid（网格）→ Block（线程块）→ Thread（线程）。启动内核时，程序员需要指定Grid和Block的维度，每个维度最多3层（x, y, z）。例如：

// 定义网格和块的大小
dim3 grid(2, 2);      // 2×2的网格，共4个Block
dim3 block(4, 4);     // 每个Block包含4×4=16个线程
kernel<<<grid, block>>>(args);  // 启动内核

上述代码共启动4个Block × 16个线程 = 64个线程。每个线程在核函数内部通过内置变量threadIdx、blockIdx、blockDim、gridDim来定位自己的全局唯一ID。对于一维情况，全局线程ID的计算公式为：

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

这个公式是CUDA编程中最基础、最常用的索引方式。对于二维或三维情况，需要分别计算各维度的索引，再组合成线性偏移。理解这个映射关系是编写正确内核的第一步。

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

CUDA的内存层次与硬件架构紧密对应，不同内存的速度、容量、作用域差异巨大。下表总结了各类内存的关键特性：

在编写内核时，应优先使用寄存器存储频繁访问的局部变量；当Block内需要数据共享时，使用共享内存可以避免重复访问全局内存；而全局内存的访问应尽量合并（coalesced）——即让Warp内32个线程访问连续的内存地址，这样硬件可以将多次访问合并为一次事务，大幅提升带宽利用率。

2.4 核函数编写与执行：一个完整的向量加法案例

下面通过一个经典的向量加法案例，完整展示核函数的编写、编译和执行流程。假设我们要计算 C = A + B，其中A、B、C均为长度为N的浮点数组。

// 核函数定义：每个线程处理一个元素
__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) {  // 边界检查，防止越界
        C[tid] = A[tid] + B[tid];
    }
}

// 主机端调用
int main() {
    int N = 1 << 20;  // 1,048,576 个元素
    size_t bytes = N * sizeof(float);

    // 1. 分配主机内存并初始化
    float *h_A = (float*)malloc(bytes);
    float *h_B = (float*)malloc(bytes);
    float *h_C = (float*)malloc(bytes);
    // ... 初始化 h_A, h_B ...

    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);

    // 3. 将数据从主机拷贝到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

    // 4. 配置内核启动参数并启动
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;  // 向上取整
    vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

    // 5. 将结果拷贝回主机
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

    // 6. 释放设备内存
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);
    return 0;
}

执行流程可以归纳为六个步骤：分配主机内存 → 分配设备内存 → 主机到设备拷贝 → 启动内核 → 设备到主机拷贝 → 释放内存。其中第4步的<<<gridSize, blockSize>>>语法是CUDA特有的内核启动配置，它告诉运行时需要创建多少个Block、每个Block有多少线程。

2.5 常见误区与调试建议

初学者在编写CUDA程序时，最容易犯以下几类错误：

• 忘记边界检查：当N不能被blockSize整除时，最后一个Block的部分线程会越界访问。必须像示例中那样加上 if (tid < N) 判断。
• 忽略错误检查：CUDA API调用（如cudaMalloc、cudaMemcpy）和内核启动都可能失败，但默认不会抛出异常。建议使用 cudaGetLastError() 或定义宏来检查每个API调用的返回值。
• 主机与设备内存混淆：在核函数中访问主机指针（如h_A）会导致程序崩溃。必须通过cudaMemcpy在主机和设备之间显式传输数据。
• 过度使用同步：在Block内使用__syncthreads()是必要的，但如果在条件分支中使用不当（如只有部分线程执行到同步点），会导致死锁。
• 忽视内存合并访问：如果线程访问的地址是跳跃的（如 A[tid * 2]），会导致内存带宽利用率下降数倍。应尽量让相邻线程访问相邻地址。

调试CUDA程序时，推荐使用 cuda-memcheck 工具检测越界和内存泄漏，使用 Nsight Systems 和 Nsight Compute 进行性能分析。在开发阶段，可以在核函数中临时加入 printf() 输出（CUDA支持设备端printf），但要注意大量输出会严重影响性能。

掌握GPU架构与CUDA编程模型是后续学习TensorRT的基础。TensorRT的层融合、张量并行等优化策略，本质上都是在利用GPU的SM架构和内存层次特性。理解Warp的执行方式、共享内存的容量限制、全局内存的带宽瓶颈，将帮助你更好地理解TensorRT为何能实现数倍甚至数十倍的推理加速。