Hacker News 의견들
  • 정말 대단함. 오래전부터 커스텀 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++에서는 둘 이상의 스레드가 같은 iout[i]에 쓰는 코드도 컴파일되지만, DisjointSliceThreadIndex에는 공개 생성자가 없고 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...
  • “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가 보통은 무시해도 될 정도지만 여기서는 중요할 수 있는 약간의 오버헤드를 추가한다는 점임
    예를 들어 배열 경계 검사가 추가 레지스터 사용을 유발해서 커널 동시성을 낮출 수 있는지 궁금함