CUDA: synchronisation des threads

Presque n'importe où, j'ai lu sur la programmation avec CUDA il y a une mention de l'importance que tous les threads d'un warp faire la même chose.
Dans mon code j'ai une situation où je ne peux pas éviter une certaine condition. Il ressemble à ceci:

//some math code, calculating d1, d2
if (d1 < 0.5)
{
    buffer[x1] += 1;  //buffer is in the global memory
}
if (d2 < 0.5)
{
    buffer[x2] += 1;
}
//some more math code.

Des threads peut entrer dans une pour la des conditions, on pourrait en conclure deux et les autres peuvent ne pas entrer en soit.

Maintenant, pour faire de tous les thread revenir à "faire la même chose" de nouveau après les conditions, devrais-je synchroniser après les conditions de l'aide __syncthreads() ? Ou est-ce en quelque sorte se produit automatiquement?
Deux threads être pas faire la même chose, en raison de l'un d'eux étant une opération de derrière, donc ruiner pour tout le monde? Ou est-il derrière les coulisses effort pour les amener à faire la même chose à nouveau après une branche?

source d'informationauteur shoosh

35

Dans une chaine, pas de fils va prendre de l'avance sur les autres. Si il y a une branche conditionnelle et il est pris par des fils dans la chaîne, mais pas d'autres (un.k.un. warp "divergence"), les autres threads juste ralenti jusqu'à ce que la branche est terminée et tous, ils "convergent" de retour ensemble sur un même enseignement. Donc, si vous avez seulement besoin de l'intérieur de la courbure de la synchronisation de threads, ce qui se passe "automagiquement."

Mais des chaines différentes ne sont pas synchronisés de cette façon. Donc, si votre algorithme exige que certaines opérations soient complètes sur de nombreux croisements, alors vous aurez besoin d'utiliser explicitement la synchronisation des appels (voir le Guide de Programmation CUDA, Section 5.4).

EDIT: réorganisé les prochains paragraphes, afin de clarifier certaines choses.

Il y a vraiment deux questions se posent ici: Instruction de la synchronisation et de la mémoire de la visibilité.
- __syncthreads() applique l'instruction de synchronisation et s'assure de la mémoire de la visibilité, mais seulement à l'intérieur d'un bloc, et non pas à travers des blocs (de Programmation CUDA Guide de l'Annexe B. 6). Il est utile pour écrire puis lire de la mémoire partagée, mais n'est pas approprié pour la synchronisation globale d'accès à la mémoire.
- __threadfence() garantit la mémoire de la visibilité, mais ne pas faire d'une instruction de synchronisation, de sorte que dans mon expérience, il est d'usage limité (mais voir le code d'exemple dans l'Annexe B. 5).
- Mondial de l'instruction de synchronisation n'est pas possible à l'intérieur d'un noyau. Si vous avez besoin d' f() fait sur tous les fils avant de les appeler g() sur n'importe quel thread, split f() et g() entre les deux types de grains et de les appeler en série à partir de l'hôte.
- Si vous avez juste besoin d'incrémenter partagé ou compteurs globaux, l'utilisation de la atomique en fonction increment atomicInc() (Annexe B. 10). Dans le cas de votre code ci-dessus, si x1 et x2 ne sont pas uniques au monde (dans tous les threads dans votre grille), non-atomique incréments entraînera dans une course à condition, semblable au dernier paragraphe de l'Annexe B. 2.4.
Enfin, gardez à l'esprit que toutes les opérations sur la mémoire globale, et les fonctions de synchronisation en particulier (y compris les atomics) sont mauvais pour la performance.

Sans connaître le problème que vous résolvez il est difficile de spéculer, mais peut-être que vous pouvez revoir la conception de votre algorithme à utiliser la mémoire partagée au lieu de la mémoire globale dans certains endroits. Cela permettra de réduire la nécessité pour la synchronisation et de vous donner un boost de performance.
2

De l'article 6.1 de la CUDA Guide de bonnes Pratiques:

Quel que soit le débit instruction de contrôle (if, switch, do, for, while) peuvent affecter de manière significative
l'instruction de débit en causant des threads d'un même warp à diverger; c'est,
pour suivre les différents chemins d'exécution. Si cela se produit, les différents chemins d'exécution
doit être sérialisé, en augmentant le nombre total d'instructions exécutées pour le compte de cette
de la chaine. Lorsque tous les chemins d'exécution différents ont terminé, le fils convergent
de retour sur le même chemin d'exécution.

Donc, vous n'avez pas besoin de faire quelque chose de spécial.
2

Gabriel réponse:

"Mondial de l'instruction de synchronisation n'est pas possible à l'intérieur d'un noyau. Si vous avez besoin de f (), sur tous les fils avant de les appeler g() sur n'importe quel thread, split f() et g() entre les deux types de grains et de les appeler en série à partir de l'hôte."

Que si la raison vous avez besoin de f() et g() dans le même thread est parce que vous êtes en utilisant un registre de la mémoire, et vous souhaitez vous inscrire ou de données partagée à partir de f pour obtenir de g?
C'est, pour mon problème, la raison de la synchronisation à travers des blocs est parce que les données à partir de f est nécessaire dans le g - et l'éclatement d'un noyau faut une grande quantité supplémentaire de la mémoire globale pour le transfert des données du registre de f à g, que je voudrais éviter
1

La réponse à votre question est non. Vous n'avez pas besoin de faire quelque chose de spécial.
De toute façon, vous pouvez résoudre ce problème, au lieu de votre code, vous pouvez faire quelque chose comme ceci:
```
buffer[x1] += (d1 < 0.5);
buffer[x2] += (d2 < 0.5);
```
Vous devriez vérifier si vous pouvez utiliser la mémoire partagée et l'accès à la mémoire globale dans un coalisées modèle. Assurez-vous également que vous NE voulez PAS écrire pour le même indice dans plus de 1 thread.

Vous devez vous connecter pour publier un commentaire.