说实话,看到这个 GitHub 仓库的时候,我第一反应是:这哥们有点 6

在显卡领域,Nvidia 筑起了一道名为 CUDA 的高墙,墙内是生态,墙外是荒原。

AMD 这么多年想破墙,搞了个 HIP 转换层,像是在做“翻译”,总隔着一层。

结果,有个叫 Zane 的程序员,直接把墙凿了个洞。

他写了个编译器,叫 BarraCUDA。

不需要 LLVM,不需要中间商,直接把 CUDA 代码变成 AMD 显卡能听懂的指令。

15,000 行代码的暴力美学

先看一段代码,非常基础的向量加法,这是 GPU 编程界的 "Hello World"。

image

__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];
}

通常情况下,你要跑这段代码,得装 Nvidia 的工具链,或者用 AMD 的 HIP 转换一下,还得依赖庞大的 LLVM 库。但在 BarraCUDA 里,命令行只有这一行:

$ ./barracuda --amdgpu-bin vector_add.cu -o vector_add.hsaco
wrote vector_add.hsaco (528 bytes code, 1 kernels)

这就完了?

No LLVM required :-)

image

作者最后加了个笑脸,但这背后全是科技与狠活。

为了绕过 LLVM 这个庞然大物,他手写了一个完整的编译器前端和后端。词法分析、语法树、类型检查、中间表示(BIR)、指令选择、寄存器分配……全是一行行敲出来的。

总共 15,117 行代码。更有意思的是他的代码风格。所有数据结构都用预分配的固定大小数组。热路径里没有 malloc,没有递归,循环都有界。

作者自己吐槽说,这代码风格严谨得让 NASA JPL的标准委员会都点头。

这种“反人类”的 C 语言写法,换来的是极致的掌控力。

AMD 官方文档在撒谎

写个编译器难吗?

难。

给 AMD 显卡写编译器,更是难上加难。

GitHub原文里有一段专门讲 GFX11 Encoding Notes,大家有兴趣可以去看看,反正我觉得脑壳疼。

如果你想给 AMD GPU 写后端,官方手册那 500 页文档里全是坑。

比如 SOP1 前缀,文档写的是一个样,实际是 0xBE800000;VOP3 的目标寄存器位宽,正常人会猜是 [15:8],结果它偏偏在 [7:0]。

更别提 RDNA 3 默认是 Wave32,而不是老架构的 Wave64。

这 1,735 行的 amdgpu_emit.c,就是作者踩完所有坑后留下的“防坑指南”。

他甚至专门准备了一个测试文件叫 notgpt.cu。

为什么叫这个名字?因为现在的 AI 写 CUDA 代码喜欢瞎堆砌功能,什么 SGEMM、直方图、前缀扫描全往上塞,还带着极其讽刺的注释。

BarraCUDA 拿来练手,居然也能跑通。

不是“护城河杀手”,但足够疯狂

评论区里有人问:这玩意儿能打破 Nvidia 的软件护城河吗?

老实讲,目前还不能。

BarraCUDA 现在只支持 GFX11,也就是 AMD 最新的 RDNA 3 架构。

这对拥有 RX 7000 系列显卡的游戏玩家是个好消息,但对用 CDNA 架构做数据中心的企业用户来说,暂时还没戏。

而且,它的局限性也很明显。

你想用 unsigned?不行,得写全 unsigned int。
你想用 +=?不行,得展开写成 a = a + b。
你想用 const?不好意思,不支持。甚至连二维数组声明都不行,得自己拍扁成一维。

但这不重要。

重要的是,它展示了另一种可能性:当一个人不再依赖庞大的工业级标准,而是回归代码本质时,他能做到什么程度。

结尾有点任性

image

项目作者是新西兰人,他在文档里留了个联系方式,顺便写了一句:

"Open an issue if theres anything you want to discuss. Or don't. I'm not your mum."(想讨论就提 issue,不想提就不提。我又不是你妈。)

这种大洋洲式的幽默,在如今全是 AI 生成内容、充满了“废话文学”的互联网上,简直像一股清流。

许可证是 Apache 2.0,他说如果这编译器真的投入生产了,记得告诉他。

“主要是为了让我能在领英上更新点更有趣的东西,总比写着玩强。”


参考链接:
https://github.com/Zaneham/BarraCUDA