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

3.1 核心原理：异构计算与CUDA编程模型

CUDA（Compute Unified Device Architecture）是英伟达提出的通用并行计算平台与编程模型，其本质是异构计算——将CPU作为主机（Host），负责逻辑控制与串行任务；将GPU作为设备（Device），负责大规模数据并行计算。CUDA编程模型的核心思想是：把数据并行问题映射到GPU的海量线程上，让数千个线程同时执行相同的指令，处理不同的数据元素。

与传统的CPU多线程编程不同，CUDA线程的创建、调度与切换由硬件高效完成，线程开销极低。程序员只需描述“做什么”（核函数），而无需关心底层线程如何被调度到流多处理器（SM）上。这种“单指令多线程”（SIMT）的执行模式，使得GPU能够以极高的吞吐量处理图形渲染、科学计算、深度学习等任务。

3.2 核函数与线程层次

核函数（Kernel）是CUDA编程的核心，使用 __global__ 修饰符声明，在设备端执行，由主机端调用。核函数必须返回 void，且不能作为类的成员函数。调用核函数时使用特殊的执行配置语法 <<<gridDim, blockDim>>> 来指定线程组织方式。

CUDA的线程层次采用两级结构：线程（Thread）组成线程块（Block），线程块组成网格（Grid）。每个线程有唯一的全局ID，可通过内置变量 threadIdx、blockIdx、blockDim 计算得出。这种层次化设计使得线程管理既灵活又高效——线程块内部的线程可以通过共享内存协作，而不同线程块之间则相互独立。

// 向量加法示例：C = A + B
__global__ void vecAdd(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 main() {
    int n = 1024;
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;
    
    // 分配设备内存并拷贝数据...
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, n);
    // 等待GPU执行完成并拷贝结果...
    return 0;
}

在上述代码中，gridSize 表示网格中线程块的数量，blockSize 表示每个线程块中的线程数。当 n = 1024、blockSize = 256 时，gridSize = 4，共启动 4×256 = 1024 个线程，每个线程处理一个元素。这种“一个线程对应一个数据元素”的映射模式是CUDA编程中最基本的并行范式。

3.3 内存管理与数据搬运

CUDA设备拥有独立于主机的显存（Global Memory），所有数据必须显式地在主机与设备之间拷贝。内存管理是CUDA编程中最容易出错的环节，也是性能优化的关键突破口。下表总结了CUDA中常见的内存类型及其特点：

标准的内存管理流程包含四个步骤：分配设备内存 → 主机数据拷贝至设备 → 执行核函数 → 结果拷贝回主机。CUDA提供了 cudaMalloc、cudaMemcpy、cudaFree 等API来完成这些操作。此外，cudaMallocManaged 创建的托管内存（Unified Memory）允许CPU和GPU通过指针直接访问同一块内存，由驱动程序自动管理数据迁移，大幅简化了编程复杂度。

3.4 案例实战：向量加法的完整实现

下面通过一个完整的向量加法案例，展示CUDA编程的标准流程。该案例包含错误检查、内存管理、核函数调用和结果验证等关键环节。

#include <cstdio>
#include <cuda_runtime.h>

#define CHECK_CUDA(call) do { \
    cudaError_t err = (call); \
    if (err != cudaSuccess) { \
        printf("CUDA错误: %s (行 %d)\n", cudaGetErrorString(err), __LINE__); \
        return -1; \
    } \
} while(0)

__global__ void vecAdd(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 main() {
    const int n = 1 << 20;  // 1,048,576 个元素
    size_t bytes = n * sizeof(float);
    
    // 1. 分配主机内存并初始化
    float *h_A = (float*)malloc(bytes);
    float *h_B = (float*)malloc(bytes);
    float *h_C = (float*)malloc(bytes);
    for (int i = 0; i < n; i++) { h_A[i] = i * 1.0f; h_B[i] = i * 2.0f; }
    
    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    CHECK_CUDA(cudaMalloc(&d_A, bytes));
    CHECK_CUDA(cudaMalloc(&d_B, bytes));
    CHECK_CUDA(cudaMalloc(&d_C, bytes));
    
    // 3. 主机数据拷贝到设备
    CHECK_CUDA(cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice));
    CHECK_CUDA(cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice));
    
    // 4. 配置并启动核函数
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, n);
    CHECK_CUDA(cudaGetLastError());  // 检查核函数启动错误
    
    // 5. 结果拷贝回主机
    CHECK_CUDA(cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost));
    
    // 6. 验证结果
    bool ok = true;
    for (int i = 0; i < n; i++) {
        if (h_C[i] != h_A[i] + h_B[i]) { ok = false; break; }
    }
    printf("向量加法%s\n", ok ? "验证通过" : "验证失败");
    
    // 7. 释放内存
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);
    return 0;
}

编译时使用 nvcc -o vecAdd vecAdd.cu 命令。运行程序后，GPU会启动 4096 个线程块（每个块256线程），共 1,048,576 个线程并行执行加法操作。通过 CHECK_CUDA 宏可以捕获CUDA API调用中的错误，这是生产级代码中不可或缺的防御性编程手段。

3.5 误区与复盘：CUDA编程的常见陷阱

• 误区一：忽视边界检查。当数据总量不是线程块大小的整数倍时，最后一个线程块的部分线程会越界访问。必须在核函数中添加 if (i < n) 这样的边界判断。
• 误区二：过度使用同步。 __syncthreads() 是线程块内同步的屏障，但滥用会导致线程等待，降低并行效率。不同线程块之间无法同步，也不应依赖执行顺序。
• 误区三：频繁的CPU-GPU数据传输。PCIe带宽远低于显存带宽，每次 cudaMemcpy 都有固定开销。应尽量减少传输次数，或在GPU上完成尽可能多的计算。
• 误区四：忽略错误检查。CUDA API调用和核函数启动都可能失败（如显存不足、配置错误），不检查返回值会让程序在后续运行中产生难以排查的异常。
• 误区五：线程块大小选择不当。线程块大小通常设为 128～512 的倍数，过小无法隐藏内存延迟，过大会导致寄存器溢出或占用率下降。