BarraCUDA - AMD GPU용 오픈소스 CUDA 컴파일러

 (github.com/Zaneham)
5P by GN⁺ | ★ favorite | 댓글 1개
  • CUDA C 소스(.cu) 를 AMD RDNA3(GFX11) 용 기계어로 직접 변환하는 독립형 오픈소스 컴파일러
  • LLVM이나 HIP 계층 없이 자체 어휘 분석기, 파서, 중간 표현(BIR) 을 통해 ELF .hsaco 바이너리를 생성
  • 15,000여 줄의 C99 코드로 작성되었으며, 단일 make 명령으로 빌드 가능
  • CUDA의 스레드 내장 변수, 공유 메모리, 원자 연산, 워프 연산, 협동 그룹 등 주요 기능을 지원
  • Apache 2.0 라이선스로 공개되어 있으며, 향후 Tenstorrent, Intel Arc, RISC-V 등 추가 아키텍처 확장을 목표로 함

BarraCUDA 개요

  • BarraCUDA는 AMD GPU용 CUDA 컴파일러로, .cu 파일을 GFX11 기계어 코드로 변환
    • 결과물은 AMD GPU에서 실행 가능한 ELF .hsaco 바이너리 형태
    • LLVM 의존성 없이 완전 독립적으로 동작
  • 전체 코드는 C99로 작성된 약 15,000줄이며, 단일 Makefile로 빌드 가능
  • 프로젝트는 뉴질랜드 기반 개발자가 개인적으로 개발

작동 방식

  • 입력된 .cu 파일을 전처리 → 어휘 분석 → 파싱 → 의미 분석 → BIR 생성 → 명령 선택 → 레지스터 할당 → 바이너리 인코딩 → ELF 출력 순서로 처리
  • BIR(BarraCUDA IR) 은 SSA 형태의 내부 표현으로, 아키텍처 독립적 설계
  • 모든 인코딩은 llvm-objdump를 통해 검증되어 디코드 오류 0건

지원 기능

  • CUDA 핵심 문법: __global__, __device__, __host__, threadIdx, blockIdx 등
  • CUDA 기능: __shared__ 메모리, __syncthreads(), 원자 연산, 워프 셔플/투표, 벡터 타입, half 정밀도, 협동 그룹 등
  • 컴파일러 기능: 완전한 C 전처리기, 오류 복구, 소스 위치 추적, 구조체 값 전달 지원

미지원 항목

  • unsigned 단독 사용, 복합 대입 연산자(+=, -= 등), const, __constant__ 메모리, 2D 공유 배열, 텍스처·서피스, 동적 병렬 실행 등은 아직 미구현
  • 다중 번역 단위 및 호스트 코드 생성은 지원하지 않음

테스트 및 로드맵

  • 14개 테스트 파일, 35개 이상 커널, 약 27KB의 기계어 코드로 검증
  • 단기 목표: 구문 보완 및 실사용 .cu 파일 호환성 강화
  • 중기 목표: 명령 스케줄링, 레지스터 할당 개선, 상수 폴딩, 루프 불변 코드 이동 등 최적화
  • 장기 목표: Tenstorrent, Intel Arc, RISC-V Vector Extension 등 새로운 백엔드 추가

라이선스

  • Apache 2.0 라이선스

함께 보면 좋은 글 β

GN⁺   [-]
Hacker News 의견들

  • 프로젝트의 README에 적힌 뉴질랜드식 유머가 인상적이었음
    LLVM 의존 없이 자체 인스트럭션 인코딩을 구현한다는 점이 신선함
    이런 프로젝트를 시작하려면 엄청난 저수준 지식이 필요함을 보여줌
    AMD 진영에서 CUDA 미지원이 NVIDIA 독점의 핑계가 되곤 하는데, 이런 시도가 시장 균형에 도움이 될 것 같음

    • 원문에서 말한 건 LLM이 아니라 LLVM임을 지적함
    • 이 프로젝트에도 분명 AI 관련 기여가 포함되어 있다고 생각함. 하지만 잘만 쓰면 AI는 훌륭한 도구가 될 수 있음
    • “Oceania식 유머” 얘기에 Beached Whale 애니메이션이 떠올랐음
    • 소스코드의 주석 “/* 80 keywords walk into a sorted bar */”가 재치 있었음 (lexer.c 링크)
    • “AI slope”라는 표현을 보고 과학적으로는 gradient descent라고 부른다는 농담을 덧붙임

  • 첫 번째 외부 이슈가 geohot이 올린 것이라 놀라웠음 (이슈 링크)
    이런 인물들이 힘을 합쳐 NVIDIA의 GPU 시장 독점을 깨는 걸 보고 싶음

  • AI 추론 워크로드를 NVIDIA GPU에서 돌리는 건 비용 부담이 큼
    이런 프로젝트는 스타트업들이 감당 가능한 대안을 만드는 데 중요함
    conv2d나 attention 같은 연산에서의 성능이 궁금함

  • “# It’s C99. It builds with gcc. There are no dependencies.”
    make 한 줄로 끝나는 단순함이 정말 아름다움

    • 이런 단순하고 직관적인 접근이 너무 마음에 듦

  • 게시물 제목의 대문자 표기를 보고, 회사 GPU 팜 이름이 “barracuda”인 이유를 이제야 깨달았음. 꽤 웃겼음

  • 만약 열정적인 개발자들이 AMD가 못한 걸 해낸다면, 웃기면서도 슬픈 일일 것 같음

    • AMD가 CUDA를 지원하지 않은 건 “못해서”가 아니라 전략적 선택이었다고 생각함
      CUDA를 지원하면 오히려 NVIDIA 생태계를 강화하는 꼴이 되기 때문임
      CUDA 대체를 원한다면 ZLUDA가 더 실용적일 수 있음
    • 오픈소스는 주주 눈치를 볼 필요가 없어서 종종 상용보다 뛰어난 결과를 냄
      하지만 커지면 결국 대기업에 인수되어 사라지는 게 아쉬움
      Linus와 git의 사례처럼, 의지와 지식만 있으면 벽을 깰 수 있음
    • AMD가 “못한” 게 아니라 “안 한” 경우가 많음
      예를 들어 FSR4를 구형 카드에서 공식 지원하지 않는 것도 그런 사례임
    • “우리에겐 HIP이 있다”는 농담으로 마무리함

  • 구형 AMD 아키텍처(GFX1010 등)까지 지원할 수 있을지 궁금함

    • 어렵지만 충분히 가능하다고 생각함. 예전엔 ROCm이 엉망이라 직접 패치했지만 지금은 훨씬 나아졌음
    • 나도 RX5700m을 써서 구형 AMD 지원 작업을 직접 하고 있음
      ISA 문서를 읽으며 바이너리 인코딩을 조정 중임
      다만 이 영역은 LLM이 잘 못하는 분야라, 직접 이해하고 수정해야 함

  • CUDA가 C++을 지원하는데, Clang/LLVM 없이 순수 C로 가는 게 제한적이지 않을까 궁금했음

    • 실제로는 CUDA가 사용하는 C++ 기능의 일부만 파싱 중임. 컴파일러 자체는 C99로 작성됨
    • 테스트를 보면 템플릿 같은 일부 C++ 기능도 처리 가능해 보임
    • LLVM의 AMD GX11 지원은 매우 한정적이라 의존하지 않는 게 나음
    • 진짜 개발자는 AI에 의존하지 않고 코드를 짬
      요즘 오픈소스에 AI가 만든 PR이 넘쳐나는데, 이 프로젝트는 그런 의존을 피한 점이 인상적임

  • AMD가 이런 프로젝트를 공식 후원해야 한다고 생각함
    세상은 NVIDIA의 독점에서 벗어날 필요가 있음

  • “OpenCL은 이제 끝난 건가?”라는 의문이 들었음
    잘 모르겠지만, 그래도 이 프로젝트는 인상적임

    • CUDA가 훨씬 더 시장 점유율을 가져갔고, OpenCL은 기반이 약함
      예전엔 Apple도 지원했지만 지금은 아닌 듯함
      마치 Unix와 Windows의 관계처럼, 기술적으로는 좋았지만 시장은 한쪽으로 쏠렸음
답변달기