CUDA编程基础：线程层次、内存模型、核函数编写与编译

在进入TensorRT多流并发实战之前，必须先夯实CUDA编程的核心地基。多流并发的本质是让多个CUDA核函数在不同流中交错或并行执行，而这一切都建立在线程组织方式、内存访问模式以及核函数编写规范之上。本节将围绕这三个核心维度展开，帮助你建立编写高效CUDA代码所需的完整心智模型。

核心原理：线程层次与内存模型

CUDA采用单指令多线程（SIMT）执行模型，GPU以网格（Grid）、线程块（Block）、线程（Thread）三级层次组织并行任务。一个核函数启动时对应一个网格，网格由若干线程块组成，每个线程块包含若干线程。硬件层面，线程块被调度到流多处理器（SM）上执行，块内线程进一步划分为32个一组的线程束（Warp），这是硬件调度的最小单位。

理解线程层次的关键在于：块内线程通过共享内存和栅栏同步协作，块间线程仅能通过全局内存通信。这意味着算法设计时，应尽量将需要频繁协作的数据放在同一线程块内处理，减少跨块通信开销。

内存模型方面，CUDA提供多级存储体系，从寄存器、共享内存、全局内存、常量内存到纹理内存，每级内存的容量、延迟和可见性各不相同。下表总结了各级内存的关键属性：

内存访问的合并（Coalescing）是性能关键：同一线程束内32个线程访问连续内存地址时，硬件将其合并为少数几次内存事务，大幅降低访存延迟。反之，若线程访问分散地址，则产生多次事务，严重拖慢性能。

核函数编写规范与编译流程

核函数使用__global__修饰符声明，必须返回void，且只能从主机端调用。编写核函数时，通过内置变量threadIdx、blockIdx、blockDim、gridDim获取线程索引信息。一个典型的向量加法核函数如下：

__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 threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);

编译流程方面，nvcc编译器将.cu文件中的主机代码与设备代码分离：设备代码经ptxas编译为PTX中间表示，再生成SASS机器码；主机代码则交由宿主编译器（如gcc/clang）处理。编译命令通常为：

nvcc -arch=sm_80 -O3 -o vector_add vector_add.cu

其中-arch指定目标GPU架构（如sm_80对应Ampere架构），-O3开启优化。实际项目中，常使用-gencode选项为多种架构生成兼容代码，或使用compute_XX指定PTX版本以实现向前兼容。

案例：矩阵乘法中的线程与内存协同

以经典的分块矩阵乘法为例，展示线程层次与内存模型的协同设计。设矩阵尺寸为M×N×K，每个线程块计算16×16的输出子矩阵。块内线程首先协作将A和B的对应子块加载到共享内存，然后通过__syncthreads()同步，确保所有线程看到完整数据后再进行计算。

#define TILE_SIZE 16

__global__ void matMulShared(const float* A, const float* B, float* C, 
                             int M, int N, int K) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    for (int t = 0; t < (K + TILE_SIZE - 1) / TILE_SIZE; ++t) {
        // 协作加载到共享内存
        if (row < M && t * TILE_SIZE + threadIdx.x < K)
            As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE_SIZE + threadIdx.x];
        else
            As[threadIdx.y][threadIdx.x] = 0.0f;
            
        if (col < N && t * TILE_SIZE + threadIdx.y < K)
            Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            Bs[threadIdx.y][threadIdx.x] = 0.0f;
            
        __syncthreads();  // 确保加载完成
        
        for (int k = 0; k < TILE_SIZE; ++k)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
            
        __syncthreads();  // 防止覆盖共享内存
    }
    
    if (row < M && col < N)
        C[row * N + col] = sum;
}

该实现通过共享内存将全局内存访问次数从2×M×N×K降低到2×M×N×K/TILE_SIZE，大幅减少高延迟访存。同时，共享内存的bank冲突需要关注：当多个线程同时访问同一bank的不同地址时会产生冲突，可通过填充（padding）或调整访问模式规避。

常见误区与性能复盘

• 误区一：忽略边界检查。当数据总量不是线程块大小的整数倍时，未加if (idx < n)判断会导致越界访问，产生未定义行为甚至系统崩溃。
• 误区二：过度使用__syncthreads()。该函数是块级屏障，若在条件分支中使用不当（如部分线程进入分支而部分未进入），会导致死锁。
• 误区三：忽视内存合并。使用row-major存储时，若按列访问（如A[j * M + i]），线程束内访问地址不连续，性能可能下降数倍。
• 误区四：寄存器溢出。每个线程使用的寄存器超过硬件限制时，数据被压入局部内存（实际在全局内存中），性能急剧恶化。可通过__launch_bounds__或maxrregcount控制。
• 误区五：主机与设备间频繁拷贝。每次cudaMemcpy都有固定开销，应尽量批量传输，或使用零拷贝内存（cudaHostAllocMapped）和统一虚拟内存（UVM）减少拷贝。