Multiprocesseurs, blocs et Threads en Streaming (CUDA)
Quelle est la relation entre un noyau CUDA, un multiprocesseur en streaming et le modèle CUDA de blocs et de threads?
Qu'est-ce qui est mappé à quoi et qu'est-ce qui est parallélisé et comment? et quoi de plus efficace, de maximiser le nombre de blocs ou le nombre de threads?
Ma compréhension actuelle est qu'il y a 8 cœurs cuda par multiprocesseur. et que chaque noyau cuda sera capable d'exécuter un bloc cuda à la fois. et tous les threads de ce bloc sont exécutés en série dans ce noyau particulier.
Est-ce correct?
3 réponses
La disposition des threads / blocs est décrite en détail dans le guide de programmation CUDA . En particulier, le chapitre 4 stipule:
L'architecture CUDA est construite autour d'un tableau évolutif de multiprocesseurs de Streaming multithread (SMs). Lorsqu'un programme CUDA sur le processeur hôte appelle une grille du noyau, les blocs de la grille sont énumérés et distribués aux multiprocesseurs avec une capacité d'exécution disponible. Les threads d'un bloc de threads 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 thread 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 chaîne unique de 32 threads - il faut donc 4 cycles d'horloge pour émettre une seule instruction pour toute la chaîne. Vous pouvez supposer que les threads dans n'importe quel warp donné s'exécutent en lock-step, mais pour synchroniser entre les chaînes, vous devez utiliser __syncthreads().
Pour la GTX 970, il existe 13 multiprocesseurs de Streaming (SM) avec 128 cœurs Cuda chacun. Les cœurs Cuda sont également appelés processeurs de flux (SP).
Vous pouvez définir des grilles qui mappent des blocs au GPU.
Vous pouvez définir des blocs qui mappent les threads aux processeurs de flux (les 128 cœurs Cuda 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 SM A SPs. Pour chaque capacité de calcul, il existe un certain nombre de threads qui peuvent résider dans un SM à la fois. Tous les blocs que vous définissez sont mis en file d'attente et attendent QU'un SM ait les ressources (nombre de SPs gratuit), puis il est chargé. Le SM commence à exécuter des déformations. Comme une chaîne n'a que 32 Threads et QU'un SM A par exemple 128 SPs, un SM peut exécuter 4 Warps à un moment donné. La chose est que si les threads accèdent à la mémoire, le thread bloquera jusqu'à ce que sa demande de mémoire soit satisfaite. Dans nombres: un calcul arithmétique sur le SP a une latence de 18-22 cycles tandis qu'un accès mémoire global non mis en cache peut prendre jusqu'à 300-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 planificateur commute pour exécuter un autre warp si disponible. Et si cette chaîne bloque, elle exécute la suivante et ainsi de suite. Ce concept est appelé cache de latence. Le nombre de déformations et la taille du bloc déterminent l'occupation (à partir du nombre de déformations SM peut choisir d'exécuter). Si l'occupation est élevée, il est plus improbable qu'il n'y ait pas de travail pour le SPs.
Votre déclaration selon laquelle chaque noyau cuda exécutera un bloc à la fois est fausse. Si vous parlez de multiprocesseurs en Streaming, ils peuvent exécuter des chaînes à partir de tous les threads qui résident 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 aurait 8 blocs résidant à partir desquels le SM peut choisir des chaînes à exécuter. Tous les threads de l' les chaînes exécutées sont exécutées en parallèle.
Vous trouverez ici des numéros pour les différentes capacités de calcul et Architectures GPU: https://en.wikipedia.org/wiki/CUDA#Limitations
Vous pouvez télécharger une feuille de calcul d'occupation à partir de NVIDIA feuille de calcul D'occupation (par Nvidia) .
Le distributeur de travail de calcul planifie un bloc de thread (CTA) sur un SM uniquement si le SM dispose de ressources suffisantes pour le bloc de thread (mémoire partagée, warps, registres, barrières, ...). Les ressources de niveau de bloc de Thread telles que la mémoire partagée sont allouées. L'allocation crée des chaînes suffisantes pour tous les threads du bloc thread. Le gestionnaire de ressources alloue le Round robin des chaînes aux sous-partitions SM. Chaque sous-partie SM contient un planificateur warp, un fichier de registre et des unités d'exécution. Une fois une distorsion est alloué à une sous-partie, il restera sur la sous-partie jusqu'à ce qu'elle soit terminée ou préemptée par un commutateur de contexte (architecture Pascal). Lors de la restauration du commutateur de contexte, le warp sera restauré sur le même SM même warp-id.
Lorsque tous les threads de warp sont terminés, le planificateur warp attend que toutes les instructions en suspens é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 dans un bloc de thread complet, les ressources au niveau du bloc sont libérées et le SM informe le distributeur de travail de calcul que le bloc est terminé.
Une fois qu'une chaîne est allouée à une sous-partie 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. Sur chaque cycle, le planificateur warp détermine quelles chaînes actives sont bloquées et lesquelles sont éligibles pour émettre une instruction. Le planificateur warp choisit le plus élevé distorsion admissible prioritaire et émet 1-2 instructions consécutives de la distorsion. Les règles pour Dual-issue sont spécifiques à chaque architecture. Si un warp émet une charge de mémoire, il peut continuer à exécuter des instructions indépendantes jusqu'à ce qu'il atteigne une instruction dépendante. La chaîne signalera ensuite calé jusqu'à la fin de la charge. La même chose est vraie pour les instructions mathématiques dépendantes. L'architecture SM est conçue pour masquer à la fois la latence ALU et la latence de la mémoire en basculant par cycle entre les warps.
Ce la réponse n'utilise pas le terme Cuda core car cela introduit un modèle mental incorrect. Les cœurs CUDA sont des unités d'exécution à virgule flottante/entière de précision unique pipelined. Le taux d'émission et la latence de dépendance sont spécifiques à chaque architecture. Chaque sous-partie SM et SM A d'autres unités d'exécution, y compris les unités de charge/magasin, les unités à virgule flottante à double précision, les unités à virgule flottante à demi-précision, les unités de branche, etc.
Afin de maximiser les performances, le développeur doit comprendre compromis entre les blocs et les chaînes et les registres/threads.
Le terme occupation est le rapport entre les déformations actives et les déformations maximales sur un SM. L'architecture Kepler-Pascal (sauf GP100) a 4 programmateurs warp par SM. Le nombre minimal de chaînes par SM doit au moins être égal au nombre d'ordonnanceurs warp. Si l'architecture a une latence d'exécution dépendante de 6 cycles (Maxwell et Pascal), vous aurez besoin d'au moins 6 warps par planificateur, soit 24 par SM (24 / 64 = 37,5% d'occupation) pour couvrir le temps de latence. Si les threads ont un parallélisme de niveau d'instruction, cela pourrait être réduit. Presque tous les noyaux émettent des instructions de latence variable telles que des charges de mémoire pouvant prendre de 80 à 1000 cycles. Cela nécessite des warps plus actifs par planificateur warp pour masquer la latence. Pour chaque noyau, il y a un compromis entre le nombre de warps et d'autres ressources telles que la mémoire partagée ou les registres, donc l'optimisation pour une occupation à 100% n'est pas conseillée car d'autres sacrifices seront probablement faits. Le CUDA profiler peut aider à identifier le taux de problème 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 thread peut avoir un impact sur les performances. Si le noyau a de grands blocs et utilise des barrières de synchronisation, les stalles de barrière peuvent être des raisons de décrochage. Cela peut être atténué en réduisant les déformations par bloc de fil.