把 CUDA C++ 从语法学到 硬件直觉
__global__ 和 threadIdx 半小时就能学完，真正的门槛是看到一段 kernel 就能判断它卡在哪。这份教程按「写对 → 写快 → 看懂编译产物 → 接入工程」的顺序展开，每一个优化都要说清楚它消除了哪个具体瓶颈。
没有显卡也能学完全部内容：四个实验室在浏览器里模拟 GPU 的执行模型。
__global__ void add(const float* a, const float* b, float* c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i];} add<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);学习路线
按顺序学，每一阶段都建立在上一阶段之上起步
写出、编译并跑通第一个 kernel3 课 · 60 分钟第一个 kernel：让 CPU 指挥 GPU 干活
host 与 device 的分工、__global__ 的含义、kernel launch 为什么是异步的
线程层级：grid / block / thread 的索引算术
把一维、二维数据映射到线程上，以及为什么每个 kernel 都要写边界检查
显存管理与错误检查：写出不骗人的 CUDA 程序
cudaMalloc / cudaMemcpy 的正确姿势、统一内存、以及那个你必须抄一遍的宏
执行模型与内存
warp、内存层级、合并访问、共享内存：性能的全部来源4 课 · 92 分钟Warp 与 SM：GPU 真正的执行单位
SIMT 模型、warp divergence、延迟隐藏，以及「32」这个数字为什么无处不在
内存层级：决定性能上限的那张表
寄存器、共享内存、L1/L2、HBM 的延迟与带宽，以及计算强度这个判断工具
内存合并：同一段代码，5 倍性能差距
warp 如何发起访存事务、什么是 stride 访问、AoS 与 SoA 的抉择
共享内存与 Bank Conflict：片上存储的正确用法
tiled 矩阵乘完整推导，以及那个神秘的 [TILE][TILE + 1]
优化实战
归约七步、occupancy 权衡、流水线重叠3 课 · 76 分钟深入底层
warp 原语、PTX 与 SASS，看清编译器做了什么2 课 · 46 分钟工程化与 AI
profiling 方法论、PyTorch 算子、Tensor Core 与 Triton3 课 · 66 分钟交互实验室
把抽象概念变成能拖动的东西线程索引可视化
grid / block / thread 到底怎么映射到数据
拖动 gridDim 与 blockDim，点任意线程看它负责哪个元素。支持一维与二维网格，并实时显示越界线程。
内存合并模拟器
一个 warp 的访存会变成几个 sector 事务
调节 stride、起始偏移和元素宽度，实时看到 32 个线程的地址分布、产生的 sector 数量和带宽利用率。
Bank Conflict 模拟器
共享内存的 32 个 bank 被怎么撞上的
输入访问模式与 padding，看到每个 bank 上落了几个线程、最大冲突路数以及需要多少个周期。
Occupancy 计算器
到底是寄存器、共享内存还是线程数在卡你
选择架构并设置 block 尺寸、每线程寄存器数、每 block 共享内存，算出理论占用率并指出瓶颈资源。
课程里出现的每个关键 kernel 都在 cuda/ 目录下有完整可运行的版本，带计时、结果校验和多版本性能对比。有 GPU 的机器上 make 即可全部构建；没有也不影响学习。
cd cudamake # 用检测到的架构编译全部示例make run # 依次跑一遍./bin/06_reduction # 归约 7 个版本的性能对比