3、CUDA编程入门：GPU架构概览、CUDA线程层次、内存层次、第一个CUDA程序（向量加法）

CUDA（Compute Unified Device Architecture）是NVIDIA推出的通用并行计算平台，它允许开发者利用GPU的大规模并行计算能力来加速通用计算任务。要真正掌握CUDA编程，首先需要理解GPU的硬件架构设计逻辑，以及CUDA如何通过线程层次和内存层次将硬件能力抽象为可编程模型。本章将从GPU架构的核心设计出发，逐步拆解CUDA的线程组织方式与内存访问模型，并通过一个完整的向量加法案例，带领你完成从编写、编译到运行、调试的第一个CUDA程序。

3.1 GPU架构核心原理：从SIMT到流式多处理器

GPU与CPU在设计哲学上存在根本差异。CPU采用控制逻辑复杂、缓存容量大、计算单元相对较少的架构，目标是最大化单线程执行速度；而GPU则采用大量计算单元（ALU）配合轻量级控制逻辑，通过极高的线程并行度来隐藏访存延迟。NVIDIA GPU的基本计算单元是流式多处理器（Streaming Multiprocessor, SM），每个SM内部包含数十个CUDA核心（即SP，Streaming Processor），以及共享内存、寄存器文件、调度器等资源。

CUDA采用SIMT（Single Instruction, Multiple Thread）执行模型。在SIMT模型中，32个线程组成一个线程束（Warp），这是硬件调度的最小单位。一个Warp中的所有线程在同一时刻执行同一条指令，但可以处理不同的数据。如果线程束内出现分支发散（即部分线程走if分支，部分走else分支），硬件会串行执行两个分支路径，导致性能下降。因此，编写CUDA代码时应尽量避免线程束内的分支发散。

现代GPU（如Ampere架构的A100）通常包含108个SM，每个SM拥有64个FP32 CUDA核心，总计约6912个核心。但实际可同时驻留的线程数远大于物理核心数——每个SM最多可驻留2048个线程（即64个Warp）。这种"物理核心少、逻辑线程多"的设计，使得GPU能够通过线程切换来掩盖内存访问延迟：当一个Warp因等待数据而阻塞时，调度器立即切换到另一个就绪的Warp执行。

3.2 CUDA线程层次：Grid、Block与Thread的协作模型

CUDA将线程组织为三层结构：网格（Grid）→ 线程块（Block）→ 线程（Thread）。这种层次化设计不仅便于编程抽象，更与硬件调度机制紧密对应。

• Thread（线程）：CUDA程序中的最小执行单元，每个线程拥有独立的程序计数器、寄存器状态和局部变量。线程通过内置变量 threadIdx.x 获取自己在块内的索引。
• Block（线程块）：由多个线程组成（通常为128~512个），块内线程可以通过共享内存（Shared Memory）协作，并通过栅栏同步（__syncthreads()）实现同步。块通过 blockIdx.x 和 blockDim.x 定位。
• Grid（网格）：由多个Block组成，一个Grid对应一个内核函数（Kernel）的启动。Grid中的Block之间无法直接通信，只能通过全局内存间接交互。

在硬件映射层面，一个Block会被调度到某个SM上执行，Block内的Warp由SM的Warp调度器分配。同一Block内的线程共享SM上的共享内存和寄存器资源，因此Block的尺寸和资源占用直接决定了SM能同时驻留多少个Block。例如，若每个Block使用32KB共享内存，而SM仅有100KB共享内存，则最多同时驻留3个Block。

在实际编程中，通常采用"网格-块-线程"的索引映射公式来计算全局线程ID：int tid = blockIdx.x * blockDim.x + threadIdx.x;。这个公式是CUDA编程中最基础、最常用的索引计算方式，务必熟练掌握。

3.3 内存层次：寄存器、共享内存与全局内存的权衡

CUDA的内存模型与硬件存储结构一一对应，理解各层级的容量、速度和访问权限是写出高性能代码的前提。下表总结了主要内存类型的关键特性：

在性能优化中，最核心的原则是减少全局内存访问次数。全局内存的延迟高达数百个时钟周期，而共享内存的延迟仅为其十分之一。因此，常见的优化策略是：将数据从全局内存加载到共享内存，在块内完成计算和复用，最后将结果写回全局内存。此外，内存合并（Memory Coalescing）也是关键——当同一Warp的32个线程访问连续的内存地址时，硬件会将这32次访问合并为一次或少数几次事务，大幅提升带宽利用率。

3.4 第一个CUDA程序：向量加法的完整实现与拆解

向量加法是CUDA的"Hello World"，它虽然简单，却完整展示了CUDA编程的核心流程：内存分配、数据拷贝、内核启动、结果回传。下面我们通过一个具体案例来演示。

场景设定：假设我们需要计算两个长度为N=1024的浮点向量A和B的和，得到向量C。在CPU上，这是一个简单的for循环；在GPU上，我们让每个线程负责计算一个元素。

完整代码（vector_add.cu）：

#include <stdio.h>
#include <cuda_runtime.h>

// 内核函数：每个线程计算一个元素
__global__ void vectorAdd(const float *A, const float *B, float *C, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {  // 边界检查，防止越界
        C[tid] = A[tid] + B[tid];
    }
}

int main() {
    int n = 1024;
    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;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);
    
    // 3. 将数据从主机拷贝到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
    
    // 4. 配置内核启动参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    // 5. 启动内核
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
    
    // 6. 将结果拷贝回主机
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
    
    // 7. 验证结果
    bool correct = true;
    for (int i = 0; i < n; i++) {
        if (h_C[i] != h_A[i] + h_B[i]) {
            correct = false;
            break;
        }
    }
    printf("结果: %s\n", correct ? "正确" : "错误");
    
    // 8. 释放内存
    free(h_A); free(h_B); free(h_C);
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    
    return 0;
}

编译与运行：使用NVIDIA提供的nvcc编译器：nvcc -o vector_add vector_add.cu，然后执行 ./vector_add。在编译时建议加上 -arch=sm_80（针对Ampere架构）或根据你的GPU型号调整计算能力参数。

案例拆解与关键点分析：

在这个案例中，我们使用了256个线程/块，因此需要 1024/256 = 4 个Block。如果N改为1000，则 blocksPerGrid = (1000+255)/256 = 4，但最后一个Block只有 1000 - 3*256 = 232 个有效线程，其余线程通过 if (tid < n) 跳过。这种"网格-块-线程"的映射模式是CUDA编程的通用范式。

3.5 常见误区与性能陷阱

在CUDA入门阶段，有几个高频错误值得特别警惕：

• 误区一：忽略错误检查。CUDA API大多返回错误码，但很多初学者不检查。建议封装一个 CHECK_CUDA_ERROR() 宏，在每次API调用后检查并打印错误信息。
• 误区二：线程数设置不合理。Block大小通常取32的倍数（如128、256、512），因为Warp大小为32。过小的Block（如8）会导致SM利用率低，过大的Block（如1024）可能超出硬件限制（每个Block最大1024线程）。
• 误区三：忽视内存合并。如果线程访问的全局内存地址不连续（如按列访问二维数组），会导致大量内存事务浪费。应尽量让相邻线程访问相邻地址。
• 误区四：在循环中频繁分配/释放设备内存。cudaMalloc/cudaFree开销较大，应尽量复用已分配的内存。
• 误区五：忘记同步。内核启动是异步的，如果后续代码依赖内核结果，需要显式调用 cudaDeviceSynchronize() 或使用 cudaMemcpy（它会隐式同步）。

最后，建议读者在完成向量加法后，尝试以下练习来巩固理解：①将N改为1000000，观察性能变化；②使用 cudaEvent 测量内核执行时间；③尝试用二维Grid/Block处理矩阵加法。这些练习将帮助你建立对CUDA执行模型更直观的感知。