第二章：CUDA编程基础回顾——线程层次结构、内存层次结构、核函数定义与调用

各位同学，咱们今天聊点实在的。CUDA编程，说白了就是让GPU帮你干活。但GPU怎么干活？你得先搞清楚它的“组织架构”和“记忆系统”。这一章，我带你回顾一下CUDA编程最核心的三个东西：线程怎么排兵布阵、内存怎么分层管理、核函数怎么写怎么调。

嗯，先别急着写代码。把底层逻辑理清了，后面调度优化才能游刃有余。

2.1 线程层次结构：GPU的“军队编制”

GPU里有成千上万个线程在同时跑。这么多线程，总得有个管理方式吧？CUDA的设计很巧妙，它把线程分成了三层：线程（Thread）、线程块（Block）、网格（Grid）。

我习惯这么理解：

线程：最基层的士兵，负责执行一条指令。
线程块：一个排的士兵，共享一块“战地食堂”（共享内存），可以互相通信。
网格：整个军队，由若干个排组成，排与排之间不能直接通信。

你想想看，为什么这么设计？说白了，是为了适配GPU的硬件架构。一个线程块里的线程，会被调度到同一个流多处理器（SM）上执行。这样它们访问共享内存就快，同步也方便。而不同线程块，可能被分配到不同的SM上，天然隔离，互不干扰。

关键点：线程块内的线程可以通过 __syncthreads() 同步，但不同线程块之间无法直接同步。这是设计上的硬约束，写代码时一定要记住。

我记得刚学CUDA那会儿，总想把所有线程都塞到一个块里。结果呢？一个块最多1024个线程（现代GPU），数据一多就崩了。后来才明白，网格和线程块是给程序员的一种“抽象”，让你能灵活地组织大规模并行计算。

来看一个典型的线程索引计算方式：

// 假设我们有一个一维的网格和一维的线程块
int tid = blockIdx.x * blockDim.x + threadIdx.x;

这个公式，我建议你刻在脑子里。它把全局线程ID映射到了具体的数据元素上。二维、三维的情况类似，只是多几个维度而已。

2.2 内存层次结构：GPU的“存储金字塔”

GPU的内存，跟CPU的内存不一样。它是一套层次化的存储体系。从快到慢、从贵到便宜，大致是这样：

内存类型	位置	访问速度	作用域	生命周期
寄存器	SM内部	最快（1周期）	单个线程	线程执行期间
共享内存	SM内部	很快（~5周期）	线程块内	线程块执行期间
L1缓存	SM内部	快（~10周期）	SM内所有线程	自动管理
L2缓存	芯片全局	中等（~100周期）	所有SM	自动管理
全局内存	显存	慢（~400周期）	所有线程+主机	程序运行期间
本地内存	显存	慢（~400周期）	单个线程	线程执行期间

这里我要多说一句：共享内存是性能优化的关键。我在项目中遇到过好几次，全局内存访问太慢，导致整个核函数被“内存墙”卡死。后来把热点数据搬到共享内存里，性能直接翻倍。

我的经验：写核函数时，先问自己三个问题：

这个数据会被多个线程反复用吗？ → 放共享内存
这个数据只属于当前线程？ → 放寄存器
这个数据太大，共享内存放不下？ → 只能放全局内存，但要注意合并访问

另外，全局内存的合并访问是个坑。我曾经写过一段代码，线程访问全局内存时是“跳跃式”的，结果带宽利用率不到10%。后来改成连续访问，性能提升了8倍。嗯，这里要注意：相邻线程最好访问相邻的地址。

2.3 核函数定义与调用：让GPU动起来

核函数（Kernel）就是跑在GPU上的函数。定义它很简单，加个 __global__ 修饰符就行。但调用方式跟CPU函数不太一样。

先看一个最基础的例子：

// 核函数定义
__global__ void vecAdd(float *A, 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() {
    // ... 分配内存、拷贝数据 ...
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    // ... 拷贝结果回主机 ...
}

这里有几个细节，我当年踩过坑：

三括号语法 <<<grid, block>>>：第一个参数是网格维度，第二个是线程块维度。可以是一维、二维、三维。
边界检查：if (i < N) 这句不能省。因为线程总数可能比数据量多，多余的线程必须“空转”退出。
异步执行：核函数调用是异步的，主机不会等它跑完。如果需要同步，得调用 cudaDeviceSynchronize()。

警告：我曾经犯过一个低级错误——在核函数里调用了 printf 打印调试信息。结果呢？GPU上的 printf 是缓冲的，程序跑完才输出，而且顺序完全乱掉。调试核函数，建议用 cuda-gdb 或者把结果写回主机再打印。

核函数的参数传递也有讲究。所有参数都是传值的，不能传引用。如果你传一个结构体，它会被复制到每个线程的寄存器或本地内存里。结构体太大？那性能就崩了。我建议：传指针，别传大结构体。

2.4 本章知识体系总览

说了这么多，咱们用一张图把核心逻辑串起来。下面这张SVG图，展示了线程层次、内存层次和核函数调用之间的关系：

这张图里，上面是线程层次（网格→线程块→线程），下面是内存层次（寄存器→共享内存→缓存→全局内存），中间是核函数调用。三者是紧密耦合的：线程块决定了你能用多少共享内存，网格大小决定了你能处理多少数据，核函数调用方式决定了性能上限。

好了，这一章的内容就这些。线程层次、内存层次、核函数调用，这三板斧是CUDA编程的基石。后面讲任务调度时，你会发现所有优化策略都绕不开它们。

想要更多资料，联系可可