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

[GN⁺]

Commentaires sur Hacker News

• L’humour néo-zélandais dans le README du projet a marqué les esprits
  Le fait d’implémenter son propre encodage d’instructions sans dépendre de LLVM est rafraîchissant
  Cela montre à quel point il faut une énorme maîtrise du bas niveau pour lancer un projet comme celui-ci
  Dans l’écosystème AMD, l’absence de prise en charge de CUDA sert souvent d’argument pour justifier le monopole de NVIDIA, donc ce genre d’initiative pourrait aider à rééquilibrer le marché

  • Quelqu’un a signalé que le texte original parlait de LLVM, pas de LLM
  • Il y a probablement aussi des contributions liées à l’IA dans ce projet. Mais bien utilisée, l’IA peut être un excellent outil
  • La mention de « l’humour océanien » a rappelé à quelqu’un l’animation Beached Whale
  • Le commentaire dans le code source « /* 80 keywords walk into a sorted bar */ » a été jugé très spirituel (lien vers lexer.c)
  • En voyant l’expression « AI slope », quelqu’un a ajouté la blague que, scientifiquement, on appelle ça une descente de gradient

• Surprise de voir que le premier ticket externe a été ouvert par geohot (lien vers l’issue)
  Certains aimeraient voir ce type de profils unir leurs forces pour briser le monopole de NVIDIA sur le marché des GPU

• Faire tourner des charges d’inférence IA sur des GPU NVIDIA coûte très cher
  Des projets comme celui-ci sont importants pour permettre aux startups de construire des alternatives abordables
  Certains se demandent quelles sont les performances sur des opérations comme conv2d ou l’attention

• « # It’s C99. It builds with gcc. There are no dependencies. »
  Cette simplicité, où un simple make suffit, est vraiment magnifique

  • Cette approche simple et directe plaît énormément

• En voyant les majuscules dans le titre du post, quelqu’un a enfin compris pourquoi la ferme de GPU de son entreprise s’appelle « barracuda ». Ça l’a bien fait rire

• Si des développeurs passionnés réussissaient là où AMD n’a pas réussi, ce serait drôle et triste à la fois

  • Selon un avis, si AMD n’a pas pris en charge CUDA, ce n’est pas parce qu’ils ne pouvaient pas, mais parce que c’était un choix stratégique
    Prendre en charge CUDA reviendrait à renforcer davantage l’écosystème de NVIDIA
    Si l’objectif est un substitut à CUDA, ZLUDA pourrait être plus pratique
  • L’open source n’a pas à se soucier des actionnaires, ce qui lui permet souvent d’obtenir des résultats supérieurs au commercial
    Mais il est regrettable que, lorsqu’un projet grossit, il finisse souvent racheté par un grand groupe puis disparaisse
    Comme le montrent Linus et git, avec de la volonté et des connaissances, on peut faire tomber des barrières
  • Souvent, AMD n’a pas « échoué », ils ont surtout choisi de ne pas le faire
    Le fait de ne pas officiellement prendre en charge FSR4 sur les anciennes cartes en serait un exemple
  • Le tout se conclut par la blague : « Nous, on a HIP »

• Certains se demandent si le support pourrait aller jusqu’aux anciennes architectures AMD (GFX1010, etc.)

  • Ce serait difficile, mais tout à fait possible. Avant, ROCm était dans un sale état et il fallait le patcher soi-même, mais la situation s’est beaucoup améliorée
  • Quelqu’un utilise aussi une RX5700m et travaille lui-même sur la prise en charge des anciens GPU AMD
    Il ajuste l’encodage binaire en lisant la documentation ISA
    Mais c’est un domaine où les LLM ne sont pas très bons, donc il faut comprendre et corriger les choses soi-même

• Comme CUDA prend en charge le C++, certains se demandaient si partir sur du pur C sans Clang/LLVM n’était pas trop limitant

  • En pratique, seul un sous-ensemble des fonctionnalités C++ utilisées par CUDA est actuellement analysé. Le compilateur lui-même est écrit en C99
  • D’après les tests, certaines fonctionnalités C++ comme les templates semblent tout de même prises en charge
  • La prise en charge d’AMD GX11 dans LLVM est très limitée, donc mieux vaut ne pas en dépendre
  • Un vrai développeur écrit du code sans dépendre de l’IA
    En ce moment, l’open source est inondé de PR générées par l’IA, et le fait que ce projet évite ce type de dépendance a été apprécié

• AMD devrait, selon certains, soutenir officiellement ce projet
  Le monde a besoin d’échapper au monopole de NVIDIA

• Certains se sont demandé : « OpenCL, c’est fini maintenant ? »
  Ils n’en sont pas sûrs, mais trouvent quand même ce projet impressionnant

  • CUDA a pris une part de marché bien plus importante, tandis qu’OpenCL a des bases fragiles
    Apple le prenait en charge autrefois, mais cela ne semble plus être le cas aujourd’hui
    Un peu comme la relation entre Unix et Windows : techniquement c’était bon, mais le marché a basculé d’un seul côté