cuda-oxide: 순수 Rust로 CUDA GPU 커널을 작성해 PTX로 컴파일하는 NVIDIA의 실험적 컴파일러

9bow · 6월 12, 2026, 7:30오전

cuda-oxide 소개

cuda-oxide는 GPU 커널을 순수 Rust로 작성해 CUDA PTX로 컴파일하는 커스텀 rustc 백엔드입니다. 기존에 CUDA 커널을 쓰려면 CUDA C++ 같은 별도의 언어나 DSL, 외부 언어 바인딩이 필요했지만, cuda-oxide는 이런 중간 장치 없이 표준 Rust 코드를 그대로 PTX로 내려보냅니다. NVIDIA 연구 조직(NVlabs)이 공개했으며, Rust 커뮤니티가 GPU 컴퓨팅을 네이티브하게 다룰 수 있도록 하는 것을 목표로 합니다.

cuda-oxide의 한 가지 특징은 단일 소스 컴파일(single-source compilation)입니다. 호스트 코드와 디바이스 코드가 같은 파일 안에 공존하고, cargo oxide build 한 번으로 함께 빌드됩니다. #[kernel] 어트리뷰트를 붙인 함수는 rustc 코드젠 백엔드를 거쳐 CUDA PTX로 컴파일되고, #[cuda_module] 매크로가 생성된 디바이스 산출물을 호스트 바이너리에 그대로 임베드합니다. 컴파일 과정은 Rust 소스에서 시작해 rustc의 MIR, MLIR과 유사한 Rust IR 프레임워크인 Pliron 기반의 IR, LLVM IR을 차례로 거쳐 최종 PTX로 이어집니다.

cuda-oxide는 아직 초기 단계(alpha)의 실험적 컴파일러로, 활발히 개발 중입니다. 저자들은 버그와 미완성 기능, API 변경이 있을 수 있다고 명시하고 있으며, 실제 작업에 써 보고 피드백을 통해 방향을 함께 잡아 가기를 권하고 있습니다. 디바이스 측에서는 타입 안전한 인덱싱, 공유 메모리, 스코프 단위 원자 연산(scoped atomics), 배리어, TMA, 워프 및 클러스터 연산을 제공하고, 호스트 측에서는 메모리 관리와 커널 실행을 담당하는 런타임(cuda-core, cuda-async)을 갖추고 있습니다.

cuda-oxide의 핵심 기능

단일 소스 컴파일 은 호스트와 디바이스 코드를 한 파일에 두고 하나의 빌드 명령으로 처리하는 방식입니다. 제네릭 함수는 모노모피제이션(monomorphization)을 통해 구체 타입으로 특수화되며, 캡처를 가진 클로저도 GPU로 전달됩니다. 예제의 map 커널처럼 Fn(T) -> T 클로저를 호스트에서 넘기면, 클로저가 캡처한 값이 스칼라화되어 커널 매개변수로 자동 전달됩니다.

타입 안전한 디바이스 추상화 는 GPU 프로그래밍에서 흔한 인덱싱 오류나 메모리 접근 실수를 Rust의 타입 시스템으로 줄이려는 시도입니다. 스레드 인덱싱, 공유 메모리, 배리어, 원자 연산, 그리고 Hopper 이상에서 쓰이는 TMA와 스레드 블록 클러스터까지 디바이스 인트린식으로 노출됩니다.

비동기 호스트 런타임 은 cuda-async 크레이트로 제공됩니다. 여기서는 {kernel}_async 형태의 호출이 즉시 실행되지 않고 지연된(lazy) DeviceOperation 을 반환하며, .sync() 나 .await 를 호출하는 시점에 실제 실행이 일어납니다. 이를 통해 여러 스트림에 걸친 GPU 작업을 조합할 수 있습니다.

성능 면에서는 GEMM 예제가 대표적입니다. 저자들의 측정에 따르면 gemm_sol 예제 는 B200(148 SM)에서 868 TFLOPS를 기록했고, 이는 cublasLt FP16 SoL의 58% 수준입니다. 이 수치는 예제가 bench/cublaslt_bench 로 cublasLt 기준선을 직접 측정해 비교한 값이며, 절대 TFLOPS는 SM 수에 따라 달라집니다.

cuda-oxide의 사용 예시

아래는 각 원소에 임의의 함수를 적용하는 제네릭 커널을 정의하고, 호스트에서 클로저를 넘겨 실행하는 예제입니다.

use cuda_device::{cuda_module, kernel, thread, DisjointSlice};
use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig};

#[cuda_module]
mod kernels {
    use super::*;

    #[kernel]
    pub fn map<T: Copy, F: Fn(T) -> T + Copy>(f: F, input: &[T], mut out: DisjointSlice<T>) {
        let idx = thread::index_1d();
        let i = idx.get();
        if let Some(out_elem) = out.get_mut(idx) {
            *out_elem = f(input[i]);
        }
    }
}

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 = kernels::load(&ctx).unwrap();

    let factor = 2.5f32;
    module
        .map::<f32, _>(
            &stream,
            LaunchConfig::for_num_elems(1024),
            move |x: f32| x * factor,
            &input,
            &mut output,
        )
        .unwrap();

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

저장소 밖에서 cargo-oxide를 설치하려면 고정된 nightly 툴체인과 함께 설치합니다. 설치 후 첫 실행 시 코드젠 백엔드가 자동으로 빌드됩니다.

# cargo-oxide 설치 (저장소 외부에서 사용할 때)
cargo +nightly-2026-04-03 install --git https://github.com/NVlabs/cuda-oxide.git cargo-oxide

# 사전 요구사항 점검 후 예제 실행
cargo oxide doctor
cargo oxide run vecadd

cargo oxide doctor 는 Rust 툴체인, CUDA 툴킷, LLVM, 코드젠 백엔드가 갖춰졌는지 확인합니다. 실행에 필요한 환경은 Rust nightly와 rust-src·rustc-dev·llvm-tools 컴포넌트, CUDA Toolkit 12.x 이상, Clang, 그리고 Linux(Ubuntu 24.04에서 테스트됨)입니다. Nix를 사용한다면 nix develop 로 재현 가능한 개발 셸을 바로 얻을 수 있습니다. 이 밖에 저장소의 examples 디렉토리 에는 벡터 덧셈부터 Blackwell 텐서 코어(tcgen05), 스레드 블록 클러스터, 비동기 MLP 파이프라인까지 60개가 넘는 예제가 들어 있습니다.

cuda-oxide의 라이선스

cuda-oxide는 대부분의 크레이트가 Apache License 2.0으로 공개되어 있어 개인 및 상업적 목적으로 자유롭게 사용할 수 있습니다. 다만 cuda-bindings 크레이트만은 NVIDIA Software License를 따르므로, 해당 크레이트를 사용할 때는 별도 라이선스 조건을 확인해야 합니다.

cuda-oxide 공식 문서 (cuda-oxide book)

cuda-oxide 프로젝트 GitHub 저장소

더 읽어보기

이 글은 GPT 모델로 정리한 글을 바탕으로 한 것으로, 원문의 내용 또는 의도와 다르게 정리된 내용이 있을 수 있습니다. 관심있는 내용이시라면 원문도 함께 참고해주세요! 읽으시면서 어색하거나 잘못된 내용을 발견하시면 덧글로 알려주시기를 부탁드립니다.

파이토치 한국 사용자 모임이 정리한 이 글이 유용하셨나요? 회원으로 가입하시면 주요 글들을 이메일로 보내드립니다!
텔레그램(Telegram)이나 Slack/Discord/Teams/Dooray/GoogleChat 등으로도 새 글 알림을 받으실 수 있습니다.

아래쪽에 좋아요를 눌러주시면 새로운 소식들을 정리하고 공유하는데 힘이 됩니다~