上一篇(host runtime)结束在
cuLaunchKernel异步返回——host 把命令塞进 stream 队列就走了。这一篇 GPU 接管:256 个线程在 SM 上并行跑 SASS、写 printf buffer、写 device 内存。然后 host 通过cuStreamSynchronize等 GPU 跑完、cuMemcpyDtoH把结果拷回——Output: 42终于出现。
0. 几个名词先说清楚
| 缩写 / 术语 | 英文全称 | 中文 | 含义 |
|---|---|---|---|
| SM | Streaming Multiprocessor | 流式多处理器 | GPU 上一个独立的并行执行单元,包含 warp scheduler、寄存器、shared memory 等 |
| warp | warp(经线 / 线束) | 线束 | 32 个 lane(线程)组成的最小调度单位,锁步执行 |
| lane | lane(车道) | 通道 | warp 里的一个执行通道,跑一个 “thread” |
| lockstep | lockstep | 锁步 | warp 内 32 个 lane 严格同步执行同一条指令 |
| benign race | benign race condition | 良性竞争 | 多个线程同时写同一地址,但因为写的都是同样的值,结果”恰好正确” |
| printf buffer | printf buffer | printf 缓冲区 | device 端循环缓冲区,vprintf 写入,driver 在同步时拷回 host |
| DtoH | Device to Host | 设备→主机 | 数据从 GPU 显存拷回 host 内存 |
| stream sync | stream synchronization | 流同步 | 阻塞 host,等 stream 队列里所有命令完成 |
1. GPU 接管之后发生了什么
cuLaunchKernel 调完,host 这边继续往下跑。GPU 这边收到 launch 命令,做的事:
1. 找一个空闲 SM 分配这个 block
2. 把 block 的 256 个 thread 切成 8 个 warp(256 ÷ 32 = 8)
3. 给每个 warp 分配寄存器(从 SM 的寄存器文件)
4. warp 排队等 warp scheduler 发射指令
5. 每个 cycle scheduler 选一个 ready warp,发射一条指令
→ warp 内 32 个 lane 锁步执行2. SM 内部结构
SM(Streaming Multiprocessor)
├── warp scheduler × 4 ← 每 cycle 选一个 ready warp 发指令
├── 寄存器文件 ← 所有活跃 warp 的寄存器都常驻
│ (65536 个 32-bit reg,大约)
├── shared memory(SRAM) ← 100KB+,block 内 thread 共享
├── L1 cache
└── 大量执行单元(ALU、FPU、tensor core 等)8 个 warp 排队:
warp 0 (lane 0..31) ← 256 个 thread 切成 8 个 warp
warp 1 (lane 32..63)
warp 2 (lane 64..95)
...
warp 7 (lane 224..255)每个 lane 都从 SASS 第一条指令开始执行——执行的是同一段代码,只是各自的寄存器状态不同。
3. 每个 lane 跑的 SASS(对应 PTX)
简化的 PTX 流程:
ld.param.b64 %rd2, [hello_constant_param_0]; ; 读 out 参数(指针)
cvta.to.global.u64 %rd1, %rd2; ; 通用 → global 地址
mov.u32 %r1, 42; ; 你的 inline asm
cvt.u64.u32 %rd3, %r1; ; xxx as i64
... 构造 printf 参数 + 调 vprintf ...
st.global.b32 [%rd1], 42; ; *out = 42
ret;256 个 lane 各自走一遍这段代码。
4. 关于 256 路 race
st.global.b32 [%rd1], 42; ← 所有 256 个线程都执行这条256 个线程同时写同一个地址(都写 42)。理论上这是经典的 race condition(数据竞争)——多个写者同时访问同一内存。
但这里race 是良性的(benign race),因为:
- 所有线程写一样的值 —— 不会出现”谁覆盖谁”的不确定
- 32-bit store 是原子的(一条 PTX 指令对应一次 4-byte 总线事务)
最终结果恰好正确——*out 收到 42。
但如果你写的是 *out = tid(每个线程写自己的 id),结果就不确定了 —— 取决于最后一个完成的线程是哪个。CUDA 里写 kernel 必须自己保证内存访问语义,否则会得到不可预测的结果。
5. printf 链路
gpu_printf!("thread xxx: {}", xxx) 在 device 端展开成调用 vprintf 这个 CUDA runtime 提供的 device 函数。vprintf 做的事:
1. 把格式串 + 参数序列化成一段二进制 payload
2. 写到 device 端一个 printf buffer(环形缓冲区)
── buffer 默认 1MB,可以用 cudaDeviceSetLimit 调整
3. ... 继续执行 ...
4. (kernel 结束 或 buffer 快满时)driver 把 buffer 拷回 host
5. host 端解码 + puts 到 stdout所以你看到的 256 行 42 <> 42 <> ... 输出不是 kernel 实时打出来的 —— 是 256 个线程各自往 printf buffer 里塞了一条记录,host 端在同步时统一打印。
这就是为什么 printf 输出经常跟普通 host stdout 顺序错乱 —— 它们走完全不同的路径。
(注:格式串是 %lld <> \0,所以输出是 42 <> 42 <> 42 <> ... 连成一串,而不是带换行的多行。)
6. to_host_vec — 同步 + 拷贝回传
let result = out_dev.to_host_vec(&stream).expect("Failed to copy result");
println!("Output: {}", result[0]);底层做两件事:
cuStreamSynchronize(stream); // ① 等 stream 上所有命令完成
cuMemcpyDtoH_v2(host_buf, dptr, 4); // ② device → host 复制 4 字节6.1 cuStreamSynchronize 的角色
之前 cuLaunchKernel 是异步的,host 这边继续往下跑。cuMemcpyDtoH 之前必须先等 stream 上所有命令完成——否则可能拷到”还没写完”的内存。
cuStreamSynchronize 会阻塞 host 线程,直到 GPU 把 stream 队列上的所有命令(包括 kernel + 之前的 memset)都干完。
6.2 printf flush 也在这一步
cuStreamSynchronize 是 driver 触发 printf buffer flush 的时机之一。256 条 42 <> ... 记录在同步时从 device 拷回 host、解码、写到 stdout——所以 printf 输出都在 Output: 42 之前出现(因为 sync 早于 memcpy 早于 println!)。
7. host 端剩余执行
println!("Output: {}", result[0]); // 输出 42到这里整个 hello_constant example 就跑完了。后续:
CudaContextdrop → 自动cuCtxDestroyDeviceBufferdrop → 自动cuMemFree- 进程退出
8. 完整时间线
host 端 GPU 端
────── ──────
[PHASE 9/9 a] 启动
[PHASE 9/9 b] cuInit/cuCtxCreate
cuMemAlloc
[PHASE 9/9 c] cuModuleLoad ←── JIT:PTX → SASS,SASS 拷到 GPU
[PHASE 9/9 d] cuLaunchKernel ──────────→ block 分配到 SM
(异步,立刻返回) ↓
↓ warp 0..7 排队
继续往下 ↓
↓ 每个 lane 跑 SASS:
to_host_vec(&stream) ├── 读参数
↓ ├── inline asm (mov.u32 %r1, 42)
cuStreamSynchronize ◀──── 等 ├── 写 printf buffer
(阻塞 host) └── st.global.b32 [%rd1], 42
↓ ↓
driver flush printf buffer GPU 完成
↓
(stdout 出现 42 <> 42 <> ...)
↓
cuMemcpyDtoH(host_buf, dptr, 4)
↓
println!("Output: {}", result[0])
↓
(stdout 出现 Output: 42)
↓
进程退出实际运行时:
INFO hello_constant: [PHASE 9/9 e] copying device buffer back to host
42 <> 42 <> 42 <> 42 <> ... (256 次) ... Output: 429. 整套 cuda-oxide 流程总览
到这里 11 步走完了。整条链路串起来:
源码 *out = 42;
│
▼ rustc 前端 (Step 03)
MIR Statement::Assign(*_1, const 42)
│
▼ cargo-oxide / rustc 启动 (Step 02)
▼ codegen_crate 拦截 (Step 04)
▼ collector + stable MIR (Step 05)
│
▼ mir-importer(statement) (Step 06)
dialect-mir mir.store %ptr, 42 : i32
│
▼ mir-importer(块尾控制流 + intrinsic) (Step 07)
dialect-nvvm nvvm.read_ptx_xxx (for xxx())
│
▼ mir-lower(DialectConversion) (Step 08)
dialect-llvm llvm.store / llvm.inline_asm
│
▼ llvm-export (Step 09)
LLVM IR store i32 42, ptr %v1
│
▼ llc / NVPTX backend (Step 09)
PTX st.global.b32 [%rd1], 42
│
▼ CUDA driver JIT (Step 10)
SASS (sm_61 真实机器码)
│
▼ cuLaunchKernel + GPU 执行 (Step 11)
GPU 256 线程各自写 *out = 42
│
▼ cuMemcpyDtoH (Step 11)
Host println!("Output: {}", 42)10. 一句话总结
cuLaunchKernel返回后 GPU 接管,256 个 thread 在 SM 上被切成 8 个 warp,每个 warp 32 个 lane 锁步执行 SASS——所有 lane 都做同一件事:读参数、跑 inline asm 拿到 42、写 printf buffer、st.global [out], 42。256 个线程同时写同一地址是良性 race(写的都是 42)。host 端to_host_vec触发cuStreamSynchronize阻塞等待 GPU 完成(同时 driver flush printf buffer 把 256 条42 <> ...输出到 stdout),然后cuMemcpyDtoH把 4 字节拷回 host,println!("Output: {}", 42)完成整个 example。整套 cuda-oxide 流程到此结束——从 Rust 源码*out = 42一路追到 GPU 上一条 SASS 指令。
下一步推荐:用本系列开头给的 dump 命令亲自跑一遍 hello_constant,对照每一站的真实输出复盘。这样能把 11 篇文章的知识真正变成你自己的。
系列上一篇: cuda-oxide:hello-constant 拆解 10——host binary 启动 + CUDA driver 加载 PTX