Skip to content
Search

BarraCUDA

lastmod 2026-02-19created 2026-02-191 edits

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

Open-source CUDA compiler for AMD GPUs

About

  • Standalone open-source compiler that translates CUDA C source (.cu) directly into machine code for AMD RDNA3 (GFX11)
  • Generates ELF .hsaco binaries through its own lexer, parser, and intermediate representation (BIR), without LLVM or HIP layers
  • Written in about 15,000 lines of C99 code, and can be built with a single make command
  • Supports major CUDA features such as thread built-in variables, shared memory, atomic operations, warp operations, and cooperative groups
  • Released under the Apache 2.0 license, with the goal of extending to additional architectures such as Tenstorrent, Intel Arc, and RISC-V in the future

BarraCUDA overview

  • BarraCUDA is a CUDA compiler for AMD GPUs that converts .cu files into GFX11 machine code
    • The output is an ELF .hsaco binary that can run on AMD GPUs
    • Works fully standalone without an LLVM dependency
  • The entire codebase is about 15,000 lines written in C99, and can be built with a single Makefile
  • The project is developed personally by a New Zealand-based developer

How it works

  • Processes the input .cu file in the order preprocessing → lexical analysis → parsing → semantic analysis → BIR generation → instruction selection → register allocation → binary encoding → ELF output
  • BIR (BarraCUDA IR) is an internal representation in SSA form, with an architecture-independent design
  • All encodings are verified with llvm-objdump, with 0 decode errors

Supported features

  • Core CUDA syntax: global, device, host, threadIdx, blockIdx, etc.
  • CUDA features: shared memory, __syncthreads(), atomic operations, warp shuffle/vote, vector types, half precision, cooperative groups, etc.
  • Compiler features: full C preprocessor, error recovery, source location tracking, passing structs by value

Unsupported items

  • Standalone unsigned, compound assignment operators (+=, -=, etc.), const, constant memory, 2D shared arrays, textures and surfaces, dynamic parallelism, etc. are not yet implemented
  • Multiple translation units and host code generation are not supported

See also

  • AMD
  • Radeon (Graphic Card)
  • CUDA
  • SCALE - Run CUDA on AMD GPUs without modification
  • CubeCL - Rust-based GPU kernels for CUDA, ROCm, and WGPU
  • BarraCUDA - Open-source CUDA compiler for AMD GPUs
  • ROCm

Favorite site

author Zer0 sitemap