Zaneham/BarraCUDA

# BarraCUDA 一个从零开始用 C99 编写的开源 CUDA C++ 编译器，它接收 `.cu` 文件并将其编译为 AMD GPU 机器码、NVIDIA PTX 和 Tenstorrent Tensix C++，计划支持更多架构。不依赖 LLVM，无依赖，无需许可。 当你审视 NVIDIA 的封闭生态并想“这能有多难？”时，这就是会发生的一切。答案是：其实挺难，但我还是做到了。 查看 [更新日志](#changelog) 了解最新动态。 ## 功能 将 CUDA C 源代码（与传给 `nvcc` 的 `.cu` 文件相同）编译为 AMD RDNA 2/3/4 二进制文件、NVIDIA PTX 或 Tenstorrent Tensix Metalium C++。 ``` ┌───────────────────────────────────────────────────────────────────────────┐ │ BarraCUDA Pipeline │ ├───────────────────────────────────────────────────────────────────────────┤ │ Source (.cu) │ │ ↓ │ │ Preprocessor → #include, #define, macros, conditionals │ │ ↓ │ │ Lexer → Tokens │ │ ↓ │ │ Parser (Recursive Descent) → AST │ │ ↓ │ │ Semantic Analysis → Type checking, scope resolution │ │ ↓ │ │ BIR (BarraCUDA IR) → SSA form, typed instructions │ │ ↓ │ │ mem2reg → Promotes allocas to SSA registers │ │ ↓ │ │ Instruction Selection │ │ ├──────────────────┬──────────────────┬────────────────────┤ │ │ ↓ AMD ↓ NVIDIA ↓ Tenstorrent │ │ │ VGPR/SGPR regalloc PTX isel + emit Tensix SFPU isel │ │ │ ↓ ↓ ↓ │ │ │ GFX9/10/11/12 .ptx text Metalium C++ │ │ │ binary encoding (driver JIT) compute/reader/writer │ │ │ ↓ ↓ ↓ │ │ │ .hsaco ELF Runs on NVIDIA Runs on Tenstorrent │ │ │ ↓ hardware hardware │ │ │ Runs on AMD │ │ │ hardware │ │ └───────────────────────────────────────────────────────────────────────────┘ ``` ## 构建 ``` # 它是 C99。它使用 gcc 构建。没有依赖项。 make # 仅此而已。没有 cmake。没有 autoconf。没有 47 步构建过程。 # 如果这不起作用，是你的 gcc 坏了，而不是 Makefile。 ``` ### 要求 - 一个 C99 编译器（gcc、clang 或你已有的任何） - 求生的意志（可选但推荐） - 不需要 LLVM。BarraCUDA 像成年人一样自行完成指令编码。 ## 使用 ``` # 编译为 AMD GPU 二进制文件（RDNA 3，默认） ./barracuda --amdgpu-bin kernel.cu -o kernel.hsaco # 编译为 RDNA 2 ./barracuda --amdgpu-bin --gfx1030 kernel.cu -o kernel.hsaco # 编译为 RDNA 4 ./barracuda --amdgpu-bin --gfx1200 kernel.cu -o kernel.hsaco # 编译为 NVIDIA PTX ./barracuda --nvidia-ptx kernel.cu -o kernel.ptx # 编译为 Tenstorrent Metalium C++ ./barracuda --tensix kernel.cu -o kernel_compute.cpp # 转储 IR（用于调试或好奇） ./barracuda --ir kernel.cu # 仅解析并转储 AST ./barracuda --ast kernel.cu # 运行语义分析 ./barracuda --sema kernel.cu # 使用毛利语（或任何带有翻译文件的语言）的错误信息 ./barracuda --lang lang/mi.txt --amdgpu-bin kernel.cu -o kernel.hsaco ``` ## 运行时启动器 BarraCUDA 包含一个极简的 HSA 运行时（`src/runtime/`），用于在真实 AMD 硬件上分发已编译的内核。零编译时依赖 ROCm —— 通过 `dlopen` 在运行时加载 `libhsa-runtime64.so`。 ``` # 同时编译运行时和示例 gcc -std=c99 -O2 -I src/runtime \ examples/launch_saxpy.c src/runtime/bc_runtime.c \ -ldl -lm -o launch_saxpy # 编译一个内核并运行它 ./barracuda --amdgpu-bin -o test.hsaco tests/canonical.cu ./launch_saxpy test.hsaco ``` 需要安装 ROCm 的 Linux 系统。完整的示例请参见 `examples/launch_saxpy.c`。 ## 支持的功能 以下 CUDA 特性可编译为可用的 GFX9/GFX10/GFX11/GFX12 机器码、NVIDIA PTX 和 Tensix Metalium C++： ### 核心语言 - `__global__`、`__device__`、`__host__` 函数限定符 - `threadIdx`、`blockIdx`、`blockDim`、`gridDim` 内置变量 - 结构体（命名与匿名内联）、枚举、类型定义、命名空间 - 指针、数组、指针运算 - 全部 C 控制流：`if`/`else`、`for`、`while`、`do-while`、`switch`/`case`、`goto`/`label` - 短路求值的 `&&` 和 `||` - 条件运算符 - 模板（基础实例化） - 多返回路径、`continue`、`break` ### CUDA 特性 - `__shared__` 内存（从 LDS 分配并正确跟踪） - `__syncthreads()` → `s_barrier` - 原子操作：`atomicAdd`、`atomicSub`、`atomicMin`、`atomicMax`、`atomicExch`、`atomicCAS`、`atomicAnd`、`atomicOr`、`atomicXor` - 战争期原语：`__shfl_sync`、`__shfl_up_sync`、`__shfl_down_sync`、`__shfl_xor_sync` - 战争期投票：`__ballot_sync`、`__any_sync`、`__all_sync` - 向量类型：`float2`、`float3`、`float4`、`int2`、`int3`、`int4`，支持 `.x`/`.y`/`.z`/`.w` 访问 - 半精度：`__half`、`__float2half()`、`__half2float()`、`__nv_bfloat16` - `__launch_bounds__`（解析、传递并强制 VGPR 上限） - 协同程序组：`cooperative_groups::this_thread_block()`，支持 `.sync()`、`.thread_rank()`、`.size()` - 运算符重载 - 数学内置函数：`sqrtf`、`rsqrtf`、`expf`、`exp2f`、`logf`、`log2f`、`log10f`、`sinf`、`cosf`、`tanf`、`tanhf`、`powf`、`fabsf`、`floorf`、`ceilf`、`truncf`、`roundf`、`rintf`、`fmaxf`、`fminf`、`fmodf`、`copysignf` - `__constant__` 内存、`__device__` 全局变量 ### 编译器特性 - 完整的 C 预处理器：`#include`、`#define`/`#undef`、函数式宏、`#ifdef`/`#ifndef`/`#if`/`#elif`/`#else`/`#endif`、`#pragma`、`#error`、`-I`/`-D` 参数 - 错误恢复（报告多个错误而不会挂起） - 多语言错误消息（`--lang <文件>`）与语言中立的 E 码 - IR 转储中的源码位置跟踪 - 按值传递结构体 ## 示例 ``` __global__ void vector_add(float *c, float *a, float *b, int n) { int idx = threadIdx.x + blockIdx.x * blockDim.x; if (idx < n) c[idx] = a[idx] + b[idx]; } ``` ``` $ ./barracuda --amdgpu-bin vector_add.cu -o vector_add.hsaco wrote vector_add.hsaco (528 bytes code, 1 kernels) ``` 无需 LLVM ✅ ## 验证过的硬件 BarraCUDA 编译的内核已在真实硅片上验证并产生正确结果： - **AMD MI300X（CDNA3，GFX942）** — 8/8 测试内核通过。蒙特卡洛中子输运产生正确的物理结果（k_eff = 0.995，与参考一致）。 - **AMD RDNA3（GFX1100）** — 通过 RDNA3 模拟器 CI 的完整测试套件。 - **NVIDIA RTX 4060 Ti** — PTX 后端，通过 CUDA Driver API 加载，由 NVIDIA 驱动 JIT 编译。蒙特卡洛中子输运基准测试产生正确结果，相比单线程 CPU 提速 3.8 倍。流程中完全不涉及 NVCC。 - **Tenstorrent Blackhole** — 编译为有效的 Metalium C++。硬件验证待开发套件接入。 ## 尚未支持的功能（目前） 诚实地说明限制很重要。以下是当前缺失的功能： - `__device__` 函数中的参数重新赋值（改用局部变量） - 纹理与表面 - 动态并行化（设备端内核启动） - 多翻译单元 - 主机代码生成（仅编译设备代码） 这些都不是架构障碍，只是“尚未实现”的项目。 ## 测试套件 14 个测试文件，35+ 个内核，约 1,700 条 BIR 指令，约 27,000 字节机器码： - `vector_add.cu` — GPU 计算的“Hello World” - `cuda_features.cu` — 原子操作、战争期原语、屏障、goto、switch、短路 - `test_tier12.cu` — 向量、共享内存、运算符重载 - `notgpt.cu` — AI 生成的 CUDA 代码，包含极富讽刺意味的注释（平铺 SGEMM、规约、直方图、前缀扫描、模板、半精度、协同程序组及“万能”内核） - `stress.cu` — N 体模拟、嵌套控制流、位操作、按值传递结构体、链式函数调用 - `canonical.cu` — 来自 NVIDIA 示例并适配解析器的规范模式 - `test_errors.cu` — 故意编写的语法错误以验证错误恢复 - `test_launch_bounds.cu` — `__launch_bounds__` 解析与 VGPR 上限强制 - `test_coop_groups.cu` — 协同程序组降级 - `mymathhomework.cu` — 三角恒等式、指数增长、牛顿-拉夫森、对数定律、双曲函数、向上/向下取整/取整、幂律、限幅 - 还包括预处理器测试、模板测试、无符号整数测试 ## 路线图 ### 近期：强化 修复已知缺口：整数字面量后缀、`const`、参数重新赋值。这些都是小型解析器/降级器的变更。目标是无需修改即可编译真实的 `.cu` 文件。 ### 中期：优化 生成的代码可用但尚未赢得任何基准。目前已完成：指令调度、常量折叠、死代码消除、基于分歧的 SSA 寄存器分配。优先级： - 循环不变代码外提 - 基于寄存器压力的占用率调优 ### 长期：更多架构 IR（BIR）是目标无关的，后端清晰分离。添加新目标编写新的 `isel` + `emit` 对。 - **NVIDIA PTX** — 已完成。编译 CUDA 为 PTX 文本，通过 CUDA Driver API 加载并由 NVIDIA 驱动 JIT 编译。参数：`--nvidia-ptx` - **Tenstorrent Tensix** — 已完成。编译 CUDA 为 TT-Metalium C++ 以用于 Blackhole。参数：`--tensix` - **Intel Arc** — Xe 架构。若能覆盖所有四大 GPU 厂商，BarraCUDA 将更完整。 - **RISC-V 向量扩展** — 当 GPU 过于主流而你希望在软核上运行 CUDA 时。 ## 贡献 **任何语言的问题与 PR 都欢迎** —— 只需附带英文翻译即可。完整风格、命名及协助指南请参见 [CONTRIBUTING.md](CONTRIBUTING.md)。 HLASM 风格的短标识符（`ra_gc`、`mk_hash`、`enc_vop3`）偶然具有文化中立性，5 个字符的标签并无“英文”属性。如果你发现漏洞或有想法，请用你习惯的语言写下它。 ## 更新日志 **2026-03-18** — NVIDIA PTX 后端（`--nvidia-ptx`）。将 CUDA 编译为 PTX 文本，通过 CUDA Driver API 加载并由 NVIDIA 驱动 JIT 编译。验证于运行蒙特卡洛中子输运基准的 RTX 4060 Ti，物理结果正确。无需 NVCC。还包括解析器、语义分析器和降级器中的匿名结构体/联合体支持（`struct { float f; int i; } cvt;` 模式）。 **2026-03-14** — 发散感知 SSA 寄存器分配器（`--ssa-ra`）。在一个 654 行的蒙特卡洛传输内核中消除全部 186 个 VGPR 溢出——临时流量减少 78%，总指令数减少 28%。利用了 Wave64 硬件上发散与统一 VGPR 溢出的 64:1 成本不对称性：统一值通过 `v_readfirstlane` 以每项 4 字节溢出，而发散值保留在寄存器中。灵感来自 Sampaio 等（2013）的发散分析。约 1,300 行 C99 代码，全部静态内存，无 malloc。 **2026-03-09** — 后 isel 验证通过（`bc_vfy`）。编码器曾信任 isel 生成有效机器指令，这并不正确。`bc_vfy` 在两次运行中（后 isel、后 RA）捕获 5 类编码违规，防止二进制文件离开编译器。其首次运行立即发现了 7 个 isel 错误，覆盖 GFX10 与 GFX942 —— 每个都是静默错误编译，会在硬件上触发“未知原因”故障。全数修复。还新增了 `bc_abend` 运行时崩溃诊断，因为若 IBM 能在 1964 年做事后转储，我们也能在 2026 年为 GPU 做此事。 **2026-03-08** — 错误定位基础设施。每个诊断现在拥有语言中立的 ID（`E001`–`E111`）。通过 `--lang <文件>` 提供外部翻译文件。英文参考位于 `lang/en.txt`，毛利语位于 `lang/mi.txt`。统一错误结构体。降级错误现已显示。 **2026-03-05** — CDNA 3 新增功能：GFX942 后端强化、MFMA、Wave64 发散、tinygrad 兼容。MI300X 上 8/8 测试通过（[PR#56](https://github.com/Zaneham/BarraCUDA/pull/56)）。 **2026-03-05** — 指令调度（[PR#52](https://github.com/Zaneham/BarraCUDA/pull/52)）。 **2026-03-03** — CDNA 2 支持（`--gfx90a`，MI250）。tinygrad 兼容。 **2026-02-28** — Tenstorrent Tensix 后端（`--tensix`）。编译 CUDA 为 TT-Metalium C++ 以用于 Blackhole。常量折叠（[PR#51](https://github.com/Zaneham/BarraCUDA/pull/51)）。死代码消除（[PR#48](https://github.com/Zaneham/BarraCUDA/pull/48)）。 **2026-02-25** — HSA 运行时启动器（[PR#40](https://github.com/Zaneham/BarraCUDA/pull/40)）。RDNA 2 支持（`--gfx1030`，[PR#38](https://github.com/Zaneham/BarraCUDA/pull/38)）。测试套件（[PR#41](https://github.com/Zaneham/BarraCUDA/pull/41)）。 **2026-02-20** — RDNA 4 支持（`--gfx1200`，[PR#32](https://github.com/Zaneham/BarraCUDA/pull/32)）。 **2026-02-16** — 初始发布。面向 AMD RDNA 3（gfx1100）的 CUDA 编译器。 ## 联系 发现漏洞？想讨论 AMDGPU 指令编码的细节？需要有人一起吐槽 GPU 计算现状？ **zanehambly@gmail.com** 如有任何问题，欢迎创建议题。或者不。我不是你妈。 位于新西兰，那里已经快到了，而 GPU 和其他地方一样困惑。 ## 许可证 Apache 2.0。随意使用。如果这个编译器 somehow 被用于生产环境，我很乐意听到，主要是为了能在 LinkedIn 上更新比“写了一个 CUDA 编译器当乐子”更有趣的内容。 ## 致谢 - **Fernando Magno Quintão Pereira** 与 **UFMG 的编译原理实验室**（Universidade Federal de Minas Gerais）。Fernando 在看到项目后联系了我，向我介绍了分歧分析论文并提供指导。SSA 寄存器分配器的存在源于那次对话。 - **学术界** — Cooper、Harvey 与 Kennedy 提供了支配树；Braun 与 Hack 提供了 SSA 溢出；Sampaio、Souza、Collange 与 Pereira 提供了分歧分析。我只是个读论文写 C 的业余爱好者。真正艰苦的工作由研究人员完成。 - **Steven Muchnick** 的《Advanced Compiler Design and Implementation》。如果这个编译器做对了一些事，都要归功于这本书。 - **Low Level** 的《Zero to Hero C》课程和 YouTube 频道。我在那里学会了 C。 - **Abe Kornelis**，一位极棒的老师。他在 [z390 Portable Mainframe Assembler](https://github.com/z390development/z390) 项目上的工作值得你花时间。 - 感谢所有发送善意与批评信息的人，感谢你们永远的学生与快乐的业余爱好者。 - 我的奶奶、爷爷、奶奶与 Baka。爱你们 x *He aha te mea nui o te ao. He tāngata, he tāngata, he tāngata.* 世界上最重要的事物是什么？是人，是人，是人。

标签：AMD GPU, C99, CUDA, GPU, GPU 编译, Hakrawler, NVIDIA PTX, RDNA, Tensix, Tenstorrent, Vectored Exception Handling, 图形编程, 客户端加密, 并行计算, 开源, 异构计算, 技术挑战, 无依赖, 源码编译, 编译器, 自研, 跨架构, 高性能计算