GPU 优化的主线，是让工作与数据都少等

Kernel 只是队列中的一项工作；Launch 返回不等于 GPU 已经算完

Host（主机 CPU）分配内存、提交 Host-to-Device（H2D，主机到设备）拷贝、Kernel launch（内核启动）和 Device-to-Host（D2H，设备到主机）拷贝。Kernel launch 和明确的异步 API 可以入队后返回；普通 memcpy、allocation 和其他 CUDA API 是否阻塞，则取决于具体 API、方向、参数与 Host memory 类型。若计时、读回结果和错误处理忽略这份逐 API 合同，程序可以“看起来很快”，甚至在错误发生后继续运行。

↔ 左右滑动查看图中细节

my_kernel<<<grid, block, 0, stream>>>(out, in, n);
CUDA_CHECK(cudaGetLastError());          // 配置、参数或此前尚未清理的错误
CUDA_CHECK(cudaStreamSynchronize(stream)); // 调试/真正依赖处：暴露异步执行错误

你定义 Grid 与 Block；硬件把 32 个 Thread 组成 Warp，再让 Block 驻留 SM

SIMT（Single Instruction, Multiple Threads，单指令多线程）允许每个 Thread 有自己的索引、寄存器状态和分支，但 SM（Streaming Multiprocessor，流式多处理器）以 Warp 为核心调度粒度。一个 Thread Block（线程块）完整驻留在一个 SM；它消耗的寄存器与 Shared Memory（共享内存）决定同一 SM 还能同时容纳多少 Blocks 和 Warps。

CUDA Tile 与 SIMT 可以出现在同一 CUDA 程序里，但当前 Tile function 与 SIMT device function 不能直接互调。CUDA Tile C++ 和 cuTile Python 共享 Tile 思路，不共享一条简单的“13.3+”安装合同；它们是并列编程入口，也不是“SIMT 已被淘汰”。

先分清作用域、管理者与物理 backing，再谈“哪层更快”

Register、Shared、Local、Global 描述的是 CUDA 编程空间；L1/TEX、L2、Device Memory 则更接近物理资源。两组名字并不是一一对应的楼梯。尤其 Local Memory 只表示每 Thread 私有地址空间，寄存器 spill（溢出）或动态索引数组可能落到片外 Device Memory，并经 L1/L2 缓存。

先用时间线找“空洞在哪”，再钻进热点 Kernel 解释“为什么等”

Nsight Systems 回答整机问题：CPU 是否及时提交、CUDA API 是否阻塞、拷贝和计算是否重叠、Stream 之间是否存在空洞；Nsight Compute 回答单 Kernel 问题：真实 memory traffic、cache、指令、occupancy（占用率）、scheduler 和 stall reason（停顿原因）。工具顺序反了，很容易精调一个并不重要的 Kernel。

大模型算子反复组合四件事：独立映射、归约、广播与分块复用

练习不是为了背模板，而是识别输出依赖：每个输出是否独立？一整行是否共享统计量？多个 Threads 是否必须交换部分和？同一片 A / B 数据能为多少输出复用？这些问题决定 Warp shuffle、Shared Memory、融合和多阶段 Kernel 的边界。

最好的实现层级不是最低，而是证据允许的最高层

高层库和编译器把成熟调度、架构适配与维护成本打包；低层代码换来更细控制，也把编译矩阵、错误面和升级负担交给你。正确顺序是先建立可信 Eager / 库基线，再依据热点选择 torch.compile、Triton、CUDA Tile 或 CUDA C++ SIMT。

↔ 左右滑动查看图中细节

每个阶段都交付“能算对、能测量、能解释、能复跑”的证据包

NVIDIA Best Practices 的 APOD（Assess, Parallelize, Optimize, Deploy，评估—并行化—优化—部署）不是一次性流程，而是每个热点的短循环。不要先通读整本手册；用一张明确 GPU 和一组固定 Shape，让每个概念都落到代码、测试、profiler capture 和复现实验清单。

学会 CUDA 的标志，是知道何时不该再写 CUDA

若成熟库或编译器已经满足端到端 Service-Level Objective（SLO，服务级目标），更低层实现只会增加维护面；若热点占比低、瓶颈在 CPU / 网络 / 排队，Kernel 再快也不会改变用户体验。下潜之前必须同时满足“值得、可测、可守住”三项。