performance - rama - poner tag en git

¿La divergencia de las ramas es realmente tan mala?

Question

He visto muchas preguntas dispersas en Internet sobre la divergencia de sucursales y cómo evitarla . Sin embargo, incluso después de leer docenas de artículos sobre cómo funciona CUDA, parece que no veo cómo evitar la divergencia de ramas ayuda en la mayoría de los casos . Antes de que nadie salte sobre mí con las garras extendidas, permítame describir lo que considero que es "la mayoría de los casos".

Me parece que la mayoría de los casos de divergencia de ramas involucran una cantidad de bloques de código verdaderamente distintos. Por ejemplo, tenemos el siguiente escenario:

if (A):
  foo(A)
else:
  bar(B)

Si tenemos dos hilos que encuentran esta divergencia, el hilo 1 se ejecutará primero, tomando el camino A. A continuación, el hilo 2 tomará el camino B. Para eliminar la divergencia, podríamos cambiar el bloque anterior para que se lea así:

foo(A)
bar(B)

Suponiendo que es seguro llamar a foo(A) en el subproceso 2 y la bar(B) en el subproceso 1, se puede esperar que mejore el rendimiento. Sin embargo, así es como lo veo:

En el primer caso, los hilos 1 y 2 se ejecutan en serie. Llamar a estos dos ciclos de reloj.

En el segundo caso, los hilos 1 y 2 ejecutan foo(A) en paralelo, luego ejecutan la bar(B) en paralelo. Esto todavía me parece como dos ciclos de reloj, la diferencia es que en el primer caso, si foo(A) implica una lectura desde la memoria, imagino que el subproceso 2 puede comenzar a ejecutarse durante esa latencia, lo que da como resultado la latencia oculta. Si este es el caso, el código divergente de la sucursal es más rápido.

Answer 1

Estás asumiendo (al menos es el ejemplo que das y la única referencia que haces) que la única forma de evitar la divergencia de ramas es permitir que todos los hilos ejecuten todo el código.

En ese caso, estoy de acuerdo en que no hay mucha diferencia.

Pero evitar la divergencia de ramas probablemente tiene más que ver con la reestructuración del algoritmo en un nivel más alto que simplemente la adición o eliminación de algunas sentencias if y hacer que el código sea "seguro" para ejecutar en todos los hilos.

Ofreceré un ejemplo. Supongamos que sé que los hilos impares necesitarán manejar el componente azul de un píxel e incluso los hilos necesitarán manejar el componente verde:

#define N 2 // number of pixel components
#define BLUE 0
#define GREEN 1
// pixel order: px0BL px0GR px1BL px1GR ...


if (threadIdx.x & 1)  foo(pixel(N*threadIdx.x+BLUE));
else                  bar(pixel(N*threadIdx.x+GREEN));

Esto significa que cada hilo alternativo toma una ruta determinada, ya sea foo o bar . Entonces ahora mi Warp tarda el doble de tiempo en ejecutarse.

Sin embargo, si reorganizo mis datos de píxeles para que los componentes de color sean contiguos, tal vez en fragmentos de 32 píxeles: BL0 BL1 BL2 ... GR0 GR1 GR2 ...

Puedo escribir un código similar:

if (threadIdx.x & 32)  foo(pixel(threadIdx.x));
else                   bar(pixel(threadIdx.x));

Todavía parece que tengo la posibilidad de divergencia. Pero dado que la divergencia ocurre en los límites de warp, un warp de give ejecuta ya sea la ruta if o la ruta else , por lo que no ocurre ninguna divergencia real.

Este es un ejemplo trivial, y probablemente estúpido, pero ilustra que puede haber formas de evitar la divergencia de la distorsión que no implique ejecutar todo el código de todas las rutas divergentes.