在为 NVIDIA GPU 编译程序时，CUDA 通常被视为参考编译框架。但编译器究竟是如何工作的？编译器在为 GPU 创建“可执行文件”时需要生成哪种指令集架构 (ISA)？编译器究竟是如何工作的？它使用哪种架构来生成可以从 CPU 启动并在 GPU 上执行的代码？
这份周日早间备忘录旨在对NVIDIA GPU编译器架构进行高层次概述。内容涵盖GPU指令集、CUDA框架中的NVCC编译器，以及最终面向GPU指令集的其他架构。
当针对 GPU 等设备进行编程时，我们需要一套指令集。以 NVIDIA 为例，该设备是流式多处理器 (SM)，如右图所示。图中值得注意的是，它包含许多并发单元，例如 IN32 加法器、特殊功能单元、GEMM+++ 张量核心等，编译器需要确保以最高效的方式（流水线式）使用这些单元（效率定义为最小化空闲时间）。
以下是所使用的术语：
机器代码：指令的紧凑二进制表示。
NVIDIA GPU 有两种指令集架构 (ISA)：
假设我们要编译这个内核：
__global__ void normalizeVector(float* x, float* y, float* z, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
float vx = x[i], vy = y[i];
float len2 = vx*vx + vy*vy;
float invLen = rsqrtf(len2); // Use SFU fast reciprocal square root
z[i] = invLen; // Normalized length
}
这里的神奇之处在于，rsqrtf 应该被分发到 SFU 管道。它是如何工作的呢？_float invLen = rsqrtf(len2);_ 这行代码会被转换成 PTX 格式的 rsqrt.approx.ftz.f32，然后转换成 SASS 格式的 MUFU.RSQ，如下图所示（来自 [godbolt](https://godbolt.org/#g:!((g:!((g:!((h:codeEditor,i:(filename:'1',fontScale:14,fontUsePx:'0',j:1,lang:cuda,selection:(endColumn:2,endLineNumber:15,positionColumn:2,positionLineNumber:15,selectionStartColumn:2,selectionStartL ineNumber:15,startColumn:2,startLineNumber:15),source:’global+void+normalizeVector(float*+x,+float*+y,+float*+z,+int+n)%0A%7B%0A++++int+i+%3D+blockIdx.x++blockDim.x+%2B+threadIdx.x%3B%0A%0A++++float+vx+%3D+x%5Bi%5D%3B%0A++++float +vy+%3D+y%5Bi%5D%3B%0A++++float+len2+%3D+vxvx+%2B+vyvy%3B%0A%0A++++//+使用+SFU+快速+倒数+平方+根%0A++++flo at+invLen+%3D+rsqrtf(len2)%3B+%0A%0A++++//+标准化+向量%0A++++x%5Bi%5D+%3D+vx++invLen%3B%0A++++y%5Bi%5D+%3D+ vy+*+invLen%3B%0A%7D’),l:‘5’,n:‘0’,o:‘CUDA+C%2B%2B+source+%231’,t:‘0’)),k:33.333333333333336,l:‘4’,n:‘0’,o:’’,s:0,t:‘0’),(g:!((h:compiler,i:(compiler:nvcc130u2,filters:(b:‘0’,binary:‘1’,binaryObject:‘1’,commentOnly:‘0’,debugCalls:‘1’, demangle:‘0’,directives:‘0’,execute:‘1’,intel:‘0’,libraryCode:‘0’,trim:‘1’,verboseDemangling:‘0’),flagsViewOpen:‘1’,fontScale:14,fontUsePx:‘0’,j:1,lang:cuda,libs:!(),options:’-arch+sm_90+-use_fast_math+-O3’,overrides:!(),selection:(en dColumn:1,endLineNumber:1,positionColumn:1,positionLineNumber:1,selectionStartColumn:1,selectionStartLineNumber:1,startColumn:1,startLineNumber:1),source:1),l:‘5’,n:‘0’,o:’+NVCC+13.0.2+(Editor+%231)’,t:‘0’)),k:33.333333333333336,l:‘4’ ,n:‘0’,o:’’,s:0,t:‘0’),(g:!((h:device,i:(compilerName:‘NVCC+13.0.2’,device:‘SASS+(sm_90)’,editorid:1,fontScale:14,fontUsePx:‘0’,j:1,selection:(endColumn:6,endLineNumber:21,positionColumn:1,positionLineNumber:1,selectionStartColumn:6,s electionStartLineNumber:21,startColumn:1,startLineNumber:1),treeid:0),l:‘5’,n:‘0’,o:‘Device+Viewer+NVCC+13.0.2+(Editor+%231,+Compiler+%231)’,t:‘0’)),k:33.33333333333333,l:‘4’,n:‘0’,o:’’,s:0,t:‘0’)),l:‘2’,n:‘0’,o:’’,t:‘0’)),version:4).
MUFU 代表多功能单元，也称为 SFU，即特殊功能单元。MUFU.RSQ 将寄存器发送到 SFU 单元的平方根“流水线”。SASS 汇编中还有许多其他指令，例如 FMUL（融合乘法）或 FFMA（融合乘加法），这些指令会被分派到张量核心单元。分派单元负责将指令发送到正确的流水线单元。
有三点值得注意：
张量比通用矩阵乘法器 (GEMM) 更核心：例如，它可以通过张量核心内的 MMA (矩阵乘加) 单元执行归约操作（求和、最小值/最大值、点积）。
INT32 单元和寄存器的区别：该单元负责计算（例如，ADD、SHIFT、XOR 等，或更一般的算术运算，或 ALU），而寄存器负责存储值。
虽然只有 16 个 INT32 单元，但却有 32 个线程。这不是错误：原因是线程束调度器会在每个周期切换两组线程。因此，在第 0 个周期，这些单元将为线程 0 到 15 执行 ALU 操作；在第 1 个周期，则为线程 16 到 31 执行 ALU 操作。
异步 Warpgroup 级矩阵乘加指令 (WGMMA) 非常强大。您可以阅读 NVIDIA 文档 以及 Aleksa Gordić 的精彩 博客 了解更多信息。
NVIDIA CUDA 编译器架构
链接到标题
现在我们对GPU底层指令集架构（ISA）有了更多了解，接下来让我们看看如何将高级C++内核编译成SASS。我们知道，用于将CUDA内核（.cu文件）转换为PTX的编译器叫做NVCC。
那么，NVCC 和 LLVM 有什么区别呢？NVCC 是一个专用的编译器驱动程序，它使用开源的 LLVM 编译器基础架构作为后端。相比之下，LLVM 是一个灵活的编译器框架，而不是像 NVCC 那样即插即用的编译器。
CUDA 编译器 (NVCC)
链接到标题
驱动程序：NVCC 是一个驱动程序，用于管理 CUDA C/C++ 代码的编译过程。它将主机（CPU）代码与设备（GPU）代码分离。
工具链：NVCC 是一个工具链，它使用主机 C++ 编译器（例如 Linux 上的 GCC/Clang）来编译 CPU 代码，并使用 NVIDIA 专用的内部编译器（例如 CICC）来编译 GPU 代码（PTX 和 SASS）。
胖二进制文件：NVCC 通常会生成“胖二进制文件”，其中包含主机可执行文件和编译后的 GPU 代码，并且通常会有多个版本以适应不同的 GPU 架构。
框架：LLVM 是一套免费、开放、模块化和可重用的编译器技术、库和工具的集合，旨在构建各种编译器和语言工具链。
中间表示 (IR)：LLVM IR 是一个与机器无关的“中间层”，它充当前端（C++、Go…）和 ISA 后端（x86、ARM…）之间的通用连接。
优化：LLVM 优化器会对 IR 应用各种性能增强（如循环展开和死代码消除），而不管源语言或目标架构如何。
值得注意的是，LLVM 可以编译 CUDA：Clang 前端（LLVM 项目的一部分）可以使用 LLVM PTX 后端 处理 CUDA 代码，而无需 nvcc 驱动程序。
Triton 编译器架构
链接到标题
Triton 语言是一种功能强大的基于 Python 的领域特定语言 (DSL)，旨在帮助创建高级的_向量化_算法，这些算法可以编译成高效的_并行化_代码，以用于 GPU 目标。
这是一个最简单的内核示例，它从一个内存块加载数据并将其写入另一个内存块。offsets 的值为 [pid*BLOCK_SIZE….(pid+1)*BLOCK_SIZE-1]。此代码片段中未显示的“魔法”在于，Triton 可以高效地操作向量 v，例如执行数学运算。
import triton.language as tl
@triton.jit
def move_kernel(in_ptr, out_ptr, BLOCK_SIZE: tl.constexpr):
pid = tl.program_id(axis=0)
offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
v = tl.load(x_ptr + offsets)
tl.store(output_ptr + offsets, v)
Triton 不仅限于一维向量，它还可以使用张量处理更高维度的数据。例如：
import triton.language as tl
@triton.jit
def load_2d_tensor( input_ptr, output_ptr,
stride_row, stride_col, # Stride of the input tensor (elements per row, elements per column)
M , N , # Number of rows and columns
BLOCK_SIZE_M,BLOCK_SIZE_N, # Block size for rows and columns
):
# X and Y tile position
pid_m, pid_n = tl.program_id(axis=0), tl.program_id(axis=1)
# Ptr in memory based on the stride
base_ptr = input_ptr + (pid_m * BLOCK_SIZE_M * stride_row) + (pid_n * BLOCK_SIZE_N * stride_col)
# Offsets for X and Y indices
offs_m, offs_n = tl.arange(0, BLOCK_SIZE_M), tl.arange(0, BLOCK_SIZE_N)
offsets_2d = (offs_m[:, None] * stride_row) + (offs_n[None, :] * stride_col)
# 2D loading ...that's magic!
loaded_2d_tensor = tl.load(base_ptr + offsets_2d)
return loaded_2d_tensor
与 CUDA 相比，Triton 可以被视为更高层次的抽象，专门用于构建神经网络内核（得益于张量块方法），并且可以生成高效的向量化并行 GPU 代码。Triton 还支持多个后端 GPU：
Triton 不使用自己的语言解析器。相反，它依赖于 Python 的运行时和抽象语法树 (AST)，然后将该 AST 转换为其自身的中间表示形式 TIR 和 TGIR，最后生成 GPU 代码。CUDA 后端编译器和 Triton 前端编译器之间的连接是通用的 LLVMIR，对于 Intel GPU，它基于 SPIRV（尽管我不确定 TGIR 是否直接降维到 SPIRV，而不是先经过中间的 LLVMIR）。
我写的是“向量化算法”，尽管其核心概念更侧重于分块的、块级模型：
分块是主要机制。Triton 的思维模型围绕着对小型、静态形状的多维子数组（或称“瓦片”）进行编程操作展开，编译器会自动将这些操作并行化到 GPU 核心上。
向量化，即在一条指令中加载/存储多个数据元素，是一种优化：编译器将其作为一项关键的内部优化策略，以在底层硬件上高效地实现这些数据块。这一概念与嵌套分块策略（见右图）相一致，后者旨在将数据块分解为微块，最终分解为纳米块，以尽可能紧密地适应机器的计算能力和内存层次结构。
Triton 语言将阶段降低到流式汇编
链接到标题
让我们来看看 Triton 语言的高级指令 tt.load 是如何被降级为 SASS 的。（注：以下代码为伪代码，实际运行效果可能有所不同。）
Triton IR (TIR)
链接到标题
在解析 Python AST 之后，初始阶段是机器无关的 Triton IR（TIR，一种 MLIR 方言）。在这个阶段，对 tile 定义了一个 tt.load 操作。tt.load 操作是一个高级的、未经优化的指令，用于将整个 tile 加载到内存中。tile 可以是单个数组，也可以是二维数组，甚至是任意维度的数组。
// simplified, illustrative TTIR fragment for 2 dimensionnal array
// load of 32 bits floating point values (without masks)
%ptrs = ... : tt.tensor<MxN, ptr<global f32>>
%vals = "tt.load"(%ptrs) : (tt.tensor<MxN, ptr<global f32>>) -> tt.tensor<MxN, f32>
对张量指针和数据块（瓦片）执行高级 tt.load 操作。第一个输入参数 %ptrs 是由 tensor 表示的 n 维张量数组。 <MxN, ptr它返回一个相同维度的张量，但返回的是值而不是指向值的指针。
实际上，tt.load 还接受一个掩码和一个默认值参数，这是从内存边界读取数据所必需的（以防向量维度未按步长对齐）。
%ptrs = ... : tensor<MxN, ptr<global f32>>
%mask = ... : tensor<MxN, i1> // in-bounds mask (optional)
%fallback = constant 0.0 : f32
%vals = "tt.load"(%ptrs, %mask, %fallback) : (tensor<MxN, ptr<global f32>>, tensor<MxN, i1>, f32) -> tensor<MxN, f32>
Triton-GPU IR（TTGIR / TGIT）
链接到标题
在下一阶段，Triton-GPU IR 用于将通用张量布局转换为特定于硬件的布局，尤其是在内存布局以及跨线程和线程束的分配方面。换句话说，TGIT 将 n 维加载拆分为更易于降维到 LLVM 的每线程向量加载操作。
// TGIR (illustrative)
%vec_ptr = tt.get_contiguous_vector_ptr %ptrs, vec_len=4
%vec_vals = ttg.load.global.v4f32 %vec_ptr // a vectorized 4-wide global load
%vec_vals_cast = f32x4 -> f32 // unpack to tensor<MxN, f32> shape
LLVM IR (LLIR)
链接到标题
Triton-GPU IR 随后被转换为以 LLVM IR 表示的线程级指令。这种底层 IR 能够很好地表示最终的机器指令，并使用标准的 LLVM 操作，例如指针运算、内存加载和同步原语。张量抽象基本被移除，取而代之的是每个线程的独立内存访问操作。
// Example LLVM IR snippet (conceptual)
%thread_ptr = getelementptr inbounds float, ptr %base_ptr, i64 %thread_offset
%value = load float, ptr %thread_ptr, align 4
表示：图块加载被分解成一系列单独的加载指令，每个线程的特定数据元素都使用指针运算，如果使用共享内存作为暂存区，则可能还会进行共享内存操作。
当 LLVM 后端面向 NVIDIA GPU 时，下一步是将其降级到 PTX。LLVM 加载指令直接映射到相应的 PTX 加载指令。
// Example PTX snippet
ld.global.f32 %r1, [%thread_ptr]; // Load from global memory
// Or if staged via shared memory:
ld.shared.f32 %r1, [%shared_ptr];
最后，NVIDIA 驱动程序编译器 (ptxas) 将 PTX 汇编代码转换为 SASS（流式汇编），即 GPU 核心实际执行的本地机器代码。如前所述，此步骤通常采用即时编译 (JIT) 方式完成。
// Example SASS snippet (conceptual, might involve specific opcodes)
MOV R1, ...
LDG.E.32 R1, [R0]; // Load from global memory
DeepSeek 的方法
链接到标题
值得注意的是，CUDA 是为 NVIDIA GPU 生成机器代码的一种非常便捷的方式。Triton 在将 PTX 降级为 SASS 时会使用 CUDA。降级到 PTX 的过程是通过 LLVM IR 完成的，但需要 CUDA 工具包提供的库（例如 libdevice.bc）。
DeepSeek 与 Triton：手写代码 vs 自动生成代码。DeepSeek 采用了另一种方法，选择手动优化 PTX：它手工编写了自定义的 PTX 内核辅助程序，将 PTX 视为 GPU 的汇编语言。DeepSeek 没有依赖通用编译器来生成这种高度非标准的专用代码。相反，Triton 作为一种通用语言，使用标准的 LLVM 编译器后端和 NVPTX 目标，自动将其中间表示（Triton-IR/MLIR）转换为优化的 PTX 汇编代码。
DeepSeek 封装器位于 csrc/kernels/utils.cuh，使用了 mbarrier.try_wait 指令
这为什么重要？引用DeepSeek技术报告中的一段话：“具体来说，我们（DeepSeek）采用了定制的PTX（并行线程执行）指令，并自动调整通信块大小，这显著降低了L2缓存的使用率以及对其他SM的干扰。”这是否意味着LLVM NVPTX后端效率并非最优？在我看来，这意味着还有进一步优化的空间，不仅在LLVM层面，而且在更好地利用裸机ISA方面也是如此！这绝对是双赢！
＃ 结论
瞧，这份备忘录又一次比预期花费了更多时间。但这是介绍 Cuda-Q 架构之前非常必要的前期研究工作。完成后我会更新这篇文章。
（图片来源：LLVM MOS）
＃ 参考
本备忘录中使用的DrawIO图表：