运行 CUDA Kernel 时到底发生了什么？
这是一个简单的 CUDA 向量加法程序：
__global__ void vadd(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];
}
在 RTX 4090 上编译运行，正确输出 c[0]=2.0 c[n-1]=2.0。但让这行代码工作，涉及数千万条 CPU 指令、多个设备文件、九百多个 ioctl 调用和一个内存映射的门铃寄存器。
本文追踪一个 CUDA kernel 从代码到 warp 并返回结果的全过程。
一、编译管线
nvcc 是一个驱动程序，运行多个编译器并组合输出：
$ nvcc --keep -arch=sm_89 -o vadd vadd.cu
输出文件：
- vadd.ptx — 设备代码的 PTX（虚拟 ISA，无限类型化寄存器）
- vadd.sm_89.cubin — PTX 转 SASS（真实硬件指令）
- vadd.fatbin — 捆绑的 cubin + PTX
- vadd.cudafe1.stub.c — 主机启动桩 + kernel 注册
- vadd.o — 最后的主机目标文件（内含 fatbin）
管线：nvcc → cicc（LLVM 后端，输出 PTX）→ ptxas（PTX → SASS）
PTX 示例（向量加法 kernel 体）：
mad.lo.s32 %r1, %r3, %r4, %r5; // r1 = ctaid*ntid + tid
setp.ge.s32 %p1, %r1, %r2; // 如果 i >= n 设置谓词 p1
@%p1 bra $L__BB0_2; // 越界则跳过
cvta.to.global.u64 %rd4, %rd1; // 转换全局地址
mul.wide.s32 %rd5, %r1, 4; // i * 4（字节偏移）
ld.global.f32 %f2, [%rd6]; // 加载 a[i]
add.f32 %f3, %f2, %f1; // a[i] + b[i]
st.global.f32 [%rd10], %f3; // 写入 c[i]
虚拟寄存器前缀：%r = 32 位整数，%rd = 64 位整数，%f = 32 位浮点，%p = 1 位谓词。
一个地址需要三条 PTX 指令，因为 PTX 不知道最终硬件——它需要先断言指针所在内存窗口，再做乘法和加法。
二、硬件细节
RTX 4090 有 128 个 SM（流式多处理器），每个 SM 有 128 个 CUDA Core。CUDA 将 thread block 分配给 SM——一个 block 运行在一个 SM 上直到完成。4096 个 thread × 256 个 block = 1M 线程。
Warp：32 个线程为一组，SM 一次执行一组 warp。warp 内所有线程执行同一指令，不同指令则发生分支发散。
Block 调度：GPC（图形处理簇）包含 TPC（纹理处理簇），TPC 包含 SM。RTX 4090 有 8 个 GPC，每个含 2 个 TPC，每个 TPC 含 4 个 SM（部分含 2 个），共 128 个 SM。硬件调度器负责将 block 分配到 SM。
三、内存之旅
内存层次：全局内存（大但慢）→ L2 缓存 → SM 内部 L1/共享内存 → 寄存器。
全局内存：24GB GDDR6X，带宽 1008 GB/s。每个 SM 有 128KB 寄存器文件和 128KB L1/共享内存（可配置）。
向量加法的一次内存访问延迟：全局内存约 200-400 周期 → L2 缓存约 60 周期 → L1 缓存约 20 周期 → 寄存器 0 周期。
在 warp 层面：ld.global.f32 触发一个 4 字节加载，如果所有 32 线程命中同一缓存行（128 字节），coalescing 只需一次内存事务。如果访问连续地址（如 a[0..31]），正好一次事务完成所有 32 次加载。
四、从 PCIe 到设备
主机到设备的旅程：
- cudaMalloc：调用
cuMemAlloc_v2→ 驱动创建虚拟内存映射 - cudaMemcpy H2D：调用
cuMemcpyHtoD_v2→ NVRM（NVIDIA 内核模块）分配 GPU 物理内存，映射到 CPU 地址空间 - 驱动通过 900 多个 ioctl 调用与内核交互
- 设置 DMA 引擎，PCIe Gen4 ×16（约 32 GB/s）
- 通过门铃寄存器通知 GPU：写入 GPC 门铃映射的 MMIO 地址
Kernel 启动过程：
- 驱动用
cuStreamCreate创建 CUDA stream cuLaunchKernel触发 12 个 ioctl（验证参数、分配 grid 等）- 最终调用
NV_ESC_REGISTER_FD注册 GPU 上下文 - GPC 的微控制器读取门铃 → 通知 SM 调度器
- SM 调度器分配 block → warp → 开始执行
五、Warp 执行
每个 SM 有 4 个 warp 调度器。每个 cycle 选择一个 ready warp 发射指令。
Warp 交错：当一个 warp 等待内存时（约 400 周期），调度器切换到另一个 ready warp。
向量加法 kernel 的指令流：
- Integer ops：计算线程 ID（
i = blockIdx * blockDim + threadIdx） - Compare + branch：越界检查
- Address calculation：指针 + 偏移
- Load：读取
a[i]、b[i]（约 200-400 周期延迟） - Multiply-add：FMAD 指令
- Store：写入
c[i]
Warp 的 occupancy（占用率）是隐藏延迟的关键——更多 warp 意味着调度器有更多候选可用。RTX 4090 每个 SM 最多 64 个 warp（2048 线程），但受寄存器文件大小限制。
六、返回结果
cudaMemcpy D2H：驱动将 GPU 内存映射到 CPU 页面，DMA 引擎通过 PCIe 传回结果。打印 c[0]=2.0 c[n-1]=2.0。
后记
一次 CUDA kernel 运行涉及：
- 编译管线：PTX → SASS 的重叠阶段
- 操作系统：ioctl 驱动调用
- 硬件微架构：GPC → TPC → SM → warp 调度 → CUDA core
- 内存层次：global → L2 → L1 → register
- DMA 与 PCIe 数据传输
"GPU 编程很简单"这句话背后是从编译器到底层硬件微架构的深厚技术栈。理解这些底层机制有助于写出高性能的 CUDA 程序。
原文出处：https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/
