- Eigenständiger Open-Source-Compiler, der CUDA-C-Quellcode (.cu) direkt in Maschinencode für AMD RDNA3 (GFX11) umwandelt
- Erzeugt ELF-Binärdateien im Format
.hsaco über einen eigenen Lexer, Parser und eine Zwischenrepräsentation (BIR) – ohne LLVM- oder HIP-Schicht
- Geschrieben in rund 15.000 Zeilen C99-Code und mit einem einzigen
make-Befehl buildbar
- Unterstützt wichtige CUDA-Funktionen wie Thread-Built-ins, Shared Memory, atomare Operationen, Warp-Operationen und Cooperative Groups
- Wird unter der Apache-2.0-Lizenz veröffentlicht und zielt künftig auf Erweiterungen für zusätzliche Architekturen wie Tenstorrent, Intel Arc und RISC-V ab
Überblick über BarraCUDA
- BarraCUDA ist ein CUDA-Compiler für AMD-GPUs, der
.cu-Dateien in GFX11-Maschinencode umwandelt
- Das Ergebnis liegt als auf AMD-GPUs ausführbare ELF-Binärdatei im Format
.hsaco vor
- Arbeitet vollständig eigenständig ohne LLVM-Abhängigkeit
- Der gesamte Code umfasst etwa 15.000 in C99 geschriebene Zeilen und lässt sich mit einem einzigen Makefile bauen
- Das Projekt wird privat von einem in Neuseeland ansässigen Entwickler entwickelt
Funktionsweise
- Eine eingegebene
.cu-Datei wird in der Reihenfolge Präprozessor → Lexing → Parsing → semantische Analyse → BIR-Erzeugung → Instruktionsauswahl → Registerallokation → Binärcodierung → ELF-Ausgabe verarbeitet
- BIR (BarraCUDA IR) ist eine interne Repräsentation in SSA-Form und architekturunabhängig ausgelegt
- Alle Encodings wurden mit
llvm-objdump verifiziert; es gab 0 Dekodierfehler
Unterstützte Funktionen
- Zentrale CUDA-Syntax:
__global__, __device__, __host__, threadIdx, blockIdx usw.
- CUDA-Funktionen:
__shared__-Memory, __syncthreads(), atomare Operationen, Warp-Shuffle/-Voting, Vektortypen, Half-Precision, Cooperative Groups usw.
- Compiler-Funktionen: vollständiger C-Präprozessor, Fehlerbehebung, Nachverfolgung von Quellpositionen, Unterstützung für die Übergabe von Strukturwerten
Noch nicht unterstützt
- Alleinstehendes
unsigned, zusammengesetzte Zuweisungsoperatoren (+=, -= usw.), const, __constant__-Memory, 2D-Shared-Arrays, Texturen und Surfaces sowie dynamische parallele Ausführung sind noch nicht implementiert
- Mehrere Translation Units und die Erzeugung von Host-Code werden nicht unterstützt
Tests und Roadmap
- Verifiziert mit 14 Testdateien, mehr als 35 Kernels und rund 27 KB Maschinencode
- Kurzfristiges Ziel: Syntax vervollständigen und die Kompatibilität mit realen
.cu-Dateien verbessern
- Mittelfristiges Ziel: Optimierungen wie Instruktionsplanung, bessere Registerallokation, Constant Folding und Loop-Invariant Code Motion
- Langfristiges Ziel: neue Backends für Tenstorrent, Intel Arc und RISC-V Vector Extension hinzufügen
Lizenz
1 Kommentare
Hacker-News-Kommentare
Der neuseeländische Humor im README des Projekts war bemerkenswert.
Es wirkte erfrischend, dass die Instruktionskodierung ohne LLVM-Abhängigkeit selbst implementiert wird.
Das zeigt, wie viel Low-Level-Wissen man braucht, um so ein Projekt zu starten.
Im AMD-Lager wird die fehlende CUDA-Unterstützung oft als Ausrede für NVIDIAs Monopol benutzt, und so ein Versuch könnte helfen, das Marktgleichgewicht wiederherzustellen.
Ich war überrascht, dass das erste externe Issue von geohot stammt (Link zum Issue).
Ich würde gern sehen, wie sich solche Leute zusammentun und NVIDIAs Monopol auf dem GPU-Markt aufbrechen.
KI-Inferenz-Workloads auf NVIDIA-GPUs laufen zu lassen, ist kostspielig.
Solche Projekte sind wichtig, damit Startups bezahlbare Alternativen schaffen können.
Ich bin neugierig auf die Performance bei Operationen wie conv2d oder attention.
„# It’s C99. It builds with gcc. There are no dependencies.”
Diese Einfachheit, bei der ein einziges
makegenügt, ist wirklich schön.Als ich die Großschreibung im Titel des Beitrags sah, habe ich erst jetzt verstanden, warum die GPU-Farm der Firma „barracuda“ heißt. Das war ziemlich lustig.
Wenn leidenschaftliche Entwickler das schaffen, was AMD nicht geschafft hat, wäre das komisch und traurig zugleich.
CUDA zu unterstützen würde am Ende nur das NVIDIA-Ökosystem weiter stärken.
Wenn man eine CUDA-Alternative will, könnte ZLUDA praktischer sein.
Schade ist nur, dass erfolgreiche Projekte am Ende oft von Großunternehmen übernommen werden und verschwinden.
Wie das Beispiel von Linus und git zeigt, kann man mit Willen und Wissen Mauern einreißen.
Dass FSR4 auf älteren Karten nicht offiziell unterstützt wird, ist zum Beispiel so ein Fall.
Ich frage mich, ob auch ältere AMD-Architekturen (GFX1010 usw.) unterstützt werden könnten.
Ich lese die ISA-Dokumentation und passe die Binärkodierung an.
Allerdings ist das ein Bereich, in dem LLMs nicht besonders gut sind, daher muss man ihn selbst verstehen und anpassen.
CUDA unterstützt C++, daher fragte ich mich, ob ein reiner-C-Ansatz ohne Clang/LLVM nicht zu eingeschränkt ist.
Heutzutage ist Open Source voller von KI erzeugter PRs, daher war es beeindruckend, dass dieses Projekt eine solche Abhängigkeit vermeidet.
Ich finde, AMD sollte so ein Projekt offiziell fördern.
Die Welt muss sich aus NVIDIAs Monopol lösen.
Ich fragte mich: „Ist OpenCL jetzt vorbei?“
Ich weiß es nicht genau, aber dieses Projekt ist trotzdem beeindruckend.
Früher wurde es auch von Apple unterstützt, heute offenbar nicht mehr.
Es ist ein bisschen wie das Verhältnis zwischen Unix und Windows: technisch gut, aber der Markt hat sich auf eine Seite verlagert.