SCALE - exécuter CUDA sur des GPU AMD sans modification

• SCALE est une boîte à outils de programmation GPGPU qui permet de compiler nativement des applications CUDA pour les GPU AMD
• Il n’est pas nécessaire de modifier les programmes CUDA ni le système de build, et la prise en charge d’un plus grand nombre de fournisseurs de GPU ainsi que des API CUDA est en cours de développement

Comment cela fonctionne-t-il ?

• SCALE apporte plusieurs innovations majeures par rapport aux autres solutions GPGPU multiplateformes
  • Il accepte les programmes CUDA tels quels. Aucun portage vers un autre langage n’est nécessaire. Cela vaut même si le programme utilise du PTX asm inline
  • Le compilateur SCALE accepte les mêmes options en ligne de commande et le même dialecte CUDA que nvcc, ce qui lui permet de fonctionner comme un remplaçant direct
  • Il « imite » une installation du NVIDIA CUDA Toolkit afin que les outils et scripts de build existants continuent de fonctionner tels quels

Quels projets ont été testés ?

• SCALE a été validé en compilant des projets CUDA open source et en exécutant leurs tests
• Les projets open source suivants sont actuellement inclus dans les tests automatisés nocturnes et les passent intégralement
  • NVIDIA Thrust, Blender Cycles, AMGX, llama-cpp, faiss, xgboost, GOMC, stdgpu, hashcat

Quels GPU sont pris en charge ?

• Les cibles GPU suivantes sont prises en charge et incluses dans les tests Nightly
  • AMD gfx1030 (Navi 21, RDNA 2.0)
  • AMD gfx1100 (Navi 31, RDNA 3.0)
• Les cibles GPU suivantes ont fait l’objet de tests manuels ponctuels et « semblent fonctionner »
  • AMD gfx1010
  • AMD gfx1101
• Le support des GPU suivants est en cours de développement
  • AMD gfx900 (Vega 10, GCN 5.0)
• Si vous souhaitez une prise en charge rapide d’une architecture GPU AMD spécifique, veuillez les contacter

Les composants de SCALE

• Un compilateur compatible nvcc capable de compiler du CUDA dialecte nvcc pour les GPU AMD, y compris le PTX asm
• Une implémentation des API CUDA Runtime et Driver pour les GPU AMD
• Une bibliothèque wrapper open source qui fournit les API « CUDA-X » en déléguant aux bibliothèques ROCm. Des bibliothèques comme cuBLAS et cuSOLVER sont gérées de cette manière

En quoi SCALE diffère-t-il des autres solutions ?

• Au lieu de proposer une nouvelle manière d’écrire des logiciels GPGPU, SCALE permet de compiler directement pour les GPU AMD des programmes écrits dans le langage CUDA largement utilisé
• SCALE vise une compatibilité complète avec NVIDIA CUDA. Le projet considère que les utilisateurs doivent pouvoir prendre en charge plusieurs fournisseurs de GPU sans maintenir plusieurs bases de code ni faire de compromis sur les performances
• Le langage de SCALE est un superset de NVIDIA CUDA, avec des extensions de langage optionnelles qui rendent l’écriture de code GPU plus simple et plus efficace pour les utilisateurs qui souhaitent s’éloigner de nvcc
• SCALE est un travail en cours. Si une API manquante empêche son utilisation, veuillez les contacter. Ils ajusteront les priorités de développement

Le résumé de GN⁺

• SCALE est une boîte à outils importante qui permet de compiler nativement des applications CUDA pour les GPU AMD
• Il offre un avantage majeur aux développeurs en permettant d’exécuter des programmes CUDA existants sur des GPU AMD sans modification
• Son objectif de compatibilité complète avec NVIDIA CUDA en fait une solution intéressante pour prendre en charge plusieurs fournisseurs de GPU
• Le projet est encore en cours de développement, et si une API nécessaire manque, il est possible de contacter l’équipe afin d’ajuster les priorités
• Parmi les projets aux fonctionnalités similaires figurent ROCm et HIP

[GN⁺]

Avis Hacker News

• Beaucoup estiment qu’AMD devrait prendre en charge une couche de traduction, mais certains pensent que c’est une mauvaise idée

  • CUDA n’a pas été conçu de manière neutre vis-à-vis des fournisseurs, et Nvidia pourrait créer des obstacles techniques et juridiques
  • Par exemple, faire tourner cuDNN ou cuBLAS par-dessus pourrait enfreindre les accords de licence
  • Ces bibliothèques Nvidia feraient alors partie de la frontière d’API qu’AMD devrait réimplémenter et prendre en charge

• Certains estiment que rechercher la compatibilité avec les bugs est une absurdité

  • Les utilisateurs importants de CUDA sont open source
  • AMD pourrait implémenter directement le support dans des projets upstream comme pytorch ou llama.cpp
  • Avec ce support, la communauté pourrait assurer la maintenance

• Certains disent ne pas comprendre comment du code fortement dépendant du matériel pourrait « simplement fonctionner » sur AMD

  • La plupart des codes CUDA sérieux prennent en compte la taille du fichier de registres et de la mémoire partagée, les instructions wgmma, les dispositions mémoire et registres optimales pour les tensor cores, les instructions du tensor memory accelerator, etc.

• Si c’est vrai, c’est impressionnant, mais certains notent que ce n’est pas open source et qu’il manque des détails précis sur son fonctionnement

  • On ne sait pas très bien pourquoi on s’attend aujourd’hui à ce qu’un projet soit open source, ou au moins source-available

• Certains estiment que la principale raison de la valorisation élevée de Nvidia est qu’AMD n’investit pas pour rendre ses GPU vraiment utiles pour le ML

  • AMD craint peut-être des mesures antitrust, ou il existe peut-être quelque chose dans son approche matérielle qui limite sa compétitivité
  • L’entreprise semble être passée à côté de milliards de dollars pendant le pic de demande des GPU pour le minage de cryptomonnaies et l’actuel boom de l’IA

• Certains ont tellement l’impression qu’AMD se trompe qu’ils ont envie de féliciter ce genre de projet

  • C’est particulièrement frustrant sur Linux, où des fonctions des ordinateurs portables existent physiquement mais ne peuvent pas être utilisées

• Une personne dit avoir travaillé chez Spectral Compute il y a quelques années

  • L’équipe technique était très intelligente et très compétente
  • À l’époque, elle ne ciblait pas seulement AMD, elle surpassait aussi le backend LLVM PTX de base et NVCC

• Une personne dit avoir déjà un peu écrit en CUDA

  • Elle se demande quelle est la configuration de base pour écrire du code destiné aux cartes AMD

• Certains trouvent ce projet excellent

  • Ils espèrent qu’AMD finira par concurrencer Nvidia directement

• C’est bien qu’il existe une page sur les limitations actuelles, mais certains soulignent que ce que la plupart des gens appellent « CUDA » ne représente qu’une petite partie des capacités réelles de CUDA

  • Un tableau comparatif des fonctionnalités avancées comme warp shuffle, atomic operations, DPX, TMA, MMA, etc., serait utile
  • L’idéal serait un tableau faisant correspondre les instructions PTX à leurs équivalents RDNA, ou à la liste des instructions qui les émulent

• Certains pensent que cela pourrait être réel, puisque c’est techniquement possible

  • Analyser du PTX inline et le mapper vers AMDGPU serait extrêmement pénible
  • Faire fonctionner un code source CUDA sans PTX inline en ciblant AMDGPU reviendrait à peu près à le remplacer par HIP
  • Certains détails peuvent toutefois poser problème, par exemple si le modèle atomique ne correspond pas ou si Volta utilise un modèle de pointeur d’instruction différent
  • Mais cela peut être fait correctement
  • AMD ne le fera probablement pas
  • CUDA n’est généralement pas quelque chose de si formidable, et l’équipe juridique risque de créer des problèmes
  • Mais d’autres acteurs peuvent tout à fait le faire