Devrais-je unifier deux noyaux similaires avec une déclaration 'if', risquant une perte de performance?

Question

J'ai 2 fonctions de noyau très similaires, dans le sens où le code est presque le même, mais avec une légère différence. Actuellement, j'ai 2 options:

• Écrire 2 méthodes différentes (mais très similaires)
• Écrire un seul noyau et mettre les blocs de code qui diffèrent dans une déclaration if / else

Dans quelle mesure une déclaration if affectera-t-elle les performances de mon algorithme?
Je sais qu'il n'y a pas de branchement, car tous les threads dans tous les blocs entreront dans le if, ou le else.
Est-ce qu'une seule déclaration if diminuera mes performances si la fonction de noyau est appelée de nombreuses fois?

Answer 1

Vous avez une troisième alternative, qui consiste à utiliser le templating C++ et à faire du variable utilisée dans l'instruction if/switch un paramètre de modèle. Instanciez chaque version du kernel dont vous avez besoin, et vous aurez ensuite plusieurs kernels effectuant des tâches différentes sans vous soucier de la divergence de branches ou de l'évaluation conditionnelle, car le compilateur optimisera le code mort et la branching avec.

Peut-être quelque chose comme ceci :

template
__global__ void kernel()
{
    switch(action) {
       case 1:
       // Premier code
       break;

       case 2:
       // Deuxième code
       break;
    }
}

template void kernel<1>();
template void kernel<2>();

Answer 2

Il diminuera légèrement vos performances, surtout s'il se trouve dans une boucle interne, car vous gaspillez un slot d'émission d'instructions de temps en temps, mais ce n'est pas du tout aussi grave que si un warp était divergent.

Si c'est vraiment important, il peut être utile de déplacer la condition en dehors de la boucle. Cependant, si le warp est vraiment divergent, pensez à comment supprimer le branching : par exemple, au lieu de

if (i>0) {
    x = 3;
} else {
    x = y;
}

Essayez

x = ((i>0)*3) | ((i<3)*y);