BarraCUDA - compilateur CUDA open source pour GPU AMD

• Compilateur open source autonome qui convertit directement le code source CUDA C (.cu) en code machine pour AMD RDNA3 (GFX11)
• Génère des binaires ELF .hsaco via son propre analyseur lexical, parseur et représentation intermédiaire (BIR), sans LLVM ni couche HIP
• Écrit en environ 15 000 lignes de code C99 et peut être compilé avec une unique commande make
• Prend en charge les principales fonctionnalités de CUDA, notamment les variables intégrées de thread, la mémoire partagée, les opérations atomiques, les opérations de warp et les groupes coopératifs
• Publié sous licence Apache 2.0, avec pour objectif d'étendre à l'avenir la prise en charge à Tenstorrent, Intel Arc, RISC-V et d'autres architectures

---

Présentation de BarraCUDA

• BarraCUDA est un compilateur CUDA pour GPU AMD qui convertit les fichiers .cu en code machine GFX11
  • Le résultat est un binaire ELF .hsaco exécutable sur les GPU AMD
  • Il fonctionne de manière totalement autonome, sans dépendance à LLVM
• L'ensemble du code représente environ 15 000 lignes écrites en C99 et peut être compilé avec un unique Makefile
• Le projet est développé à titre personnel par un développeur basé en Nouvelle-Zélande

Fonctionnement

• Le fichier .cu en entrée est traité selon la chaîne prétraitement → analyse lexicale → parsing → analyse sémantique → génération de BIR → sélection d'instructions → allocation de registres → encodage binaire → sortie ELF
• BIR (BarraCUDA IR) est une représentation interne sous forme SSA, conçue pour être indépendante de l'architecture
• Tous les encodages ont été vérifiés via llvm-objdump, avec zéro erreur de décodage

Fonctionnalités prises en charge

• Syntaxe de base de CUDA : __global__, __device__, __host__, threadIdx, blockIdx, etc.
• Fonctionnalités CUDA : mémoire __shared__, __syncthreads(), opérations atomiques, shuffle/vote de warp, types vectoriels, précision half, groupes coopératifs, etc.
• Fonctionnalités du compilateur : préprocesseur C complet, récupération après erreur, suivi des positions dans le code source, prise en charge du passage de structures par valeur

Éléments non pris en charge

• L'utilisation de unsigned seul, les opérateurs d'affectation composés (+=, -=, etc.), const, la mémoire __constant__, les tableaux partagés 2D, les textures et surfaces, ainsi que l'exécution parallèle dynamique ne sont pas encore implémentés
• Les unités de traduction multiples et la génération de code hôte ne sont pas prises en charge

Tests et feuille de route

• Validation effectuée sur 14 fichiers de test, plus de 35 kernels et environ 27 Ko de code machine
• Objectif à court terme : compléter la syntaxe et renforcer la compatibilité avec des fichiers .cu réels
• Objectif à moyen terme : optimisations comme l'ordonnancement des instructions, l'amélioration de l'allocation de registres, le constant folding et le déplacement de code invariant de boucle
• Objectif à long terme : ajout de nouveaux backends comme Tenstorrent, Intel Arc, RISC-V Vector Extension

Licence

• Licence Apache 2.0