Ce qui se passe en interne quand on exécute un kernel CUDA

 (fergusfinn.com)
2 points par GN⁺ 3 시간 전 | 1 commentaires | Partager sur WhatsApp
  • Même un simple programme CUDA d’addition de vecteurs passe, avant d’obtenir le résultat 2.000000, par un pipeline de compilation, des appels au driver, une file de commandes GPU, l’ordonnancement des warps, la hiérarchie mémoire et un sémaphore de complétion
  • nvcc sépare le code hôte et le code device, produit du PTX avec cicc, du SASS avec ptxas, puis regroupe le cubin et le PTX dans un fatbin intégré à l’exécutable Linux
  • La syntaxe de lancement vadd<<<4096, 256>>> est transformée en stub de lancement côté hôte, et les arguments da, db, dc, n sont transmis au driver via le runtime CUDA et libcuda.so.1
  • L’exécution GPU démarre avec QMD, pushbuffer, GPFIFO, GP_PUT et une écriture MMIO dans la doorbell ; les 128 SM de la RTX 4090 exécutent une configuration de 4096 blocs et 256 threads par blocs au niveau des warps
  • À cause de sa faible intensité arithmétique, qui nécessite 12 octets transférés pour une addition de float, ce kernel est limité par la bande passante mémoire dans Nsight Compute : 10,78 μs, 79,65 % du pic DRAM et 5,17 % de warp issue

Kernel d’exemple et périmètre observé

  • Le programme d’exemple utilise le kernel CUDA vadd, qui additionne deux tableaux de float et stocke le résultat dans un troisième tableau
    • Avec n = 1 << 20, il traite 1 048 576 float
    • La configuration de lancement est vadd<<<4096, 256>>>(da, db, dc, n) et utilise 4096 * 256 = n threads
  • Compilé pour une RTX 4090 avec nvcc -arch=sm_89 puis exécuté, il affiche c[0]=2.000000 c[n-1]=2.000000
  • Même ce résultat d’une seule ligne implique des dizaines de millions d’instructions CPU, un device file, environ 900 ioctl et un registre doorbell mappé en mémoire

Comment nvcc fabrique l’exécutable

  • Avec nvcc --keep, on peut inspecter directement les artefacts du pipeline de compilation
    • vadd.ptx : le PTX du code device produit par cicc
    • vadd.sm_89.cubin : le SASS du code device produit par ptxas
    • vadd.fatbin : le fatbin regroupant le cubin et le PTX
    • vadd.cudafe1.stub.c : le stub de lancement hôte et le code d’enregistrement du kernel
    • vadd.o : l’objet hôte final contenant le fatbin
  • Le code hôte est traité par le compilateur hôte, tandis que le kernel device vadd passe par les étapes cicc et ptxas
  • Le PTX est une ISA virtuelle, qui utilise un nombre infini de registres virtuels typés et ne reflète pas directement le nombre de registres matériels réels
    • Le PTX de l’exemple contient le calcul blockIdx.x * blockDim.x + threadIdx.x, le contrôle de borne, le chargement global, l’addition float et le stockage global
    • Les pointeurs CUDA sont par défaut des generic pointers ; ils sont donc convertis en adresses globales avec cvta.to.global avant l’utilisation de ld.global
    • mul.wide.s32 convertit l’index en offset par unités de 4 octets, soit sizeof(float), et l’étend de 32 à 64 bits
  • Le SASS correspond aux vraies instructions propres à l’architecture ; dans la sortie ciblant la RTX 4090, il apparaît sous une forme plus compacte que le PTX
    • S2R copie des registres spéciaux comme SR_CTAID.X et SR_TID.X vers des registres généraux
    • La combinaison PTX de mul.wide et add est fusionnée en IMAD.WIDE dans le SASS
    • La conversion cvta est absorbée dans le processus d’adressage
  • Les opérandes c[0x0][...] pointent vers la constant bank 0 gérée par le driver
    • Les pointeurs a, b, c se trouvent à 0x160, 0x168, 0x170
    • n se trouve à 0x178
    • La géométrie de lancement, comme blockDim.x, et les valeurs ABI se trouvent aussi dans la même bank
  • Le cubin est un fichier ELF, le même format conteneur que les exécutables Linux
    • Le fatbinary regroupe le cubin et le PTX
    • Sur cette RTX 4090, c’est le SASS qui est réellement exécuté, mais le PTX est inclus comme fallback que le driver peut compiler en JIT pour d’autres architectures
    • Le PTX étant du texte brut verbeux, nvcc le compresse par défaut

Comment le code hôte prépare le lancement

  • Le frontend de compilation cudafe++ insère un constructeur caché exécuté avant main
    • Ce constructeur enregistre le fatbinary embarqué auprès du runtime CUDA
    • Il associe le pointeur de fonction côté hôte vadd au nom manglé du kernel device dans le fatbin
  • La syntaxe vadd<<<4096, 256>>>(da, db, dc, n) est transformée en host launch stub généré
    • da, db, dc, n sont placés dans l’argument buffer en mémoire hôte avec les offsets alignés 0, 8, 16, 24
    • Ces offsets correspondent aux emplacements 0x160, 0x168, 0x170, 0x178 que le SASS lit dans la constant bank 0
  • Le stub appelle __cudaLaunch en lui passant l’adresse de la fonction dummy vadd côté hôte
    • Cette adresse n’est pas une adresse de fonction à exécuter sur le CPU, mais sert de clé pour interroger la table d’enregistrement du runtime
    • Le runtime trouve le nom de symbole device correspondant, puis passe la main au driver en mode utilisateur closed-source libcuda.so.1
  • Lors du premier appel GPU, le runtime CUDA ouvre dynamiquement libcuda.so.1 et crée un contexte
    • Avec strace, on peut voir l’ouverture de /lib/x86_64-linux-gnu/libcuda.so.1
    • Le contexte contient un channel par lequel le CPU communique avec le GPU
  • Depuis CUDA 12.2, le chargement des modules est lazy par défaut
    • L’upload du cubin SASS est différé jusqu’au premier lancement d’un kernel donné
    • Il est contrôlable via CUDA_MODULE_LOADING

La file de commandes qui transmet le travail au GPU

  • Le GPU ne reçoit pas un appel de fonction comme un CPU et ne saute pas vers un entry point
    • Il lit, au-delà du bus PCIe, un flux de commandes du driver situé dans la mémoire hôte
    • cuLaunchKernel place la commande de lancement finalisée dans ce flux et notifie le GPU
  • Lors de la première exécution, le driver copie le SASS du kernel en mémoire GPU
    • Il alloue un code buffer et y copie le SASS
  • Le channel contient deux structures clés en RAM hôte
    • pushbuffer : zone mémoire dans laquelle le driver écrit des methods, c’est-à-dire des commandes GPU
    • GPFIFO : ring buffer de pointeurs vers des spans du pushbuffer
  • Une entrée GPFIFO est constituée de deux mots de 32 bits représentant le (base, length) d’un span de pushbuffer
  • Le GPU et le driver suivent les positions de consommation et de production au moyen de deux curseurs
    • GP_GET : indique jusqu’où le GPU a consommé
    • GP_PUT : indique jusqu’où le driver a produit
    • Les deux se trouvent dans une structure par channel appelée USERD
  • Lors du lancement du kernel, le driver écrit des methods dans un span du pushbuffer, fait pointer une entrée GPFIFO vers celui-ci, puis avance GP_PUT
  • Sur les GPU modernes, le host engine ne surveille pas continuellement le curseur, d’où la nécessité d’une doorbell
    • Le GPU mappe dans le process une petite fenêtre de registres
    • Le driver écrit le token de soumission de travail du channel dans le registre doorbell
    • Après réception de la doorbell, le host engine lit GP_PUT et récupère par DMA l’entrée GPFIFO et le span du pushbuffer

Les informations d’exécution contenues dans le QMD

  • Le lancement commence par une rafale de methods SET_INLINE_QMD_ADDRESS_A/B et LOAD_INLINE_QMD_DATA
  • Le QMD (Queue Meta Data) est le descripteur de lancement du compute grid
    • Il contient les dimensions de grille et de bloc, 4096 et 256
    • Il contient le nombre de registres par thread et les besoins en shared memory
    • Il contient l’adresse de début du programme et l’adresse de la constant bank contenant les arguments du kernel
    • Il contient aussi l’emplacement où signaler la complétion
  • Les arguments packés par le host stub sont copiés par le driver dans la constant bank, dont l’adresse est enregistrée dans le QMD
  • Le QMD indique au GPU l’emplacement du SASS, la manière de configurer le programme parallèle et l’emplacement du signal de complétion
  • cuLaunchKernel retourne au moment où la doorbell est déclenchée
    • L’appel est asynchrone, donc le CPU peut continuer à s’exécuter pendant que le travail GPU progresse

SM, warps et occupation

  • Le host engine transmet le QMD au compute work distributor
    • Ce composant est unique à l’échelle du GPU
    • Il distribue le flux linéaire d’instructions SASS aux SM afin de l’exécuter comme un programme parallèle
  • Le GPU cible, la GeForce RTX 4090, utilise 128 SM
    • Le lancement est composé de 4096 blocs et de 256 threads par bloc
  • Chaque SM possède son cache d’instructions local, et chaque warp actif conserve son compteur ordinal
    • Depuis Volta, il existe un modèle d’Independent Thread Scheduling avec compteur ordinal et call stack par thread
    • L’issue reste toutefois réalisée au niveau du warp
  • Dans le kernel d’exemple, la limite de ressources détermine la résidence des blocs
    • 256 threads = 8 warps par bloc
    • ptxas réserve 16 registres par thread
    • Du point de vue des registres, 16 blocs par SM seraient possibles
    • La capacité en threads est de 1 536 threads actifs par SM, donc 1536 / 256 = 6 blocs seulement sont possibles
    • Ainsi, au maximum 6 blocs par SM, soit 48 warps, sont résidents
  • Un SM est divisé en 4 processing blocks, ou sub-partitions
    • Les 48 warps résidents sont répartis uniformément entre les 4 sub-partitions
    • Chaque warp scheduler gère 12 warps actifs lorsqu’il est plein
    • À chaque cycle, il choisit un warp éligible et dispatch l’instruction suivante vers 32 lanes

Conditions pour qu’un warp soit éligible

  • Le GPU n’extrait pas beaucoup de dépendances dynamiques dans un thread unique comme le ferait une exécution out-of-order sur CPU
    • Il garde de nombreux warps résidents et, lorsqu’un stall survient, bascule vers un autre warp afin de masquer la latence
    • Le compilateur planifie les timings prévisibles, tandis qu’un hardware scoreboard traite les parties difficiles à prédire
  • Une instruction SASS de 128 bits contient un control-code payload écrit par ptxas
    • Les instructions à latence fixe contiennent un static stall count
    • Le yield hint indique s’il faut céder la priorité au scheduler
    • Les opérations à latence variable utilisent 6 barrières de scoreboard physiques par warp
  • Dans la section SASS de l’exemple, les deux LDG.E définissent la même barrière de scoreboard B2
    • FADD a B2 comme wait-on
    • Tant que les deux loads ne sont pas revenus et que la barrière n’est pas levée, ce warp est inéligible
    • Pendant ce temps, le scheduler choisit un autre warp de la même sub-partition
  • Le passage de FADD à STG.E est traité comme une latence fixe
    • FADD a stall=5 et park le warp pendant quelques cycles, jusqu’à ce que le résultat R9 soit prêt
    • Aucune barrière séparée n’est nécessaire
  • Ce control payload est masqué dans la sortie par défaut de nvdisasm
    • Dans l’encodage brut 128 bits de cuobjdump -sass, il est inclus dans le deuxième mot de 64 bits
    • Son layout n’est pas documenté ; il a été reconstruit par microbenchmarking

Accès mémoire et mesure de performance

  • Lorsqu’un warp exécute LDG.E, ses 32 threads calculent chacun une adresse
    • L’exemple accède à des tableaux de float consécutifs, donc l’ensemble du warp demande un bloc contigu de 32 * 4 = 128 bytes
  • L’unité load/store du SM effectue du request coalescing
    • Elle fusionne les 32 requêtes de 4 octets en 4 requêtes de secteur de 32 octets
    • Si l’accès n’avait pas été contigu, elle aurait pu lire plus de données que nécessaire
  • Une requête coalescée consulte d’abord le L1 Data Cache local du SM
    • En cas de miss, elle passe par l’interconnexion crossbar vers une slice du L2 Cache de 72 MB
    • En cas de miss en L2 également, elle passe par le memory controller et le bus mémoire vers la VRAM GDDR6X
  • Le store STG.E suit en principe le même chemin en sens inverse
  • Les mesures Nsight Compute montrent que ce kernel est memory-bound
    • 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
  • Le kernel a une intensité arithmétique très faible
    • Il effectue une addition float pour deux loads de 4 octets et un store de 4 octets, soit 12 octets transférés au total
    • Côté lectures DRAM, il lit 8,4 MB en 10,78 μs, soit environ 780 GB/s, environ 4/5 du pic
    • La sortie c de 4 MB tient dans les 72 MB du L2, donc elle n’est pas flushée vers la DRAM avant d’être lue par la copie device-to-host

Comment le résultat revient au CPU

  • Comme le kernel launch retourne au CPU au moment où la doorbell est déclenchée, le GPU doit signaler séparément sa complétion
  • Lorsque les 4096 blocs ont tous été retirés, le GPU poste le completion semaphore contenu dans le QMD
    • Le champ fence du QMD se trouve dans les mots 23–24
  • Dans le default stream, cudaMemcpy(c, dc, ...) est placé après le kernel
    • Le moteur de copie GPU reste gated jusqu’à ce que le sémaphore soit levé
    • Comme c est encore dirty dans les 72 MB du L2, la lecture du moteur de copie est servie depuis le L2 sans aller-retour DRAM
    • Les données traversent PCIe vers la mémoire hôte
  • Une fois la copie terminée, le moteur de copie poste son propre sémaphore
    • L’attente de cudaMemcpy côté hôte prend fin
    • c redevient de la mémoire hôte ordinaire
    • printf lit c[0] et c[n-1] depuis la RAM et les écrit sur stdout

Comment observer l’intérieur d’un lancement

  • Lire les open kernel modules ne suffit pas à vérifier directement certains comportements, car libcuda est closed-source
  • Les écritures de methods ne passent pas par un syscall : elles sont écrites directement dans un buffer write-combined déjà mappé ; pour voir le pushbuffer, il faut donc lire la mémoire
  • Un shim LD_PRELOAD peut envelopper mmap afin d’enregistrer les zones mappées depuis /dev/nvidia*
    • Si le programme de test appelle la fonction de dump du shim juste après le lancement, il peut afficher le pushbuffer mappé
    • Le dump recherche la rafale de methods correspondant à SET_INLINE_QMD_ADDRESS_A
  • Le header d’une method du pushbuffer contient, sous forme de bit fields, l’opcode, le nombre de payloads, l’index de subchannel et l’offset de registre
    • 0x0318 correspond à SET_INLINE_QMD_ADDRESS_A
    • 0x0320 + i * 4 correspond à LOAD_INLINE_QMD_DATA(i)
    • Dans le dump, on observe une rafale d’increasing-method de count 66 ; elle embarque inline deux mots d’adresse et 64 mots QMD, soit un QMD total de 256 octets
    • Dans le QMD, le mot 12 vaut 0x1000 et le mot 18 vaut 0x100, ce qui correspond aux valeurs de lancement 4096 et 256
  • La configuration du driver passe par des ioctl
    • Sur un programme à un seul kernel, strace enregistre 948 ioctl
    • La plupart correspondent à du setup effectué une seule fois
    • Les principaux file descriptors sont /dev/nvidiactl et /dev/nvidia-uvm
    • Le magic byte des ioctl du NVIDIA resource manager est 0x46, soit 'F'
    • Le numéro de commande 0x2A est interprété comme NV_ESC_RM_CONTROL, et 0x2B comme NV_ESC_RM_ALLOC
  • Dans vadd.cudafe1.stub.c, généré avec nvcc --keep, on peut aussi voir le code d’enregistrement au démarrage
    • La fonction marquée __attribute__((__constructor__)) s’exécute avant main
    • __cudaRegisterBinary et __cudaRegisterEntry associent le pointeur de fonction hôte vadd à l’entry point device _Z4vaddPKfS0_Pfi

À lire aussi

1 commentaires

 
GN⁺ 3 시간 전
Avis sur Hacker News

  • Article intéressant, et l’explication du sémaphore du flux par défaut était également amusante.
    J’aime le fait que CUDA gère implicitement la synchronisation des commandes, tout en permettant d’utiliser sélectivement les commandes parallèles via les flux.
    Cela contraste avec Vulkan, qui fait porter toute la complexité de la synchronisation à l’utilisateur dès le départ.

  • Côté matériel, il existe une partie de la documentation publique.
    Pas besoin de lire absolument le source du kernel pour trouver la documentation des méthodes ou le format QMD.
    Voir https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...

  • Très utile.
    La partie sur doorbell et QMD, en particulier, a été la plus utile, car elle montre comment la syntaxe d’exécution CUDA se relie concrètement à ce qui est soumis au GPU.
    La plupart des explications s’arrêtent à peu près aux kernels, blocs et warps, mais cet article rend le chemin CPU→pilote→GPU beaucoup plus facile à suivre.

  • Le code de contrôle est un peu plus complexe que ce qui est décrit dans l’article.
    En réalité, c’est plus proche d’une consultation de table que de bits dans un mot de contrôle.

  • Aujourd’hui, certaines entreprises ont pour activité principale d’optimiser des kernels pour les faire tourner plus vite.
    Je me demande si ces entreprises finiront un jour par être supplantées par une bibliothèque open source qui ferait ça très bien.
    Nvidia pourrait très bien sortir quelque chose de ce genre à tout moment.
    Ou bien les grands fournisseurs pourraient racheter ces entreprises pour en faire un moat visant à accélérer l’inférence, et les choses pourraient encore mieux tourner pour elles.

    • À court terme, les acquisitions-acquihires semblent assez plausibles.
      Cela dit, quand on voit les progrès des modèles sur des benchmarks liés comme kernelbench, je pense que des solutions plus généralisées finiront forcément par arriver.
      Le problème est qu’à chaque nouvelle génération de matériel apparaissent souvent des contraintes ou fonctionnalités que les modèles existants n’ont jamais vues.
      Par exemple, le tcgen05 de Blackwell a été, à un moment, un cas hors distribution.
      Si les modèles commencent à mieux généraliser, ce ne sera peut-être pas un obstacle fatal, mais pour l’instant cela reste au moins un frein.
      [1] https://kernelbench.com/
    • Faire tourner CUDA à grande échelle demande une quantité écœurante de temps d’ingénierie pour gérer les bugs des pilotes et bibliothèques Nvidia.
      Je n’ai pas vu grand monde souhaiter dépendre encore davantage des bibliothèques Nvidia.
    • Je pense que ce ne sera probablement pas le cas.
      Les détails de la charge de travail — les paramètres exacts, la représentation des données en mémoire, les plages de valeurs, etc. — font énormément diverger les stratégies d’optimisation.

  • Je viens de terminer un master en HPC et j’ai suivi des cours sur CUDA, MPI+CUDA et OpenCL ; lire ce genre d’article avant les cours m’aurait probablement beaucoup aidé.
    J’ai particulièrement apprécié le contexte autour de la partie expliquant ce que signifie qu’un warp soit exécutable.

  • D’abord, c’est un bon article, qui explore bien de nombreux recoins.
    Cela dit, si l’on ne passe pas par la runtime API de CUDA, une grande partie du vaudou en espace utilisateur disparaît.
    En utilisant l’API driver, en prenant le code source du kernel sous forme de chaîne et en le compilant avec le compilateur runtime de NVIDIA, on voit mieux ce qui se passe.
    Tout ne devient pas transparent, mais une bonne partie, oui.
    Une version plus « primitive » se trouve ici :
    https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
    Pour voir la même chose sous la forme d’une API C++ moderne, beaucoup plus lisible tout en restant entièrement transparente, voir ceci :
    https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
    C’est un programme d’exemple de ma bibliothèque header-only CUDA API wrappers.

    • L’API driver est agréable parce qu’elle permet de traiter les kernels CUDA comme des shaders rechargeables à chaud.
      Pouvoir développer en modifiant le code pendant l’exécution est amusant.

  • Sur du bare metal ?