Programmation GPU

Les codes (1) et (2) ci-dessus font tous deux exécuter A à la moitié des threads et B à l'autre moitié. Quelle explication est correcte du point de vue de l'efficacité des warps (32 threads par warp) ?

if (threadIdx.x % 2 == 0) A(); else B();   // (1)
if ((threadIdx.x / 32) % 2 == 0) A(); else B();   // (2)
  1. A(1) est plus efficace car les branchements y sont uniformes
  2. B(1) fait diverger les branchements dans tous les warps, alors que (2) les rend uniformes par warp ; (2) est donc plus efficace
  3. CLes deux ont exactement la même efficacité
Afficher la réponse et l'explication

Réponse : B. (1) fait diverger les branchements dans tous les warps, alors que (2) les rend uniformes par warp ; (2) est donc plus efficace

Dans (1), la séparation se fait entre threads pairs et impairs ; chaque warp contient des threads A et B, et tous les warps exécutent successivement A et B. Dans (2), la séparation se fait sur l'indice divisé par 32 ; les 32 threads d'un warp sont tous du même côté, et chaque warp n'exécute que A ou que B. Pour une même tâche, aligner la condition sur l'unité du warp supprime la divergence.

S'entraîner sur ce thème Retour au thème

Autres questions sur Programmation GPU

  1. Quelle est la définition de l'accès mémoire coalescé (coalescing) sur GPU ?
  2. Avec i l'indice de chaque thread, lire le tableau tous les 32 éléments comme ci-dessus est plus lent que lire a[i]. Quelle en est la raison principale ?
  3. Pour traiter sur GPU des données de particules ayant des positions x, y et z, quelle disposition est avantageuse du point de vue de l'accès mémoire coalescé ?
  4. Quelle est la caractéristique de la mémoire partagée (shared memory) en CUDA ?
  5. Quel est le rôle de __syncthreads() en CUDA ?

Comparatifs sur ce thème