📚 课程目录第1章 CUDA编程模型回顾

1. CUDA编程模型回顾

各位同学好，我是你们的老朋友。今天咱们聊聊CUDA编程模型的基础。说实话，这部分内容看起来简单，但很多人在实际项目中栽跟头，就是因为基础没打牢。我当年刚接触GPU编程时，也踩过不少坑，今天把这些经验分享给大家。

1.1 GPU架构基础

GPU和CPU的设计哲学完全不同。CPU是「全能选手」，什么活都能干，但一次只能干几件事。GPU呢？它是「偏科生」，擅长同时干成千上万件简单的事。

你想想看，一个典型的GPU有几千个计算核心。我习惯把GPU想象成一个大型工厂，里面有几千个工人。每个工人能力一般，但人多力量大。CPU则像几个顶尖专家，单兵作战能力极强。

具体到NVIDIA的GPU架构，有几个关键概念：

SM（Streaming Multiprocessor）：流式多处理器，是GPU的核心计算单元。每个SM包含多个CUDA核心、共享内存、寄存器等资源。
Warp：线程束，32个线程为一组。这是GPU执行指令的基本单位。嗯，这里要注意，warp里的线程执行的是同一指令。
CUDA Core：CUDA核心，执行浮点运算和整数运算的基本单元。

重要概念：GPU的并行度来自两个方面：一是多个SM之间的并行，二是SM内部多个warp的并行。我见过不少新手只关注前者，忽略了后者，结果性能上不去。

下面这张图展示了GPU架构的核心层次关系：

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

CUDA的线程组织方式，说白了就是三级结构：Grid（网格）→ Block（线程块）→ Thread（线程）。

我个人习惯这样理解：

Thread：单个工人，执行具体计算任务
Block：一个班组，工人们可以互相协作（通过共享内存）
Grid：整个工厂，包含所有班组

为什么会设计成三级结构？因为GPU的硬件资源有限。一个SM能同时处理的block数量是有限的，block里的线程数也有限制。我记得有一次做项目，把block设成1024个线程，结果发现SM只能同时跑两个block，性能反而下降了。

经验之谈：block大小建议设为32的倍数（warp大小），常用128或256。太小了SM利用率低，太大了寄存器压力大。我一般从256开始调。

代码里怎么用？看这个例子：

// 定义grid和block维度
dim3 gridDim(4, 4);    // 4x4 = 16个block
dim3 blockDim(16, 16); // 16x16 = 256个线程/block

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

// kernel内部获取线程索引
__global__ void myKernel(float* data, int N) {
    int bx = blockIdx.x;  // block在grid中的x索引
    int by = blockIdx.y;  // block在grid中的y索引
    int tx = threadIdx.x; // 线程在block中的x索引
    int ty = threadIdx.y; // 线程在block中的y索引
    
    // 计算全局线程ID
    int idx = (by * gridDim.x + bx) * blockDim.x * blockDim.y 
            + ty * blockDim.x + tx;
    
    if (idx < N) {
        data[idx] = data[idx] * 2.0f;
    }
}

1.3 内存层次结构

CUDA的内存模型，是性能优化的关键。说白了，就是「距离计算核心越近，速度越快，容量越小」。

内存类型	位置	访问速度	生命周期	作用域
寄存器	SM内部	~1周期	线程	单个线程
共享内存	SM内部	~5周期	Block	同一block内线程
全局内存	显存	~400周期	程序	所有线程
常量内存	显存（带缓存）	~100周期	程序	所有线程（只读）
纹理内存	显存（带缓存）	~100周期	程序	所有线程（只读）

避坑指南：我曾经在项目里把大量数据存在局部变量里，以为这样快。结果每个线程用了太多寄存器，导致SM能同时运行的线程数大幅减少，性能反而下降了。记住：寄存器是有限资源！

共享内存是性能优化的利器。同一block内的线程可以通过共享内存交换数据，避免反复访问慢速的全局内存。我建议你养成一个习惯：只要block内线程需要协作，优先考虑共享内存。

// 共享内存使用示例
__global__ void sharedMemExample(float* input, float* output, int N) {
    __shared__ float shared[256];  // 声明共享内存数组
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tid;
    
    // 从全局内存加载到共享内存
    if (idx < N) {
        shared[tid] = input[idx];
    }
    __syncthreads();  // 同步：确保所有线程都加载完毕
    
    // 在共享内存上做计算
    if (tid < 128) {
        shared[tid] += shared[tid + 128];
    }
    __syncthreads();
    
    // 写回全局内存
    if (idx < N) {
        output[idx] = shared[tid];
    }
}

1.4 CUDA程序编译与执行流程

CUDA程序的编译流程，和普通C程序不太一样。你想想看，代码里既有CPU代码（host），又有GPU代码（device），编译器得分别处理。

具体流程是这样的：

预处理：处理宏定义、头文件包含等
编译分离：nvcc编译器将.cu文件中的host代码和device代码分离
device代码编译：device代码被编译成PTX（并行线程执行）中间代码
PTX优化：PTX代码经过优化，生成cubin（CUDA二进制）文件
host代码编译：host代码由gcc/clang等编译成目标文件
链接：将所有目标文件链接成可执行文件

执行流程更直观：

小技巧：使用cudaMemcpyAsync配合流（stream）可以实现数据传输和计算的重叠，这是提升性能的重要手段。我习惯把数据分成几块，一块在计算时，另一块在传输，流水线作业。

最后说一句，CUDA编程模型的核心就这些内容。但真正用好它，需要你在实践中不断体会。我见过太多人把GPU当成CPU用，结果性能还不如CPU。记住：理解硬件，才能写出高效的代码。