Tiny GPU : un GPU minimal implémenté en Verilog

 (github.com/adam-maj)
2 points par GN⁺ 2024-04-27 | 1 commentaires | Partager sur WhatsApp
  • tiny-gpu est une implémentation minimale de GPU en Verilog destinée à apprendre, depuis les bases, comment un GPU fonctionne au niveau matériel ; elle se concentre sur les principes communs aux GPGPU et aux accélérateurs de ML plutôt que sur du matériel dédié au graphisme
  • L’implémentation se compose de moins de 15 fichiers Verilog documentés, de documents sur l’architecture et l’ISA, de kernels d’addition et de multiplication de matrices, ainsi que d’un support pour la simulation de kernels et le traçage d’exécution
  • Le GPU exécute un seul kernel à la fois : il charge la mémoire programme et la mémoire de données, configure thread_count, puis lance le kernel en activant le signal start
  • Pour simplifier, chaque cœur traite un block à la fois ; chaque thread possède son ALU, son LSU, son PC et son fichier de registres, mais on suppose que tous les threads convergent vers le même PC après chaque instruction
  • La plupart des fonctions des GPU modernes — caches multi-niveaux, mémoire partagée, coalescence mémoire, pipeline, ordonnancement des warps, divergence de branchement, barrières — sont exclues afin de privilégier une structure pédagogique

Le problème que tiny-gpu cherche à résoudre

  • Il existe beaucoup de ressources pour apprendre les CPU, de l’architecture jusqu’aux signaux de contrôle, mais les détails techniques bas niveau des GPU modernes restent pour la plupart propriétaires en raison d’un marché très concurrentiel
  • Il existe de nombreuses ressources sur la programmation GPU, mais très peu pour comprendre comment un GPU fonctionne au niveau matériel
  • Les implémentations open source de GPU comme Miaow et VeriGPU visent l’exhaustivité fonctionnelle et l’exécution, ce qui rend leur structure complexe
  • tiny-gpu supprime une grande partie de la complexité d’une carte graphique de production et se concentre sur les éléments essentiels communs aux accélérateurs matériels modernes
    • Les composants importants de l’architecture GPU
    • La manière dont le modèle de programmation SIMD est implémenté dans le matériel
    • La manière dont un GPU gère une bande passante mémoire limitée

Architecture globale

  • tiny-gpu est conçu pour n’exécuter qu’un seul kernel à la fois
  • La procédure d’exécution d’un kernel est la suivante
    • Charger le code du kernel dans la mémoire programme globale
    • Charger les données nécessaires dans la mémoire de données
    • Indiquer le nombre de threads à exécuter dans le registre de contrôle du périphérique
    • Mettre le signal start à high pour lancer l’exécution du kernel
  • Le GPU se compose des unités suivantes
    • Registre de contrôle du périphérique
    • Dispatcher
    • Nombre variable de cœurs de calcul
    • Contrôleurs mémoire pour la mémoire de données et la mémoire programme
    • Cache

Exécution des kernels et répartition des threads

  • Le registre de contrôle du périphérique sert à stocker les métadonnées d’exécution du kernel ; dans tiny-gpu, il ne stocke que thread_count, le nombre total de threads à exécuter
  • Le dispatcher répartit les threads entre plusieurs cœurs de calcul lorsque le kernel démarre
    • Il regroupe les threads exécutables en parallèle en blocks
    • Il envoie les blocks aux cœurs disponibles pour traitement
    • Il signale la fin de l’exécution du kernel une fois tous les blocks traités
  • Un cœur simplifié traite un seul block à la fois
  • Chaque thread possède une ALU, un LSU, un PC et un fichier de registres dédiés
  • La gestion de l’exécution des instructions des threads sur ces ressources est l’un des problèmes difficiles des GPU

Structure mémoire et contrôleurs

  • Le GPU est conçu pour s’interfacer avec une mémoire globale externe ; pour simplifier, la mémoire de données et la mémoire programme sont séparées
  • Spécifications de la mémoire de données
    • Adressage sur 8 bits
    • 256 lignes au total
    • Données sur 8 bits
    • Chaque ligne stocke une valeur inférieure à 256
  • Spécifications de la mémoire programme
    • Adressage sur 8 bits
    • 256 lignes au total
    • Données sur 16 bits
    • Selon l’ISA, chaque instruction occupe 16 bits
  • Le contrôleur mémoire suit les requêtes mémoire provenant des cœurs, limite les requêtes en fonction de la bande passante réelle de la mémoire externe et transmet les réponses aux bonnes ressources
  • Chaque contrôleur mémoire possède un nombre fixe de canaux selon la bande passante de la mémoire globale
  • Le cache est une fonctionnalité en cours de développement : il stocke dans la SRAM du périphérique les données récupérées depuis la mémoire externe, afin de les récupérer plus rapidement lors des requêtes suivantes et de réserver la bande passante mémoire aux nouvelles données

Composition interne des cœurs

  • Chaque cœur dispose d’un scheduler unique qui gère l’exécution des threads
  • Le scheduler de tiny-gpu exécute jusqu’au bout les instructions d’un block, puis récupère un nouveau block, et exécute les instructions de tous les threads dans un ordre synchronisé
  • Des schedulers plus avancés peuvent améliorer l’utilisation des ressources grâce au pipeline et à l’ordonnancement des warps
  • La principale contrainte du scheduler vient de la latence lors du chargement et du stockage de données en mémoire globale
    • La plupart des instructions peuvent être exécutées de manière synchrone
    • Les opérations load-store comme LDR et STR sont asynchrones, il faut donc organiser l’exécution des instructions autour de longues attentes
  • Le Fetcher récupère de manière asynchrone dans la mémoire programme l’instruction correspondant au compteur de programme courant
  • Le Decoder décode l’instruction récupérée en signaux de contrôle pour l’exécution des threads
  • Le fichier de registres de chaque thread conserve les données en cours de calcul et permet le motif SIMD
    • Les registres en lecture seule contiennent %blockIdx, %blockDim et %threadIdx
    • Le kernel peut s’exécuter sur des données différentes selon l’ID local du thread
  • L’ALU de chaque thread traite les instructions arithmétiques ADD, SUB, MUL et DIV
  • CMP indique si le résultat de la différence entre deux registres est négatif, nul ou positif, et stocke le résultat dans le registre NZP de l’unité PC
  • Le LSU de chaque thread accède à la mémoire de données globale et gère LDR, STR ainsi que la latence mémoire asynchrone
  • Le PC de chaque thread détermine la prochaine instruction à exécuter
    • Par défaut, il augmente de 1 à chaque instruction
    • BRnzp branche vers une ligne spécifique de la mémoire programme si la condition du registre NZP définie par le CMP précédent est satisfaite
    • Les boucles et les conditions sont implémentées de cette manière
  • Pour simplifier, tiny-gpu suppose que tous les threads convergent vers le même PC après chaque instruction
  • Dans un vrai GPU, des threads individuels peuvent brancher vers des PC différents ; le groupe de threads traité ensemble se divise alors en plusieurs flux d’exécution, ce qui provoque une divergence de branchement

ISA

  • tiny-gpu implémente une ISA de 11 instructions pour exécuter de simples kernels de preuve de concept, comme l’addition et la multiplication de matrices
  • Instructions prises en charge
    • BRnzp : saute vers une autre ligne de la mémoire programme si la condition NZP est satisfaite
    • CMP : compare deux valeurs de registre et stocke le résultat dans le registre NZP
    • ADD, SUB, MUL, DIV : opérations arithmétiques de base pour les calculs tensoriels
    • LDR : charge des données depuis la mémoire globale
    • STR : stocke des données dans la mémoire globale
    • CONST : charge une valeur constante dans un registre
    • RET : signale la fin d’exécution du thread courant
  • Chaque registre est désigné sur 4 bits, pour un total de 16 registres
    • Les 13 registres de R0 à R12 sont des registres libres en lecture-écriture
    • Les 3 derniers sont des registres spéciaux en lecture seule qui fournissent %blockIdx, %blockDim et %threadIdx, nécessaires au SIMD

Flux d’exécution

  • Lorsqu’il exécute une instruction, chaque cœur suit le flux de contrôle suivant
    • FETCH : récupérer l’instruction suivante du PC courant
    • DECODE : décoder l’instruction en signaux de contrôle
    • REQUEST : si LDR ou STR est nécessaire, demander les données à la mémoire globale
    • WAIT : si nécessaire, attendre la réponse de la mémoire globale
    • EXECUTE : effectuer le calcul sur les données
    • UPDATE : mettre à jour le fichier de registres et le registre NZP
  • Ce flux de contrôle est conçu pour la simplicité et la compréhension
  • Dans une implémentation réelle, certaines étapes peuvent être fusionnées afin d’optimiser le temps de traitement, ou bien un pipeline peut coordonner l’exécution de plusieurs instructions sur les ressources du cœur
  • Chaque thread suit le même chemin d’exécution pour calculer sur les données de son fichier de registres dédié
  • La structure ressemble à un diagramme de CPU, mais se distingue par la présence de %blockIdx, %blockDim et %threadIdx dans des registres en lecture seule, ce qui rend le SIMD possible

Exemples de kernels

  • Des kernels d’addition et de multiplication de matrices ont été écrits comme preuves de concept de l’ISA
  • Les fichiers de test du dépôt peuvent simuler entièrement ces kernels sur le GPU, et générer l’état de la mémoire de données ainsi qu’une trace complète d’exécution

  • Addition de matrices

    • matadd.asm additionne deux matrices 1 x 8
    • Les 8 additions élément par élément sont chacune effectuées dans un thread séparé
    • Les registres %blockIdx, %blockDim et %threadIdx démontrent la programmation SIMD
    • Les instructions LDR et STR incluent la gestion mémoire asynchrone

  • Multiplication de matrices

    • matmul.asm multiplie deux matrices 2 x 2
    • Il calcule élément par élément le produit scalaire des lignes et colonnes concernées
    • CMP et BRnzp démontrent les branchements à l’intérieur d’un thread
    • Comme tous les branchements reconvergent, cela fonctionne avec l’implémentation actuelle de tiny-gpu

Simulation

  • Pour exécuter la simulation d’un kernel, iverilog et cocotb sont nécessaires
  • Procédure de préparation
    • Installer le compilateur Verilog et cocotb avec brew install icarus-verilog et pip3 install cocotb
    • Télécharger la dernière version de sv2v, décompresser l’archive et ajouter le binaire à $PATH
    • Exécuter mkdir build à la racine du dépôt
  • Les simulations de kernels s’exécutent avec make test_matadd et make test_matmul
  • Les résultats d’exécution sont écrits sous forme de fichiers de logs dans test/logs
    • État initial de la mémoire de données
    • Trace complète d’exécution du kernel
    • État final de la mémoire de données
  • Au début de chaque fichier de log, on voit les matrices d’entrée ; à la fin, la mémoire de données finale contient la matrice résultat
  • La trace d’exécution inclut, à chaque cycle, l’état d’exécution de tous les threads de tous les cœurs
    • Instruction courante
    • PC
    • Valeurs des registres
    • Informations d’état

Fonctionnalités GPU avancées volontairement absentes

  • Pour simplifier, tiny-gpu exclut la plupart des améliorations de performance et de fonctionnalités des GPU modernes

  • Caches multi-niveaux et mémoire partagée

    • Les GPU modernes utilisent plusieurs niveaux de cache afin de réduire les accès à la mémoire globale
    • tiny-gpu n’implémente qu’un seul niveau de cache entre la ressource demandeuse et le contrôleur mémoire, pour stocker les données récentes
    • Les caches multi-niveaux réduisent le temps de chargement en mettant en cache les données fréquemment utilisées plus près de leur lieu d’utilisation
    • Les GPU utilisent aussi parfois de la mémoire partagée afin que les threads d’un même block puissent échanger des résultats intermédiaires

  • Coalescence mémoire

    • Plusieurs threads exécutés en parallèle accèdent souvent à des adresses contiguës, par exemple à des éléments adjacents d’une matrice
    • La coalescence mémoire analyse les requêtes mémoire en file d’attente et fusionne les requêtes adjacentes en une seule transaction
    • L’objectif est de réduire le temps consacré à l’adressage et de traiter les requêtes ensemble

  • Pipeline

    • Les cœurs de tiny-gpu ne commencent l’instruction suivante qu’une fois terminée l’exécution d’une instruction pour un groupe de threads
    • Les GPU modernes streament l’exécution de plusieurs instructions séquentielles tout en garantissant l’exécution dans l’ordre pour les instructions dépendantes
    • Cela améliore l’utilisation des ressources afin d’éviter que les ressources du cœur restent inactives dans des situations comme l’attente de requêtes mémoire asynchrones

  • Ordonnancement des warps

    • Un block est divisé en warps, des lots de threads pouvant être exécutés ensemble
    • Lorsqu’un warp est en attente, les instructions d’un autre warp sont exécutées, ce qui permet de traiter plusieurs warps simultanément sur un même cœur
    • C’est similaire au pipeline, mais appliqué aux instructions de threads différents

  • Divergence de branchement

    • tiny-gpu suppose que tous les threads d’un même lot se trouvent au même PC après chaque instruction
    • En pratique, selon les données, des threads individuels peuvent brancher vers des lignes différentes
    • Les threads ayant des PC différents sont séparés en flux d’exécution distincts, et le moment où ils reconvergent doit également être géré

  • Synchronisation et barrières

    • Les GPU modernes peuvent définir des barrières pour faire attendre les groupes de threads d’un même block jusqu’à ce que tous aient atteint un point donné
    • C’est utile lorsque les threads doivent échanger des données partagées, afin de garantir que le traitement des données est terminé

Prochaines étapes

  • Les améliorations envisagées sont les suivantes
    • Ajouter un cache d’instructions simple
    • Construire un adaptateur permettant d’utiliser le GPU dans Tiny Tapeout 7
    • Ajouter une divergence de branchement de base
    • Ajouter une coalescence mémoire de base
    • Ajouter un pipeline de base
    • Optimiser le flux de contrôle et l’utilisation des registres afin d’améliorer le temps de cycle
    • Écrire un kernel graphique de base ou ajouter un matériel graphique simple pour démontrer les capacités graphiques
  • Les utilisateurs qui souhaitent améliorer le dépôt peuvent contribuer via des PR

À lire aussi

1 commentaires

 
GN⁺ 2024-04-27
Commentaires Hacker News
  • Le marché des GPU est tellement concurrentiel que les détails techniques de bas niveau des architectures modernes restent pour la plupart non publics
    Intel fait figure d’exception en publiant beaucoup de documentation technique sur ses GPU : https://kiwitree.net/~lina/intel-gfx-docs/prm/
    On peut aussi trouver en ligne les manuels des i810/815, et à part l’étrange trou avant les 965, où les 855/910/915/945 manquent, la documentation a été assez régulièrement maintenue
    • AMD publie aussi pas mal de documentation : https://www.amd.com/en/developer/browse-by-resource-type/documentation.html
      Cela inclut même la documentation sur l’architecture du jeu d’instructions pour des produits actuels et anciens, mais cela semble davantage destiné aux implémenteurs qu’à des explications de haut niveau pour passionnés
    • Les drivers Linux d’Intel sont aussi de bonne qualité et intégrés au mainline
      J’aimerais que toutes les entreprises suivent cette approche
    • C’est une ressource de 2018, mais elle reste en partie pertinente : The Thirty Million Line Problem - Casey Muratori
  • Projet vraiment génial, et c’est agréable de voir ce type de projet matériel se développer publiquement
    Cela dit, je vois ça davantage comme un coprocesseur SIMD
    Pour l’appeler un GPU, je pense qu’il devrait au moins avoir une forme quelconque de sortie d’affichage
    Je sais bien que le terme est devenu assez souple récemment, avec Nvidia et d’autres qui vendent aussi comme GPU des variantes d’architectures graphiques réservées aux serveurs, mais dans la conception d’un GPU, la partie graphique représente encore aujourd’hui une part importante de la complexité
    • S’il traite des graphismes, je pense qu’on peut le considérer comme un GPU même sans sortie
      Un GPU qui n’affiche rien peut rester utile
      À mon travail, il y a environ 75 stations de travail équipées de Quadro milieu de gamme, mais les cartes n’ont que du mini-DisplayPort et l’entreprise n’achète que des câbles HDMI, donc elles sont toutes branchées sur la partie graphique intégrée
      Malgré ça, ces cartes accélèrent toujours les logiciels et traitent les graphismes ; elles ne font simplement pas l’affichage à l’écran
  • Très bien. Je soutiens pleinement les travaux sur les GPU open core
    Il existe aussi d’autres exemples : https://github.com/jbush001/NyuziProcessor
    • Ce serait bien d’avoir une implémentation minimale de CUDA sur l’un de ces processeurs open core
      Quel volume faudrait-il pour faire fabriquer économiquement ce type de processeur chez TSMC ou une autre fonderie ?
  • Projet vraiment remarquable
    J’aimerais me mettre aux FPGA, mais honnêtement il est difficile de savoir par où commencer, et l’ensemble du domaine paraît assez intimidant
    Mon objectif final serait de créer une carte d’accélération pour les LLM ; même si c’est un objectif fixé un peu arbitrairement, j’ai l’impression qu’il y aurait beaucoup de recoupements avec ce projet, avec peut-être surtout la partie offload mémoire qui différerait pour charger des modèles plus grands
    • Il faut changer de cadre mental
      Débuter avec les FPGA doit être découpé en plusieurs sous-compétences, et il faut aussi ajuster ses attentes
      On n’attendrait pas d’un ingénieur logiciel qu’il construise d’emblée un ordinateur complet à partir des principes de base, écrive une architecture de jeu d’instructions, comprenne le code machine, le transforme en assembleur, puis développe encore un langage de programmation pour créer des applications en Python
      Il faut commencer par le haut et descendre dans la pile
      Si on abstrait la complexité et qu’on se concentre sur la construction de systèmes avec des IP déjà prêtes, la conception FPGA est en réalité assez accessible
      Je recommande souvent quelque chose comme MATLAB, car avec une DevKit disposant d’un design de référence, on peut créer une première application avec HDL Coder
      Sinon, on se retrouve avec l’énorme charge d’apprendre l’architecture du calcul numérique, Verilog, le timing, les transceivers/entrées-sorties, le pin planning, Quartus/Vivado, la simulation/vérification, les systèmes embarqués, etc.
      En résumé, il faut commencer par la conception au niveau système, apprendre à prendre des IP plug-and-play et à les assembler au niveau supérieur, puis insérer ces modules dans un design de référence déjà prêt
      Ensuite, on peut enlever progressivement les couches pour faire apparaître la complexité sous-jacente
    • Je suis dans la même situation, et voici mon plan
      1. Lire Harris, Harris, Digital Design and Computer Architecture. (2022). Elsevier : https://doi.org/10.1016/c2019-0-00213-0
      2. Suivre le cours RVFpga de l’auteur pour construire un vrai CPU RISC-V sur FPGA : https://www.youtube.com/watch?v=ePv3xD3ZmnY
    • Je recommande ce parcours
      1. Cloner le dépôt pédagogique https://github.com/yuri-panchul/basics-graphics-music ; c’est une collection de petits exercices pour les personnes qui apprennent Verilog depuis le début, rédigée par Yuri Panchul, qui a travaillé chez Imagination pour le développement de GPU
      2. Se procurer l’une des dizaines de cartes FPGA prises en charge ainsi que des accessoires comme des boutons et des LED
      3. Installer Yosys et les outils associés
      4. Commencer par lab01 DeMorgan et faire autant d’exercices du dépôt que possible
        On peut faire les exercices en parallèle de la lecture de Harris&Harris
        Une fois les exercices et le livre terminés, ce sera le moment de démarrer son propre projet
        À noter qu’il y a aussi des rencontres hebdomadaires chez HackerMojo, auxquelles on peut participer via Zoom même sans être dans la Valley
    • Je ne sais pas à quel stade tu en es, mais ces ressources m’ont aidé à mieux comprendre la logique numérique et les architectures CPU/GPU
      1. https://learn.saylor.org/course/CS301
      2. https://www.coursera.org/learn/comparch
  1. https://hdlbits.01xz.net/wiki/Main_Page
  • Si vous voulez accélérer les LLM, il faut d’abord comprendre l’architecture
    C’est par là qu’il faut commencer
    Le matériel est en fait la partie facile, et aussi la partie difficile du point de vue de la fabrication
  • Y a-t-il une raison de mélanger les opérateurs d’assignation non bloquante et bloquante dans ce bloc always séquentiel ?
    • Ça ressemble à une variable locale
    • Si on ne s’obsède pas trop sur la correspondance entre simulation et synthèse, ça passe
  • Il y a longtemps, j’ai fait quelque chose de similaire en VHDL
    Il y avait un site appelé opencores qui rassemblait plusieurs projets HDL open source
    Je me demande s’il existe aujourd’hui de bons simulateurs HDL distribués à grande échelle, de niveau HPC
    Utiliser des GPU modernes pour la simulation au niveau RTL semble pertinent
  • L’ALU implémente vraiment une instruction DIV directement au niveau matériel ?
    Sur des CUDA cores modernes, est-ce courant d’avoir une division comme véritable instruction, ou bien est-ce généralement émulé en logiciel ?
    Un vrai circuit matériel de division prend énormément de place, donc je ne m’attendais pas à le voir dans l’ALU d’un GPU
    En Verilog, écrire une seule ligne comme DIV: begin alu_out_reg <= rs / rt; end est trop facile, mais cette seule ligne consomme beaucoup de silicium
    Si on se contente de simuler le Verilog, on risque de ne pas s’en rendre compte
    • Ce n’est qu’un projet pour apprendre le Verilog
      Le projet s’arrête à la simulation, et il faudrait beaucoup plus de travail pour en faire un vrai matériel
  • C’est encore un « GPU » sans capacités graphiques
    Personnellement, je pense qu’il faudrait appeler ça autrement
    • La première question, c’est déjà pourquoi CPU et GPU ont été séparés au départ
      L’écart entre les deux se réduit et chacun ajoute des fonctions de l’autre, mais il reste encore une différence importante
      À mon avis, cela a à voir avec la loi d’Amdahl [0]
      Dans ce sens, on peut appeler le CPU un processeur optimisé pour la latence, et le GPU un processeur optimisé pour le débit
      Plus précisément, on pourrait aussi appeler [1] le CPU un processeur à dépendances de données longues et profondes, et le GPU un processeur à dépendances de données larges et plates
      [0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
      [1]: https://en.wikipedia.org/wiki/Data_dependency
    • On pourrait appeler ça un TPU, c’est-à-dire une unité de traitement tensoriel
      Un tenseur n’est qu’un tableau à n dimensions
      On peut ajouter du logiciel ou du firmware par-dessus pour le faire fonctionner comme un GPU
    • J’ai déjà envisagé de lancer un projet de « display adapter », mais je me suis retrouvé bloqué avant même de commencer, incapable de comprendre le protocole de communication entre le driver GOP de l’UEFI et l’adaptateur d’affichage
      J’ai essayé de reconstituer les morceaux à partir du code source d’EDK2, mais il n’est pas clair dans quelle mesure cela est spécifique à QEMU
    • On peut simplement appeler ça un MPU, c’est-à-dire une unité de traitement matriciel
    • Le terme qui semble s’imposer est AIA, c’est-à-dire accélérateur IA
  • L’hypothèse de tiny-gpu selon laquelle tous les threads « convergent » vers le même compteur ordinal après chaque instruction est une simplification beaucoup trop naïve
    Sur un vrai GPU, des threads individuels peuvent bifurquer vers des PC différents, et il se produit une divergence de branchement où un groupe de threads d’abord traité ensemble se sépare en exécutions distinctes
    Il aurait sans doute mieux valu commencer par faire de la programmation GPU avant de fabriquer un GPU en silicium
    Et en plus, ça ne correspond pas vraiment à du SIMD
    C’est la même personne qui avait déjà dit avoir fabriqué un CPU en assemblant les circuits d’autres gens pour faire clignoter une LED
    • Le premier point, ce n’est pas en gros comme appeler __syncthreads() à chaque exécution ?