第一章：CUDA编程模型回顾

各位同学好，我是你们的CUDA讲师。今天咱们来聊聊CUDA编程模型的基础。说实话，这部分内容看起来简单，但很多坑都藏在这些基础里。我见过不少项目，就是因为对模型理解不透彻，性能死活上不去。

GPU架构概览：从硬件到软件

先说说GPU的硬件架构。你想想看，CPU和GPU最大的区别是什么？CPU是「大核少核」，GPU是「小核多核」。NVIDIA的GPU核心叫CUDA Core，一个SM（Streaming Multiprocessor）里通常有几十到上百个CUDA Core。

我个人习惯把GPU想象成一个「工厂」：

• SM 就是车间，每个车间有自己的生产线
• CUDA Core 就是工人，负责具体干活
• Warp 是32人一组的小队，一起行动

嗯，这里要注意：Warp是GPU调度的基本单位。32个线程组成一个Warp，它们执行相同的指令。如果遇到分支，那就麻烦了——部分线程执行，部分线程等待，这就是所谓的「Warp Divergence」。我在项目中遇到过这种情况，一个简单的if-else语句，性能直接腰斩。

线程层次结构：Grid-Block-Thread

CUDA的线程组织方式，说白了就是三级结构：

• Grid：一个Kernel启动对应一个Grid
• Block：Grid由多个Block组成，Block之间不能通信
• Thread：Block由多个Thread组成，Thread之间可以共享数据

为什么会这样设计？我刚开始学的时候也纳闷。后来做项目才明白：这是为了适配GPU的硬件架构。Block对应SM，Thread对应CUDA Core。一个Block里的线程跑在同一个SM上，所以可以共享数据。不同Block可能跑在不同SM上，通信代价太大。

来看个简单的例子：

// 定义线程块和网格大小
dim3 blockDim(16, 16);  // 每个Block有16x16=256个线程
dim3 gridDim(32, 32);   // 网格有32x32=1024个Block

// Kernel启动
myKernel<<<gridDim, blockDim>>>(d_data, N);

// Kernel内部获取线程索引
__global__ void myKernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idy = blockIdx.y * blockDim.y + threadIdx.y;
    int globalIdx = idy * gridDim.x * blockDim.x + idx;
    
    if (globalIdx < N) {
        data[globalIdx] = data[globalIdx] * 2.0f;
    }
}

内存层次结构：Global/Shared/Register/Local

GPU的内存层次，说白了就是「速度越快的，容量越小」。我画了个图，方便你理解：

这张图我每次讲课都会放。你看，从寄存器到全局内存，延迟差了两个数量级。所以优化的核心思路就是：尽量用快的，少用慢的。

具体来说：

• 寄存器：每个线程私有，速度最快。但数量有限，一个线程最多255个寄存器。用多了会 spill 到本地内存，性能暴跌。
• 共享内存：同一个Block内的线程共享。我经常用它来做数据复用，比如矩阵分块计算。
• 全局内存：所有线程都能访问，但速度最慢。访问时要尽量合并访问（Coalesced Access）。
• 本地内存：寄存器不够用时，编译器自动把数据放到本地内存。说白了就是全局内存，只是逻辑上是私有的。

CUDA程序编译与执行流程

CUDA程序的编译，说白了就是「两步走」：

1. 主机代码（CPU部分）：用普通C++编译器编译
2. 设备代码（GPU部分）：用NVCC编译器编译成PTX或cubin

NVCC的工作流程是这样的：

源文件(.cu) 
    → 分离主机代码和设备代码
    → 设备代码编译成PTX（中间表示）
    → PTX再编译成cubin（二进制）
    → 最终打包成可执行文件

嗯，这里有个细节：PTX是中间表示，类似于Java的字节码。不同架构的GPU（比如Turing、Ampere、Hopper）有不同的cubin格式。所以NVCC会生成多个版本的cubin，运行时根据GPU型号选择加载。

执行流程就更直观了：

1. CPU把数据从主机内存拷贝到GPU全局内存
2. CPU启动Kernel，GPU开始执行
3. GPU执行完毕后，CPU把结果拷贝回来

最后，给你个编译命令的参考：

nvcc -arch=sm_80 -O3 -o my_program my_program.cu

# 常用选项说明：
# -arch=sm_80：针对Ampere架构优化
# -O3：最高优化级别
# -G：调试模式（会降低性能）
# --ptxas-options=-v：查看寄存器使用情况

好了，这一章的内容就到这里。记住：CUDA编程模型是基础中的基础，理解透了才能写出高性能的算子。下一章我们会深入实战，到时候见。