BarraCUDA - compilador CUDA de código abierto para GPU AMD

• Compilador independiente de código abierto que convierte directamente código fuente CUDA C (.cu) a código máquina para AMD RDNA3 (GFX11)
• Genera binarios ELF .hsaco mediante su propio analizador léxico, parser y representación intermedia (BIR), sin capas LLVM ni HIP
• Está escrito en más de 15,000 líneas de código C99 y puede compilarse con un solo comando make
• Soporta funciones clave de CUDA como variables integradas de hilo, memoria compartida, operaciones atómicas, operaciones warp y grupos cooperativos
• Se publica bajo licencia Apache 2.0 y apunta a expandirse en el futuro a arquitecturas adicionales como Tenstorrent, Intel Arc y RISC-V

---

Resumen de BarraCUDA

• BarraCUDA es un compilador CUDA para GPU AMD que convierte archivos .cu en código máquina GFX11
  • El resultado tiene formato de binario ELF .hsaco ejecutable en GPU AMD
  • Funciona de forma totalmente independiente sin dependencia de LLVM
• Todo el código suma unas 15,000 líneas escritas en C99 y puede compilarse con un único Makefile
• El proyecto fue desarrollado de manera personal por un desarrollador radicado en Nueva Zelanda

Cómo funciona

• El archivo .cu de entrada se procesa en este orden: preprocesamiento → análisis léxico → parsing → análisis semántico → generación de BIR → selección de instrucciones → asignación de registros → codificación binaria → salida ELF
• BIR (BarraCUDA IR) es una representación interna en forma SSA con un diseño independiente de la arquitectura
• Toda la codificación fue verificada con llvm-objdump, con 0 errores de decodificación

Funciones compatibles

• Sintaxis principal de CUDA: __global__, __device__, __host__, threadIdx, blockIdx, etc.
• Funciones de CUDA: memoria __shared__, __syncthreads(), operaciones atómicas, shuffle/votación warp, tipos vectoriales, precisión half, grupos cooperativos, etc.
• Funciones del compilador: preprocesador de C completo, recuperación de errores, seguimiento de ubicación en el código fuente y soporte para paso de valores de estructuras

Elementos aún no compatibles

• Aún no implementa el uso de unsigned por sí solo, operadores de asignación compuesta (+=, -=, etc.), const, memoria __constant__, arreglos compartidos 2D, texturas y superficies, ni paralelismo dinámico
• No soporta múltiples unidades de traducción ni generación de código host

Pruebas y hoja de ruta

• Validado con 14 archivos de prueba, más de 35 kernels y alrededor de 27 KB de código máquina
• Objetivo a corto plazo: completar la sintaxis y mejorar la compatibilidad con archivos .cu de uso real
• Objetivo a mediano plazo: optimizaciones como scheduling de instrucciones, mejor asignación de registros, constant folding y movimiento de código invariante de bucle
• Objetivo a largo plazo: añadir nuevos backends como Tenstorrent, Intel Arc y RISC-V Vector Extension

Licencia

• Licencia Apache 2.0

[GN⁺]

Opiniones de Hacker News

• Le llamó la atención el humor neozelandés escrito en el README del proyecto
  Le parece novedoso que implemente su propia codificación de instrucciones sin depender de LLVM
  Esto demuestra que para iniciar un proyecto así se necesita un enorme conocimiento de bajo nivel
  En el lado de AMD, la falta de soporte para CUDA suele usarse como excusa para el dominio de NVIDIA, y este tipo de intentos podría ayudar al equilibrio del mercado

  • Señala que en el texto original no hablaban de un LLM, sino de LLVM
  • Cree que este proyecto también incluye claramente aportes relacionados con IA. Aun así, si se usa bien, la IA puede ser una gran herramienta
  • La mención del “humor de Oceanía” le recordó la animación Beached Whale
  • Le pareció ingenioso el comentario en el código fuente: “/* 80 keywords walk into a sorted bar */” (enlace a lexer.c)
  • Al ver la expresión “AI slope”, añadió la broma de que científicamente eso se llama gradient descent

• Le sorprendió que el primer issue externo lo publicara geohot (enlace al issue)
  Le gustaría ver a figuras así uniendo fuerzas para romper el monopolio del mercado de GPU de NVIDIA

• Ejecutar cargas de trabajo de inferencia de IA en GPU de NVIDIA tiene un costo elevado
  Proyectos como este son importantes para que las startups construyan alternativas accesibles
  Le da curiosidad el rendimiento en operaciones como conv2d o attention

• “# It’s C99. It builds with gcc. There are no dependencies.”
  La simplicidad de resolver todo con una sola línea de make es realmente hermosa

  • Le encanta este enfoque tan simple e intuitivo

• Al ver el uso de mayúsculas en el título de la publicación, recién entendió por qué el nombre del GPU farm de su empresa es “barracuda”. Le dio bastante risa

• Si desarrolladores apasionados terminan haciendo lo que AMD no pudo hacer, sería gracioso y triste al mismo tiempo

  • Cree que AMD no dejó de soportar CUDA porque “no pudiera”, sino por una decisión estratégica
    Soportar CUDA terminaría reforzando el ecosistema de NVIDIA
    Si lo que se quiere es una alternativa a CUDA, ZLUDA podría ser más práctico
  • El software de código abierto no necesita quedar bien con los accionistas, así que a menudo logra resultados superiores a los comerciales
    Pero da pena que, cuando crece, termine siendo absorbido por una gran empresa y desaparezca
    Como muestran los casos de Linus y git, con voluntad y conocimiento se pueden derribar muros
  • Hay muchos casos en los que AMD no es que “no pudiera”, sino que “no quiso”
    Por ejemplo, el hecho de que FSR4 no tenga soporte oficial en tarjetas antiguas también sería un caso así
  • Cierra con la broma: “nosotros tenemos HIP”

• Le da curiosidad si podrá dar soporte incluso a arquitecturas AMD antiguas (como GFX1010)

  • Cree que es difícil, pero totalmente posible. Antes ROCm era un desastre y tenía que parchearlo por su cuenta, pero ahora está mucho mejor
  • Él también usa una RX5700m, así que está trabajando directamente en el soporte para AMD antiguo
    Está ajustando la codificación binaria mientras lee documentos de ISA
    Pero esta es un área en la que los LLM no son muy buenos, así que hay que entenderla y corregirla personalmente

• Como CUDA soporta C++, se preguntaba si ir por C puro sin Clang/LLVM no sería limitante

  • En la práctica, solo está parseando una parte de las funciones de C++ que usa CUDA. El compilador en sí está escrito en C99
  • Por las pruebas, parece que también puede manejar algunas funciones de C++ como templates
  • El soporte de LLVM para AMD GX11 es muy limitado, así que es mejor no depender de él
  • Un desarrollador de verdad no depende de la IA para escribir código
    Últimamente el open source está lleno de PR hechas por IA, y en este proyecto impresionó que se evitara esa dependencia

• Cree que AMD debería patrocinar oficialmente este proyecto
  El mundo necesita salir del monopolio de NVIDIA

• Se preguntó: “¿OpenCL ya está acabado?”
  No está seguro, pero aun así el proyecto le parece impresionante

  • CUDA se quedó con una cuota de mercado mucho mayor, y OpenCL tiene una base débil
    Antes Apple también lo soportaba, pero ahora ya no parece ser así
    Como la relación entre Unix y Windows, técnicamente era bueno, pero el mercado se inclinó hacia un solo lado