Skip to content

BarraCUDA

Open-source CUDA compiler targeting AMD GPUs (and more in the future!). Compiles .cu to GFX11 machine code.

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

About

  • 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 공유 배열, 텍스처·서피스, 동적 병렬 실행 등은 아직 미구현
  • 다중 번역 단위 및 호스트 코드 생성은 지원하지 않음

See also

  • AMD
  • Radeon (Graphic Card)
  • CUDA
  • SCALE - AMD GPU에서 수정 없이 CUDA 실행
  • CubeCL - CUDA, ROCm, WGPU를 위한 Rust 기반 GPU 커널
  • BarraCUDA - AMD GPU용 오픈소스 CUDA 컴파일러
  • ROCm

Favorite site