4. CUDA编程模型入门

CUDA（Compute Unified Device Architecture）编程模型是GPU硬件加速与CPU协同设计的核心抽象层。它并非简单地将GPU视为一个协处理器，而是通过一套层次化的执行模型和内存模型，将GPU的大规模并行硬件能力（数千个CUDA核心、多级缓存、高带宽显存）映射为程序员可操作的逻辑结构。理解这一模型，是编写高性能异构程序、定位性能瓶颈、以及进行软硬件协同设计的基础。本章将深入剖析核函数（Kernel）的定义与生命周期、线程层次（Grid/Block/Warp）的硬件映射关系、以及四种关键内存（Global/Shared/Local/Constant）的特性与适用场景，并介绍主流的编译与调试工具链。

上图展示了CUDA编程模型的核心映射关系：CPU负责串行逻辑与调度，GPU通过Grid-Block-Warp三层结构组织海量线程，并配合Global/Shared/Local/Constant四种内存空间实现数据交换。理解这张图，是掌握后续所有性能优化技巧的前提。

4.1 核函数（Kernel）的定义与生命周期

核心原理：核函数是运行在GPU上的并行函数，使用 __global__ 修饰符声明。它由CPU（Host）端发起调用，但在GPU（Device）端执行。核函数必须返回 void，且不能是类的成员函数（除非使用静态成员）。当CPU调用核函数时，GPU会启动一个庞大的线程网格（Grid）来执行该函数体，每个线程执行相同的代码路径，但通过内置变量（如 threadIdx、blockIdx）区分不同的数据。

操作流程：一个典型的核函数生命周期包含以下步骤：

1. 数据准备：在Host端分配内存（malloc）并初始化数据。
2. 设备内存分配：使用 cudaMalloc 在GPU上分配显存空间。
3. 数据拷贝：通过 cudaMemcpy 将输入数据从Host内存复制到Device显存（方向为 cudaMemcpyHostToDevice）。
4. 核函数启动：使用三括号语法 kernel<<>>(args) 指定网格和块维度，异步启动核函数。
5. 同步与回拷：调用 cudaDeviceSynchronize() 等待核函数执行完毕，再将结果从Device拷贝回Host（cudaMemcpyDeviceToHost）。
6. 资源释放：使用 cudaFree 释放显存，使用 free 释放Host内存。

// 核函数定义：向量加法
__global__ void vectorAdd(const float *A, const float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

// Host 端调用
int main() {
    int N = 1 << 20;
    size_t bytes = N * sizeof(float);
    float *h_A, *h_B, *h_C;
    float *d_A, *d_B, *d_C;
    
    // 1. 分配Host内存并初始化
    h_A = (float*)malloc(bytes); h_B = (float*)malloc(bytes); h_C = (float*)malloc(bytes);
    // ... 初始化 h_A, h_B ...
    
    // 2. 分配Device显存
    cudaMalloc(&d_A, bytes); cudaMalloc(&d_B, bytes); cudaMalloc(&d_C, bytes);
    
    // 3. 拷贝数据到Device
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
    
    // 4. 启动核函数：1个Grid，256个线程/Block
    int threads = 256;
    int blocks = (N + threads - 1) / threads;
    vectorAdd<<<blocks, threads>>>(d_A, d_B, d_C, N);
    
    // 5. 同步并拷贝结果回Host
    cudaDeviceSynchronize();
    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.2 线程层次：Grid、Block与Warp的硬件映射

核心原理：CUDA的线程层次是理解GPU并行执行的关键。一个Grid包含多个Block，一个Block包含多个Thread。硬件上，GPU的调度器（SM，Streaming Multiprocessor）以Warp（32个线程）为最小调度单位。这意味着，一个Block中的线程会被划分为若干个Warp，每个Warp内的线程在同一个指令周期内执行相同的指令（SIMT，单指令多线程）。如果Warp内线程发生分支分歧（if-else），则所有分支路径都会被串行执行，导致性能下降。

判断方法：选择合适的Grid和Block维度是性能调优的第一步。以下是经验法则：

案例拆解：假设我们需要处理一个 1024×1024 的矩阵加法。如果采用 16×16 的Block（共256线程），则Grid维度为 64×64。每个Block负责一个 16×16 的数据块。在硬件上，这256个线程被划分为8个Warp（256/32）。SM会将这些Warp分配给不同的CUDA核心执行。如果矩阵维度不是Block维度的整数倍，则需要添加边界检查（如 if (row < height && col < width)），防止越界访问。

4.3 内存模型：Global、Shared、Local与Constant

核心原理：CUDA的内存模型决定了数据访问的速度和范围。理解各内存空间的物理位置和生命周期，是优化带宽和延迟的关键。

• Global Memory（全局内存）：位于显存（VRAM）中，容量最大（数GB），但延迟最高（数百个时钟周期）。所有线程（包括Host）均可访问。它用于存储需要长期保存或跨Block共享的数据。访问Global Memory时，应尽量满足合并访问（Coalesced Access）条件，即同一Warp的32个线程访问连续的内存地址，这样硬件可以将多次访问合并为一次或少数几次内存事务。
• Shared Memory（共享内存）：位于SM内部，容量小（通常几十KB），但延迟极低（约1-2个时钟周期）。它被同一个Block内的所有线程共享，常用于Block内部的数据复用和线程间通信。使用Shared Memory需要手动管理数据拷贝和同步（__syncthreads()），但能显著减少对Global Memory的访问次数。
• Local Memory（局部内存）：每个线程私有的内存空间。虽然逻辑上独立，但物理上位于显存中。当寄存器溢出（寄存器不够用）或访问动态索引的局部数组时，编译器会将数据放入Local Memory。频繁使用Local Memory会导致性能下降，应尽量避免。
• Constant Memory（常量内存）：只读内存空间，位于显存中，但具有专用的缓存。当所有线程访问同一个常量地址时，缓存会以广播方式发送给所有线程，速度极快。适合存储系数、权重等只读参数。

实战场景：在矩阵乘法中，如果直接使用Global Memory，每个输出元素需要读取A的一行和B的一列，会产生大量重复的Global访问。优化方法是：将A和B的子矩阵加载到Shared Memory中，然后从Shared Memory中读取数据进行计算。这样，每个数据只从Global读取一次，后续计算全部在Shared Memory中完成，计算访存比大幅提升。

// 使用 Shared Memory 优化矩阵乘法（片段）
__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 bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = by * BLOCK_SIZE + ty;
    int col = bx * BLOCK_SIZE + tx;
    
    float sum = 0.0f;
    for (int k = 0; k < N; k += BLOCK_SIZE) {
        // 协作加载 A 和 B 的子矩阵到 Shared Memory
        As[ty][tx] = A[row * N + k + tx];
        Bs[ty][tx] = B[(k + ty) * N + col];
        __syncthreads(); // 确保所有线程加载完成
        
        for (int i = 0; i < BLOCK_SIZE; i++) {
            sum += As[ty][i] * Bs[i][tx];
        }
        __syncthreads(); // 防止下一轮覆盖前，所有线程用完当前数据
    }
    C[row * N + col] = sum;
}

4.4 编译与调试工具链

核心原理：CUDA程序的编译和调试比传统CPU程序更复杂，因为它涉及Host和Device两套代码。NVIDIA提供了完整的工具链来支持开发。

操作流程：

1. 编译：使用 nvcc 编译器。它会将 .cu 文件中的Host代码（如 main）交给主机编译器（如 g++）编译，将Device代码（核函数）编译为PTX（并行线程执行）或SASS（流式汇编）代码。常用编译选项包括：-arch=sm_80（指定GPU架构）、-O3（优化级别）、-G（生成调试信息，会禁用优化）。
2. 调试：使用 cuda-gdb 或 Nsight 系列工具。cuda-gdb 是GDB的扩展，支持设置断点、查看GPU变量、检查Warp状态。Nsight Systems用于分析CPU/GPU时间线，定位同步瓶颈；Nsight Compute用于分析核函数内部的性能指标，如占用率、内存吞吐量、指令混合等。
3. 性能分析：使用 nvprof（旧）或 Nsight Systems（新）获取核函数执行时间、内存拷贝时间、SM利用率等数据。

案例拆解：假设我们运行一个向量加法程序，发现结果错误。使用 cuda-gdb 的步骤如下：

1. 使用 nvcc -G -g vectorAdd.cu -o vectorAdd 编译生成调试版本。
2. 启动 cuda-gdb ./vectorAdd。
3. 在核函数入口处设置断点：(cuda-gdb) break vectorAdd。
4. 运行程序：(cuda-gdb) run。
5. 当程序停在断点时，使用 info threads 查看当前活跃的线程，使用 print threadIdx.x 查看线程索引，使用 print A[0] 查看变量值。
6. 如果发现 i 的值超出范围，检查 blockIdx.x 和 blockDim.x 的计算逻辑。

通过上述方法，可以快速定位是索引计算错误、数据未同步还是内存越界等问题。