BarraCUDA: LLVM 없이 맨땅에서 깎아 만든 광기의 AMD용 CUDA 컴파일러

엔지니어링의 세계에서 “바퀴를 다시 발명하지 마라(Don’t reinvent the wheel)“는 격언은 거의 십계명과도 같습니다. 특히 컴파일러처럼 복잡한 영역에서는 더욱 그렇습니다. 보통 새로운 언어나 타겟을 지원하려면 LLVM이라는 거인의 어깨 위에 올라타는 것이 정석입니다. 그런데 최근 Hacker News를 뜨겁게 달군 프로젝트 하나가 이 불문율을 정면으로 거스르고 나타났습니다.

바로 BarraCUDA 입니다.

이 프로젝트는 이름부터 도발적입니다. Nvidia의 CUDA를 AMD GPU에서 돌리겠다는 야심도 크지만, 더 놀라운 건 구현 방식입니다. LLVM 의존성 ‘Zero’. 15,000줄의 순수 C99 코드. make 한 번으로 빌드 완료. 이건 단순히 기술적인 도전이 아니라, 거의 ‘차력쇼’에 가까운 엔지니어링 퍼포먼스입니다.

오늘은 이 미친(?) 프로젝트가 어떻게 Nvidia의 Walled Garden에 균열을 내고 있는지, 그리고 시니어 엔지니어 관점에서 이 코드가 왜 경이로운지 뜯어보겠습니다.

1. LLVM 없이 컴파일러를 만든다는 것의 의미

보통 우리가 GPU 컴파일러를 만든다면 흐름은 뻔합니다. Clang 기반의 프론트엔드를 쓰고, LLVM IR로 변환한 뒤, AMDGPU 백엔드를 태우는 식이죠. 하지만 BarraCUDA의 개발자는 이 과정을 거부했습니다. 이유는 단순합니다. “얼마나 어렵겠어? (How hard can it be?)”

결과적으로 꽤 어려웠다고 고백하지만, 그는 결국 해냈습니다. 이 컴파일러는 .cu 파일을 읽어서 토큰화(Lexing), 파싱(Parsing), 의미 분석(Semantic Analysis)을 거쳐 BIR (BarraCUDA IR) 이라는 독자적인 중간 표현으로 변환합니다. 그 후 직접 레지스터 할당(Register Allocation)을 수행하고, AMD RDNA 3 (GFX11) 기계어 코드(ELF .hsaco)를 뱉어냅니다.

이게 얼마나 대단한 일이냐면, 개발자가 AMD의 ISA(Instruction Set Architecture) 매뉴얼 500페이지를 정독하고, 비트 단위의 인코딩을 직접 C 코드로 구현했다는 뜻입니다. amdgpu_emit.c 파일 1,735줄은 그 고통과 인내의 결정체입니다.

2. 파이프라인 딥다이브: 교과서가 울고 갈 정석 구현

BarraCUDA의 파이프라인은 컴파일러 이론의 정석을 따르지만, 그 구현체는 극도로 가볍고 실용적입니다.

  • C99 & No Dependencies: malloc을 핫 패스(hot path)에서 쓰지 않습니다. 고정 크기 배열을 미리 할당해두고 씁니다. 재귀 호출도 없고, 모든 루프는 유한합니다. 마치 NASA나 JPL의 임베디드 코딩 스탠다드를 보는 듯한 느낌입니다.
  • BIR (BarraCUDA IR): 자체적인 SSA(Static Single Assignment) 형태의 IR을 가집니다. 여기서 mem2reg 패스를 통해 스택 할당을 레지스터로 승격시키는 최적화까지 수행합니다.
  • Instruction Selection: 약 1,700줄의 핸드라이팅 코드로 BIR을 AMD 기계어 명령어로 매핑합니다.

제가 특히 인상 깊게 본 부분은 검증 방식 입니다. LLVM을 컴파일에 쓰진 않지만, 자신이 생성한 바이너리가 맞는지 확인하기 위해 llvm-objdump를 사용하여 “숙제 검사”를 맡겼다는 점입니다. 디코딩 실패가 0건이라니, 얼마나 집요하게 디버깅했는지 알 수 있는 대목입니다.

3. 실제로 무엇이 돌아가는가?

“그래서 이거 Hello World만 되는 거 아니야?”라고 생각할 수 있지만, 지원 범위가 꽤 넓습니다.

  • CUDA 핵심 기능: threadIdx, blockIdx, __shared__ 메모리, __syncthreads() 등.
  • 제어 흐름: if/else, for, while, goto, switch 등 C의 제어 구문을 대부분 지원합니다.
  • 고급 기능: 워프(Warp) 인트린식(__shfl_sync), 아토믹 연산(atomicAdd), 심지어 템플릿의 기본적인 인스턴스화까지 지원합니다.

물론 한계는 명확합니다. C++의 복잡한 기능(클래스, 상속 등)이나 unsigned 단독 사용 같은 사소한 문법들이 아직 지원되지 않습니다. 하지만 vector_add.cu 같은 커널은 문제없이 GFX11 바이너리로 떨어집니다.

__global__ void vector_add(float *c, float *a, float *b, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n)
        c[idx] = a[idx] + b[idx];
}

이 코드를 ./barracuda --amdgpu-bin으로 돌리면 528바이트짜리 .hsaco 파일이 튀어나옵니다. 링커도, 복잡한 빌드 시스템도 필요 없습니다.

4. Hacker News의 반응: LLVM vs LLM

재미있는 점은 HN 댓글 창에서 벌어진 해프닝입니다. 프로젝트 작성자가 “No LLVM dependency”를 강조했는데, 일부 유저들이 이를 “No LLM(Large Language Model) dependency”로 오독하고 “AI 시대에 AI를 안 쓰다니 신선하다!”라며 칭찬을 쏟아낸 것이죠.

작성자는 솔직하게 등판해서 “아니요, 저 테스트 요약이나 보일러플레이트 코드 짤 때 Ollama랑 Claude 썼는데요?”라고 이실직고했습니다. 하지만 이 오해와 별개로, 커뮤니티는 이 프로젝트의 “야생성” 에 열광했습니다.

“A will to live (optional but recommended)”

README에 적힌 요구사항 중 하나입니다. 이 유머러스함과 동시에, AMD가 수십억 달러를 들여서라도 했어야 할 일을 개인 개발자가 취미로 해냈다는 점에 많은 엔지니어들이 씁쓸함과 경의를 동시에 표하고 있습니다. ZLUDA가 실용적인 접근이라면, BarraCUDA는 로맨틱한 접근이랄까요.

5. 총평: 프로덕션 레벨인가? 아니, 그 이상이다.

냉정하게 말해서, 당장 내일 회사 프로덕션 코드에 BarraCUDA를 도입할 수는 없습니다. RDNA 3(GFX11)만 지원하는 점, C++ 표준 라이브러리 부재, 최적화 패스의 부족(Instruction Scheduling 등) 때문에 성능이나 범용성 면에서 갈 길이 멉니다. 실무에서 CUDA 코드를 AMD에서 돌려야 한다면 여전히 HIP이나 ZLUDA가 맞습니다.

하지만 BarraCUDA의 가치는 ‘실용성’에 있지 않습니다.

이 프로젝트는 “소프트웨어의 복잡도에 대한 저항” 입니다. 수십 기가바이트짜리 SDK와 복잡한 툴체인 없이도, 원리만 이해하고 있다면 맨손으로 GPU를 제어할 수 있다는 것을 증명했습니다. 주니어 엔지니어들이 “이건 라이브러리가 없어서 못 해요”라고 할 때, 시니어로서 보여주고 싶은 “근성”의 결정체입니다.

Nvidia의 독점이 견고한 이유는 하드웨어가 아니라 CUDA라는 소프트웨어 생태계 때문입니다. BarraCUDA 같은 시도들이 쌓여서 “CUDA 코드는 꼭 Nvidia GPU가 아니어도 된다”는 인식이 퍼질 때, 비로소 그 철옹성에 진짜 균열이 생길 것입니다.

만약 여러분이 컴파일러 백엔드나 GPU 아키텍처에 관심이 있다면, 이 15,000줄의 코드는 그 어떤 전공 서적보다 훌륭한 교재가 될 것입니다. 일독을 권합니다.