Multiprocesseurs, blocs et threads en streaming (CUDA)

Quelle est la relation entre un cœur CUDA, un multiprocesseur en continu et le modèle CUDA de blocs et de threads?

Qu’est-ce qui est mis en correspondance avec quoi et qu’est-ce qui est parallélisé et comment? et ce qui est le plus efficace, maximiser le nombre de blocs ou le nombre de threads?


Selon ma compréhension actuelle, il existe 8 cuda cores par multiprocesseur. et que chaque kernel cuda pourra exécuter un bloc cuda à la fois. et tous les threads de ce bloc sont exécutés en série dans ce kernel particulier.

Est-ce correct?

La disposition des threads / blocs est décrite en détail dans le guide de programmation CUDA . En particulier, le chapitre 4 indique:

L’architecture CUDA est construite autour d’un ensemble évolutif de multiprocesseurs multithreads (SM). Lorsqu’un programme CUDA sur le processeur hôte appelle une grid du kernel, les blocs de la grid sont énumérés et dissortingbués aux multiprocesseurs avec une capacité d’exécution disponible. Les threads d’un bloc de thread s’exécutent simultanément sur un multiprocesseur et plusieurs blocs de thread peuvent s’exécuter simultanément sur un multiprocesseur. Lorsque les blocs de threads se terminent, de nouveaux blocs sont lancés sur les multiprocesseurs libérés.

Chaque SM contient 8 cœurs CUDA, et à tout moment, ils exécutent une seule chaîne de 32 threads – il faut donc 4 cycles pour émettre une seule instruction pour toute la chaîne. Vous pouvez supposer que les threads d’un war donné s’exécutent en lock-step, mais pour synchroniser les warps, vous devez utiliser __syncthreads() .

Pour la GTX 970, il existe 13 multiprocesseurs de streaming (SM) avec 128 Cuda Cores chacun. Cuda Cores sont également appelés processeurs de stream (SP).

Vous pouvez définir des grids qui mappent des blocs sur le GPU.

Vous pouvez définir des blocs mappant les threads aux processeurs de stream (les 128 Cuda Cores par SM).

Une chaîne est toujours formée de 32 fils et tous les fils d’une chaîne sont exécutés simultanément.

Pour utiliser toute la puissance possible d’un GPU, vous avez besoin de beaucoup plus de threads par SM que le SP possède des SP. Pour chaque capacité de calcul, il existe un certain nombre de threads pouvant résider dans un SM à la fois. Tous les blocs que vous définissez sont en queue et attendent qu’un SM ait les ressources (nombre de SP libres), puis il est chargé. Le SM commence à exécuter des Warps. Comme un Warp ne possède que 32 threads et qu’un SM a par exemple 128 SP, un SM peut exécuter 4 Warps à un moment donné. La chose est que si les threads accèdent à la mémoire, le thread va bloquer jusqu’à ce que sa demande de mémoire soit satisfaite. En chiffres: un calcul arithmétique sur le processeur de service a une latence de 18 à 22 cycles, tandis qu’un access à la mémoire globale non mis en cache peut prendre entre 300 et 400 cycles. Cela signifie que si les threads d’une chaîne attendent des données, seul un sous-ensemble des 128 SP fonctionnerait. Par conséquent, le programmateur bascule pour exécuter une autre chaîne si disponible. Et si cette chaîne bloque, elle exécute la suivante et ainsi de suite. Ce concept est appelé masquage de latence. Le nombre de chaînes et la taille des blocs déterminent l’occupation (à partir du nombre de chaînes que le SM peut choisir d’exécuter). Si l’occupation est élevée, il est plus improbable que les SP ne travaillent pas.

Votre déclaration selon laquelle chaque kernel cuda exécutera un bloc à la fois est incorrecte. Si vous parlez de Streaming Multiprocessors, ils peuvent exécuter des warps de tous les threads résidant dans le SM. Si un bloc a une taille de 256 threads et que votre GPU autorise 2048 threads à résider par SM, chaque SM aura 8 blocs résidant à partir desquels le SM pourra choisir les warps à exécuter. Tous les threads des warps exécutés sont exécutés en parallèle.

Vous trouverez des chiffres pour les différentes capacités de calcul et architectures GPU ici: https://en.wikipedia.org/wiki/CUDA#Limitations

Vous pouvez télécharger une feuille de calcul d’occupation à partir de la feuille de calcul de l’ occupation Nvidia (fournie par Nvidia) .

Compute Work Dissortingbutor planifie un bloc de threads (CTA) sur un SM uniquement si le SM dispose de ressources suffisantes pour le bloc de threads (mémoire partagée, distorsions, registres, barrières, …). Les ressources au niveau du bloc de thread, telles que la mémoire partagée, sont allouées. L’allocation crée suffisamment de chaînes pour tous les threads du bloc de threads. Le gestionnaire de ressources alloue les warp round robin aux sous-partitions SM. Chaque sous-partition SM contient un planificateur de warp, un fichier de registre et des unités d’exécution. Une fois qu’un warp est alloué à une sous-partition, il restra sur la sous-partition jusqu’à ce qu’il soit terminé ou préempté par un changement de contexte (architecture Pascal). Lors du changement de contexte, la restauration du warp sera restaurée sur le même identifiant de warp du même SM.

Lorsque tous les threads de Warp ont terminé, le planificateur warp attend que toutes les instructions émises par le Warp soient terminées, puis le gestionnaire de ressources libère les ressources de niveau Warp qui incluent warp-id et register file.

Lorsque toutes les déformations d’un bloc de thread sont terminées, les ressources de niveau bloc sont libérées et le SM avertit le dissortingbuteur de travaux de calcul que le bloc est terminé.

Une fois qu’une chaîne est allouée à une sous-partition et que toutes les ressources sont allouées, la chaîne est considérée comme active, ce qui signifie que le planificateur de chaîne suit activement l’état de la chaîne. À chaque cycle, le programmateur de chaîne détermine quelles chaînes actives sont bloquées et lesquelles sont éligibles pour émettre une instruction. Le programmateur de chaîne sélectionne la chaîne éligible la plus prioritaire et émet 1 à 2 instructions consécutives à partir du Warp. Les règles de double édition sont spécifiques à chaque architecture. Si une chaîne émet une charge de mémoire, elle peut continuer à exécuter des instructions indépendantes jusqu’à ce qu’elle atteigne une instruction dépendante. La chaîne sera alors signalée jusqu’à ce que le chargement soit terminé. La même chose est vraie pour les instructions mathématiques dépendantes. L’architecture SM est conçue pour masquer la latence de l’ALU et de la mémoire en commutant par cycle entre les chaînes.

Cette réponse n’utilise pas le terme de kernel CUDA car cela introduit un modèle mental incorrect. Les cœurs CUDA sont des unités d’exécution à virgule flottante et à nombre entier en simple précision. Le taux d’émission et la latence des dépendances sont spécifiques à chaque architecture. Chaque sous-partition SM et SM possède d’autres unités d’exécution, y compris des unités de chargement / stockage, des unités à virgule flottante double précision, des unités à virgule flottante demi-précision, des unités de twig, etc.

Afin d’optimiser les performances, le développeur doit comprendre le compromis entre blocs et warps vs. registres / thread.

Le terme «occupation» est le rapport des chaînes actives aux chaînes maximales sur un SM. L’architecture Kepler-Pascal (sauf GP100) possède 4 ordonnanceurs de chaîne par SM. Le nombre minimal de chaînes par SM devrait au moins être égal au nombre de programmateurs de chaîne. Si l’architecture a une latence d’exécution dépendante de 6 cycles (Maxwell et Pascal), il vous faudra au moins 6 warps par planificateur soit 24 par SM (24/64 = 37,5% d’occupation) pour couvrir la latence. Si les threads ont un parallélisme au niveau des instructions, cela peut être réduit. Presque tous les kernelx émettent des instructions de latence variables telles que des charges de mémoire pouvant prendre 80 à 1000 cycles. Cela nécessite des warps plus actifs par planificateur de chaîne pour masquer la latence. Pour chaque kernel, il existe un compromis entre le nombre de chaînes et d’autres ressources, telles que la mémoire partagée ou les registres. Il est donc déconseillé d’optimiser l’occupation à 100% étant donné que d’autres sacrifices sont possibles. Le profileur CUDA peut aider à identifier le taux d’émission d’instruction, l’occupation et les raisons de décrochage afin d’aider le développeur à déterminer cet équilibre.

La taille d’un bloc de fil peut avoir un impact sur les performances. Si le kernel a de gros blocs et utilise des barrières de synchronisation, alors les blocages peuvent être une raison majeure. Cela peut être atténué en réduisant les chaînes par bloc de fil.