3. CUDA编程入门：核函数、线程层次、内存模型

CUDA（Compute Unified Device Architecture）是NVIDIA提供的并行计算平台与编程模型。要真正理解CUDA程序如何运行，必须抓住三个核心概念：核函数（Kernel）、线程层次（Thread Hierarchy）和内存模型（Memory Model）。这三者共同决定了GPU上并行任务的组织方式、执行效率以及数据访问的带宽表现。

3.1 核函数：运行在GPU上的入口函数

核函数是使用__global__修饰符声明的函数，它由CPU（主机端）调用，但在GPU（设备端）上执行。核函数具有以下关键特征：

必须返回void类型，不能直接返回数据到主机端。
调用时使用三括号语法<<<gridDim, blockDim>>>指定执行配置。
核函数内部可以访问线程索引（threadIdx、blockIdx）来区分不同线程的工作。
核函数中不能调用主机端函数（如printf在部分架构受限，malloc需使用设备端版本）。

// 一个简单的向量加法核函数
__global__ void vectorAdd(float *A, float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

// 主机端调用
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);

上述代码中，每个线程处理一个元素的加法。当n不是threadsPerBlock的整数倍时，通过if (i < n)进行边界检查，避免越界访问。

易错点：核函数是异步执行的。调用kernel<<<...>>>()后，CPU会立即继续执行后续代码，不会等待GPU完成。如需同步，必须调用cudaDeviceSynchronize()或使用cudaMemcpy（该操作隐式同步）。

3.2 线程层次：Grid-Block-Thread三级结构

CUDA采用网格（Grid）→ 线程块（Block）→ 线程（Thread）的三级层次结构。这种设计既保证了硬件可扩展性（不同GPU的SM数量不同），又提供了逻辑上的灵活性。

Thread：最小的执行单元，每个线程拥有独立的程序计数器和寄存器状态。
Block：一组线程的集合，块内线程可以通过共享内存（Shared Memory）协作，并通过栅栏同步（__syncthreads()）保持一致性。
Grid：由多个Block组成，Grid内的Block之间没有直接通信机制（除非通过全局内存和原子操作）。

线程索引的计算公式为：全局线程ID = blockIdx.x * blockDim.x + threadIdx.x。对于二维或三维网格，需要分别计算各维度索引。

  硬件映射关系：Block会被调度到GPU的流多处理器（SM）上执行。一个SM可以同时驻留多个Block，但Block内的线程数不能超过SM的最大线程容量（通常为1024或2048）。因此，Block大小通常设置为128、256或512，以平衡占用率和资源消耗。

3.3 内存模型：从寄存器到全局内存的层级

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

内存类型	作用域	生命周期	访问速度	典型用途
寄存器（Register）	单个线程	线程执行期间	最快（~1周期）	局部变量、中间计算结果
共享内存（Shared Memory）	单个Block	Block执行期间	快（~20周期）	Block内线程协作、数据复用
全局内存（Global Memory）	所有线程 + 主机	程序运行期间	慢（~400-800周期）	主数据存储、CPU-GPU通信
常量内存（Constant）	所有线程（只读）	程序运行期间	快（有缓存时）	核函数参数、查找表
纹理内存（Texture）	所有线程（只读）	程序运行期间	快（有缓存）	图像处理、空间局部性访问

全局内存访问是性能瓶颈的主要来源。当线程访问连续的内存地址时，硬件会将多个访问合并为一次事务（内存合并）。因此，应尽量让相邻线程访问相邻地址。

// 合并访问（推荐）：threadIdx.x 对应连续地址
float val = data[blockIdx.x * blockDim.x + threadIdx.x];

// 非合并访问（不推荐）：线程间地址跳跃
float val = data[threadIdx.x * blockDim.x + blockIdx.x];

3.4 案例：矩阵乘法中的内存优化

以经典的矩阵乘法C = A × B为例，展示内存模型的实际应用。朴素实现中，每个线程直接读取全局内存中的A和B，会产生大量重复访问。

// 朴素版本：每个线程计算C[i][j]
__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;
    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;
}

优化思路：利用共享内存分块（Tiling）。将矩阵划分为子块，每个Block加载一个子块到共享内存，然后反复复用，大幅减少全局内存访问次数。

  分块矩阵乘法流程：
  将A和B划分为TILE_SIZE × TILE_SIZE的子块。
每个Block负责计算一个输出子块，协作加载对应的A子块和B子块到共享内存。
使用__syncthreads()确保共享内存写入完成后再读取。
循环累加所有子块的乘积结果。

#define TILE_SIZE 16

__global__ void matMulTiled(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = by * TILE_SIZE + ty;
    int col = bx * TILE_SIZE + tx;
    float sum = 0.0f;
    
    for (int tile = 0; tile < N / TILE_SIZE; tile++) {
        // 协作加载到共享内存
        As[ty][tx] = A[row * N + tile * TILE_SIZE + tx];
        Bs[ty][tx] = B[(tile * TILE_SIZE + ty) * N + col];
        __syncthreads();
        
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[ty][k] * Bs[k][tx];
        }
        __syncthreads(); // 防止下一轮覆盖前被读取
    }
    C[row * N + col] = sum;
}

该优化将全局内存访问次数从O(N³)降低到O(N³ / TILE_SIZE)，在N=1024时，访问量减少64倍，性能提升显著。

3.5 常见误区与调试建议

误区1：忽略边界检查。当数据长度不是Block大小的整数倍时，未加if (i < n)会导致越界访问，可能引发未定义行为或程序崩溃。务必在核函数中检查索引范围。

误区2：过度使用共享内存。共享内存容量有限（通常48KB~164KB），过大的TILE_SIZE会导致每个SM驻留的Block数减少，降低占用率。应根据具体GPU调整TILE_SIZE（常用16或32）。

误区3：忘记同步。在共享内存读写之间、以及Block内线程协作时，必须调用__syncthreads()。遗漏同步会导致数据竞争，且错误难以复现。

调试建议：使用cuda-memcheck（或compute-sanitizer）检测越界和未初始化内存访问；使用cudaError_t err = cudaGetLastError()检查核函数启动是否成功；在开发阶段可先用小数据量（如N=32）验证正确性，再扩展到全量数据。

掌握核函数、线程层次和内存模型是CUDA编程的基石。核函数定义了并行任务的入口，线程层次提供了灵活的组织方式，而内存模型则决定了数据访问的效率。在实际开发中，三者需要协同考虑：根据算法特性选择合适的Block大小，利用共享内存减少全局访问，并通过边界检查保证安全性。这些基础概念将直接影响到后续TensorRT推理引擎中自定义算子的实现效率。