企业官网建设流程全解析

Booth 是一个开源 GPU 编译器，它直接把 CUDA C、HIP 和 Triton 源码编译成 AMD、NVIDIA、Tenstorrent 甚至 x86 CPU 的可执行机器码——全程不依赖 LLVM，不依赖 NVCC，整个编译器只用纯 C99 写成，一条make命令就能构建。对于想读懂 GPU 编译原理、或需要一条轻量可控编译链路的开发者来说，这是一个难得的完整样本。

🧭 一句话认识 Booth

Booth 的工作方式和你熟悉的nvcc、ROCm 工具链很像：给它.cu或.hip文件，它吐出目标平台能直接加载的二进制。区别在于中间没有任何"黑盒"：

• 输入：CUDA C、HIP、Triton Python 内核（也支持 Fortran、OCaml、MLIR 文本）
• 中间表示：项目自研的 BIR（Booth Intermediate Representation），指令集完整文档见 docs/INSTRUCTIONS.md
• 输出：AMD RDNA 2/3/4 二进制、NVIDIA PTX/cubin、Tenstorrent Metalium C++ 与原生 RV32IM ELF、x86-64 主机目标文件，甚至可以在没有 GPU 的笔记本上直接跑 Triton 矩阵乘法

项目以汇编语言鼻祖 Kathleen Booth 命名，二进制叫kath，细节和来龙去脉都写在 README.md 里。

🛠 为什么值得"纯 C99"做编译器

多数 GPU 工具链建立在 LLVM 之上，规模庞大、构建复杂。Booth 证明了另一条路：

• 构建依赖只有一个 C99 编译器，Makefile 中-std=c99加上严格的告警开关就是全部构建配置
• 前端（词法、解析、语义分析）全部手写，位于 src/fe/，支持完整的 C 预处理器、模板、错误恢复
• 优化 Pass（常量折叠、死代码消除、SROA、mem2reg、设备函数内联、感知分支发散的 SSA 寄存器分配）全部可见可改，位于 src/ir/
• 官方态度非常诚实：docs/features.md 明确列出了"什么能编译、什么还不行"，没有夸大

⚙️ 一条流水线：前端 → BIR → 多后端

所有前端（CUDA/HIP/Triton/MLIR…）都汇入同一个 BIR，之后走完全相同的管线，这在 src/main.c 的共享后端分发器中有清晰注释。新增目标架构时只需实现isel（指令选择）和emit（发射）两个环节，接口契约在 src/backend/backend.h，官方还准备了可直接复制的骨架 src/backend/skeleton/，说明文档见 docs/backends.md。

🚀 三步快速上手：零依赖构建与运行

1. 获取源码：git clone https://gitcode.com/gh_mirrors/bar/Booth（也可直接下载官方预编译包，支持 Linux / macOS / Windows）
2. 构建：在项目根目录执行make，唯一前提是系统里有 gcc 或 clang
3. 运行：

./kath run kernel.cu # 自动检测本机设备，编译并运行 ./kath doctor # 报告可用设备并自检 ./kath --ir kernel.cu # 查看编译产物（BIR）

完整命令、每个后端和每个开关的说明都在 docs/usage.md。

🌐 不只有 CUDA：Triton、HIP 与更多前端

Booth 的前端阵容比名字里"GPU 编译器"听起来更广：

• Triton：--triton直接解析@triton.jitPython 内核，tl.dot矩阵乘法可在纯 CPU 后端跑通
• HIP：.hip文件自动启用，与任意后端自由组合
• Fortran / OCaml / MLIR 文本：分别经由 LFortran、kcomp和--mlir标志接入同一条 BIR 管线
• 多翻译单元：一次传多个.cu文件，独立编译后链接进同一模块

✅ 真实硬件验证，不只是玩具

docs/hardware.md 列出了已通过实硅验证的平台：AMD MI300X（8/8 测试内核通过）、RDNA3、NVIDIA RTX 4060 Ti（PTX 路径由驱动 JIT 加载，全程无 NVCC 参与，蒙卡基准获得 3.8 倍加速）。测试语料包含 14 个源文件、35+ 内核、约 27KB 机器码，源码在 tests/。

📚 新手阅读路线：从哪些文件开始

• 想看前端：src/fe/lexer.c、src/fe/parser.c
• 想看中间表示与优化：src/ir/bir.h、src/ir/bir_cfold.c
• 想看机器码发射：src/amdgpu/encode.c、src/nvidia/emit.c
• 想看编译到运行的闭环：src/exec/ 与运行时 src/runtime/
• 想了解后续方向：docs/roadmap.md

Booth 用纯 C99 证明了：从 CUDA 源码到多架构机器码的完整路径，可以小到一个人读得完。无论是学习 GPU 编译原理，还是为特殊硬件搭建自己的编译后端，它都是一份干净、诚实、可以逐行读懂的开源 GPU 编译器参考实现。