Diffusion multiprocesseurs, Les blocs et les fils (CUDA)

Question

Quelle est la relation entre un noyau CUDA, un multiprocesseur flux et le modèle CUDA de blocs et de fils?

Ce qui est mis en correspondance à quoi et ce qui est parallélisé et comment? et ce qui est plus efficace, de maximiser le nombre de blocs ou le nombre de threads?

---

Ma compréhension actuelle est qu'il ya 8 cœurs de cuda par multiprocesseur. et que chaque noyau de cuda sera en mesure d'exécuter un bloc de cuda à la fois. et tous les fils de ce bloc sont exécutées en série dans ce noyau particulier.

Est-ce correct?

Answer 1

La disposition fil / bloc est décrit en détail dans la section Guide de programmation CUDA . En particulier, le chapitre 4 états:

  

L'architecture CUDA est construit autour d'une gamme évolutive de streaming multiprocesseurs (SM) multithread. Lorsqu'un programme CUDA sur la CPU hôte invoque une grille du noyau, les blocs de la grille sont dénombrés et distribués aux multiprocesseurs avec une capacité d'exécution disponible. Les fils d'un bloc de fil d'exécuter simultanément sur une multiprocesseur, et plusieurs blocs de filetage peuvent exécuter simultanément sur une multiprocesseur. Comme les blocs de fil se terminent, de nouveaux blocs sont lancés sur les multiprocesseurs laissés vacants.

Chaque SM contient 8 cœurs CUDA, et à la fois ils sont l'exécution d'une seule chaîne de 32 fils - de sorte qu'il prend 4 cycles d'horloge pour délivrer une instruction unique pour la chaîne entière. Vous pouvez supposer que les discussions dans une chaîne donnée exécuter en lock-step, mais à travers Synchronize funes, vous devez utiliser __syncthreads().

Answer 2

Pour la GTX 970 il y a 13 streaming multiprocesseurs (SM) avec 128 Cuda Cores chacun. Cuda Cores sont aussi appelés Stream Processors (SP).

Vous pouvez définir des grilles qui cartes blocs au GPU.

Vous pouvez définir des blocs qui tracent des fils de Stream Processors (128 Cuda Cores par SM).

Une chaîne est toujours formé par des fils 32 et tous les fils de chaîne sont exécutées une simulaneously.

Pour utiliser toute la puissance possible d'un GPU dont vous avez besoin beaucoup plus de fils par SM que le SM a SPs. Pour chaque capacité Compute il y a un certain nombre de fils qui peuvent résider dans une SM à la fois. Tous les blocs que vous définissez sont mis en attente et attendre SM d'avoir les ressources (nombre de SPs gratuit), il est chargé. Le SM commence à exécuter Déformations. Depuis une chaîne a seulement 32 fils et un SM a par exemple 128 un SM SPs peut exécuter 4 Déformations à un moment donné. La chose est si les threads ne l'accès mémoire le fil bloque jusqu'à ce que sa demande de mémoire est satisfaite. En chiffres: Un calcul arithmétique sur la SP a une latence de cycles tandis qu'un 18-22 accès à la mémoire globale non mis en cache peut prendre jusqu'à 300-400 cycles. Cela signifie que si les fils d'une chaîne sont en attente de données qu'un sous-ensemble des 128 SPs travailleraient. À cet effet le planificateur bascule pour exécuter une autre chaîne si elle est disponible. Et si cela bloque de chaîne exécute la suivante et ainsi de suite. Ce concept est appelé latence de dissimulation. Le nombre de funes et la taille du bloc déterminent l'occupation (de combien de SM le funes peut choisir d'exécuter). Si l'occupation est élevé, il est plus improbable qu'il n'y a pas de travail pour les SPs.

Votre déclaration que chaque noyau de cuda exécutera un bloc à la fois est erroné. Si vous parlez de streaming multiprocesseurs ils peuvent exécuter funes de tous les fils qui résident dans le SM. Si un bloc a une taille de 256 fils et vos allowes GPU 2048 fils à SM résident par chaque SM aurait 8 blocs résidant à partir de laquelle le SM peut choisir funes à exécuter. Tous les fils des chaînes exécutées sont exécutées en parallèle.

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

Vous pouvez télécharger une feuille de calcul d'occupation de Nvidia Occupation feuille de calcul (par Nvidia) .

Answer 3

Le distributeur travail Compute planifiera un bloc de fil (CTA) sur une SM que si le SM dispose de ressources suffisantes pour le bloc de fil (mémoire partagée, funes, registres, barrières, ...). ressources au niveau de blocs de threads comme mémoire partagée sont attribués. L'allouer crée chaînes suffisantes pour tous les fils dans le bloc de filetage. Le gestionnaire de ressources en utilisant round robin aux sous-partitions SM. Chaque sous-partition de SM contient un ordonnanceur de fils de chaîne, le fichier de registre et les unités d'exécution. Une fois qu'une chaîne est affectée à un sous-partition, il restera sur la sous-partition jusqu'à ce qu'elle se termine ou est préempté par un changement de contexte (architecture Pascal). Le changement de contexte restaurer la chaîne sera restaurée à la même SM-même chaîne id.

Lorsque tous les fils en chaîne ont terminé le temps d'attente du planificateur de chaîne pour toutes les instructions émis par la chaîne à remplir et puis le gestionnaire de ressources libère les ressources au niveau de la chaîne qui comprennent la chaîne-id et fichier de registres.

Lorsque tous les funes dans un bloc de fil ressources au niveau des blocs complets sont ensuite libérés et les SM informe le Distributeur travail Compute que le bloc a terminé.

Une fois la chaîne est affectée à un sous-partition et toutes les ressources sont attribuées la chaîne est considérée comme active dans le sens que le planificateur de chaîne est suivi activement l'état de la chaîne. Sur chaque cycle la chaîne planificateur déterminer quels actifs funes sont bloquées et qui sont admissibles à émettre une instruction. Le planificateur de chaîne prend la plus haute priorité chaîne de droit et les questions 1-2 instructions consécutives de la chaîne. Les règles de double question sont propres à chaque architecture. Si une chaîne émet une charge de la mémoire, il peut continuer à des instructions indépendantes exécutées jusqu'à ce qu'il atteigne une instruction à charge. La chaîne fera rapport au point mort jusqu'à ce que finalise la charge. La même chose est vraie pour les instructions de mathématiques à charge. L'architecture de SM est conçu pour masquer la fois ALU et la latence de mémoire en passant par cycle entre les fils de chaîne.

Cette réponse ne pas utiliser le terme noyau CUDA comme il introduit un modèle mental incorrect. cœurs CUDA sont des unités virgule flottante simple précision en pipeline / exécution entiers. Le taux d'émission et la latence dépendance est spécifique à chaque architecture. Chaque SM et SM a sous-partition autres unités d'exécution, y compris les unités de chargement / stockage, les unités double précision en virgule flottante, la moitié des unités de précision en virgule flottante, des unités de dérivation, etc.

Afin de maximiser la performance du développeur doit comprendre le commerce hors des blocs par rapport funes par rapport aux registres / fil.

Le taux d'occupation du terme est le rapport des fils de chaîne actives pour fils de chaîne maximale sur un SM. Kepler - l'architecture Pascal (sauf GP100) ont 4 ordonnanceurs de chaîne par SM. Le nombre minimal de fils de chaîne par SM doit au moins être égal au nombre d'ordonnanceurs de chaîne. Si l'architecture a un temps d'attente d'exécution en fonction de 6 cycles (Maxwell et Pascal), alors vous auriez besoin d'au moins 6 par funes programmateur qui est 24 par SM (24/64 = 37,5% d'occupation) pour couvrir le temps d'attente. Si les fils ont le parallélisme de niveau d'instruction, cela pourrait être réduit. Presque tous les grains émettent des instructions de latence variables telles que les charges de mémoire qui peuvent prendre 80-1000 cycles. Cela nécessite plus funes actifs par programmateur de chaîne à la latence du cache. Pour chaque noyau il y a un compromis entre le nombre de points et d'autres ressources funes telles que la mémoire partagée ou registres afin d'optimiser les occupation de 100% ne sont pas informés que sera probablement fait un autre sacrifice. Le profileur CUDA peut aider à identifier les taux d'émission de l'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 influer sur les performances. Si le noyau a de grands blocs et utilise des barrières de synchronisation, puis les stalles de barrière peut être une des raisons de venir décrochage. Cela peut être atténué en réduisant lafunes par bloc de fil.

Answer 4

Il y a multiflux multiprocesseur sur un seul appareil.
Un SM peut contenir plusieurs blocs. Chaque bloc peut contenir plusieurs threads.
Un SM ont plusieurs cœurs CUDA (en tant que développeur, vous ne devriez pas se soucier de cela parce qu'il est extrait par la chaîne), qui travaillera sur fil. SM travaille toujours sur la chaîne de fils (toujours 32). Une chaîne ne travaillant sur un fil de même bloc.
SM et de bloquer les deux ont des limites sur le nombre de fils, le nombre de registre et la mémoire partagée.