4、GPU架构与CUDA编程入门：GPU硬件架构（SM、线程束）、CUDA编程模型（线程层次、内存层次）、第一个CUDA程序

GPU（Graphics Processing Unit）最初为图形渲染而设计，但其高度并行的计算能力使其成为通用高性能计算（GPGPU）的核心载体。理解GPU的硬件架构是编写高效CUDA程序的前提，而CUDA编程模型则提供了将硬件并行能力映射为软件线程层次的抽象机制。本章将深入剖析GPU的流式多处理器（SM）与线程束（Warp）机制，系统讲解CUDA的线程层次与内存层次，并通过一个完整的向量加法案例，带领读者完成从环境配置到性能分析的首个CUDA程序闭环。

上图清晰地展示了GPU硬件架构与CUDA软件模型之间的映射关系。左侧为物理硬件层，包含多个SM，每个SM内部有若干SP（Streaming Processor，流处理器）以及共享内存；右侧为CUDA编程模型，由Grid、Block、Thread三级层次构成，其中Warp是硬件调度的基本单位。理解这种映射关系，是后续进行性能优化和异构协同设计的基础。

4.1 GPU硬件架构核心原理：SM与线程束

GPU采用SIMT（Single Instruction, Multiple Thread，单指令多线程）执行模型，其硬件核心由大量SM组成。每个SM包含数十至上百个CUDA核心（即SP）、共享内存、寄存器文件、调度器以及特殊功能单元（SFU）。以NVIDIA Ampere架构的A100为例，每个SM包含64个FP32 CUDA核心、4个第三代Tensor Core、192KB可配置的L1/共享内存，以及4个调度器。SM是GPU执行计算任务的基本物理单元，一个线程块（Block）被调度到某个SM上执行后，该SM负责管理该Block内所有线程的生命周期。

线程束（Warp）是GPU硬件调度的最小单位，由32个线程组成。当一个Block被分配到SM后，SM将其中的线程划分为多个Warp。例如，一个包含256个线程的Block会被划分为8个Warp。硬件调度器以Warp为单位进行指令分发，这意味着同一个Warp内的32个线程在同一时刻执行相同的指令，但操作不同的数据。这种设计极大地简化了控制逻辑，将更多晶体管用于计算单元。然而，这也带来了一个关键问题：如果Warp内线程发生分支发散（即不同线程执行不同路径），则所有路径会被串行执行，导致性能下降。

在异构协同计算场景中，GPU的SM架构决定了其适合处理数据并行任务。例如，在图像处理中，每个像素的计算相互独立，可以映射到不同的线程上。而在CPU+GPU协同工作中，CPU负责逻辑控制和数据预处理，GPU则专注于大规模并行计算。理解SM的资源限制（如寄存器数量、共享内存大小）对于优化CUDA程序至关重要，因为资源使用量直接决定了SM上能同时驻留的Block数量，进而影响GPU的占用率（Occupancy）。

4.2 CUDA编程模型：线程层次与内存层次

CUDA编程模型采用三层线程层次结构：Grid（网格）、Block（线程块）、Thread（线程）。一个Grid包含多个Block，一个Block包含多个Thread。这种层次设计不仅与硬件架构（Block映射到SM，Thread映射到SP）完美对应，还提供了灵活的并行粒度控制。在编写内核函数（Kernel）时，开发者通过gridDim、blockDim、blockIdx和threadIdx四个内置变量来索引线程的全局位置。例如，要计算全局线程ID，可以使用公式：int tid = blockIdx.x * blockDim.x + threadIdx.x;。这种索引方式使得线程可以精确地访问对应的数据元素。

CUDA的内存层次是性能优化的核心。从速度最快到最慢依次为：寄存器（Register）、共享内存（Shared Memory）、全局内存（Global Memory）、常量内存（Constant Memory）和纹理内存（Texture Memory）。寄存器是每个线程私有的，访问延迟为1个时钟周期；共享内存是Block内线程共享的，延迟约为20-30个时钟周期，容量有限（通常为48KB-164KB）；全局内存是所有线程共享的，容量大（可达数十GB），但延迟高达400-600个时钟周期。因此，减少全局内存访问、充分利用共享内存是CUDA优化的关键策略。

在CPU与GPU协同的异构系统中，内存管理尤为重要。主机（CPU）端内存与设备（GPU）端内存是物理隔离的，需要通过PCIe总线进行数据拷贝。CUDA提供了cudaMalloc、cudaMemcpy、cudaFree等API来管理设备内存。此外，统一虚拟地址空间（UVA）和零拷贝内存（Zero-Copy Memory）技术可以简化数据共享，但需要权衡性能。对于多核异构系统，合理划分数据在CPU内存、GPU全局内存和共享内存中的分布，是提升整体性能的关键。

4.3 第一个CUDA程序：向量加法实战

向量加法是CUDA的"Hello World"，虽然简单，但完整展示了CUDA编程的核心流程。我们以一个包含N=1024个元素的浮点向量加法为例，演示从主机端代码编写、内核函数设计到编译运行的全过程。该案例的目标是：在GPU上并行计算C[i] = A[i] + B[i]，并与CPU串行结果进行正确性验证。

首先，我们需要编写内核函数（Kernel Function）。内核函数使用__global__修饰符声明，在设备端执行。每个线程处理一个元素，通过线程索引定位数据位置：

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

接下来是主机端代码。主机端负责分配内存、初始化数据、配置内核执行参数并启动内核。关键步骤包括：使用cudaMalloc为设备端数组分配显存，使用cudaMemcpy将数据从主机拷贝到设备，配置Grid和Block的维度（例如blockSize = 256，gridSize = (N + blockSize - 1) / blockSize），然后调用内核。最后，将结果从设备拷贝回主机，并释放显存。

int main() {
    int numElements = 1024;
    size_t size = numElements * sizeof(float);
    
    // 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);
    
    // 初始化数据
    for (int i = 0; i < numElements; ++i) {
        h_A[i] = rand() / (float)RAND_MAX;
        h_B[i] = rand() / (float)RAND_MAX;
    }
    
    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    
    // 拷贝数据到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    
    // 配置内核执行参数
    int blockSize = 256;
    int gridSize = (numElements + blockSize - 1) / blockSize;
    
    // 启动内核
    vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, numElements);
    
    // 拷贝结果回主机
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    
    // 验证结果
    bool correct = true;
    for (int i = 0; i < numElements; ++i) {
        if (fabs(h_C[i] - (h_A[i] + h_B[i])) > 1e-5) {
            correct = false;
            break;
        }
    }
    printf("结果验证: %s\n", correct ? "通过" : "失败");
    
    // 释放内存
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);
    return 0;
}

编译时，需要使用nvcc编译器：nvcc -o vectorAdd vectorAdd.cu。运行程序后，输出"结果验证: 通过"。这个简单案例揭示了CUDA编程的几个关键点：第一，内核函数中的边界检查if (i < numElements)是必要的，因为Grid大小可能大于实际数据长度；第二，cudaMemcpy是同步操作，会阻塞主机直到拷贝完成；第三，内核启动是异步的，主机代码会继续执行，因此需要cudaMemcpy来确保内核执行完毕。

4.4 常见误区与总结复盘

在CUDA编程入门阶段，开发者常陷入以下误区，导致程序运行错误或性能不佳。

• 误区一：忽略错误检查。 CUDA API调用和内核启动都可能失败（如显存不足、非法内存访问），但默认情况下错误是异步的，不会立即中断程序。正确做法是使用cudaGetLastError()检查每个API调用的返回值，并使用cudaDeviceSynchronize()同步等待内核执行完毕后再检查错误。
• 误区二：线程索引计算错误。 在多维Grid/Block场景下，混淆blockIdx.x、blockIdx.y、threadIdx.x等索引，导致数据访问越界或重复计算。建议在编写内核时，先画出索引映射图，明确每个线程负责的数据范围。
• 误区三：过度使用全局内存。 频繁访问全局内存会因高延迟严重拖慢性能。应尽量将数据加载到共享内存或寄存器中复用。例如，在矩阵乘法中，通过分块（Tiling）技术将子矩阵加载到共享内存，可减少全局内存访问次数达数十倍。
• 误区四：忽视Warp发散。 当Warp内线程因条件分支执行不同指令时，所有分支路径会被串行执行，导致效率下降。应尽量使分支条件基于threadIdx的连续范围，或使用预计算和查找表避免发散。
• 误区五：主机与设备间频繁数据传输。 PCIe传输带宽远低于GPU内部带宽，频繁的小批量数据传输会严重降低整体性能。应尽量在设备端完成数据处理，减少主机-设备交互次数。

总结复盘本章内容，我们首先剖析了GPU的SM与Warp硬件架构，理解了SIMT执行模型和资源限制；其次系统学习了CUDA的线程层次（Grid-Block-Thread）和内存层次（寄存器-共享-全局等），掌握了索引计算和内存选择的原则；最后通过向量加法案例，完整实践了CUDA程序的开发流程。在异构协同计算的大背景下，GPU并非万能，其优势在于大规模数据并行，而CPU擅长逻辑控制和复杂分支。实际项目中，需要根据任务特性合理划分CPU与GPU的工作负载，并通过异步操作、内存复用等技术手段，实现计算与通信的重叠，从而最大化系统吞吐量。下一阶段的学习应聚焦于共享内存的深入使用、原子操作、CUDA流与事件，以及性能分析工具（如NVIDIA Nsight）的应用，逐步构建高性能异构计算的核心能力。