SCALE - CUDA auf AMD-GPUs ohne Änderungen ausführen

• SCALE ist ein GPGPU-Programmier-Toolkit, das es ermöglicht, CUDA-Anwendungen nativ für AMD-GPUs zu kompilieren
• Weder CUDA-Programme noch Build-Systeme müssen angepasst werden, und Unterstützung für weitere GPU-Anbieter sowie zusätzliche CUDA-APIs ist in Entwicklung

Wie funktioniert es?

• SCALE bringt im Vergleich zu anderen plattformübergreifenden GPGPU-Lösungen einige wichtige Innovationen mit
  • Es akzeptiert CUDA-Programme unverändert. Ein Port auf eine andere Sprache ist nicht nötig. Das gilt auch dann, wenn das Programm Inline-PTX-asm verwendet
  • Der SCALE-Compiler akzeptiert dieselben Kommandozeilenoptionen und denselben CUDA-Dialekt wie nvcc und funktioniert so als Drop-in-Ersatz
  • Es gibt sich als installierte NVIDIA CUDA Toolkit-Umgebung aus, sodass bestehende Build-Tools und Skripte unverändert weiter funktionieren

Welche Projekte wurden getestet?

• SCALE wurde validiert, indem Open-Source-CUDA-Projekte kompiliert und deren Tests ausgeführt wurden
• Derzeit sind die folgenden Open-Source-Projekte in den nächtlichen automatisierten Tests enthalten und bestehen vollständig
  • NVIDIA Thrust, Blender Cycles, AMGX, llama-cpp, faiss, xgboost, GOMC, stdgpu, hashcat

Welche GPUs werden unterstützt?

• Die folgenden GPU-Ziele werden unterstützt und sind in den Nightly-Tests enthalten
  • AMD gfx1030 (Navi 21, RDNA 2.0)
  • AMD gfx1100 (Navi 31, RDNA 3.0)
• Die folgenden GPU-Ziele wurden in vorläufigen manuellen Tests geprüft und „scheinen zu funktionieren“
  • AMD gfx1010
  • AMD gfx1101
• An Unterstützung für die folgenden GPUs wird gearbeitet
  • AMD gfx900 (Vega 10, GCN 5.0)
• Wer zügig Unterstützung für eine bestimmte AMD-GPU-Architektur benötigt, soll sich melden

Bestandteile von SCALE

• Ein mit nvcc kompatibler Compiler, der den nvcc-Dialekt von CUDA für AMD-GPUs kompilieren kann, einschließlich PTX-asm
• Eine Implementierung der CUDA-Runtime- und Driver-API für AMD-GPUs
• Eine Open-Source-Wrapper-Bibliothek, die durch Delegation an ROCm-Bibliotheken „CUDA-X“-APIs bereitstellt. Bibliotheken wie cuBLAS und cuSOLVER werden auf diese Weise behandelt

Unterschiede zwischen SCALE und anderen Lösungen

• Statt eine neue Art zu bieten, GPGPU-Software zu schreiben, ermöglicht SCALE die direkte Kompilierung von Programmen, die in der weit verbreiteten CUDA-Sprache geschrieben sind, für AMD-GPUs
• SCALE strebt vollständige Kompatibilität mit NVIDIA CUDA an. Dahinter steht die Überzeugung, dass Nutzer mehrere GPU-Anbieter unterstützen können sollten, ohne mehrere Codebasen pflegen oder Kompromisse bei der Performance eingehen zu müssen
• Die Sprache von SCALE ist eine Obermenge von NVIDIA CUDA und bietet optionale Spracherweiterungen, die das Schreiben von GPU-Code für Nutzer erleichtern und effizienter machen sollen, die sich von nvcc lösen möchten
• SCALE ist ein laufendes Projekt. Falls eine API fehlt, die die Nutzung verhindert, soll man sich melden. Dann werden die Entwicklungsprioritäten angepasst

Zusammenfassung von GN⁺

• SCALE ist ein wichtiges Toolkit, das die native Kompilierung von CUDA-Anwendungen für AMD-GPUs ermöglicht
• Dass bestehende CUDA-Programme ohne Änderungen auf AMD-GPUs laufen können, ist ein großer Vorteil für Entwickler
• Das Ziel vollständiger Kompatibilität mit NVIDIA CUDA ist vorteilhaft für die Unterstützung mehrerer GPU-Anbieter
• Das Projekt ist noch in Arbeit; wenn eine benötigte API fehlt, kann das Entwicklerteam kontaktiert werden, um Prioritäten anzupassen
• Ähnliche Projekte mit vergleichbarer Funktionalität sind ROCm und HIP

[GN⁺]

Hacker-News-Kommentare

• Viele sind der Meinung, AMD sollte eine Übersetzungsschicht unterstützen, aber es gibt auch die Ansicht, dass das eine schlechte Idee ist

  • CUDA ist nicht herstellerneutral entworfen, und Nvidia könnte technisch und rechtlich Hürden aufbauen
  • Zum Beispiel könnte das Ausführen von cuDNN oder cuBLAS darauf gegen Lizenzvereinbarungen verstoßen
  • Diese Nvidia-Bibliotheken würden dann Teil der API-Grenze, die AMD nachimplementieren und unterstützen müsste

• Es gibt die Meinung, dass es töricht ist, Bug-Kompatibilität anzustreben

  • Wichtige CUDA-Nutzer sind Open Source
  • AMD könnte Unterstützung direkt in Upstream-Projekten wie pytorch oder llama.cpp implementieren
  • Mit entsprechender Unterstützung könnte die Community die Wartung übernehmen

• Es gibt die Ansicht, dass nicht nachvollziehbar ist, wie stark hardwareabhängiger Code auf AMD „einfach so funktionieren“ soll

  • Der meiste ernsthafte CUDA-Code berücksichtigt Registerdatei und Shared-Memory-Größe, wgmma-Instruktionen, optimale Speicher- und Register-Layouts für Tensor Cores, Tensor Memory Accelerator-Instruktionen usw.

• Falls es stimmt, wäre es beeindruckend, aber es ist nicht Open Source, und es fehlen genaue Details zur Funktionsweise

  • Ich verstehe nicht ganz, warum man heute erwartet, dass ein Projekt Open Source oder zumindest source-available ist

• Es gibt die Meinung, dass der Hauptgrund für Nvidias hohe Bewertung darin liegt, dass AMD nicht darin investiert, GPUs für ML nützlich zu machen

  • Möglicherweise fürchtet AMD kartellrechtliche Maßnahmen, oder es gibt etwas an ihrem Hardware-Ansatz, das die Wettbewerbsfähigkeit einschränkt
  • Das Unternehmen scheint während des Nachfragebooms nach GPUs fürs Krypto-Mining und jetzt des aktuellen AI-Booms Milliarden an Dollar verpasst zu haben

• Es gibt die Meinung, dass AMD so viel falsch macht, dass man ein Projekt wie dieses fast feiern möchte

  • Besonders unter Linux ist es extrem frustrierend, wenn Funktionen eines Laptops physisch vorhanden sind, aber nicht genutzt werden können

• Ich habe vor ein paar Jahren bei Spectral Compute gearbeitet

  • Das technische Team war sehr klug und äußerst kompetent
  • Damals zielte man nicht nur auf AMD, sondern übertraf auch das grundlegende LLVM-ptx-Backend und NVCC

• Ich habe ein wenig CUDA geschrieben

  • Ich frage mich, wie das grundlegende Setup aussieht, um Code für AMD-Karten zu schreiben

• Es gibt die Meinung, dass dieses Projekt großartig ist

  • Ich freue mich darauf, dass AMD in direkten Wettbewerb mit Nvidia tritt

• Es ist gut, dass es eine Seite zu den aktuellen Einschränkungen gibt, aber vieles von dem, was die meisten Leute als „CUDA“ bezeichnen, ist nur ein kleiner Teil der tatsächlichen CUDA-Funktionalität

  • Eine Vergleichstabelle für fortgeschrittene Funktionen wie Warp Shuffle, atomare Operationen, DPX, TMA, MMA usw. wäre hilfreich
  • Ideal wäre eine Tabelle, die PTX-Instruktionen auf entsprechende RDNA-Instruktionen oder auf Instruktionsfolgen abbildet, die diese emulieren

• Es gibt die Meinung, dass es real sein könnte, weil es technisch möglich ist

  • Inline-PTX zu parsen und auf AMDGPU abzubilden, wäre äußerst schmerzhaft
  • Bei CUDA-Quellcode ohne Inline-PTX auf AMDGPU zu zielen, wäre ähnlich wie ihn durch HIP zu ersetzen
  • Manche Details könnten fragwürdig sein, etwa wenn das atomare Modell nicht übereinstimmt oder Volta ein anderes Instruction-Pointer-Modell hat
  • Es kann jedoch korrekt umgesetzt werden
  • AMD wird das nicht tun
  • CUDA ist im Allgemeinen nicht gerade besonders gut, und die Rechtsabteilung würde Probleme verursachen
  • Andere könnten es jedoch durchaus schaffen