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)
- A(1) est plus efficace car les branchements y sont uniformes
- B(1) fait diverger les branchements dans tous les warps, alors que (2) les rend uniformes par warp ; (2) est donc plus efficace
- 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.