有同学课上问我，他看到AMD设备上面能够有ZLUDA，那么 macOS 上能不能跑 CUDA。

答案基本上是：不行。而且这个"不行"背后的事情，看来还比我原本想的复杂得多。

咱们今天聊聊 ZLUDA 这个项目，顺便说说在 macOS 上做类似事情的难度。

ZLUDA 是什么

ZLUDA 是一个让你在非 NVDIA 显卡上跑 CUDA 程序的项目。不用改代码，不用重新编译，直接把 CUDA 二进制丢过去就能跑，性能还接近原生。

项目作者 Andrzej Janik，一开始是自己搞的开源项目。后来 AMD 好像是资助了让他把 ZLUDA 适配到 AMD 显卡上。然后，好像是投资方把代码强行撤回了，这事儿当初还上的新闻。后来 Janik 宣布找到了新的投资方，项目活了过来。现在主攻的是机器学习方向——llama.cpp 已经完整支持，PyTorch 和 vLLM 正在搞。

这东西怎么实现的

ZLUDA 分两层。

第一层，API 拦截。CUDA 程序启动时会加载 libcuda.so（Linux）或 nvcuda.dll（Windows）。ZLUDA 用自己的动态库替换这些库，导出一模一样的函数名。程序调用 cuMemAlloc，实际上执行的是 ZLUDA 的 cuMemAlloc，然后 ZLUDA 转手调用 AMD 的 hipMalloc。在代码里，你会看到 zluda_common/src/lib.rs:46-47 里的 FromCuda trait——这是所有 CUDA 到 HIP 类型转换的核心。

目前 ZLUDA 导出了大约 140 多个 CUDA Driver API 函数，覆盖了上下文管理、内存分配、kernel 启动、stream、event、graph 这些核心功能。用到了 HIP 的 FFI 绑定（hip_runtime-sys），cuBLAS 映射到 rocBLAS，cuDNN 映射到 MIOpen，cuSPARSE 映射到 rocSPARSE。

第二层，PTX 编译器。这才是真正的硬骨头。

CUDA 程序里的 GPU 代码不是直接编译成机器码的，而是先编译成 PTX（Parallel Thread Execution），一种类似汇编的中间表示。NVIDIA 的驱动会在运行时把 PTX 编译成 GPU 能执行的机器码。

ZLUDA 要做的事情是：解析 PTX，把它转换成 LLVM IR，然后用 AMD 的 LLVM 后端编译成 AMD GPU 的机器码（HSACO 格式）。

可以在 ptx/src/pass/mod.rs:61 看到整个编译流水线——整整 18 个 pass：

normalize_identifiers：把所有字符串标识符换成数字 ID
replace_known_functions：模式匹配，把已知的 PTX 指令替换成函数调用
normalize_predicates：处理 PTX 的谓词寄存器——NVIDIA GPU 的谓词机制，AMD GPU 上没有直接对应
optimize_function_arguments：函数参数优化
resolve_function_pointers：解析函数指针
fix_special_registers：这是重头戏。PTX 里有一堆特殊寄存器，比如 %tid（线程 ID）、%ntid（线程块大小）、%ctaid（线程块 ID），每个都有特定语义。得映射到 AMD 的等价物上
expand_operands：操作数展开
insert_post_saturation：饱和运算支持——PTX 的 add.sat 等指令
deparamize_functions：去参数化
replace_instructions_with_functions_fp_required：需要精确浮点行为的指令降级
normalize_basic_blocks：基本块规范化
remove_unreachable_basic_blocks：死代码消除
instruction_mode_to_global_mode：指令模式规范化——这个 pass 曾经占编译时间的 60%
insert_explicit_load_store：显式内存操作
insert_implicit_conversions：隐式类型转换——PTX 允许在整数和浮点数之间隐式转换，但 LLVM IR 不行
replace_instructions_with_functions：最终指令降级
hoist_globals：全局变量提升
llvm::emit：发射 LLVM IR

然后 LLVM 后端（llvm_zluda/src/compile.rs）接管，链接 AMD 的设备库（ockl.bc、ocml.bc），运行 LLVM 优化 passes，最终生成 AMD GPU 能跑的代码。

光是这一堆 pass 的名字，你就知道这事多复杂了。

说说编译原理层面的事

我承认，我看这东西的时候头大了，编译原理就我学过那点皮毛。但读代码的过程中，有些事情还是挺容易搞明白的。

PTX 和 AMD GPU 的指令集之间，差异比看起来大得多。

NVIDIA GPU 用的是 SIMT（Single Instruction, Multiple Thread）模型，一个 warp 里有 32 个线程同时执行同一条指令。AMD GPU 也有类似的 wavefront，但细节不一样。wavefront 可以是 32 或 64 个线程，AMD 不同架构用的还不一致——RDNA 系列用 wave32，老的 GCN 和 CDNA 用 wave64。ZLUDA 得处理这个差异，这感觉好坑啊。

更深层的麻烦在于，PTX 有一些 AMD GPU 根本没有对应物的东西：

谓词寄存器：PTX 每条指令前面可以加 @p0 这种谓词，条件执行。AMD 没有这个，得转换成分支和 select 指令
饱和运算：add.sat 这类指令，得用多个基础指令模拟
取整模式：NVIDIA 支持 RTN（向最近偶数取整）、RZ（向零取整）、RP（向正无穷取整）、RM（向负无穷取整）。FAQ 里明确说了，Vulkan 和 OpenCL 都暴露不了这些
非规格化数模式：PTX 可以控制 flush-to-zero 等浮点行为
内联汇编：PTX 可以直接嵌入 asm 语句
任意虚函数调用：PTX 支持通过函数指针做虚调用
bindless image：无绑定的纹理/图像访问模式

ZLUDA 还在 FAQ 里写了一个很坦诚的列表，列出了哪些特性永远无法通过 Vulkan 或 OpenCL 暴露，只能通过原生的 AMD 后端来做。

Linux/Windows 上的注入方式

ZLUDA 要让自己被加载进 CUDA 程序，在不同平台上有完全不同的技术方案。

Windows 上，用的是一个叫 Detours 的库（在 detours-sys/ 和 zluda_inject/ 里）。原理是：以 SUSPENDED 状态启动目标进程，在它还没开始跑之前，注入 ZLUDA 的 DLL，hook 掉 LoadLibraryA/LoadLibraryW 等 API 函数。当程序试图加载 nvcuda.dll 时，ZLUDA 在中间拦一道，换成自己的库。

Linux 上用的是 LD_AUDIT 机制（zluda_ld/src/lib.rs）。这个是 GNU 动态链接器的一个高级特性，允许你在库加载时拦截和替换。ZLUDA 拦截了 libcuda.so、libcublas.so、libcudnn.so、libcufft.so、libcusparse.so、libnvidia-ml.so 等一堆库的加载。

这两种注入方式，都是针对 x86/x86_64 平台的。macOS 的动态链接器机制完全不同——用的是 dyld，不是 ld.so。而且 macOS 有 SIP（System Integrity Protection），DYLD_INSERT_LIBRARIES 这个环境变量在受保护进程里是失效的。光是"让自己被加载进去"这一步，macOS 上就已经比 Linux 多了道门槛。

macOS 上为啥这么难

ZLUDA 的官方 FAQ 关于 macOS 只有一句话：

Unlikely to ever happen. There is very little non-deprecated CUDA software for macOS, and what remains will soon be unsupported.

翻译一下：不太可能支持。macOS 上没几个不废弃的 CUDA 软件，剩下的也快凉了。

这话听着挺绝的，但有它的道理。NVIDIA 自 2018 年起就不给 macOS 做驱动了——NVIDIA 和苹果彻底闹掰。最后的官方驱动还是 High Sierra（10.13）时代的，连 Metal 都不支持。所以 macOS 上压根没有能跑 CUDA 的物理环境，你要做的不是一层翻译，而是两层：先翻译成 Metal，再让 Metal 驱动 Apple Silicon 的 GPU。

这比 ZLUDA 的局面难得多。ZLUDA 最聪明的地方在于：它选了一个底层很接近的目标。AMD 的 HIP 运行时本身就是照着 CUDA 设计的，API 语义大部分能直接映射。cuMemAlloc 到 hipMalloc 的转换，就是换个函数名的事情。最难的 PTX 编译，也因为有 AMD LLVM 后端的支持，可以走 LLVM IR 这条成熟路径。

macOS 上没有 HIP，只有 Metal。

Metal 是什么呢。Metal 是苹果自己的图形和计算 API，从 2014 年推出到现在已经发展到 Metal 4 了。它的着色语言是 MSL（Metal Shading Language），基于 C++14，但跟 CUDA C++ 和 OpenCL C 都有很大差异。

Metal 有几个让 CUDA 翻译特别头疼的特性：

首先，Metal 没有独立的计算栈。它确实有计算 kernel（compute shader），但是整个 API 是为图形渲染设计的，计算作为副功能。CUDA 有一整套独立于图形之外的 API：流、事件、图、纹理对象、surface 对象、设备内存管理。这些概念在 Metal 里要么根本没有，要么得用完全不同的思路来实现。

其次，Metal 的着色语言内核模型跟 CUDA 的思路根本不同。CUDA 里你可以做指针运算、动态内存分配、虚函数调用、设备端 kernel 启动（Dynamic Parallelism），Metal 里这些都属于"想多了"的范畴。Metal Shading Language 的内存模型更接近 C++ 的 restrict 风格——严格别名、编译器可以激进优化——而 CUDA 的内存模型宽松得多。

第三，Metal 没有 SPIR-V 支持。ZLUDA 走的 PTX→LLVM IR→AMD ISA 这条路，关键是在于 LLVM IR 这个通用中间层。AMD GPU 的 LLVM 后端是开源的，你可以直接拿来用。但 Apple GPU 的指令集根本不公开，苹果也不提供 LLVM 后端。Metal 的编译器内置在 Xcode 里，通过 metal 命令行工具把 MSL 编译成 GPU 机器码——这是一个黑盒。

这意味着，如果你想在 macOS 上做 CUDA 翻译，你不能走 PTX→LLVM IR→GPU ISA 的路。你只能走 PTX→MSL 的路，然后让苹果的编译器编译 MSL。

PTX→MSL 的翻译，比 PTX→LLVM IR 难得多。

因为 MSL 是高级着色语言，不是中间表示。ZLUDA 的 18 个 pass 里，前 17 个都在降低抽象层级，最终发射的是 LLVM IR——一种接近机器码的表示。但如果你要输出 MSL，你不能直接降低到机器码级别，你得生成一个符合 MSL 语法和语义约束的高级程序。

举个例子：PTX 里的一条 atom.global.add.gpu.b32 %r1, [%r2], %r3（全局内存的原子加操作），在 LLVM IR 里可以是简单的 atomicrmw add，但到了 MSL 里，你得用 atomic_fetch_add_explicit 函数，而且要指明 memory order。MSL 的原子操作模型比 LLVM IR 更接近 C++ 的 std::atomic，语义细节一对不上就会出 bug。

再举个例子：PTX 的 shfl.sync.down.b32（warp 内线程间数据交换），在 AMD 上可以用 __shfl_down 原生指令，但 MSL 里没有直接对应的操作。Metal 用 simd_ 前缀的函数来实现类似功能，但语义和行为不完全等价。

MoltenVK 的启示

有一个项目叫 MoltenVK，https://github.com/KhronosGroup/MoltenVK 做的事从精神上挺像的——把 Vulkan 翻译到 Metal 上运行。

MoltenVK 是 Valve 和 The Brenwill Workshop 搞的，现在已经相当成熟了。Dota 2 的 Mac 版就是用 MoltenVK 跑的。但这个项目花了快 10 年，投入了大量人力，才达到现在这个质量。

MoltenVK 要做的事情其实比 ZLUDA 简单——Vulkan 和 Metal 都是图形 API，虽然设计哲学不同但目标域重叠很大。而 CUDA 和 Metal 的重叠域就小多了。

MoltenVK 的实现里，最大的挑战之一就是 SPIR-V 到 MSL 的翻译。SPIR-V 是一种中间表示，跟 LLVM IR 挺像的。MoltenVK 用开源的 SPIRV-Cross 库来做这个事。SPIRV-Cross 的作者 Bill Hollings 写了大量代码来处理各种边界情况——比如 Vulkan 的 descriptor set 怎么映射到 Metal 的 argument buffer，push constant 怎么变成 Metal buffer 绑定，等等。

但即使有 SPIR-V 这种标准化的 IR，翻译的质量和性能也不是完美的。有些 Vulkan 特性在 Metal 上只能用软件模拟，性能差好几个数量级。MoltenVK 团队花了好多年才把常见的游戏场景跑顺。

回到 CUDA→Metal，你要面对的是一个远更复杂的问题。CUDA 没有标准的 IR（PTX 是 NVIDIA 私有格式），CUDA 的 API 面积极大（140+ 个函数只是 Driver API，还有 Runtime API、cuBLAS、cuDNN、cuFFT、cuSPARSE、NPP 等等），而且计算语义比图形复杂得多。

苹果自己的尝试：Metal 上的机器学习

苹果不是没有在 Metal 上搞计算。他们为 PyTorch 和 JAX 都做了 Metal 后端，可以让训练和推理跑在 Apple Silicon 的 GPU 上。

但这个方案的思路跟 ZLUDA 完全不同。ZLUDA 追求的是"不改源码直接运行"，苹果的方案是"你把框架换成 Metal 后端的版本，重新跑"。

PyTorch 的 Metal 后端用的是 MPS（Metal Performance Shaders）框架，它直接在 PyTorch 的算子层做适配——每个 torch.matmul 调用，后端接到后翻译成 MPS 的矩阵乘法。这不是输入级的翻译，是算子级的替代。

这种方式能有效利用 GPU，但它只适合能改代码、能换框架的场景。存量的大量 CUDA 二进制——比如各种闭源的科学计算软件、渲染器、加密矿工——完全没法用。ZLUDA 做的是底层兼容，苹果做的是上层适配，两个方向完全不同。

最大的障碍其实是法律

ZLUDA 的故事里还有一个没法回避的问题：法律风险。

NVIDIA 从 2021 年开始，在 CUDA 的 EULA 里加了一条：

You may not reverse engineer, decompile or disassemble any portion of the output generated using SDK elements for the purpose of translating such output artifacts to target a non-NVIDIA platform.

翻译：你不能对 CUDA SDK 生成的输出做逆向工程，目的是把它翻译到非 NVIDIA 平台上跑。

这条条款针对的就是 ZLUDA 这类项目。ZLUDA 目前没有被 NVIDIA 找过麻烦，Janik 在 Reddit 上也专门说过"没有收到过来自 NVIDIA 的任何法律威胁"。但法律风险悬在那，任何人做类似项目都得掂量这个。

AMD 当初撤代码也是出于类似的法律顾虑。给一个"可能侵犯 NVIDIA 版权条款"的项目提供资金，对大公司来说是公关和法律双重风险。

在 macOS 上做，你还多一层：苹果的 EULA 和 App Store 审核政策。Metal 的底层细节是未公开的，对 Metal 做能力边界探索的每一步都可能踩到"疑似逆向工程"的红线。这比在 AMD 平台上麻烦得多——AMD 的 HIP 和 ROCm 是开源的，AMD LLVM 后端也是开源的，你可以光明正大地做。

再说说 macOS 上的 GPU 生态

2026 年了，macOS 的 GPU 计算生态比几年前好了一些，但也有限。

苹果在推自己的 MLX 框架（对标 PyTorch），Pure ML Framework 也可以直接用 Metal 跑。还有 llama.cpp 的 Metal 后端已经比较成熟了，可以在 M1/M2/M3/M4 的 GPU 上跑大模型推理。

但这些都不是 CUDA 兼容。它们是"我们知道 CUDA 不好搞，所以干脆做个新的生态"的路线。

如果你想运行的是 PyTorch，现在直接用 Metal 后端就行了。如果你想运行的是 llama.cpp，直接在 Mac 上编译 Metal 版就行了。但如果你有一个十年前的 CUDA 物理模拟程序，源代码早就丢了，只有二进制，那你还是得找一台带 NVIDIA 显卡的机器。

所以到底有没有希望

我不想画饼，感觉如果在 macOS 上做一个类似 ZLUDA 的 CUDA 兼容层，技术上可行。MoltenVK 证明了 A→Metal 翻译这种事是能做的。但投入会非常巨大。

ZLUDA 能做到今天这样，靠的是几个关键条件：（1）AMD 的 HIP 运行时提供了接近 CUDA 语义的 API；（2）AMD 的 LLVM 后端是开源的，能走 PTX→LLVM IR→GPU ISA 的路；（3）作者 Janik 全职干了好几年。

在 macOS 上，这三个条件一个都不满足。（1）Metal 跟 CUDA 语义差距巨大；（2）Apple GPU LLVM 后端不存在，编译路径不透明；（3）没有足够的人力。

还有一个我觉得挺关键的因素：苹果自己也不想看到 CUDA 在 Mac 上跑。苹果的策略一直是推自己的技术栈——Metal、MLX、MPS、Core ML。如果 CUDA 能在 Mac 上跑，开发者就没动力用 Metal 了。这不是技术问题，这是商业竞争。

之前 OpenAI 跟 AMD 签了 6GW 的算力协议，AMD 股价应声大涨——NVIDIA 的 CUDA 生态垄断是整个行业最头疼的问题之一。如果苹果哪天突然宣布"我们的 Mac 也能跑 CUDA 的代码了"，NVIDIA 的股价大概也得受影响。但苹果为什么要这样做呢？苹果的 Apple Silicon GPU 是给自己用的，不是卖给第三方的。帮助打破 CUDA 的垄断，对苹果的直接利益不大。

最后的想法

ZLUDA 是一个让我印象深刻的项目。不是因为它能做到什么，而是因为它让我看到，在 GPU 计算的底层，我们离真正的可移植性到底有多远。

NVIDIA 花了将近 20 年构建 CUDA 生态——2006 年发布，到现在积累了几亿行依赖 CUDA 的代码。而所有试图打破这个生态的努力——AMD 的 HIP、Intel 的 oneAPI/SYCL、Khronos 的 OpenCL 和 Vulkan Compute——没有一个真正成功到足以威胁 CUDA 的地位。

ZLUDA 的聪明之处在于"我不跟你争生态，我让你不用改代码就能跑"。但这种聪明也有局限——它依赖于目标平台跟 CUDA 在底层足够像。

macOS 上的 Apple Silicon，跟 NVIDIA GPU 从根本上就不像。不只是指令集不像，是整个编程模型、内存模型、执行模型都不像。想要在 Mac 上跑 CUDA 二进制，你需要的不是一个翻译层，而是一个完整的虚拟机。

我觉得，Mac 用户最好的选择，不是等着某天 CUDA 兼容层出现，而是直接拥抱 Metal 生态。如果你有没法迁移的 CUDA 代码，那就老老实实弄台 Linux 机器，买个带 NVIDIA 卡的服务器。毕竟，我自己所有模型的训练也都跑在 Linux 机器上。

预览时标签不可点

Name cleared

微信扫一扫赞赏作者

Like the AuthorOther Amount

赞赏后展示我的头像

作品

暂无作品

Like the Author

Other Amount

最低赞赏 ¥0

Back

Other Amount

赞赏金额

最低赞赏 ¥0

基础知识 · 目录

基础知识

上一篇学习率的前世今生：从全局单一到逐层定制下一篇电车难题？厂商的智驾会在事故前突然退出么？

搜索「」网络结果

调整当前正文文字大小

100%