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

在构建全链路推理管线时，CUDA是连接深度学习框架与底层GPU硬件的关键桥梁。理解CUDA编程模型，不仅是为了写出能运行的核函数，更是为了后续优化TensorRT引擎、设计高效数据流时，能够准确判断性能瓶颈所在。本节从GPU硬件架构出发，逐步深入到线程组织、内存使用与编译流程，为后续的算子融合与引擎部署打下坚实基础。

3.1 GPU硬件架构：从SM到Warp

现代NVIDIA GPU（如Ampere、Hopper架构）的核心计算单元是流式多处理器（Streaming Multiprocessor, SM）。每个SM包含若干CUDA核心（FP32/INT32单元）、Tensor Core（用于矩阵乘加）、共享内存（Shared Memory）、寄存器堆以及调度器。GPU由多个SM组成，例如A100拥有108个SM，H100拥有132个SM。

执行模型的关键在于Warp：一个Warp由32个线程组成，是硬件调度的最小单位。SM以Warp为单位发射指令，所有线程在同一时刻执行同一条指令（SIMT模式）。若线程分支发散，则不同路径会被串行执行，导致性能下降。因此，编写核函数时应尽量避免线程间的分支分歧。

核心硬件参数速览（以A100为例）

SM数量：108
每个SM的CUDA核心：64个FP32 + 64个INT32
每个SM的Tensor Core：4个（第三代）
最大线程数/SM：2048
最大Warp数/SM：64
共享内存大小/SM：最多164KB（可配置）
寄存器文件大小/SM：256KB

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

CUDA采用三层线程组织结构，从大到小依次为：Grid（网格）→ Block（线程块）→ Thread（线程）。核函数启动时，由程序员指定Grid和Block的维度与大小。硬件将Block调度到SM上，Block内的线程通过共享内存和同步机制协作。

Thread：最小执行单元，拥有独立的寄存器和程序计数器。
Block：一组线程，可共享共享内存，通过__syncthreads()同步。同一Block内的线程必须能在同一SM上共存。
Grid：一组Block，Grid内的Block之间无法直接同步或共享内存，只能通过全局内存通信。

线程索引通过内置变量访问：threadIdx.x、blockIdx.x、blockDim.x。对于一维数据，全局线程ID的计算公式为：

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

对于二维或三维网格，需要结合gridDim和blockDim的y、z维度进行换算。合理选择Block大小（通常为128或256）能有效隐藏内存延迟，提升SM占用率。

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

CUDA内存模型是性能优化的核心。不同内存的访问速度与容量差异巨大，理解其特性是编写高效核函数的前提。

内存类型	作用域	生命周期	访问速度	典型用途
寄存器	单个线程	线程内	最快	局部变量、中间计算
共享内存	Block内所有线程	Block生命周期	极快（~20ns）	数据复用、Block内通信
全局内存	所有线程 + 主机	程序运行期间	慢（~400ns）	输入/输出数据、跨Block通信
常量内存	所有线程（只读）	程序运行期间	有缓存加速	只读参数、权重系数
纹理内存	所有线程（只读）	程序运行期间	有缓存加速	图像处理、插值访问

优化要点：全局内存访问应尽量合并（coalesced），即同一Warp的32个线程访问连续的内存地址，这样硬件可将多次访问合并为一次事务。共享内存应避免存储体冲突（bank conflict），否则访问会被串行化。

3.4 核函数编写与编译流程

核函数使用__global__修饰符定义，必须返回void。调用时使用<<<grid, block>>>语法指定执行配置。以下是一个向量加法的完整示例：

// 核函数定义
__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);
    // ... 拷贝结果回主机、释放内存 ...
}

编译流程分为两步：首先使用nvcc将.cu文件编译为PTX（并行线程执行）中间代码，再通过ptxas将PTX编译为SASS（SM汇编指令）。nvcc会调用主机编译器（如gcc）处理主机端代码。

# 编译为可执行文件
nvcc -arch=sm_80 -O3 vector_add.cu -o vector_add

# 查看生成的PTX代码
nvcc -arch=sm_80 -ptx vector_add.cu -o vector_add.ptx

# 查看SASS代码
cuobjdump -sass vector_add

其中-arch=sm_80指定目标架构为Ampere（A100），若需兼容多种GPU可使用-gencode选项生成多个架构的二进制码。

3.5 常见误区与性能陷阱

误区1：Block越大越好。Block过大会导致寄存器溢出（spill）到局部内存，反而降低性能。应通过cudaOccupancyMaxPotentialBlockSize计算最佳Block大小。

误区2：忽略内存合并。若线程访问全局内存时地址跳跃（如按列访问二维数组），会导致大量额外事务，性能可能下降数倍。

误区3：过度使用同步。__syncthreads()会强制Block内所有线程等待，若使用不当会造成死锁或性能浪费。

误区4：未检查错误。CUDA API和核函数执行是异步的，必须使用cudaGetLastError()和cudaDeviceSynchronize()检查错误。

3.6 案例：矩阵乘法的共享内存优化

以经典的矩阵乘法C = A × B为例，朴素实现中每个线程直接读取全局内存，重复访问严重。通过共享内存分块（tiling）可大幅减少全局内存访问次数：

__global__ void matMulShared(const float* A, const 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 tile = 0; tile < N / BLOCK_SIZE; ++tile) {
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * BLOCK_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * BLOCK_SIZE + threadIdx.y) * N + col];
        __syncthreads();
        
        for (int k = 0; k < BLOCK_SIZE; ++k) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

该实现将每个Block负责的矩阵块加载到共享内存，计算时从共享内存读取，全局内存访问次数从O(N³)降低到O(N³/BLOCK_SIZE)。在A100上，配合Tensor Core的wmma API可进一步获得数量级提升，这将在后续TensorRT章节中深入探讨。

掌握CUDA编程基础后，下一步需要理解如何将多个算子融合为单个核函数，以及如何利用TensorRT的层融合技术减少内核启动开销。这些优化手段将直接决定推理管线的端到端延迟与吞吐量。