Pregunta ¿Debería unificar dos kernels similares con una declaración 'if', arriesgando la pérdida de rendimiento?


Tengo 2 funciones de kernel muy similares, en el sentido de que el código es casi el mismo, pero con una ligera diferencia. Actualmente tengo 2 opciones:

  • Escribe 2 métodos diferentes (pero muy similares)
  • Escriba un kernel único y coloque los bloques de código que difieren en una instrucción if / else

¿Cuánto afectará una declaración if a mi rendimiento de algoritmo?
Sé que no hay ramificación, ya que todos los hilos en todos los bloques entrarán en el if o el else.
Entonces, ¿una sola instrucción if disminuirá mi rendimiento si la función kernel es invocada muchas veces?


33
2018-05-30 17:45


origen


Respuestas:


Tiene una tercera alternativa, que es usar plantillas de C ++ y hacer que la variable que se utiliza en la sentencia if / switch sea un parámetro de plantilla. Crea una instancia de cada versión del kernel que necesites, y luego tienes varios kernels haciendo cosas diferentes sin divergencia de rama o evaluación condicional de la que preocuparte, porque el compilador optimizará el código muerto y la ramificación con él.

Quizás algo como esto:

template<int action>
__global__ void kernel()
{
    switch(action) {
       case 1:
       // First code
       break;

       case 2:
       // Second code
       break;
    }
}

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

92
2018-05-30 18:26



Disminuirá un poco su rendimiento, especialmente si está en un bucle interno, ya que está desperdiciando un slot de edición de instrucciones de vez en cuando, pero no es tanto como si un warp fuera divergente.

Sin embargo, si es un gran problema, puede valer la pena mover la condición fuera del circuito. Sin embargo, si la distorsión es realmente divergente, piense cómo eliminar la ramificación: por ejemplo, en lugar de

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

tratar

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

4
2018-05-30 21:41