¿Debería unificar dos kernels similares con una statement ‘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 statement 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?

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 twig o evaluación condicional de la que preocuparte, porque el comstackdor optimizará el código muerto y la ramificación con él.

Quizás algo como esto:

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

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. Si la disformidad es realmente divergente, piense en 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);