CUDA/openCL; ramas de reescritura como no ramificación expresión

Question

mayoría de las veces se requiere una rama en un programa CUDA o OpenCL, como:

for (int i=0; i<width; i++) 
{ 
    if(i % threadIdx.x == 0) 
    quantity += i*i; 
} 

el código puede siempre (o al menos, la mayor parte del tiempo) reescribirse en estilo no-ramificación:

for (int i=0; i<width; i++) 
{ 
    quantity += i*i* (i % threadIdx.x != 0); 
} 

la disyuntiva parece ser ya sea corriendo en una sola ranura urdimbre frente hacer más cálculos en todos los hilos (en el segundo caso, la suma se ejecuta siempre, sólo que a veces el valor es cero)

Suponiendo que las operaciones de bifurcación tomarán múltiples ranuras de urdimbre para cada posible bifurcación, uno esperaría que el segundo sea consistentemente mejor que el primero, ahora mi pregunta es; ¿Puedo confiar en que el compilador optimice 1) en 2) cuando tenga sentido, o no exista un criterio de aplicación amplia, lo que implica que no se puede decidir en general cuál es mejor sin intentar y perfilar?

Answer 1

No tengo muchas memorias sobre CUDA, pero ¿por qué no paralelizas tu ciclo? Debería usar operaciones atómicas [1] para agregar su cálculo. ¡Espero que esto ayude! Lo siento si no es el caso.

1. atómicos Operaciones: http://supercomputingblog.com/cuda/cuda-tutorial-4-atomic-operations/

Answer 2

En mi experiencia - es totalmente hasta el compilador-escritores para optimizar este tipo de casos extremos.

Entonces, ¿puedo pensar en algún caso donde 1) no se puede convertir en 2)? He aquí uno: escribí núcleos donde era más eficiente ejecutar ciertas partes de los cálculos cada 10 hilos o algo así, en cuyo caso no se puede inferir tal optimización aunque exista una operación matemática (una división y un resta) que pueda arroje el mismo resultado independientemente del resultado condicional en comparación con "ejecutar en todos menos arrojar cero".

Sin embargo, incluso dado que la comprobación de threadId == 0 es un escenario bastante común, no tengo idea de si está realmente optimizado para. Apostaría a que depende de la implementación E incluso del dispositivo en sí (CPU vs GPU).

Tendrás que probarlo para descubrir realmente qué funciona mejor, no solo por el motivo anterior sino también porque el programador del trabajo puede comportarse de manera diferente en función de lo caro que es programar/iniciar/detener un conjunto de hilos en lugar de tenerlos a todos corriendo (y la mayoría proporciona un resultado de cero/identidad).

Espero que esto ayude!

Answer 3

Las operaciones del módulo son razonablemente costosas: estoy razonablemente seguro de que agregar el módulo requeriría más tiempo que el de tener una sola instrucción que solo se ejecuta 1 subproceso. Su única declaración de bifurcación, if sin else, solo bloqueará los otros hilos mientras se esté ejecutando dicha declaración. Debido a que los gpus están optimizados para un cambio de contexto muy rápido, debe haber muy poco costo para eso.

Sin embargo, se recomienda no utilizar largas instrucciones de bifurcación: un cálculo en serie demasiado grande en la GPU (es decir, un hilo haciendo todo el trabajo) niega la ventaja del paralelismo.