비 오는 수요일 오후, 어느 개발자의 모니터.

한쪽에는 C++로 작성된 CUDA 커널이, 다른 쪽에는 Rust 프로젝트가 띄워져 있다. 두 언어를 잇기 위해 복잡한 FFI(서로 다른 언어 간 함수 호출 인터페이스) 바인딩을 설정하며 씨름하는 모습이 보인다.

이런 번거로운 풍경이 곧 바뀐다.

cuda-oxide의 기술적 구성과 컴파일 파이프라인

NVIDIA AI 연구팀이 Rust 코드를 PTX(NVIDIA GPU의 어셈블리 수준 중간 표현)로 직접 컴파일하는 실험적 컴파일러 백엔드인 cuda-oxide를 공개했다. 이 도구는 도메인 특화 언어나 C/C++ 코드 없이 표준 Rust 코드를 통해 SIMT(단일 명령 다중 스레드 방식) GPU 커널을 작성하게 한다. 핵심은 rustc(Rust 컴파일러)의 코드 생성 백엔드를 커스텀하여 CPU 코드가 아닌 디바이스 코드를 생성하는 것이다.

컴파일 과정은 Rust 소스에서 시작해 rustc 프론트엔드를 거쳐 rustc_public(Rust 컴파일러 내부 API를 안정적으로 제공하는 인터페이스)의 Stable MIR(버전 관리가 되는 안정적인 Rust 중간 표현)로 변환된다. 이후 Pliron(Rust로 작성된 MLIR-like 중간 표현 프레임워크)을 통해 dialect-mir, mem2reg, dialect-llvm 단계를 거쳐 LLVM IR(.ll) 파일이 생성된다. 마지막으로 외부 llc(LLVM의 정적 컴파일러 바이너리)가 이를 PTX(.ptx) 어셈블리로 변환하며, 이 파일은 런타임에 CUDA 드라이버에 의해 로드된다.

설치를 위해서는 리눅스 환경(Ubuntu 24.04 테스트 완료)과 CUDA 12.x 버전이 필요하다. Rust nightly-2026-04-03 툴체인과 함께 rust-src, rustc-dev 컴포넌트가 설치되어 있어야 하며, NVPTX 백엔드가 활성화된 LLVM 21 이상의 버전이 요구된다. 특정 llc 바이너리를 지정하려면 아래와 같이 환경변수를 설정한다.

bash
export LLC_PATH=/path/to/llc

기존 GPU 프로그래밍 방식과의 차이점

예전에는 GPU 커널을 작성하려면 C++를 사용해 CUDA 프로그래밍 모델을 직접 다루거나, Triton(GPU 커널 작성을 돕는 파이썬 기반 언어) 같은 추상화 도구에 의존해야 했다. Rust 생태계에서도 Rust-GPU(Vulkan/그래픽스 컴퓨트 타겟)나 CubeCL(JIT 런타임을 사용하는 임베디드 DSL) 같은 시도가 있었으나, cuda-oxide는 CUDA의 프로그래밍 모델 자체를 Rust 내부로 가져오는 데 집중한다.

유사한 프로젝트인 rust-cuda와 비교하면 지향점이 명확히 갈린다. rust-cuda가 async/.await 같은 Rust의 편의성을 GPU로 가져가려는 접근이라면, cuda-oxide는 C++의 __global__ 함수를 작성하는 것과 유사하게 CUDA의 SIMT 실행 모델과 디바이스 내장 함수를 Rust에서 네이티브하게 표현하는 것을 목표로 한다. 두 프로젝트는 서로 보완적인 관계이며, 연구팀은 rust-cuda 유지관리자와 협력하고 있음을 밝혔다.

개발자가 체감하는 기술적 차이는 컴파일러 최적화 단계에서 나타난다. cuda-oxide는 rustc의 JumpThreading(함수 호출을 if문 양쪽 분기로 복제하는 최적화) 기능을 디바이스 코드에서 비활성화한다. CPU에서는 안전한 최적화지만, GPU에서는 모든 스레드가 동일한 bar.sync(배리어 동기화) 지점에서 만나야 하는 배리어 시맨틱을 깨뜨리기 때문이다. 이를 위해 동기화 프리미티브를 LLVM IR에서 convergent로 표시하여 최적화 과정에서 위치가 바뀌거나 복제되지 않도록 제어한다.

실제 구현 단계에서 호스트 코드와 디바이스 코드는 하나의 .rs 파일에 공존한다. 개발자는 #[kernel] 절차적 매크로를 사용하여 커널 함수를 지정하며, cargo oxide 명령어를 통해 빌드한다.

bash
cargo oxide --trace

이 명령을 실행하면 백엔드가 cuda_oxide_kernel_<hash>_<name> 접두사가 붙은 함수를 찾아 PTX 파이프라인으로 보내고, 나머지 호스트 코드는 표준 LLVM 백엔드로 처리한다. 결과적으로 하나의 빌드로 호스트 바이너리와 .ptx 파일이 동시에 생성된다. 라이브러리 의존성에서 오는 디바이스 코드는 커널이 실제로 호출할 때만 .rlib 메타데이터에서 Stable MIR을 읽어 컴파일하는 지연 컴파일 방식을 채택했다.

이제 GPU 커널 개발의 핵심은 언어의 선택이 아니라, 하드웨어의 특성을 얼마나 정교하게 Rust의 타입 시스템으로 정의하느냐의 싸움이 된다.