Comment réduire le nombre de branches/instructions if dans mes noyaux ?

Question

Existe-t-il une manière intelligente de réduire le nombre de l if statements dans le noyau d'un CUDA ?

J'écris une application qui calculera un Hamiltonien à plusieurs corps (simulation d'un système quantique). Les calculs dépendent fortement de conditional expressions .

La raison pour laquelle je veux réduire ces instructions est qu'elles introduisent une surcharge de performance. (entier warp saisit chaque option de if(){} else if(){} si la condition n'est pas remplie, la thread pour un temps donné, reste inactif).

Question :
1. Will switch() La déclaration résout le problème ?
2. Le code ci-dessous est destiné à représenter une idée générale :

class tag_option_1 {};
class tag_option_2 {};
class tag_option_3 {};

template<typename T> __device__
int cal_something(int ab, int cd)
{
    return -12345; // error value. default case is an error.
};

template<> __device__
int cal_something<tag_option_1>(int ab, int cd)
{
    // return something
}

template<> __device__
int cal_something<tag_option_2>(int ab, int cd)
{
    // return something
}

template<> __device__
int cal_something<tag_option_3>(int ab, int cd)
{
    // return something
}

////////////////////////////////
// version #1:

__global__
void calc_hamiltonian(int * foo, int * bar)
{
    unsigned int tid = /* calce thread index*/;

    // do something... 

    if (/* condition */)
    {
        cal_something<tag_option_1>(foo[tid], bar[tid]);
    }
    else if(/* condition */)
    {
        cal_something<tag_option_2>(foo[tid], bar[tid]);
    }
    else if(/* condition */)
    {
        cal_something<tag_option_3>(foo[tid], bar[tid]);
    }
    // no default case.

    // do something... 
}

////////////////////////////////
// version #2:

// "magical" way to select a version:
// variant is meant to be something like "boost::variant" but implementented without a single "if" statement.
// This "magical" step is meant to be resolved at compile time.
__devcie__
variant <tag_option_1, tag_option_2, tag_option_3> 
version_selector(int ab, int cd)
{
    // magic happens here.
}

__global__
void calc_hamiltonian(int * foo, int * bar)
{
    unsigned int tid = /* calce thread index*/;

    // do something... 

    cal_something <version_selector(foo[tid], bar[tid])> (foo[tid], bar[tid]);

    // do something... 
}

Existe-t-il un moyen d'implémenter un version #2 de l'exemple ci-dessus, ou bien il est impossible en CUDA C/C++ ?

Answer 1

Je suis généralement d'accord avec la recommandation de @njuffa de ne pas essayer de tordre artificiellement votre style de codage naturel, et que vous devriez être après la performance (et la lisibilité et la maintenabilité) plutôt que de compter les branches dans votre code source. D'autant plus que le compilateur peut parfois les faire disparaître.

Cela dit...

Quelques méthodes "intelligentes" de sagesse commune pour réduire le nombre de branches (dans CUDA et en général) :

1. Faire en sorte que les conditions soient vérifiées à l'avance, de préférence au moment de la compilation.

Explication par l'exemple. Version 1 :

 void foo(int* a, bool cond) {
     ...
     for(int i = 0; i < lots; i++) {
         if (cond) do_stuff()
         else do_other_stuff();
     }
     ...
 }

 bool cond = check_stuff();
 foo(data, cond);

Version 2 :

 void foo(int* a, bool cond) {
     ...
     if (cond) {
         for(int i = 0; i < lots; i++) { do_stuff(); }
     }
     else {
         for(int i = 0; i < lots; i++) { do_other_stuff(); }
     }
     ...
 }

 bool cond = check_stuff();
 foo(data, cond);

Version 3 :

 template <bool Cond>
 void foo(int* a) {
     ...
     if (cond) {
         for(int i = 0; i < lots; i++) { do_stuff(); }
     }
     else {
         for(int i = 0; i < lots; i++) { do_other_stuff(); }
     }
     ...
 }

 bool cond = check_stuff();
 if (cond) foo<true>(data) else foo<false>(data);

Et ce qui est bien avec la version 3, c'est que pendant qu'elle regarde comme s'il y avait une branche, ce n'est pas le cas en réalité - le compilateur prend soit seulement l'instruction "then", soit seulement l'instruction "else", mais pas les deux dans la même fonction.

Passer de la version 1 à la version 2 est quelque chose que le compilateur peut avoir la gentillesse de faire pour vous ; mais parfois ce n'est pas aussi simple que dans l'exemple et vous devez vous en occuper vous-même. Passer de la version 2 à la version 3 est quelque chose que le compilateur ne fera jamais pour vous.

1.1 Considérer les noyaux de modélisation sur la taille des blocs

Ce n'est pas toujours - en fait, pas souvent - utile, mais il y a un exemple bien connu donné par Mark Harris dans son livre intitulé présentation sur l'optimisation des réductions parallèles avec CUDA. Jetez un œil à l'optimisation n° 6 dans les diapositives 24 à 27. Mais n'essayez pas quelque chose comme ça - qui est laid et quelque peu fragile - à moins que vous n'ayez attentivement chronométré votre exécution pour vous assurer qu'elle en vaut la peine.

2. Faire en sorte que les comportements divergent en données plutôt que dans flux de contrôle

Version 1 :

 void foo(int* a, int *b) {
     ...
     if (check(a[global_thread_index]) { b[global_thread_index]++; }
 }

Version 2 :

 void foo(int* a, int *b) {
     ...
     b[global_thread_index] += check(a[global_thread_index]);
 }

(en supposant que la vérification renvoie un booléen, ou un entier 0 en cas d'échec et 1 en cas de succès).

Ici, je ne suis pas sûr de ce que fera le compilateur CUDA ; de plus, vous payez une pénalité de lisibilité en écrivant ce code et en brisant sans doute le principe de l'utilisation de l'algorithme de l'utilisateur. "principe du moindre étonnement" . Mais vous pouvez trouver des exemples moins artificiels.

2.1 Limiter les branches aux choix d'opérations trinaires de valeur

Il y a aussi une version 3 :

 void foo(int* a, int *b) {
     ...
     b[global_thread_index] = check(a[global_thread_index]) ? 1 : 0;
 }

Maintenant, il y a toujours une branche - l'opérateur trinaire n'est qu'un raccourci pour un "if", mais si vous pouvez amener votre code à cet état, la divergence sera limitée à une seule déclaration et à chaque branche, et peut-être même moins, si le compilateur CUDA parvient à utiliser l'option de contrôle de l'opérateur trinaire. déclaration slct PTX :

slct

Sélectionnez un opérande source, en fonction du signe du troisième opérande.

cela "emballe" la sémantique d'une branche dans la logique combinatoire d'une seule instruction.

Bien sûr, slct pourrait être utilisé par le compilateur dans d'autres cas ; cela ne dépend pas de vous.

3. Éviter efficacement les branches en s'efforçant d'obtenir un consensus à l'échelle de la chaîne.

(voir également le commentaire de @RobertCrovella au même effet).

Encore une fois, l'explication par l'exemple.

 void foo(int* a, int *b) {
     ...
     if (threadIdx.x % 2 == 0) { do_stuff(); }
     else { do_other_stuff(); }
     ...
 }

Version 2 :

 void foo(int* a, int *b) {
     ...
     if (threadIdx.x >= blockDim.x / 2) { do_stuff(); }
     else { do_other_stuff(); }
     ...
 }

Cela garantit que tous les warps, sauf peut-être celui du milieu, auront soit toutes les voies qui remplissent la condition, soit toutes les voies qui ne la remplissent pas. Et cela signifie qu'aucune voie dans ces warps ne devra attendre sans rien faire pendant que d'autres voies exécutent l'autre branche.

Pour un exemple concret, consultez les diapositives 7 à 13 de l'exposé de Mark Harris intitulé présentation Je l'ai mentionné plus haut.