CUDA: __syncthreads() dentro de sentencias if

Question

Tengo una pregunta acerca de la sincronización de CUDA. Particularmente, necesito alguna aclaración sobre la sincronización en las declaraciones if. Quiero decir, si pongo un __syncthreads() bajo el alcance de una instrucción if golpeada por una fracción de los hilos dentro del bloque, ¿qué ocurre? Pensé que algunos hilos permanecerán "para siempre" esperando los otros hilos que no alcanzarán el punto de sincronización. Por lo tanto, he escrito y ejecutado un código de ejemplo para inspeccionar:

__global__ void kernel(float* vett, int n) 
{ 
    int index = blockIdx.x*blockDim.x + threadIdx.x; 
    int gridSize = blockDim.x*gridDim.x; 

    while(index < n) 
    { 
     vett[index] = 2; 
     if(threadIdx.x < 10) 
     { 
      vett[index] = 100; 
      __syncthreads(); 
     } 
     __syncthreads(); 

     index += gridSize; 
    } 
} 

Sorprendentemente, he observado que la salida era bastante "normal" (64 elementos, BLOCKSIZE 32):

100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 
100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 

Así

I modificado ligeramente el código de la siguiente manera:

__global__ void kernel(float* vett, int n) 
{ 
    int index = blockIdx.x*blockDim.x + threadIdx.x; 
    int gridSize = blockDim.x*gridDim.x; 

    while(index < n) 
    { 
     vett[index] = 2; 
     if(threadIdx.x < 10) 
     { 
      vett[index] = 100; 
      __syncthreads(); 
     } 
     __syncthreads(); 
      vett[index] = 3; 
     __syncthreads(); 

     index += gridSize; 
    } 
} 

Y la salida fue:

3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 
3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 

Una vez más, estaba equivocado: pensé que los hilos dentro de la declaración if, después de modificar el elemento del vector, permanecerían en estado de espera y nunca saldrían del alcance de if. Entonces ... ¿podría aclarar qué sucedió? ¿Un hilo que se encuentra después de un punto de sincronización desbloquea los hilos que esperan en la barrera? Si necesita reproducir mi situación, utilicé CUDA Toolkit 5.0 RC con SDK 4.2. Muchas gracias por adelantado.

Answer 1

No debe usar __syncthreads() a menos que la instrucción se alcance en todos los hilos dentro de un bloque de hilos, siempre. Desde el programming guide (B.6):

__syncthreads() está permitido en código condicional pero sólo si los evalúa condicionales idéntica en toda la secuencia de rosca, de lo contrario es probable que colgar o producir efectos secundarios no deseados de la ejecución de código.

Básicamente, su código no es un programa CUDA bien formado.

Answer 2

El modelo CUDA es MIMD pero las GPU NVIDIA actuales implementan __syncthreads() en warp granularity en lugar de thread. Significa que estos son warps inside a thread-block que están sincronizados no necesariamente threads inside a thread-block. __syncthreds() espera a que todos los 'urdidos' de bloque de hilos alcancen la barrera o salgan del programa. Consulte Henry Wong's Demistifying paper para obtener más información.

Answer 3

En resumen, el comportamiento es undefined. Por lo tanto, algunas veces puede hacer lo que usted desea, o no, o (bastante probable) simplemente bloqueará o bloqueará su kernel.

Si tiene curiosidad por saber cómo funcionan las cosas internamente, debe recordar que los hilos no se ejecutan de forma independiente, sino que se combinan (un grupo de 32 hilos) a la vez.

Esto, por supuesto, crea un problema con las ramas condicionales donde el condicional no se evalúa uniformemente a lo largo de la urdimbre. El problema se resuelve mediante la ejecución de ambas rutas, una después de otra, cada una con los subprocesos desactivados que no se supone que ejecuten esa ruta. IIRC en hardware existente la rama se toma primero, luego la ruta se ejecuta donde no se toma la rama, pero este comportamiento es indefinido y por lo tanto no se garantiza.

Esta ejecución de rutas separada continúa hasta cierto punto para que el compilador pueda determinar que se garantiza que todos los hilos de las dos rutas de ejecución separadas (el "punto de reconvergencia" o el "punto de sincronización") lo alcancen. Cuando la ejecución de la primera ruta del código llega a este punto, se detiene y en su lugar se ejecuta la segunda ruta del código. Cuando la segunda ruta alcanza el punto de sincronización, todos los hilos se habilitan nuevamente y la ejecución continúa uniformemente desde allí.

La situación se vuelve más complicada si se encuentra otra rama condicional antes de la sincronización. Este problema se resuelve con una pila de rutas que aún deben ejecutarse (por suerte, el crecimiento de la pila es limitado, ya que podemos tener, como máximo, 32 rutas de código diferentes para una urdimbre).

Donde se insertan los puntos de sincronización es undefined e incluso varía ligeramente entre las arquitecturas, por lo que de nuevo no hay garantías. El único comentario (no oficial) que recibirá de Nvidia es que el compilador es bastante bueno para encontrar puntos de sincronización óptimos. Sin embargo, a menudo hay problemas sutiles que pueden mover el punto óptimo más abajo de lo que cabría esperar, especialmente si los hilos salen antes.

Ahora, para comprender el comportamiento de la directiva __syncthreads(), (que se traduce en una instrucción bar.sync en PTX) es importante tener en cuenta que esta instrucción no se ejecuta por hilo, sino por toda la deformación a la vez (independientemente de si alguno de los hilos está desactivado o no) porque solo se deben sincronizar los urdimbres de un bloque. Los hilos de una urdimbre ya se están ejecutando en sincronización, y la sincronización adicional no tendrá ningún efecto (si todos los subprocesos están habilitados) o llevará a un interbloqueo cuando intente sincronizar los subprocesos de diferentes rutas de código condicional.

Puede trabajar desde esta descripción hasta cómo se comporta su código particular. Pero tenga en cuenta que todo esto es indefinido, no hay garantías, y confiar en un comportamiento específico puede romper su código en cualquier momento.

Es posible que desee consultar el PTX manual para obtener más detalles, en particular para la instrucción compilada por __syncthreads(). También vale la pena leer el "Demystifying GPU Microarchitecture through Microbenchmarking" paper de Henry Wong, al que ahmad hace referencia a continuación. A pesar de que, por ahora, la arquitectura obsoleta y la versión CUDA, las secciones sobre ramificación condicional y __syncthreads() parecen ser generalmente válidas.

Answer 4

__syncthreads() se utiliza para sincronizar subprocesos dentro de un bloque. Eso significa que todos los hilos en el bloque esperarán a que todos se completen antes de continuar.

Considere el caso en el que hay algunos subprocesos en un bloque, que ingresan el enunciado if y otros no. Esos hilos esperando, serán bloqueados; siempre esperando

En general, no es un buen estilo para poner sincronizar en una instrucción if-conditional. Lo mejor es evitarlo y rediseñar tu código si lo tienes. El propósito de la sincronización es asegurarse de que todos los hilos procedan juntos, ¿por qué los filtra mediante el uso de la instrucción if en primer lugar?

Para agregar, si se requiere sincronización entre los bloques. Se requiere reiniciar el kernel.