NVlabs/cuda-oxide 源码：https://github.com/NVlabs/cuda-oxide

cuda-oxide

cuda-oxide 是一个用于在纯 Rust 中编译 GPU 内核的自定义 rustc 后端。该工作空间包含：

• 单源编译 — 主机代码和设备代码位于同一文件，通过一条 cargo oxide build 命令构建
• 一个 rustc 代码生成后端，将 #[kernel] 函数编译为 CUDA PTX
• 设备端抽象（类型安全索引、共享内存、作用域原子操作、屏障、TMA、warp/cluster 操作）
• 用于内存管理和内核启动的主机端运行时（cuda-core、cuda-async）
• 基于 Pliron（https://github.com/vaivaswatha/pliron）的 Rust 原生编译流水线 —— 一个 Rust 中的类 MLIR IR 框架（Rust → Rust MIR → Pliron IR → LLVM IR → PTX）

项目状态

cuda-oxide 是一个实验性编译器，展示了如何以纯 Rust 原生编写 CUDA SIMT 内核 —— 无需 DSL，无需外部语言绑定 —— 并将其提供给更广泛的 Rust 社区。

该项目处于早期阶段（alpha），正在积极开发中：您可能会遇到 bug、功能不完整和 API 破坏性变更，因为我们正在努力改进它。尽管如此，我们希望您能在自己的工作中尝试它，并通过分享使用体验来帮助塑造其发展方向。如果您有兴趣为该项目做出贡献，请参阅 CONTRIBUTING.md。

快速开始

use cuda_device::{kernel, thread, DisjointSlice};
use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig};
use cuda_host::{cuda_launch, load_kernel_module};

// 设备端：泛型内核，对每个元素应用任意函数。
// F 可以是带捕获的闭包 —— rustc 将其单态化为具体类型。
#[kernel]
pub fn map<F: Fn(T) -> T, T: Copy>(f: F, input: &[T], mut out: DisjointSlice<T>) {
    let idx = thread::index_1d();
    if let Some(out_elem) = out.get_mut(idx) {
        *out_elem = f(input[idx.get()]);
    }
}

fn main() {
    let ctx = CudaContext::new(0).unwrap();
    let stream = ctx.default_stream();

    let data: Vec<f32> = (0..1024).map(|i| i as f32).collect();
    let input = DeviceBuffer::from_host(&stream, &data).unwrap();
    let mut output = DeviceBuffer::<f32>::zeroed(&stream, 1024).unwrap();

    let module = load_kernel_module(&ctx, "host_closure").unwrap();

    // 使用闭包启动 —— factor 被自动捕获并传递到 GPU
    let factor = 2.5f32;
    cuda_launch! {
        kernel: map::<_, f32>,
        stream: stream,
        module: module,
        config: LaunchConfig::for_num_elems(1024),
        args: [move |x: f32| x * factor, slice(input), slice_mut(output)]
    }.unwrap();

    let result = output.to_host_vec(&stream).unwrap();
    assert!((result[1] - 2.5).abs() < 1e-5);
}

上述示例定义了一个泛型 #[kernel] 函数 map，它接受任意 Fn(T) -> T 闭包。在主机端，CudaContext 和 DeviceBuffer 管理 GPU 上下文和内存，cuda_launch! 将内核分派到 GPU。闭包 move |x| x * factor 被捕获、标量化，并作为 PTX 内核参数自动传递。PTX 在单次 cargo build 调用中与主机二进制文件一起生成。

对于可组合的异步 GPU 工作，相同的启动点看起来几乎相同：stream: 消失，cuda_launch_async! 返回一个惰性 DeviceOperation，执行在调用 .sync() 或 .await 时发生。

use cuda_async::device_operation::DeviceOperation;
use cuda_host::cuda_launch_async;

// 假设 `module`、`input` 和 `output` 来自 cuda-async 设置：
let factor = 2.5f32;
cuda_launch_async! {
    kernel: map::<_, f32>,
    module: module,
    config: LaunchConfig::for_num_elems(1024),
    args: [move |x: f32| x * factor, slice(input), slice_mut(output)]
}
.sync()?; // 或：.await?;

有关完整的异步设置，请参阅 async_mlp 示例和 crates/cuda-async/README.md。

# 构建并运行示例
cargo oxide run host_closure

# 显示完整编译流水线（Rust MIR → dialect-mir → mem2reg → dialect-llvm → LLVM IR → PTX）
cargo oxide pipeline vecadd

# 使用 cuda-gdb 调试
cargo oxide debug vecadd --tui

环境配置

前置要求

• cargo-oxide —— 驱动构建流水线的 cargo 子命令（cargo oxide run、build、debug 等）
• Rust nightly，含 rust-src 和 rustc-dev 组件（在 rust-toolchain.toml 中固定）
• CUDA Toolkit（12.x+）
• LLVM 21+，含 NVPTX 后端（llc 必须在 PATH 中）
• Clang + libclang 开发头文件（clang-21 / libclang-common-21-dev）—— 构建主机 cuda-bindings crate 时 bindgen 需要
• Linux（在 Ubuntu 24.04 上测试）

为什么需要 LLVM 21？ 我们生成的 TMA / tcgen05 / WGMMA 内联函数，LLVM 20 及更早版本的 llc 无法处理。简单内核可能仍可在旧版 llc 上运行，但任何 Hopper / Blackwell 相关功能都需要 21+。

安装

cargo-oxide

在 cuda-oxide 仓库内，cargo oxide 通过工作空间别名即可直接使用。要在仓库外使用（您自己的项目）：

cargo install --git https://github.com/NVlabs/cuda-oxide.git cargo-oxide

首次运行时，cargo-oxide 将自动获取并构建代码生成后端。

Rust

# 工具链通过 rust-toolchain.toml 自动安装
# 如需手动安装：
rustup toolchain install nightly-2026-04-03
rustup component add rust-src rustc-dev --toolchain nightly-2026-04-03

CUDA

export PATH="/usr/local/cuda/bin:$PATH"
nvcc --version

LLVM

# Ubuntu/Debian
sudo apt install llvm-21

# 验证 NVPTX 支持
llc-21 --version | grep nvptx

流水线会自动在 PATH 中查找 llc-22 和 llc-21（按此顺序）。要固定特定二进制文件，请设置 CUDA_OXIDE_LLC=/usr/bin/llc-21。

Clang（主机 cuda-bindings）

主机 cuda-bindings crate 运行 bindgen，它会加载 libclang 并需要 clang 自身的 resource-dir stddef.h —— 仅 libclang1-* 运行时是不够的。

sudo apt install clang-21  # 或 libclang-common-21-dev

cargo oxide doctor 会提前检测此问题；否则症状是在主机构建期间出现晦涩的 'stddef.h' file not found 错误。

验证安装

# 检查所有前置条件是否就绪
cargo oxide doctor

# 端到端构建并运行示例
cargo oxide run vecadd

cargo oxide doctor 会验证您的 Rust 工具链、CUDA 工具包、LLVM 和代码生成后端。如果一切配置正确，cargo oxide run vecadd 会将 Rust 内核编译为 PTX，在 GPU 上启动它，并打印 ✓ SUCCESS: All 1024 elements correct!。

示例

crates/rustc-codegen-cuda/examples/ 中有 46 个示例。亮点：

cargo oxide run vecadd
cargo oxide run gemm_sol

Crate 概览

面向用户的 Crate

编译器 Crate

构建工具

文档

状态

亮点：

• 端到端 Rust -> PTX 编译
• 统一的单源编译（主机 + 设备在同一文件中）
• 带单态化的泛型函数
• 带捕获的闭包（move 和非 move，通过 HMM）
• 用户自定义结构体、枚举、模式匹配
• 完整的 GPU 内联函数支持（线程、warp、共享内存、屏障、TMA、簇、原子操作）
• 跨 crate 内核
• Blackwell+ 的 LTOIR 生成（设备端 LTO）
• 设备 FFI：通过 LTOIR 实现 Rust <-> C++/CCCL 互操作
• MathDx 集成：cuFFTDx 线程级 FFT、cuBLASDx 块级 GEMM
• 主机运行时：cuda-core（显式控制）和 cuda-async（可组合异步操作）
• GEMM SoL：在 B200 上达到 868 TFLOPS（cuBLAS SoL 的 58%），使用 cta_group::2、CLC、4 阶段流水线

文档

进行中： 🚧 cuda-oxide book (https://nvlabs.github.io/cuda-oxide/) 是该项目的主要参考。它涵盖了 Rust 中的 SIMT 内核编写、同步和异步 GPU 编程、编译器架构等内容。

要在本地构建并提供书籍服务，请参阅 cuda-oxide-book/README.md。

生态

cuda-oxide 是众多正在积极开发中的 Rust + GPU 项目之一。该领域的项目解决不同部分的问题 —— 用于图形的 Vulkan/SPIR-V、通过 LLVM 的隐式卸载、第三方 CUDA 后端、安全驱动绑定 —— 我们一直在与更广泛的 Rust GPU 社区维护者合作，探讨如何共同推动 Rust 中的 GPU 计算发展。

关于 cuda-oxide 相对于其他项目的定位，请参阅书籍的 Ecosystem 附录（https://nvlabs.github.io/cuda-oxide/appendix/ecosystem.html）。

许可证

cuda-bindings crate 基于 NVIDIA 软件许可证授权：LICENSE-NVIDIA。

所有其他 crate 基于 Apache License, Version 2.0 授权：LICENSE-APACHE。