# CUDA-oxide: Nvidia의 공식 Rust-to-CUDA 컴파일러 > Clean Markdown view of GeekNews topic #29414. Use the original source for factual precision when an external source URL is present. ## Metadata - GeekNews HTML: [https://news.hada.io/topic?id=29414](https://news.hada.io/topic?id=29414) - GeekNews Markdown: [https://news.hada.io/topic/29414.md](https://news.hada.io/topic/29414.md) - Type: GN+ - Author: [neo](https://news.hada.io/@neo) - Published: 2026-05-12T09:55:26+09:00 - Updated: 2026-05-12T09:55:26+09:00 - Original source: [nvlabs.github.io](https://nvlabs.github.io/cuda-oxide/index.html) - Points: 1 - Comments: 1 ## Topic Body - **cuda-oxide**는 안전에 가까운 관용적 Rust로 SIMT GPU 커널을 작성하고 표준 Rust 코드를 PTX로 직접 컴파일하는 실험적 컴파일러임 - DSL이나 외국어 바인딩 없이 **Rust**만 사용하며, 소유권·트레이트·제네릭 이해를 전제로 하고 async 장은 `.await` 지식도 필요함 - v0.1.0은 **초기 알파** 릴리스라 버그, 미완성 기능, API 파괴적 변경을 예상해야 함 - 예제는 `cargo oxide run vecadd`로 실행하며, `#[cuda_module]` 안의 `#[kernel]` 함수가 `thread::index_1d()`로 벡터 덧셈을 수행함 - `#[cuda_module]`은 **디바이스 아티팩트**를 호스트 바이너리에 포함하고, 타입 지정 로더와 커널별 실행 메서드를 생성함 --- ### 사용 방식과 생성 코드 - ## 빠른 시작 - 설치 전제 조건을 갖춘 뒤 `cargo oxide run vecadd`로 예제를 빌드하고 실행함 - 설치 안내는 [prerequisites](https://nvlabs.github.io/cuda-oxide/getting-started/installation.html)에 있음 - 예제는 `#[cuda_module]` 모듈 안에 `#[kernel]` 함수 `vecadd`를 정의하고, `thread::index_1d()`로 인덱스를 얻어 `a[i] + b[i]`를 `DisjointSlice<f32>`에 기록함 - 호스트 쪽에서는 `CudaContext::new(0)`, 기본 스트림, `kernels::load(&ctx)`를 사용하고, `DeviceBuffer::from_host`, `DeviceBuffer::<f32>::zeroed`, `LaunchConfig::for_num_elems(1024)`로 커널을 실행함 - 실행 결과를 `c.to_host_vec(&stream)`로 가져와 `result[0] == 3.0`을 확인함 - ## `#[cuda_module]` 동작 - `#[cuda_module]`은 생성된 디바이스 아티팩트를 호스트 바이너리에 포함함 - 타입이 지정된 `kernels::load` 함수와 커널별 실행 메서드를 생성함 - 특정 sidecar 아티팩트를 로드하거나 커스텀 실행 코드를 만들어야 할 때는 낮은 수준의 `load_kernel_module`과 `cuda_launch!` API도 계속 사용 가능함 ### 전제와 방향 - cuda-oxide는 Rust의 타입 시스템과 소유권 모델로 GPU 커널을 작성하는 것을 목표로 하며, 안전성을 1급 목표로 둠 - GPU에는 미묘한 부분이 있으므로 [the safety model](https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-model.html)을 읽을 필요가 있음 - DSL이 아니라 순수 Rust를 PTX로 컴파일하는 커스텀 `rustc` 코드 생성 백엔드임 - GPU 작업을 지연 실행되는 `DeviceOperation` 그래프로 구성하고, 스트림 풀에 스케줄링하며, `.await`로 결과를 기다리는 **비동기 실행**을 지원함 - Rust의 소유권, 트레이트, 제네릭에 익숙하다는 전제를 두며, 이후 async GPU 프로그래밍 장은 `async`/`.await`와 tokio 같은 런타임 지식도 필요함 - 참고 자료로 [The Rust Programming Language](https://doc.rust-lang.org/book/), [Rust by Example](https://doc.rust-lang.org/rust-by-example/), [Async Book](https://rust-lang.github.io/async-book/)이 제공됨 - v0.1.0 릴리스는 초기 알파 단계이며, **버그**, 미완성 기능, API 파괴적 변경을 예상해야 함 ## Comments ### Comment 57270 - Author: neo - Created: 2026-05-12T09:55:26+09:00 - Points: 1 ###### [Hacker News 의견들](https://news.ycombinator.com/item?id=48096692) - 정말 대단함. 오래전부터 커스텀 CUDA 커널과 [https://crates.io/crates/cudarc]()를 써 왔는데, 이건 거의 **드롭인 대체재**가 될 수도 있어 보임 특히 **빌드 시간**이 어떻게 비교될지 궁금함. 대부분의 Rust CUDA 크레이트는 CMake나 nvcc 호출에 의존해서 컴파일이 고통스럽게 느려질 수 있음 마침 지난주에 빌드 시간을 프로파일링하다가 sccache 같은 도구가 산출물 캐싱으로 재빌드 시간을 크게 줄일 수 있다는 걸 봤지만, 그래도 커스텀 nvcc 호출 비용은 남음. 예를 들어 Hugging Face의 candle도 커널 컴파일에서 커스텀 nvcc 명령을 호출함: [https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...]() - Cudarc는 정말 좋음 Rust CUDA 크레이트 대부분이 CMake나 nvcc를 호출해서 컴파일이 느려진다는 부분은 개인적으로 크게 겪어보지 못했음. 빌드 스크립트를 처리하려고 만든 `cuda_setup` 크레이트를 보면, 단순한 `build.rs`라 파일이 바뀔 때만 다시 컴파일하고, Rust의 CPU 쪽 코드에 비하면 컴파일 시간이 아주 작음 - 다른 사람들도 cuda-oxide가 cudarc의 거의 드롭인 대체재처럼 보인다고 생각하는지 궁금함 그렇게 되면 정말 좋겠지만, 개인적으로는 아마 **보완재**에 가깝지 않을까 봄. cuda-oxide를 차별화하는 요소가 뭔지도 궁금하고, NVIDIA가 완전히 통제한다는 점 이상으로 무엇이 있는지 궁금함 - “직접 PTX로” 간다는 게 이상함. 최근 **NVIDIA MLIR**도 꽤 좋고 빠름. 아니면 CuTile이 쓰는 더 쉽고 최근 유행하는 **Tile IR** [1]을 대상으로 할 수도 있었을 것임 Tile IR은 조금 더 높은 수준이라 타깃팅하기 훨씬 쉽고, 에필로그 융합 같은 데서만 손해를 봄 [1] [https://docs.nvidia.com/cuda/tile-ir/]() [2] [https://developer.nvidia.com/cuda/tile]() - Rust의 **메모리 모델**을 CUDA 의미론에 어떻게 맞췄는지가 꽤 궁금함. CUDA C++와 비교해 무엇이 다른지, Rust 타입 시스템이 실제로 CUDA에 더 많은 안전성을 줄 수 있는지도 알고 싶음 GPU 커널 작성은 본질적으로 안전하지 않다고 생각함. 하드웨어 동작 방식과 항상 극한 최적화를 해야 한다는 점 때문에 안전한 언어를 만들기가 너무 어려움 - 크게 보이는 차이는 4가지임. 첫째, `cudaFree`를 수동 호출하는 방식과 달리 **use-after-free**와 drop 의미론을 다룸 둘째, C++의 `void*` 인자는 포인터 배열이고 개수만 검증하지만, 여기서는 `cuda_launch!`로 커널 인자를 강제함 셋째, 가변 쓰기의 별칭 문제임. C++에서는 둘 이상의 스레드가 같은 `i`로 `out[i]`에 쓰는 코드도 컴파일되지만, `DisjointSlice`와 `ThreadIndex`에는 공개 생성자가 없고 [https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52...]() `index_1d`, `index_2d`, `index_2d_runtime` API만 쓰게 되어 있음 넷째, C++에서는 `std::string`이나 사실상 어떤 POD든 `cuda memcpy`해서 상태를 망가뜨릴 수 있지만, 여기서는 `DisjointSlice`, 스칼라, 클로저만 받음 [https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...]() 자세한 내용은 [https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo...]()와 [https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...]()에 있음. 물론 전부 잡아내는 건 아니지만, 생 raw `.cu` 파일보다 정의되지 않은 동작을 막는 **가드레일**을 훨씬 많이 주는 듯함 - 참고로 Rust의 메모리 모델은 의도적으로 C++와 거의 완전히 같음. **원자적 연산**도 같고, 출처(provenance) 같은 개념도 있음 GPU 프로그래밍에 편한 언어인지는 아직 봐야겠지만, GPU 특유의 이상한 요소들을 전부 활용하면서 안전한 코드를 쓰는 괜찮은 DSL 비슷한 API는 만들 수 있어도 놀랍지 않음. CUDA도 결국 그런 것 아닌가 싶음 - 문서에 꽤 자세히 나와 있음. **안전 계층**, 대부분 안전한 계층, 안전하지 않은 계층이 있음 Rust의 `Send`/`Sync` 모델에 쉽게 끼워 넣기 어려웠던 안전하지만 병렬적인 작업에는 약간의 투박함이 필요함 - 목표에 따라 다를 듯함. Rust로 애플리케이션을 쓰고 가끔 그 안에서 GPU 계산을 쓰고 싶은 입장에서는, 솔직히 크게 신경 쓰지 않음 메모리 모델이나 소유권 모델을 낮은 마찰로 활용할 수 있다면 좋음. 하지만 그 때문에 사용 경험이 많이 불편해진다면 그런 방식은 원하지 않음 기준선은 현재 Cudarc가 하는 방식이라고 봄. 메모리 관리가 많이 개입되지는 않고, FFI를 감싼 명령형 문법과 커널이 바뀔 때 nvcc를 호출하는 빌드 스크립트 몇 줄 정도임 - 이게 **Slang**[0]에는 어떤 의미일지 궁금함. 사람들이 더 현대적인 언어로 GPU 프로그래밍을 하고 싶어 한다는 게 핵심일 텐데, 이제 Rust를 쓰면 되는 것처럼 보임 참고로 Slang은 꽤 좋아함 [0]: [https://shader-slang.org/]() - 셰이더 작성은 적어도 현재로서는 CUDA 커널 작성과 실질적으로 다름. 셰이더는 동시에 더 높은 수준이면서 더 낮은 수준이고, 특정하고 제한된 드라이버/GPU 기능 집합을 위해 설계된 결과로 특이한 점이 많음 예를 들면 **디스크립터 세트**, 리소스 레지스터, 디스패치 제한 같은 것들임 - 대상이 다름. Slang 쪽은 AI 알고리즘보다 **그래픽스 프로그래밍**에 더 관심이 많음 셰이딩 언어는 기능 면에서도 더 사용자 친화적임. 게다가 NVIDIA는 이미 Slang을 프로덕션에서 쓰고 있고, 그 사람들이 셰이더 파이프라인을 Rust로 다시 작성하지는 않을 것임 - Rust와 “안전한” 프로그래밍 언어 얘기와 관련해서, NVIDIA가 **Spark/Ada**를 어떻게 쓰는지 더 자세히 아는 사람이 있는지 궁금함 찾을 수 있는 건 아래 내용뿐임 [https://www.adacore.com/case-studies/nvidia-adoption-of-spar...]() - 지난 DEF CON에서 자세한 발표를 했음: [https://www.youtube.com/watch?v=KhWtkZmOPn4]() - AdaCore 콘퍼런스에서 NVIDIA가 2020년에 발표한 “Securing the Future of Safety and Security of Embedded Software” 녹화가 있음 [https://www.youtube.com/watch?v=2YoPoNx3L5E]() - “DSL도 없고, 외국어 바인딩도 없고, 그냥 Rust”라는 문구부터 공식 CUDA 포트라면서 도입 문단도 제대로 신경 쓰지 않은 느낌임 그래도 무시하고 문서를 읽어보려 했는데, 커스텀 IR이 나와서 흥미로워지려는 순간 “MLIR 구현은 C++에 TableGen을 곁들였고, LLVM 전체를 컴파일해야 하는 빌드 시스템과 커리어 선택을 의심하게 만드는 디버깅 세션이 있다”는 식의 문장을 보고 더는 이 업계를 진지하게 받아들이기 어려워짐 - 전체 코드베이스가 대체로 **AI로 작성**된 것 같음 - 웹페이지에 AI를 쓰지 않았으면 “왜 NVIDIA는 자사 웹사이트와 문서를 AI로 쓰지 않나? AI 공장과 수천 에이전트를 관리하는 직원이라는 자기 이야기를 믿지 않는 건가?”라는 반응이 나왔을 것임 이건 AI 과대홍보 기업에게 기대할 만한 정확한 **자사 제품 활용**처럼 보임 - 이름도 CUDA-oxide라고 붙였는데, Rust 언어 이름의 유래가 산화가 아니라 **균류**라는 걸 모르는 티를 냄 - 정확히 뭐가 불만인지 모르겠음. 누군가 MLIR이 매우 복잡하고 LLVM에 의존적이라고 관찰한 것 때문인가? - TileLang [https://github.com/tile-ai/tilelang]()과 Tile Kernels [https://github.com/deepseek-ai/TileKernels]() 같은 것들이 언젠가 **CUDA를 구식**으로 만들 것임 - CUDA는 거의 20년 됐고, 앞으로도 수년간 사라지지 않을 것임 - 근거가 너무 적은데 꽤 큰 주장임 - [https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo...]() 문서를 보면, GPU 커널은 같은 메모리를 동시에 보는 수천 스레드에서 돌고, CPU에서 Rust는 소유권과 빌림으로 데이터 경합을 막지만 GPU에는 SM당 2048개 스레드가 같은 함수에서 시작해 같은 출력 버퍼를 가리키므로 빌림 검사기가 이를 위해 설계되지 않았다고 함 cuda-oxide는 흔한 경우인 “스레드 하나가 원소 하나를 쓴다”는 구조적으로 안전하게 만들고, 공유 메모리·워프 셔플·하드웨어 내장 함수 같은 드문 경우는 문서화된 계약이 있는 `unsafe`를 요구하며, TMA·텐서 코어·클러스터 수준 통신 같은 최전선 기능은 하드웨어 복잡성에 맞춰 완전 수동으로 둔다고 함 그런데 이건 그다지 **Rust답지 않음**. Rust에서는 기존 추상화가 문제에 잘 맞지 않으면 새로운 안전한 추상화를 만듦. Rust for Linux가 그런 예임 안전하지 않다면 Rust를 쓰는 이유가 뭔지 의문임. 마지막 성능을 짜내야 하는 사람에게 unsafe API를 제공하는 건 괜찮지만, 그게 기본값이 되어서는 안 됨 `io_uring`이나 Vulkan 같은 API의 사용자 공간 라이브러리와 비교하게 됨. 그런 것들을 위한 안전한 API 설계는 꽤 어렵고, 실제로 sound하지 않은 시도들도 있음 - 이게 **호스트와 디바이스 사이에 구조체를 공유**하게 해줄지 아는 사람이 있는지 궁금함. 기존 Rust/CUDA 워크플로에서 지금까지 빠져 있던 큰 부분이 바로 그것임 그 사이의 직렬화/바이트 장벽도 마찬가지임 - CUDA에서 Rust를 쓸 때 경계했던 부분은 Rust가 보통은 무시해도 될 정도지만 여기서는 중요할 수 있는 약간의 **오버헤드**를 추가한다는 점임 예를 들어 배열 경계 검사가 추가 레지스터 사용을 유발해서 커널 동시성을 낮출 수 있는지 궁금함