BarraCUDA - compilador CUDA open source para GPUs AMD

• Compilador open source independente que converte diretamente código-fonte CUDA C (.cu) em código de máquina para AMD RDNA3 (GFX11)
• Gera binários ELF .hsaco por meio de seu próprio analisador léxico, parser e representação intermediária (BIR), sem LLVM nem a camada HIP
• Escrito em cerca de 15 mil linhas de código C99 e pode ser compilado com um único comando make
• Suporta recursos principais do CUDA, como variáveis internas de thread, memória compartilhada, operações atômicas, operações de warp e grupos cooperativos
• Disponibilizado sob a licença Apache 2.0 e, no futuro, pretende expandir para arquiteturas adicionais como Tenstorrent, Intel Arc e RISC-V

---

Visão geral do BarraCUDA

• BarraCUDA é um compilador CUDA para GPUs AMD que converte arquivos .cu em código de máquina GFX11
  • O resultado é gerado no formato de binários ELF .hsaco executáveis em GPUs AMD
  • Funciona de forma totalmente independente, sem dependência de LLVM
• Todo o código do projeto tem aproximadamente 15 mil linhas escritas em C99 e pode ser compilado com um único Makefile
• O projeto é desenvolvido de forma independente por um desenvolvedor baseado na Nova Zelândia

Como funciona

• O arquivo .cu de entrada é processado na sequência pré-processamento → análise léxica → parsing → análise semântica → geração de BIR → seleção de instruções → alocação de registradores → codificação binária → saída ELF
• BIR (BarraCUDA IR) é uma representação interna em formato SSA, projetada para ser independente de arquitetura
• Toda a codificação foi validada com llvm-objdump, com 0 erros de decodificação

Recursos suportados

• Sintaxe principal do CUDA: __global__, __device__, __host__, threadIdx, blockIdx etc.
• Recursos do CUDA: memória __shared__, __syncthreads(), operações atômicas, shuffle/votação de warp, tipos vetoriais, precisão half, grupos cooperativos etc.
• Recursos do compilador: pré-processador C completo, recuperação de erros, rastreamento de localização no código-fonte e suporte a passagem de valores de struct

Itens ainda não suportados

• Uso isolado de unsigned, operadores de atribuição composta (+=, -= etc.), const, memória __constant__, arrays compartilhados 2D, texturas e surfaces, paralelismo dinâmico etc. ainda não foram implementados
• Não há suporte a múltiplas unidades de tradução nem à geração de código host

Testes e roadmap

• Validado com 14 arquivos de teste, mais de 35 kernels e cerca de 27 KB de código de máquina
• Meta de curto prazo: complementar a sintaxe e ampliar a compatibilidade com arquivos .cu de uso real
• Meta de médio prazo: otimizações como escalonamento de instruções, melhoria na alocação de registradores, constant folding e loop-invariant code motion
• Meta de longo prazo: adicionar novos backends como Tenstorrent, Intel Arc e RISC-V Vector Extension

Licença

• Licença Apache 2.0

[GN⁺]

Comentários do Hacker News

• O humor neozelandês escrito no README do projeto chamou atenção
  Foi interessante ver que ele implementa a própria codificação de instruções sem depender do LLVM
  Isso mostra como é preciso ter um enorme conhecimento de baixo nível para começar um projeto desses
  No lado da AMD, a falta de suporte a CUDA costuma servir de desculpa para o domínio da NVIDIA, então esse tipo de tentativa parece poder ajudar no equilíbrio do mercado

  • Apontaram que o texto original falava de LLVM, não de LLM
  • Acho que este projeto também claramente inclui contribuições relacionadas a IA. Ainda assim, se bem usada, IA pode ser uma ótima ferramenta
  • Ao ler “humor da Oceania”, lembrei da animação Beached Whale
  • O comentário no código-fonte /* 80 keywords walk into a sorted bar */ foi espirituoso (link para lexer.c)
  • Ao ver a expressão “AI slope”, acrescentaram a piada de que, cientificamente, isso se chama gradient descent

• Fiquei surpreso que a primeira issue externa tenha sido aberta pelo geohot (link da issue)
  Queria ver gente assim unindo forças para quebrar o monopólio da NVIDIA no mercado de GPUs

• Rodar workloads de inferência de IA em GPUs da NVIDIA tem um custo alto
  Projetos assim são importantes para que startups construam alternativas viáveis financeiramente
  Fiquei curioso sobre o desempenho em operações como conv2d e attention

• “# It’s C99. It builds with gcc. There are no dependencies.”
  A simplicidade de resolver tudo com uma única linha de make é realmente bonita

  • Gosto demais dessa abordagem simples e direta

• Ao ver o uso de maiúsculas no título da postagem, finalmente percebi por que o nome do GPU farm da empresa é “barracuda”. Achei bem engraçado

• Se desenvolvedores apaixonados fizerem o que a AMD não conseguiu, seria engraçado e triste ao mesmo tempo

  • Acho que a AMD não deixou de dar suporte a CUDA por “não conseguir”, e sim por uma escolha estratégica
    Dar suporte a CUDA acabaria fortalecendo ainda mais o ecossistema da NVIDIA
    Se a ideia é ter uma alternativa ao CUDA, ZLUDA talvez seja mais prático
  • O open source não precisa se preocupar com acionistas, então muitas vezes entrega resultados melhores que soluções comerciais
    Mas é triste quando cresce e acaba sendo comprado por uma grande empresa e desaparece
    Como no caso do Linus e do git, com vontade e conhecimento dá para romper barreiras
  • Há muitos casos em que a AMD não “não conseguiu”, mas simplesmente “não fez”
    Por exemplo, o fato de o FSR4 não ter suporte oficial em placas antigas também entra nisso
  • Terminaram com a piada: “temos o HIP”

• Fiquei pensando se daria para suportar arquiteturas AMD mais antigas (como GFX1010)

  • Acho difícil, mas totalmente possível. Antes o ROCm era uma bagunça e eu mesmo precisei aplicar patches, mas hoje está bem melhor
  • Eu também uso uma RX5700m, então estou trabalhando pessoalmente no suporte a AMDs antigas
    Estou lendo a documentação da ISA e ajustando a codificação binária
    Só que essa é uma área em que LLMs não costumam ir bem, então é preciso entender e corrigir por conta própria

• Como o CUDA suporta C++, fiquei me perguntando se seguir com C puro sem Clang/LLVM não seria limitante

  • Na prática, no momento ele só está fazendo parse de parte dos recursos de C++ usados pelo CUDA. O compilador em si é escrito em C99
  • Pelos testes, parece que ele consegue lidar até com alguns recursos de C++, como templates
  • O suporte do LLVM a AMD GX11 é bem limitado, então é melhor não depender dele
  • Desenvolvedor de verdade escreve código sem depender de IA
    Hoje o open source está cheio de PRs feitos por IA, e foi marcante ver que este projeto evita esse tipo de dependência

• Acho que a AMD deveria patrocinar oficialmente um projeto assim
  O mundo precisa sair do monopólio da NVIDIA

• Fiquei com a dúvida: “o OpenCL já era?”
  Não sei ao certo, mas mesmo assim o projeto é impressionante

  • O CUDA acabou ficando com uma fatia de mercado muito maior, e o OpenCL tem uma base fraca
    Antigamente até a Apple dava suporte, mas hoje aparentemente não mais
    É como a relação entre Unix e Windows: tecnicamente era bom, mas o mercado acabou pendendo para um lado