BarraCUDA - Kompiler CUDA open-source untuk GPU AMD

• Kompiler open-source mandiri yang langsung mengubah source CUDA C (.cu) menjadi kode mesin untuk AMD RDNA3 (GFX11)
• Menghasilkan biner ELF .hsaco melalui lexer, parser, dan representasi menengah (BIR) miliknya sendiri tanpa lapisan LLVM atau HIP
• Ditulis dalam sekitar 15.000 baris kode C99 dan dapat dibangun dengan satu perintah make
• Mendukung fitur utama CUDA seperti variabel bawaan thread, memori bersama, operasi atomik, operasi warp, dan cooperative groups
• Dirilis dengan lisensi Apache 2.0, dengan target perluasan arsitektur tambahan di masa depan seperti Tenstorrent, Intel Arc, dan RISC-V

---

Ikhtisar BarraCUDA

• BarraCUDA adalah kompiler CUDA untuk GPU AMD yang mengubah file .cu menjadi kode mesin GFX11
  • Hasilnya berupa biner ELF .hsaco yang dapat dijalankan di GPU AMD
  • Bekerja sepenuhnya secara mandiri tanpa dependensi LLVM
• Seluruh kode proyek terdiri dari sekitar 15.000 baris yang ditulis dalam C99, dan dapat dibangun dengan satu Makefile
• Proyek ini dikembangkan secara pribadi oleh pengembang yang berbasis di Selandia Baru

Cara kerjanya

• File .cu yang dimasukkan diproses dalam urutan preprocessing → analisis leksikal → parsing → analisis semantik → pembuatan BIR → pemilihan instruksi → alokasi register → encoding biner → output ELF
• BIR (BarraCUDA IR) adalah representasi internal berbentuk SSA yang dirancang agar independen terhadap arsitektur
• Semua encoding diverifikasi melalui llvm-objdump dan menghasilkan 0 kesalahan decode

Fitur yang didukung

• Sintaks inti CUDA: __global__, __device__, __host__, threadIdx, blockIdx, dll.
• Fitur CUDA: memori __shared__, __syncthreads(), operasi atomik, shuffle/voting warp, tipe vektor, presisi half, cooperative groups, dll.
• Fitur kompiler: preprocessor C lengkap, pemulihan kesalahan, pelacakan lokasi source, dukungan passing nilai struct

Yang belum didukung

• Penggunaan unsigned secara tunggal, operator penugasan gabungan (+=, -=, dll.), const, memori __constant__, array shared 2D, texture/surface, dynamic parallelism, dan lainnya masih belum diimplementasikan
• Tidak mendukung beberapa translation unit maupun pembuatan host code

Pengujian dan roadmap

• Diverifikasi dengan 14 file uji, lebih dari 35 kernel, dan sekitar 27KB kode mesin
• Target jangka pendek: melengkapi sintaks dan memperkuat kompatibilitas dengan file .cu untuk penggunaan nyata
• Target jangka menengah: optimasi seperti instruction scheduling, peningkatan alokasi register, constant folding, dan loop-invariant code motion
• Target jangka panjang: menambahkan backend baru seperti Tenstorrent, Intel Arc, dan RISC-V Vector Extension

Lisensi

• Lisensi Apache 2.0

[GN⁺]

Komentar Hacker News

• Humor ala Selandia Baru yang tertulis di README proyek ini cukup berkesan
  Menarik karena proyek ini mengimplementasikan encoding instruksi sendiri tanpa bergantung pada LLVM
  Ini menunjukkan bahwa memulai proyek seperti ini membutuhkan pengetahuan tingkat rendah yang sangat mendalam
  Di kubu AMD, tidak adanya dukungan CUDA sering dijadikan alasan bagi dominasi NVIDIA, jadi upaya seperti ini tampaknya bisa membantu menyeimbangkan pasar

  • Menunjukkan bahwa yang dimaksud di teks asli bukan LLM melainkan LLVM
  • Tetap merasa proyek ini jelas mencakup kontribusi terkait AI. Namun jika digunakan dengan baik, AI bisa menjadi alat yang sangat bagus
  • Saat melihat pembahasan tentang “humor Oceania”, jadi teringat animasi Beached Whale
  • Komentar di source code “/* 80 keywords walk into a sorted bar */” terasa jenaka (tautan lexer.c)
  • Melihat istilah “AI slope”, lalu menambahkan candaan bahwa secara ilmiah itu disebut gradient descent

• Kaget karena isu eksternal pertama ternyata dibuat oleh geohot (tautan isu)
  Ingin melihat orang-orang seperti ini bekerja sama untuk mematahkan monopoli pasar GPU NVIDIA

• Menjalankan workload inferensi AI di GPU NVIDIA itu mahal
  Proyek seperti ini penting agar startup bisa membangun alternatif yang lebih terjangkau
  Penasaran dengan performanya pada operasi seperti conv2d atau attention

• “# It’s C99. It builds with gcc. There are no dependencies.”
  Kesederhanaan yang selesai hanya dengan satu baris make benar-benar indah

  • Sangat suka pendekatan yang sesederhana dan sejelas ini

• Melihat kapitalisasi di judul posting, baru sadar kenapa nama GPU farm di kantor adalah “barracuda”. Cukup lucu

• Kalau para developer penuh semangat berhasil melakukan hal yang tidak dilakukan AMD, itu akan terasa lucu sekaligus menyedihkan

  • Menurutnya AMD tidak mendukung CUDA bukan karena “tidak bisa”, melainkan karena pilihan strategis
    Mendukung CUDA justru akan memperkuat ekosistem NVIDIA
    Jika menginginkan alternatif CUDA, ZLUDA mungkin lebih praktis
  • Open source tidak perlu khawatir pada pemegang saham, sehingga sering menghasilkan hasil yang lebih baik daripada produk komersial
    Namun sayangnya ketika membesar, sering kali akhirnya diakuisisi perusahaan besar lalu menghilang
    Seperti kasus Linus dan git, jika ada kemauan dan pengetahuan, tembok pun bisa ditembus
  • Sering kali AMD bukan “tidak bisa” melainkan “tidak mau”
    Misalnya, tidak memberikan dukungan resmi FSR4 pada kartu lama juga termasuk contoh seperti itu
  • Ditutup dengan candaan, “kita punya HIP”

• Penasaran apakah proyek ini bisa mendukung arsitektur AMD lama seperti GFX1010

  • Sulit, tetapi menurutnya sangat mungkin. Dulu ROCm berantakan sampai harus ditambal sendiri, tapi sekarang jauh lebih baik
  • Saya juga memakai RX5700m, jadi sedang mengerjakan dukungan AMD lama secara langsung
    Sambil membaca dokumen ISA, sedang menyesuaikan encoding biner
    Namun ini adalah area yang tidak dikuasai LLM dengan baik, jadi harus benar-benar dipahami dan diperbaiki sendiri

• CUDA mendukung C++, jadi sempat bertanya apakah pendekatan pure C tanpa Clang/LLVM tidak akan terlalu terbatas

  • Sebenarnya saat ini hanya sedang mem-parse sebagian fitur C++ yang digunakan CUDA. Kompilernya sendiri ditulis dalam C99
  • Dari pengujiannya, tampaknya beberapa fitur C++ seperti template juga bisa ditangani
  • Dukungan LLVM untuk AMD GX11 sangat terbatas, jadi lebih baik tidak bergantung padanya
  • Developer sejati menulis kode tanpa bergantung pada AI
    Belakangan ini open source dibanjiri PR buatan AI, jadi proyek ini terasa mengesankan karena menghindari ketergantungan seperti itu

• AMD seharusnya mensponsori secara resmi proyek seperti ini
  Dunia perlu keluar dari dominasi NVIDIA

• Muncul pertanyaan, “Apakah OpenCL sekarang sudah berakhir?”
  Tidak terlalu tahu, tetapi proyek ini tetap mengesankan

  • CUDA mengambil pangsa pasar yang jauh lebih besar, sementara OpenCL fondasinya lemah
    Dulu Apple juga mendukungnya, tapi sekarang sepertinya tidak lagi
    Mirip seperti hubungan Unix dan Windows, secara teknis bagus, tetapi pasar condong ke satu sisi