BarraCUDA - AMD GPU를 대상으로 하는 오픈소스 CUDA 컴파일러

• 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 등 새로운 백엔드 추가

기술적 세부 구조

• 주요 소스 구성
  • lexer.c (747줄): 토큰화
  • parser.c (1,500줄): 재귀 하강 파서
  • sema.c (1,725줄): 타입 및 스코프 해석
  • bir.c + bir_lower.c (3,032줄): SSA 기반 중간 표현
  • amdgpu_isel.c (1,788줄): 명령 선택
  • amdgpu_emit.c (1,735줄): 레지스터 할당 및 ELF 생성
• 동적 메모리 할당 없음, 재귀 호출 없음, 고정 크기 배열 기반 구조로 안정성 확보

라이선스 및 기타

• Apache 2.0 라이선스로 자유로운 사용 및 수정 가능
• 프로젝트는 611개 스타, 18개 포크를 기록
• 개발자는 Steven Muchnick의 컴파일러 설계서와 Low Level C 강좌를 주요 참고 자료로 명시