3. CUDA编程基础：编程模型、线程层次与内存管理

3.1 核心原理：从GPU架构到CUDA编程模型

CUDA（Compute Unified Device Architecture）是NVIDIA推出的通用并行计算平台。理解CUDA编程模型，首先需要建立异构计算的心智模型：CPU作为主机（Host）负责逻辑控制与串行任务，GPU作为设备（Device）负责大规模数据并行计算。两者通过PCIe总线或NVLink互联，各自拥有独立的存储空间。

CUDA编程模型的核心抽象是核函数（Kernel）——一个由 __global__ 修饰符声明、在设备端执行的函数。当主机调用核函数时，GPU会启动大量线程并行执行同一段代码，这种模式称为单指令多线程（SIMT）。与CPU多核不同，GPU线程的创建、调度与切换由硬件完成，开销极低，因此可以创建数十万个线程来充分利用计算资源。

3.2 线程层次与索引映射

CUDA将线程组织为三层结构：网格（Grid）→ 线程块（Block）→ 线程（Thread）。核函数启动时，通过执行配置 <<<gridDim, blockDim>>> 指定网格和线程块的维度。每个维度均可为一维、二维或三维，分别通过 blockIdx、threadIdx 内建变量访问。

在深度学习场景中，最常见的是一维或二维索引。例如处理一张 N×C×H×W 的特征图时，通常将 blockIdx.x 映射到输出通道或批次维度，将 threadIdx.x 映射到空间位置。计算全局线性索引的公式为：

// 一维网格 + 一维线程块
int tid = blockIdx.x * blockDim.x + threadIdx.x;

// 二维网格 + 二维线程块（处理图像）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;   // 行主序展平

线程块的最大尺寸受硬件限制（当前为1024线程/块），而网格维度上限为 2^31-1。因此，当数据量超过单块容量时，需要采用网格-步长循环（Grid-Stride Loop）模式，让每个线程处理多个元素：

for (int i = tid; i < N; i += blockDim.x * gridDim.x) {
    result[i] = input[i] * 2.0f;
}

3.3 内存管理与数据搬运

GPU拥有独立于CPU的内存空间。CUDA提供以下内存管理API：

API 函数	功能	注意事项
cudaMalloc	在设备端分配全局内存	返回设备指针，不能直接解引用
cudaMemcpy	主机与设备间数据拷贝	需指定方向（H2D / D2H / D2D）
cudaFree	释放设备内存	必须与 cudaMalloc 成对使用
cudaMallocManaged	分配统一内存（可被CPU/GPU访问）	依赖GPU的UVM硬件支持
cudaHostAlloc	分配页锁定主机内存	提高DMA传输效率，但占用物理内存

一个标准的CUDA程序流程如下：

1. 分配主机内存并初始化数据（malloc 或 new）；
2. 分配设备显存（cudaMalloc）；
3. 将输入数据从主机拷贝到设备（cudaMemcpy，方向 cudaMemcpyHostToDevice）；
4. 配置核函数执行参数并启动核函数；
5. 将计算结果从设备拷贝回主机（cudaMemcpyDeviceToHost）；
6. 释放所有分配的内存（cudaFree + free）。

3.4 核函数编写与编译流程

核函数使用 __global__ 修饰符声明，必须返回 void。编写核函数时需注意：

• 核函数不能调用主机端函数（如 printf 在计算能力3.5+可用，但仅限调试）；
• 核函数支持递归（计算能力3.5+），但深度受限且性能差，应避免；
• 核函数内不能使用 static 变量或 C++ 标准库容器；
• 参数通过值传递，结构体参数需注意对齐。

以下是一个完整的向量加法示例：

// vecAdd.cu
#include <cstdio>
#include <cuda_runtime.h>

__global__ void vecAddKernel(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];
    }
}

int main() {
    const int N = 1 << 20;  // 1M 元素
    size_t bytes = N * sizeof(float);

    float *h_A = new float[N];
    float *h_B = new float[N];
    float *h_C = new float[N];
    for (int i = 0; i < N; i++) { h_A[i] = i * 0.5f; h_B[i] = i * 1.5f; }

    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);

    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

    int threads = 256;
    int blocks = (N + threads - 1) / threads;
    vecAddKernel<<<blocks, threads>>>(d_A, d_B, d_C, N);

    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

    // 验证结果...
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    delete[] h_A; delete[] h_B; delete[] h_C;
    return 0;
}

编译时使用 nvcc 编译器。nvcc会将 .cu 文件中的主机代码交给宿主C++编译器（如g++），设备代码则编译为PTX或SASS。常用编译选项：

# 基础编译（默认计算能力）
nvcc -o vecAdd vecAdd.cu

# 指定架构（Ampere A100 = sm_80, Hopper H100 = sm_90）
nvcc -arch=sm_80 -O3 -o vecAdd vecAdd.cu

# 包含调试信息
nvcc -g -G -o vecAdd_debug vecAdd.cu

3.5 案例：矩阵乘法的朴素实现与性能瓶颈

以 C = A × B（均为 1024×1024）为例，朴素实现让每个线程计算输出矩阵的一个元素：

__global__ void matMulKernel(const float* A, const float* B, float* C, int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; k++) {
            sum += A[row * K + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

该实现虽然正确，但性能极差。原因在于内层循环每次迭代都从全局内存读取 A[row*K+k] 和 B[k*N+col]，而全局内存延迟高达数百周期。每个线程需要执行 2×K 次全局内存访问，导致计算单元大量空闲等待。

3.6 误区与复盘

误区一：将GPU当作"更快的CPU"。GPU的并行优势来源于大量线程的延迟隐藏，而非单线程执行速度。如果核函数内存在大量分支发散（warp divergence）或串行依赖，性能可能远低于CPU。

误区二：忽略主机与设备间的数据传输开销。PCIe 4.0 x16的理论带宽约32GB/s，而A100的显存带宽超过2TB/s。如果核函数计算量小但数据搬运频繁，整体性能将受限于PCIe带宽。应尽量减少H2D/D2H拷贝次数，或将多个操作融合到单个核函数中。

误区三：线程块尺寸选择随意。线程块大小应为32的倍数（一个warp的大小），通常选择128或256。过小的块无法充分隐藏延迟，过大的块可能导致寄存器溢出或占用率下降。

误区四：忽视错误检查。CUDA错误是异步的——核函数启动后，错误可能在下一次同步API调用时才被捕获。建议在每次核函数启动后调用 cudaGetLastError()，并在关键节点使用 cudaDeviceSynchronize() 同步。

复盘清单：编写CUDA程序时，按以下顺序自查：

1. 数据规模是否足够大（通常需 >10万元素）以抵消启动开销？
2. 内存访问是否合并（coalesced）？相邻线程是否访问相邻地址？
3. 是否存在bank conflict（共享内存访问冲突）？
4. 核函数是否包含不必要的 if 分支或 printf？
5. 是否使用了 -O3 和正确的 -arch 编译选项？
6. 是否用 nvprof 或 Nsight Compute 验证了实际占用率和内存吞吐？