背景
高级DSL（如Triton）本质是把高层算子描述逐级 lowering 成 CUDA 内核（PTX或cubin）。本文先梳理 NVCC 从源码到可执行的编译链路，再对照 Triton 的编译流程，说明高级 DSL 如何落到 GPU 上运行。
CUDA编译流程
CPU的基于已经发布的指令集架构，使得针对同样架构的新的CPU不需要修改编译选项即可运行，而GPU在不断的改进过程中难以保证二进制的兼容，所以nvcc利用了两段编译模型来保证这一点。
从cuda到ptx再到cubin。其中PTX可以被视作一种虚拟的GPU架构的组装。他不和实际的硬件直接对应，而是类似于hava字节码，是一种虚拟的中间体系结构。而cubin即可认为是运行在硬件上的最终的代码，与GPU的代际强相关。nvcc通常使用compute_xy来表示PTX的架构，而使用sm_xy来表示cubin的架构。
CUDA的编译过程如下：输入程序首先经过设备编译预处理，被编译为CUDA二进制文件(cubin)和PTX中间代码，这些内容会被放入一个fatbinary中。随后输入程序会再次进行主机编译预处理，通过代码合成将fatbinary嵌入其中，并将CUDA特定的C++扩展转换为标准C++结构。最后，C++主机编译器会将嵌入了fatbinary的合成主机代码编译成主机对象文件。具体步骤如图1所示。每当主机程序启动设备代码时，CUDA运行时系统会检查嵌入的fatbinary，以获取适用于当前GPU的fatbinary镜像。
C++ Preprocessor
输出：.cpp4.ii——主机分支、.cpp1.ii——设备分支
作用：处理.cu中的预编译指令（#include、#define、条件编译#if等）。拆分代码，将混合的主机代码和设备代码初步分离，为后续工具分工做准备。
nvcc -E vector_add.cu -o vector_add.cpp4.ii --keep --generate-dependencies
cicc（cuda Intermediate Compiler）
输出：.ptx（并行执行虚拟指令），PTX是GPU架构无关的虚拟指令集，不绑定硬件
作用：基于LLVM框架做优化（循环展开、内存访问优化等），生成与GPU架构解耦的中间表示IR。
// 编译ptx指令nvcc -ptx vector_add.cu -o vector_add.ptx -arch=compute_70// vector_add.ptx 内容示例.version 7.6.target sm_70.address_size 64.visible .entry vectorAdd( .param .u64 vectorAdd_param_0, .param .u64 vectorAdd_param_1, .param .u64 vectorAdd_param_2, .param .u32 vectorAdd_param_3){ .reg .pred %p<2>; .reg .b32 %r<10>; .reg .b64 %rd<11>;
// ... PTX指令 ld.param.u64 %rd1, [vectorAdd_param_0]; ld.param.u64 %rd2, [vectorAdd_param_1]; // ... 更多指令}
说明ptx是gpu无关的虚拟指令集，类似于java字节码，不同GPU架构提供统一的中间表示。
ptas（PTX汇编器）
作用：将PTX转换为特定GPU架构的机器码（如SM80、SM90对应指令）；执行寄存器分配、指令调度等硬件优化，生成CPU可执行的二进制文件。
nvcc -cubin vector_add.cu -o vector_add.cubin -arch=sm_70
fatbinary（胖二进制生成器）
输出：.fatbin.c文件（多个GPU架构的.cubin文件的集合）
1. 将PTX转换为特定GPU架构的机器码（如SM80、SM90对应指令）；执行寄存器分配、指令调度等硬件优化，生成CPU可执行的二进制文件。
2. fatbin机制通过将多版本的GPU代码打包到一个单一文件，解决了GPU架构碎片和未来设备兼容性问题。他在编译时预优化主流框架的性能，同时在运行时通过PTX JIT编译确保灵活性，是CUDA生态中实现一次编译，处处执行的关键技术。（至于PTX JIT的耗时，因为缓存等机制，只会增加启动耗时）。
nvcc -gencode arch=compute_50,code=sm_50 \ -gencode arch=compute_60,code=sm_60 \ -gencode arch=compute_70,code=sm_70 \ vector_add.cu -o vector_add.fatbin.c
cudafe++（cuda前端）
输出：.cudafe1.stub.c（主机代码中记录的设备代码的存根）、.cudafe1.cpp（最终主机代码）
作用：这一阶段将处理cuda特有的语法（如<<< >>>），将设备代码嵌入主机，生成最终的运行时主机代码，.cudafe1.stub.c是一个中间结果，存储了设备代码的符号，但没有具体实现，具体的实现放到了.fatbin.c。这一阶段可以认为是处理有的主机代码。
C++ Compiler（主机编译器）
g++ -c vector_add.cudafe1.cpp -o vector_add.o \ -I/usr/local/cuda/include \ -fPIC -O2 -std=c++11
Host Linker（主机链接器）
此外，上图中的阴影部分的nvlink和fatbinary部分是为了处理多个cuda代码，将不同的.cu文件链接到一起，并生成最终的胖二进制。
cuda编译的中间文件
注意nvcc不会区分对象文件，库文件或资源文件。在执行链接阶段时，她只是将这些类型的文件传递给连接器。
Triton编译流程
前端（Frontend）：用于将用户使用python编写的Kernel转换为Triton IR；
优化器（Optimizer）：各类pass将Triton IR逐步转换并优化为TritonGPU IR；
后端（Backend）：TritonGPU IR逐步转化为LLVM IR；
对于nvidia显卡，通过LLVM编译器转换为ptx或cubin。
以Triton官网的向量加实例01-vector-add.py在NVIDIA上的编译流程为例进行说明，运行这个kernel文件会生成在~/.cache/id中生成如下文件，这些都是triton编译中间结果和运行时缓存。
Triton IR（add_kernel.ttir）：与硬件无关的高级中间表示，表示深度学习模型的计算图，Triton IR 的特点包括:1、高级抽象：Triton IR 提供了高级抽象，使得开发者可以用接近于高级深度学习框架的方式来描述计算图。2、操作表示：它包含了一系列的操作（ops），如矩阵乘法、卷积、激活函数等，这些都是深度学习模型中常见的操作。3、优化：在 Triton IR 层面，编译器可以应用一些高级优化，如死代码消除、常量折叠等。4、转换：Triton IR 可以被转换为更接近硬件的 Triton GPU IR，以便进行进一步的优化。
TritonGPU IR（add_kernel.ttgir）：低中间表示，专门针对GPU硬件进行优化。TritonGPU IR的特点包括：1.硬件特定优化：TritonGPU IR包含了针对特定GPU架构的优化，如内存访问模式、线程布局等；2.并行性表示：线程块、网格等；3.转换为 LLVM IR：Triton GPU IR 可以被转换为 LLVM IR，以便利用 LLVM 的优化和代码生成能力。
LLVM IR（add_kernel.llir）：1.平台无关：LLVM IR 可以在不同的硬件平台上使用，LLVM 提供了多种后端来生成特定平台的机器代码。2.优化：LLVM 提供了大量的优化通道，可以在 LLVM IR 层面进行指令组合等优化。3.代码生成：LLVM IR 最终会被转换为特定硬件平台的机器代码，LLVM 提供了成熟的代码生成后端，如针对 x86、ARM、NVIDIA GPU 等。4.模块化：LLVM IR 是模块化的，可以表示程序的各个部分，如函数、全局变量、类型等。
NVPTX backend：LLVM有多种代码生成backend，通过设置好target features，调用LLVM codegen讲LLVM IR转成PTX汇编代码，也就是.ptx文件；
cubin文件：Triton 通常调用 ptxas 把 PTX 汇编成 cubin，直接把cubin bytes缓存下来，在执行时读取cubin bytes，直接用CUDA Driver API加载。
到这里triton kernel就可以调用cuda的程序完成kernel的运行了，我们发现，并没有像nvcc那样继续执行fatbinary、cudafe++等过程，因为triton并没有cuda c++ API哪些语法，也不需要被编译成可执行文件，其中host测的控制逻辑，在Triton runtime（C++扩展）里，已经提前编译好了，每次运行不需要重新编译链接。nvcc的编译可执行文件：“生成 .o → host linker → executable”过程，变成triton的运行时动态加载：“生成 cubin → driver load module → get function → launch”，这些都有专门的cuda API完成。
