Was intern passiert, wenn man einen CUDA-Kernel ausführt

 (fergusfinn.com)
2 Punkte von GN⁺ 3 시간 전 | 1 Kommentare | Auf WhatsApp teilen
  • Selbst ein einfaches CUDA-Programm zur Vektoraddition durchläuft bis zum Ergebnis 2.000000 eine Compiler-Pipeline, Treiberaufrufe, eine GPU-Befehlswarteschlange, Warp-Scheduling, Speicherhierarchien und ein Completion-Semaphore
  • nvcc trennt Host-Code und Device-Code, erzeugt mit cicc PTX, mit ptxas SASS, bündelt cubin und PTX in einem fatbin und legt dieses in die Linux-Executable
  • Die Launch-Syntax vadd<<<4096, 256>>> wird in einen Host-Launch-Stub umgewandelt; die Argumente da, db, dc, n werden über die CUDA-Runtime und libcuda.so.1 an den Treiber übergeben
  • Die GPU-Ausführung beginnt mit QMD, pushbuffer, GPFIFO, GP_PUT und einem doorbell-MMIO-Schreibzugriff; die 128 SMs der RTX 4090 führen die Konfiguration aus 4096 Blöcken und 256 Threads pro Block auf Warp-Ebene aus
  • Wegen der niedrigen arithmetischen Intensität dieses Kernels, der pro float-Addition 12 Byte übertragen muss, ist er in Nsight Compute mit 10,78 μs, 79,65 % des DRAM-Peaks und 5,17 % Warp-Issue von der Speicherbandbreite abhängig

Beispiel-Kernel und Beobachtungsumfang

  • Das Beispielprogramm addiert mit dem CUDA-Kernel vadd zwei float-Arrays und speichert das Ergebnis in einem dritten Array
    • Mit n = 1 << 20 werden 1.048.576 floats verarbeitet
    • Die Launch-Konfiguration lautet vadd<<<4096, 256>>>(da, db, dc, n) und verwendet 4096 * 256 = n Threads
  • Wenn man für die RTX 4090 mit nvcc -arch=sm_89 kompiliert und ausführt, wird c[0]=2.000000 c[n-1]=2.000000 ausgegeben
  • Selbst an dieser einen Ergebniszeile sind zig Millionen CPU-Instruktionen, device files, etwa 900 ioctls und speichergemappte doorbell-Register beteiligt

Wie nvcc die ausführbare Datei erzeugt

  • Mit nvcc --keep kann man die Artefakte der Compiler-Pipeline direkt ansehen
    • vadd.ptx: PTX des von cicc erzeugten Device-Codes
    • vadd.sm_89.cubin: SASS des von ptxas erzeugten Device-Codes
    • vadd.fatbin: ein fatbin, das cubin und PTX bündelt
    • vadd.cudafe1.stub.c: Host-Launch-Stub und Kernel-Registrierungscode
    • vadd.o: finales Host-Objekt inklusive fatbin
  • Der Host-Code wird vom Host-Compiler verarbeitet, während der Device-Kernel vadd die Schritte cicc und ptxas durchläuft
  • PTX ist eine virtuelle ISA, nutzt typisierte, unendlich viele virtuelle Register und spiegelt die tatsächliche Anzahl der Hardware-Register nicht direkt wider
    • Das Beispiel-PTX enthält die Berechnung blockIdx.x * blockDim.x + threadIdx.x, Grenzprüfung, global load, float add und global store
    • CUDA-Pointer sind standardmäßig generic pointer; daher werden sie mit cvta.to.global in global addresses umgewandelt, bevor ld.global verwendet wird
    • mul.wide.s32 wandelt den Index in einen Offset in Einheiten von 4 Byte, also sizeof(float), um und erweitert von 32 auf 64 Bit
  • SASS sind architekturspezifische echte Instruktionen und erscheinen in der Ausgabe für die RTX 4090 kompakter als PTX
    • S2R kopiert Spezialregister wie SR_CTAID.X und SR_TID.X in allgemeine Register
    • Die Kombination aus mul.wide und add in PTX wird in SASS zu IMAD.WIDE zusammengeführt
    • Die cvta-Umwandlung geht im Adressierungsprozess auf
  • Der Operand c[0x0][...] verweist auf die vom Treiber verwaltete constant bank 0
    • Die Pointer a, b, c liegen bei 0x160, 0x168, 0x170
    • n liegt bei 0x178
    • Auch Launch-Geometrie und ABI-Werte wie blockDim.x befinden sich in derselben bank
  • cubin ist eine ELF-Datei, also dasselbe Containerformat wie eine Linux-Executable
    • Die fatbinary bündelt cubin und PTX gemeinsam
    • Auf dieser RTX 4090 wird tatsächlich SASS ausgeführt, PTX ist aber als Fallback enthalten, den der Treiber auf anderen Architekturen per JIT kompilieren kann
    • Da PTX ausführlicher Klartext ist, komprimiert nvcc es standardmäßig

Wie Host-Code den Launch vorbereitet

  • Das Compiler-Frontend cudafe++ fügt einen versteckten Constructor ein, der vor main ausgeführt wird
    • Dieser Constructor registriert die eingebettete fatbinary bei der CUDA-Runtime
    • Er verknüpft den Host-seitigen Funktionspointer vadd mit dem gemangelten Device-Kernel-Namen im fatbin
  • Die Syntax vadd<<<4096, 256>>>(da, db, dc, n) wird in den generierten Host-Launch-Stub umgewandelt
    • da, db, dc, n werden jeweils an den Offsets 0, 8, 16, 24 ausgerichtet in einem Argumentpuffer im Host-Speicher abgelegt
    • Diese Offsets entsprechen den Positionen 0x160, 0x168, 0x170, 0x178, die SASS aus constant bank 0 liest
  • Der Stub ruft __cudaLaunch auf und übergibt dabei die Adresse der Host-seitigen Dummy-Funktion vadd
    • Diese Adresse ist keine auf der CPU auszuführende Funktionsadresse, sondern dient als Schlüssel für die Lookup-Tabelle der Runtime
    • Die Runtime findet den entsprechenden Device-Symbolnamen und übergibt dann an den closed-source User-Mode-Treiber libcuda.so.1
  • Beim ersten GPU-Aufruf öffnet die CUDA-Runtime dynamisch libcuda.so.1 und erzeugt einen Context
    • In strace ist zu sehen, dass /lib/x86_64-linux-gnu/libcuda.so.1 geöffnet wird
    • Der Context enthält einen channel, über den die CPU mit der GPU kommuniziert
  • Seit CUDA 12.2 ist module loading standardmäßig lazy
    • Das Hochladen des SASS-cubin wird bis zum ersten Launch eines bestimmten Kernels aufgeschoben
    • Es lässt sich über CUDA_MODULE_LOADING steuern

Die Befehlswarteschlange, die Arbeit an die GPU übergibt

  • Eine GPU nimmt keinen Funktionsaufruf wie eine CPU entgegen und springt nicht zu einem Entry Point
    • Sie liest über den PCIe-Bus hinweg einen Treiber-Command-Stream im Host-Speicher
    • cuLaunchKernel legt den fertigen Launch-Befehl in diesen Stream und benachrichtigt die GPU
  • Bei der ersten Ausführung kopiert der Treiber das Kernel-SASS in den GPU-Speicher
    • Er allokiert einen Code-Buffer und kopiert SASS hinein
  • Im channel gibt es zwei zentrale Strukturen im Host-RAM
    • pushbuffer: Speicherbereich, in den der Treiber methods schreibt, also GPU-Befehle
    • GPFIFO: ein Pointer-Ringbuffer, der auf pushbuffer-Spans zeigt
  • Ein GPFIFO-Eintrag besteht aus zwei 32-Bit-Wörtern, die (base, length) eines pushbuffer-Spans angeben
  • GPU und Treiber verfolgen mit zwei Cursors, wo Arbeit konsumiert und produziert wurde
    • GP_GET: zeigt an, bis wohin die GPU konsumiert hat
    • GP_PUT: zeigt an, bis wohin der Treiber produziert hat
    • Beide liegen in einer per-channel-Struktur namens USERD
  • Beim Kernel-Launch schreibt der Treiber methods in einen pushbuffer-Span, lässt einen GPFIFO-Eintrag darauf zeigen und schiebt GP_PUT vor
  • Auf modernen GPUs überwacht die host engine den Cursor nicht kontinuierlich, daher ist eine doorbell nötig
    • Die GPU mappt ein kleines Registerfenster in den Prozess
    • Der Treiber schreibt das Work-Submit-Token des channels in das doorbell-Register
    • Nach dem Empfang der doorbell liest die host engine GP_PUT und holt GPFIFO-Eintrag und pushbuffer-Span per DMA

Welche Ausführungsinformationen QMD enthält

  • Der Launch beginnt mit einem method burst aus SET_INLINE_QMD_ADDRESS_A/B und LOAD_INLINE_QMD_DATA
  • QMD (Queue Meta Data) ist der Launch-Descriptor des Compute-Grids
    • Enthält die Grid- und Blockgrößen 4096, 256
    • Enthält die Registerzahl pro Thread und den Bedarf an shared memory
    • Enthält die Programmstartadresse und die Adresse der constant bank mit den Kernel-Argumenten
    • Enthält auch die Position, an die Completion gemeldet wird
  • Die vom Host-Stub gepackten Argumente kopiert der Treiber in die constant bank, und die Adresse dieser bank wird im QMD eingetragen
  • QMD teilt der GPU mit, wo SASS liegt, wie das parallele Programm aufgebaut ist und wo das Completion-Signal abgelegt werden soll
  • cuLaunchKernel kehrt in dem Moment zurück, in dem die doorbell ausgelöst wird
    • Der Aufruf ist asynchron, die CPU kann also weiterlaufen, während die GPU-Arbeit stattfindet

SM, Warp und Occupancy

  • Die host engine übergibt QMD an den compute work distributor
    • Diese Komponente existiert einmal pro GPU
    • Sie verteilt den linearen SASS-Instruktionsstrom auf die SMs und lässt ihn als paralleles Programm ausführen
  • Die Ziel-GPU GeForce RTX 4090 verwendet 128 SMs
    • Der Launch besteht aus 4096 Blöcken und 256 Threads pro Block
  • Jeder SM hat einen lokalen Instruction Cache, und aktive Warps halten einen Program Counter
    • Seit Volta gibt es ein Independent-Thread-Scheduling-Modell mit Program Counter und Call Stack pro Thread
    • Das Issue erfolgt weiterhin auf Warp-Ebene
  • Im Beispiel-Kernel bestimmt das Resource Limit die Block Residency
    • 256 threads = 8 warps pro Block
    • ptxas reserviert 16 Register pro Thread
    • Nach Registern wären 16 Blöcke pro SM möglich
    • Die Thread-Kapazität beträgt 1.536 aktive Threads pro SM, also sind nur 1536 / 256 = 6 Blöcke möglich
    • Daher können maximal 6 Blöcke pro SM, also 48 Warps, resident sein
  • Ein SM ist in 4 processing blocks, also sub-partitions, unterteilt
    • Die 48 resident Warps werden gleichmäßig auf 4 sub-partitions verteilt
    • Jeder Warp-Scheduler verwaltet im vollen Zustand 12 aktive Warps
    • In jedem Cycle wählt er einen eligible Warp aus und dispatched die nächste Instruktion an 32 Lanes

Bedingungen, unter denen ein Warp eligible wird

  • Eine GPU extrahiert dynamische Abhängigkeiten aus einem einzelnen Thread nicht in großem Umfang wie bei der Out-of-Order-Ausführung einer CPU
    • Sie hält viele resident Warps bereit und wechselt bei Stalls zu einem anderen Warp, um Latenz zu verbergen
    • Der Compiler plant vorhersehbares Timing, und die Hardware-Scoreboard behandelt schwer vorhersehbare Teile
  • Eine 128-Bit-SASS-Instruktion enthält eine von ptxas geschriebene control-code payload
    • Fixed-Latency-Instruktionen enthalten eine statische Stall-Zahl
    • Ein Yield-Hint teilt mit, ob Scheduler-Priorität abgegeben werden soll
    • Operationen mit variabler Latenz verwenden sechs per-warp physical scoreboard barriers
  • Im Beispiel-SASS-Abschnitt setzen die beiden LDG.E dieselbe Scoreboard-Barrier B2
    • FADD hat B2 als wait-on
    • Bis beide Loads zurückkommen und die Barrier gelöscht ist, ist der betreffende Warp ineligible
    • Der Scheduler wählt in der Zwischenzeit andere Warps derselben sub-partition aus
  • Der Übergang von FADD zu STG.E wird als Fixed Latency behandelt
    • FADD hat stall=5 und parkt den Warp einige Cycles, bis das Ergebnis in R9 bereitsteht
    • Eine separate Barrier ist nicht nötig
  • Diese control payload ist in der Standardausgabe von nvdisasm verborgen
    • In der rohen 128-Bit-Codierung von cuobjdump -sass steckt sie im zweiten 64-Bit-Wort
    • Das Layout ist nicht dokumentiert, sondern wurde durch Microbenchmarking rekonstruiert

Speicherzugriffe und Performance-Messung

  • Wenn ein Warp LDG.E ausführt, berechnen alle 32 Threads jeweils ihre Adresse
    • Im Beispiel handelt es sich um Zugriffe auf aufeinanderfolgende float-Arrays, sodass der gesamte Warp einen zusammenhängenden Block von 32 * 4 = 128 bytes anfordert
  • Die Load/Store-Unit des SM führt request coalescing aus
    • Sie fasst die 32 Anforderungen à 4 Byte zu vier 32-Byte-Sector-Requests zusammen
    • Bei nicht aufeinanderfolgenden Zugriffen könnte mehr Datenvolumen gelesen werden als nötig
  • Eine coalesced request prüft zuerst den lokalen L1 Data Cache des SM
    • Bei einem Miss geht sie über das Crossbar-Interconnect zu einem Slice des 72 MB großen L2 Cache
    • Bei einem weiteren Miss geht sie über Memory Controller und Memory Bus in den GDDR6X-VRAM
  • Auch ein STG.E-Store folgt im Prinzip demselben Pfad in umgekehrter Richtung
  • Die Messwerte aus Nsight Compute zeigen, dass dieser Kernel memory-bound ist
    • launch__grid_size: 4.096
    • launch__block_size: 256
    • launch__registers_per_thread: 16
    • launch__waves_per_multiprocessor: 5,33
    • sm__warps_active.avg.pct_of_peak: 82,77 %
    • smsp__issue_active.avg.pct_of_peak: 5,17 %
    • dram__throughput.avg.pct_of_peak: 79,65 %
    • gpu__time_duration.sum: 10,78 μs
  • Der Kernel hat eine sehr niedrige arithmetische Intensität
    • Er führt eine float-Addition pro zwei 4-Byte-Loads und einem 4-Byte-Store aus, also pro insgesamt 12 Byte Übertragung
    • Auf DRAM-Read-Seite liest er 8,4 MB in 10,78 μs, etwa 780 GB/s und damit rund 4/5 des Peaks
    • Die 4-MB-Ausgabe c passt in den 72-MB-L2 und wird daher erst dann in den DRAM geflusht, wenn der Device-to-Host-Copy sie liest

Wie das Ergebnis zur CPU zurückkommt

  • Da der Kernel-Launch in dem Moment zur CPU zurückkehrt, in dem die doorbell ausgelöst wurde, muss die GPU den Abschluss separat melden
  • Wenn alle 4096 Blöcke retired sind, postet die GPU das im QMD enthaltene Completion-Semaphore
    • Das fence field des QMD liegt in den Wörtern 23–24
  • Im default stream liegt cudaMemcpy(c, dc, ...) hinter dem Kernel
    • Die GPU-copy engine bleibt gated, bis das Semaphore gesetzt ist
    • Da c im 72-MB-L2 noch dirty ist, wird der Lesezugriff der copy engine aus dem L2 bedient, ohne Roundtrip zum DRAM
    • Die Daten werden über PCIe in den Host-Speicher übertragen
  • Nach Abschluss des Kopiervorgangs postet die copy engine ihr eigenes Semaphore
    • Das Warten von cudaMemcpy auf dem Host endet
    • c ist wieder normaler Host-Speicher
    • printf liest c[0] und c[n-1] aus dem RAM und gibt sie auf stdout aus

Wie man in den Launch hineinblickt

  • Allein das Lesen der open kernel modules reicht nicht aus, um manche Abläufe direkt zu prüfen, weil libcuda closed-source ist
  • Method writes laufen ohne Syscall ab und schreiben direkt in einen bereits gemappten write-combined Buffer; um den pushbuffer zu sehen, muss man also Speicher auslesen
  • Mit einem LD_PRELOAD-Shim kann man mmap umhüllen und Bereiche aufzeichnen, die von /dev/nvidia* gemappt wurden
    • Wenn ein Testprogramm direkt nach dem Launch die Dump-Funktion des Shims aufruft, kann es den gemappten pushbuffer ausgeben
    • Der Dump sucht den method burst, der SET_INLINE_QMD_ADDRESS_A entspricht
  • Ein pushbuffer-method-header enthält Opcode, Payload Count, Subchannel Index und Register Offset als Bitfelder
    • 0x0318 ist SET_INLINE_QMD_ADDRESS_A
    • 0x0320 + i * 4 ist LOAD_INLINE_QMD_DATA(i)
    • Im Dump ist ein increasing-method burst mit Count 66 zu sehen; er enthält zwei Address-Wörter und 64 QMD-Wörter, also insgesamt ein 256-Byte-QMD inline
    • Wort 12 im QMD ist 0x1000, Wort 18 ist 0x100; sie entsprechen den Launch-Werten 4096 und 256
  • Das Driver Setup läuft über ioctl
    • Bei einem One-Kernel-Programm protokolliert strace 948 ioctls
    • Die meisten davon sind einmaliges Setup
    • Die wichtigsten File Descriptors sind /dev/nvidiactl und /dev/nvidia-uvm
    • Das ioctl-Magic-Byte des NVIDIA Resource Managers ist 0x46, also 'F'
    • Command Number 0x2A wird als NV_ESC_RM_CONTROL, 0x2B als NV_ESC_RM_ALLOC interpretiert
  • In der von nvcc --keep erzeugten Datei vadd.cudafe1.stub.c sieht man auch den Startup-Registrierungscode
    • Eine Funktion mit __attribute__((__constructor__)) wird vor main ausgeführt
    • Über __cudaRegisterBinary und __cudaRegisterEntry werden der Host-Funktionspointer vadd und der Device-Entry-Point _Z4vaddPKfS0_Pfi verknüpft

Verwandte Beiträge

1 Kommentare

 
GN⁺ 3 시간 전
Hacker-News-Kommentare

  • Ein interessanter Artikel, und auch die Erklärung des Semaphors des Default-Streams war spannend
    Es ist gut, dass CUDA die Befehlssynchronisation implizit übernimmt und parallele Befehle optional über Streams nutzbar macht
    Das steht im Kontrast zu Vulkan, das die gesamte Komplexität der Synchronisation von Anfang an dem Nutzer überlässt

  • Auf der Hardware-Seite gibt es einige öffentlich zugängliche Dokumente
    Man muss nicht zwingend den Kernel-Source lesen, um Methodendokumente oder das QMD-Format zu finden
    Siehe https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...

  • Sehr nützlich
    Besonders der Teil zu Doorbell und QMD war am hilfreichsten, weil er zeigt, wie die CUDA-Ausführungssyntax mit dem zusammenhängt, was tatsächlich an die GPU übermittelt wird
    Die meisten Erklärungen enden bei Kernel, Block und Warp, aber dieser Artikel macht den Pfad CPU→Treiber→GPU viel leichter nachvollziehbar

  • Der Steuerungscode ist etwas komplexer, als im Artikel beschrieben
    In der Praxis ist es eher ein Tabellen-Lookup als Bits in einem Steuerwort

  • Es gibt inzwischen Firmen, deren Hauptgeschäft darin besteht, Kernel zu optimieren, damit sie schneller laufen
    Ich frage mich, ob solche Firmen irgendwann von einer Open-Source-Bibliothek verdrängt werden, die das extrem gut kann
    Nvidia könnte so etwas jederzeit selbst herausbringen
    Oder es läuft für diese Firmen sogar noch besser, wenn große Anbieter sie übernehmen, um daraus einen moat für schnellere Inferenz zu machen

    • Kurzfristig scheinen Acqui-Hires ziemlich wahrscheinlich
      Wenn man sich aber ansieht, wie sich Modelle bei einschlägigen Benchmarks wie kernelbench verbessern, glaube ich, dass am Ende auch allgemeinere Lösungen kommen werden
      Das Problem ist, dass mit jeder neuen Hardware-Generation oft Einschränkungen oder Features auftauchen, die bestehende Modelle noch nie gesehen haben
      Zum Beispiel war Blackwells tcgen05 zeitweise ein Out-of-Distribution-Fall
      Wenn Modelle besser zu generalisieren beginnen, muss das keine fatale Hürde sein, aber zumindest im Moment ist es noch ein Stolperstein
      [1] https://kernelbench.com/
    • Wenn man CUDA in großem Maßstab betreibt, geht ekelhaft viel Engineering-Zeit für Bugs in Nvidia-Treibern und -Bibliotheken drauf
      Ich kenne nicht viele Leute, die darauf hoffen, noch stärker von Nvidia-Bibliotheken abhängig zu werden
    • Wahrscheinlich nicht
      Denn die Details des Workloads – also genaue Parameter, die Repräsentation der Daten im Speicher und die Wertebereiche – führen zu sehr unterschiedlichen Optimierungsstrategien

  • Ich habe gerade meinen HPC-Master abgeschlossen und Kurse zu CUDA, MPI+CUDA und OpenCL besucht; wenn ich vor den Kursen so einen Artikel gelesen hätte, hätte mir das viel geholfen
    Besonders der Kontext rund um den Teil, der behandelt, was es bedeutet, dass ein Warp ausführbar ist, war gut

  • Zunächst einmal: ein guter Artikel, der viele Ecken und Winkel gründlich ausleuchtet
    Allerdings verschwindet vieles von dem voodoo-artigen Kram im User-Space, wenn man nicht über die runtime API von CUDA geht
    Wenn man die Treiber-API verwendet und den Kernel-Source als String an NVIDIAs Runtime-Compiler übergibt, kann man viel besser sehen, was passiert
    Nicht alles, aber doch ziemlich viel wird transparent
    Eine „rohere“ Version gibt es hier:
    https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
    Wenn man dasselbe in Form einer viel besser lesbaren und dennoch völlig transparenten modernen C++-API sehen will, dann hier:
    https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
    Das ist ein Beispielprogramm aus meiner Header-only-Bibliothek für CUDA-API-Wrapper

    • Die Treiber-API ist nett, weil man CUDA-Kernel wie hot-reloadbare Shader behandeln kann
      Es macht Spaß, während der Laufzeit Code zu ändern und so zu entwickeln

  • Auf Bare Metal?