111 votes

Streaming multiprocesseurs, Blocs et Threads (CUDA)

Quelle est la relation entre un cœur CUDA, un multiprocesseur de streaming et le modèle CUDA de blocs et de fils?

Qu'est-ce qui est associé à quoi, et qu'est-ce qui est parallélisé et comment? Et qu'est-ce qui est plus efficace, maximiser le nombre de blocs ou le nombre de fils?


Ma compréhension actuelle est qu'il y a 8 cœurs CUDA par multiprocesseur et que chaque cœur CUDA pourra exécuter un bloc CUDA à la fois, et que tous les fils de ce bloc sont exécutés séquentiellement dans ce cœur particulier.

Est-ce correct?

79voto

Edric Points 11387

La mise en page du fil / bloc 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 tableau évolutif de multiprocesseurs de flux multithreadés (SM). Lorsqu'un programme CUDA sur l'UC hôte appelle une grille de noyaux, les blocs de la grille sont énumérés et distribués aux multiprocesseurs ayant 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 threads 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 guerre de 32 threads - il faut donc 4 cycles d'horloge pour émettre une seule instruction pour toute la guerre. Vous pouvez supposer que les threads dans une guerre donnée s'exécutent de manière synchrone, mais pour synchroniser à travers les guerres, vous devez utiliser __syncthreads().

55voto

JoeFox Points 617

Pour la GTX 970, il y a 13 multiprocesseurs de flux (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 cartographient des blocs vers le GPU.

Vous pouvez définir des blocs qui cartographient des threads vers les processeurs de flux (les 128 cœurs Cuda par SM).

Une warp est toujours formée de 32 threads et tous les threads d'une warp 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 de SP. Pour chaque capacité de calcul, il y a 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 dispose des ressources (nombre de SP libres), puis il est chargé. Le SM commence à exécuter des Warps. Puisqu'une Warp n'a que 32 Threads et qu'un SM a par exemple 128 SP, un SM peut exécuter 4 Warps en même temps. Le problème est que si les threads accèdent à la mémoire, le thread sera bloqué jusqu'à ce que sa demande mémoire soit satisfaite. En chiffres : Un calcul arithmétique sur le SP a une latence de 18-22 cycles alors qu'un accès à la mémoire globale non mis en cache peut prendre jusqu'à 300-400 cycles. Cela signifie que si les threads d'une warp attendent des données, seul un sous-ensemble des 128 SP travaillerait. Par conséquent, l'ordonnanceur passe à exécuter une autre warp si disponible. Et si cette warp bloque, il exécutera la suivante et ainsi de suite. Ce concept s'appelle masquage de la latence. Le nombre de warps et la taille du bloc déterminent l'occupation (à partir de combien de warps le SM peut choisir d'exécuter). Si l'occupation est élevée, il est moins probable qu'il n'y ait pas de travail pour les SP.

Votre affirmation selon laquelle chaque cœur Cuda exécutera un bloc à la fois est incorrecte. Si vous parlez des multiprocesseurs de flux, 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 permet à 2048 threads de résider par SM, chaque SM aurait 8 blocs résidents à partir desquels le SM peut choisir les warps à exécuter. Tous les threads des warps exécutées 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 auprès de Nvidia Feuille de calcul d'occupation (par Nvidia).

8voto

Greg Smith Points 5397

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

Lorsque tous les threads dans un warp sont terminés, le planificateur de warp attend que toutes les instructions en attente émises par le warp soient terminées, puis le gestionnaire de ressources libère les ressources au niveau du warp, y compris l'identifiant warp et le fichier de registres.

Lorsque tous les warps dans un bloc de threads sont terminés, alors les ressources au niveau du bloc sont libérées et le SM notifie le Compute Work Distributor que le bloc est terminé.

Une fois qu'un warp est alloué à une sous-partition et que toutes les ressources sont allouées, le warp est considéré comme actif, ce qui signifie que le planificateur de warp suit activement l'état du warp. À chaque cycle, le planificateur de warp détermine quels warps actifs sont bloqués et lesquels sont éligibles pour émettre une instruction. Le planificateur de warp choisit le warp éligible de priorité la plus élevée et émet 1 à 2 instructions consécutives à partir du warp. Les règles de double émission sont spécifiques à chaque architecture. Si un warp émet une lecture mémoire, il peut continuer à exécuter des instructions indépendantes jusqu'à ce qu'il atteigne une instruction dépendante. Le warp signalera alors une attente jusqu'à ce que la lecture soit terminée. Il en va de même pour les instructions mathématiques dépendantes. L'architecture du SM est conçu pour masquer à la fois la latence de l'ALU et de la mémoire en alternant par cycle entre les warps.

Cette réponse n'utilise pas le terme "cœur CUDA" car cela introduit un modèle mental incorrect. Les cœurs CUDA sont des unités d'exécution pipelinées en virgule flottante/entier simple précision. Le débit d'émission et la latence de dépendance sont spécifiques à chaque architecture. Chaque sous-partition de SM et SM dispose d'autres unités d'exécution, y compris des unités de chargement/stockage, des unités de virgule flottante double précision, des unités de virgule flottante demi-précision, des unités de branchement, etc.

Pour maximiser les performances, le développeur doit comprendre le compromis entre blocs, warps et registres/thread.

Le terme "occupation" est le ratio de warps actifs sur le nombre maximum de warps sur un SM. Les architectures Kepler - Pascal (à l'exception du GP100) ont 4 planificateurs de warps par SM. Le nombre minimal de warps par SM doit au moins être égal au nombre de planificateurs de warps. Si l'architecture a une latence d'exécution dépendante de 6 cycles (Maxwell et Pascal), alors vous aurez besoin d'au moins 6 warps par planificateur, ce qui donne 24 par SM (24 / 64 = 37,5 % d'occupation) pour couvrir la latence. Si les threads ont un parallélisme au niveau de l'instruction, cela pourrait être réduit. Presque tous les noyaux émettent des instructions à latence variable telles que des lectures mémoire pouvant prendre 80 à 1000 cycles. Cela nécessite plus de warps actifs par planificateur de warp pour masquer la latence. Pour chaque noyau, il existe un point de compromis entre le nombre de warps et d'autres ressources telles que la mémoire partagée ou les registres, donc optimiser pour une occupation de 100 % n'est pas conseillé car un autre sacrifice sera probablement fait. Le profil CUDA peut aider à identifier le débit d'émission des instructions, l'occupation et les raisons d'attente afin d'aider le développeur à déterminer cet équilibre.

La taille d'un bloc de threads peut avoir un impact sur les performances. Si le noyau a de grands blocs et utilise des barrières de synchronisation, alors les attentes de barrière peuvent être une des raisons de l'attente. Cela peut être atténué en réduisant les warps par bloc de threads.

-3voto

liu km Points 1

Il y a plusieurs multiprocesseurs de streaming sur un seul appareil.
Un SM peut contenir plusieurs blocs. Chaque bloc peut contenir plusieurs threads.
Un SM a plusieurs cœurs CUDA (en tant que développeur, vous ne devriez pas vous soucier de cela car il est abstrait par warp), qui travailleront sur un thread. SM travaille toujours sur un warp de threads (toujours 32). Un warp ne travaillera que sur des threads du même bloc.
SM et bloc ont tous deux des limites sur le nombre de threads, le nombre de registres et la mémoire partagée.

Prograide.com

Prograide est une communauté de développeurs qui cherche à élargir la connaissance de la programmation au-delà de l'anglais.
Pour cela nous avons les plus grands doutes résolus en français et vous pouvez aussi poser vos propres questions ou résoudre celles des autres.

Powered by:

X