跳转至

03 LLVM 后端代码生成与 PTX 到 SASS 编译体系

1. 现代 GPU 编译工具链层次 (以 CUDA / NVCC 为例)

[ CUDA C++ / Triton 源代码 ]
        ↓ (Clang / NVCC 前端)
[ LLVM IR (高层通用中间表示) ]
        ↓ (NVPTX 后端代码生成器)
[ PTX (Parallel Thread Execution - 虚拟机器指令集, 具备向前兼容性) ]
        ↓ (PTX JIT Compiler / ptxas 汇编器)
[ SASS (Streaming Assembler - 对应具体物理芯片微架构的真实机器二进制指令) ]

2. PTX 与 SASS 指令对照分析

语义 虚拟指令 (PTX) 真实物理微架构指令 (SASS - Hopper/Blackwell)
浮点乘加 fma.rn.f32 %f3, %f1, %f2, %f0; FFMA R3, R1, R2, R0;
矩阵张量计算 mma.sync.aligned.m16n8k16.row.col ... HMMA.16816.F32 R4, R8, R12, R16;
共享内存加载 ld.shared.f32 %f0, [%r1]; LDS.U.32 R0, [R1];
线程同步 bar.sync 0; BAR.SYNC 0x0;