GPU计算基础：GPU架构、并行计算概念与CUDA编程模型

在进入TensorRT的推理优化之前，必须先理解GPU为何能以远超CPU的吞吐量处理矩阵运算。CPU由少数几个高性能核心组成，每个核心拥有大容量缓存和复杂的分支预测单元，擅长处理串行、逻辑复杂的任务；而GPU则由数千个轻量级核心组成，它们共享控制逻辑和缓存，专为大规模数据并行计算设计。这种架构差异决定了两种处理器适合完全不同的工作负载。

GPU硬件架构：从SM到CUDA Core

NVIDIA GPU的基本计算单元是流式多处理器（Streaming Multiprocessor, SM）。每个SM内部包含若干CUDA Core（即FP32算术单元）、特殊函数单元（SFU）、共享内存（Shared Memory）和寄存器堆。以Ampere架构的A100为例，每个SM包含128个CUDA Core，整颗芯片共有108个SM，合计13824个CUDA Core。GPU执行任务时，以线程束（Warp）为最小调度单位——每个Warp包含32个线程，它们在同一时刻执行同一条指令（SIMT模式）。如果线程束内出现分支分歧，不同路径会被串行执行，因此应尽量避免Warp内的条件分支。

并行计算的核心概念

GPU并行计算遵循单指令多线程（SIMT）模型。与CPU的SIMD（单指令多数据）不同，SIMT允许每个线程拥有独立的寄存器状态和内存地址，只是执行相同的指令流。这种设计让程序员可以像编写串行代码一样编写并行逻辑，由硬件负责线程调度和资源分配。

并行计算的关键在于数据并行：将大规模数据切分为多个独立片段，每个线程处理一个片段。例如，对长度为100万的向量做逐元素加法，可以启动100万个线程，每个线程只负责一个元素的加法运算。这种“单线程单数据”的映射方式是最简单也最高效的并行模式。

内存层次结构是另一个必须理解的核心。GPU内存从快到慢、从容量小到大依次为：

内存类型	访问速度	容量	作用域	典型用途
寄存器	最快（1周期）	每个线程最多255个	单个线程	局部变量、中间计算
共享内存	约20~30周期	每个SM最多48~228KB	Block内所有线程	数据复用、Block内通信
全局内存	约400~600周期	数GB~数十GB	所有线程	主数据存储、Kernel间通信
常量/纹理内存	有缓存加速	64KB常量	所有线程（只读）	常量参数、只读查找表

共享内存是GPU性能优化的核心资源。由于它位于芯片内部，访问延迟远低于全局内存，且支持广播和合并访问。在矩阵乘法中，通过将子矩阵块加载到共享内存，可以显著减少对全局内存的重复访问，这是实现高计算吞吐的关键技术。

CUDA编程模型：从主机到设备

CUDA采用异构编程模型，将CPU视为主机（Host），GPU视为设备（Device）。主机负责逻辑控制和数据准备，设备负责大规模并行计算。一个典型的CUDA程序包含以下步骤：

1. 分配设备内存（cudaMalloc）
2. 将输入数据从主机拷贝到设备（cudaMemcpy，方向为HostToDevice）
3. 启动Kernel（使用<<<grid, block>>>语法指定线程配置）
4. 将计算结果从设备拷贝回主机（cudaMemcpy，方向为DeviceToHost）
5. 释放设备内存（cudaFree）

下面是一个向量加法的完整CUDA示例：

// 向量加法 Kernel
__global__ void vecAdd(float* a, 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 n = 1 << 20;  // 1,048,576 个元素
    size_t bytes = n * sizeof(float);
    
    // 1. 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    
    // 2. 拷贝输入数据到设备
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);
    
    // 3. 启动 Kernel：256 线程/Block，共 n/256 个 Block
    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vecAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);
    
    // 4. 拷贝结果回主机
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    
    // 5. 释放设备内存
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

Kernel启动时，blockIdx.x表示当前Block在网格中的索引，threadIdx.x表示线程在Block内的索引，blockDim.x表示Block的维度。通过公式 idx = blockIdx.x * blockDim.x + threadIdx.x 可以计算出全局唯一的线程ID。这种索引计算是CUDA编程中最基础也最重要的模式。

案例：矩阵乘法的并行化策略

矩阵乘法 C = A × B 是深度学习中最核心的运算。朴素实现中，每个线程计算C矩阵的一个元素：

__global__ void matMulNaive(float* A, float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; k++) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

这种实现虽然正确，但性能很差——每个线程需要从全局内存读取A的一行和B的一列，产生大量重复访问。优化方案是使用分块矩阵乘法（Tiled Matrix Multiplication）：将矩阵划分为小块（如16×16），每个Block负责一个小块的计算，先将子块加载到共享内存，再进行计算。这样，每个全局内存数据只被读取一次到共享内存，后续计算全部在片上完成，访存效率提升数十倍。

常见误区与复盘

误区一：线程越多性能越好。实际上，线程数量超过硬件可并行执行的极限后，多余线程只会增加调度开销。合理的线程配置应使每个SM的占用率达到50%~75%左右，留出足够的寄存器资源。

误区二：忽略内存拷贝开销。PCIe总线的带宽远低于GPU显存带宽。如果Kernel执行时间很短，而数据拷贝时间很长，整体性能可能反而不如CPU。解决方法是使用异步拷贝（cudaMemcpyAsync）和流水线技术，让数据拷贝与计算重叠执行。

误区三：在Kernel内使用printf或malloc。这些操作在GPU上代价极高，会严重拖慢执行速度。调试时应将数据拷贝回主机再输出，或使用cuda-gdb等专用调试工具。

误区四：忽视Warp发散。当Warp内32个线程执行不同的分支路径时，所有路径都会被串行执行，导致性能下降为原来的1/32。应尽量让Warp内的线程执行相同的控制流，或将分支条件设计为基于threadIdx的连续区间。

复盘一个实际案例：某团队在优化Transformer推理时，将LayerNorm的Kernel从128线程/Block改为256线程/Block，性能反而下降了15%。分析发现，LayerNorm的归约操作需要Block内线程同步，更大的Block导致同步开销增加，且共享内存使用量翻倍，降低了SM上可驻留的Block数量。最终调整为128线程/Block并配合向量化加载（float4），性能提升了40%。这说明线程配置必须结合具体算法的访存模式和同步需求来设计，不能盲目套用经验值。

理解GPU架构和CUDA编程模型是后续学习TensorRT的基础。TensorRT的层融合、精度校准、内核自动调优等核心功能，本质上都是在CUDA编程模型之上进行的自动化优化。只有深刻理解线程组织、内存层次和并行模式，才能在使用TensorRT时做出正确的性能分析和调优决策。