第一章：GPU架构与CUDA基础

1.1 GPU硬件架构概览

说实话，我第一次接触GPU编程时，最大的困惑就是——这玩意儿跟CPU到底有啥区别？

CPU像个全能选手，什么活儿都能干，但一次只能干几件。GPU呢？它是个偏科生，只会做简单计算，但一次能同时干成千上万件。我习惯把GPU想象成一个超级大食堂：一个CPU是米其林大厨，精雕细琢；而GPU是一百个打饭阿姨，动作简单但架不住人多。

具体到硬件层面，NVIDIA的GPU核心架构是这样的：

SM（Streaming Multiprocessor）：流式多处理器，GPU的基本计算单元。每个SM里有多个CUDA核心、共享内存、寄存器等资源。
CUDA Core：执行算术运算的最小单元。一个SM通常包含几十到上百个CUDA核心。
Warp：32个线程组成一个warp，这是GPU调度和执行的基本单位。嗯，这个数字32很关键，后面你会反复见到它。
Memory Hierarchy：全局内存、共享内存、寄存器、L1/L2缓存，各有各的脾气。

核心概念：GPU通过大量线程并行执行来隐藏延迟。当一个warp在等数据时，调度器立刻切换到另一个warp。这就是所谓的「零成本上下文切换」——CPU做不到这一点。

我在项目中遇到过一个问题：一个核函数在GTX 1080上跑得飞快，换到V100上反而变慢了。排查了半天，发现是warp调度策略不同导致的。你看，硬件细节真的会咬人。

1.2 CUDA编程模型

CUDA编程模型，说白了就是「主机（CPU）发号施令，设备（GPU）干活」。我个人习惯把流程拆成三步：

数据准备：在主机端分配内存，把数据从CPU搬到GPU。
启动核函数：告诉GPU「嘿，用这些线程，跑这个函数」。
结果回传：把计算结果从GPU搬回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() {
    // ... 分配内存、拷贝数据 ...
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
    // ... 拷贝结果回主机 ...
}

注意那个 <<<gridSize, blockSize>>> 语法，这是CUDA特有的。左边是网格大小（多少个线程块），右边是块大小（每个块里多少个线程）。

我的经验：刚开始写CUDA时，我总喜欢把blockSize设成256或512。后来发现，对于不同的硬件和问题规模，最优值差别很大。建议你多试试128、256、512这几个值，用profiler看看效果。

1.3 线程层次结构

CUDA的线程组织方式，我管它叫「俄罗斯套娃」——一层套一层：

Thread（线程）：最底层，每个线程执行核函数的一份拷贝。
Block（线程块）：一组线程，可以协作（通过共享内存和同步）。
Grid（网格）：一组线程块，构成一次核函数调用。

你想想看，为什么要有这种层次结构？

原因有两个。第一，硬件限制——一个SM的资源有限，不能塞太多线程。把线程分成块，每个块分配到不同的SM上，就能突破这个限制。第二，数据局部性——同一个块里的线程可以共享数据，减少对全局内存的访问。

避坑指南：我曾经把一个块里的线程数设成了1025，结果程序直接崩了。为什么？因为硬件规定每个块最多1024个线程（某些老架构是512）。记住这个上限，别踩坑。

下面这张图展示了线程层次结构的关系：

1.4 内存层次结构

GPU的内存系统，我刚开始学的时候觉得特别绕。其实记住一句话就行：离计算单元越近，越快，但越小。

内存类型	位置	访问速度	作用域	生命周期
寄存器	SM内部	~1 cycle	单个线程	核函数执行期间
共享内存	SM内部	~5 cycles	线程块内	核函数执行期间
L1缓存	SM内部	~10 cycles	SM内所有线程	自动管理
L2缓存	芯片级	~100 cycles	所有SM	自动管理
全局内存	显存	~400-800 cycles	所有线程	程序运行期间
常量内存	显存（带缓存）	~100 cycles（缓存命中）	所有线程（只读）	程序运行期间

这里有个关键点：全局内存很慢。慢到什么程度？一个全局内存访问的时间，够GPU执行几百条算术指令了。所以优化的核心思路就是——尽量减少全局内存访问，多用共享内存和寄存器。

我的习惯：写核函数时，我会先问自己三个问题：

这个数据能不能放在寄存器里？
能不能用共享内存让一个块里的线程复用？
如果必须用全局内存，能不能做到合并访问（coalesced access）？

这三个问题想清楚了，性能基本不会太差。

说到合并访问，我踩过一个坑。有一次做矩阵转置，我按行读取、按列写入，结果全局内存访问模式乱七八糟，带宽利用率只有10%。后来改成按列读取、按行写入，配合共享内存做转置，带宽利用率直接飙到80%以上。你看，访问模式真的很重要。

注意：共享内存虽然快，但容量有限。一个SM通常只有几十KB到一百多KB。如果你一个块里用了太多共享内存，SM能同时容纳的块数就会减少，反而可能降低整体性能。这是个trade-off，需要根据具体情况权衡。

好了，GPU架构和CUDA基础就聊到这儿。记住：理解硬件是优化的前提，而CUDA编程模型给了你精细控制硬件的工具。后面我们会一步步深入，把这些知识变成实实在在的性能提升。

想要更多资料，联系可可

📚 课程目录

第一章：GPU架构与CUDA基础

1.1 GPU硬件架构概览

1.2 CUDA编程模型

1.3 线程层次结构

1.4 内存层次结构