[Synchronization] ¿Puedo usar __syncthreads () después de haber soltado los hilos?


Answers

Actualización de Compute Capability 7.x (Volta):

Con la introducción de la programación independiente de subprocesos entre subprocesos en un warp, CUDA es finalmente más estricto en la práctica, y ahora coincide con el comportamiento documentado. De la guía de programación :

Aunque __syncthreads () se ha documentado consistentemente como la sincronización de todos los subprocesos en el bloque de subprocesos, Pascal y las arquitecturas anteriores solo podían imponer la sincronización en el nivel warp. En ciertos casos, esto permitió que una barrera tuviera éxito sin ser ejecutada por cada hilo, siempre que al menos un hilo en cada urdimbre alcanzara la barrera. Comenzando con Volta, las __syncthreads incorporadas en CUDA () y la instrucción PTX bar.sync (y sus derivadas) se aplican por subproceso y, por lo tanto, no tendrán éxito hasta que todas las subprocesos no salientes del bloque alcancen el éxito. Es probable que el código que explota el comportamiento anterior se estanque y deba modificarse para garantizar que todos los subprocesos no salidos lleguen a la barrera.

A continuación se encuentra la respuesta anterior, que hablaba sobre el comportamiento pre-Volta.

Actualización : esta respuesta puede no agregar nada encima de talonmies '(dependiendo de tu comprensión del tema, supongo), pero a riesgo de ser demasiado detallado, estoy presentando la información que me ayudó a entender esto mejor. Además, si no está interesado en cómo las cosas podrían funcionar "bajo el capó" o qué podría ser posible más allá de la documentación oficial, no hay nada que ver aquí. Dicho todo esto, todavía no recomiendo hacer suposiciones más allá de lo que está oficialmente documentado, especialmente en un entorno que espera soportar arquitecturas múltiples o futuras. Principalmente, quería señalar que, si bien la Guía de programación de CUDA lo menciona explícitamente como una mala práctica, el comportamiento real de __syncthreads() puede ser algo diferente de cómo se describe y para mí es interesante. Lo último que quiero es difundir la información errónea, ¡entonces estoy abierto a la discusión y a la revisión de mi respuesta!

Algunos hechos importantes

No hay TL; DR para esta respuesta, ya que hay demasiado potencial para la interpretación errónea, pero aquí hay algunos hechos relevantes para comenzar:

  • __syncthreads() comporta como una barrera para warps en un bloque en lugar de todos los hilos en un bloque, aunque cuando se usa como se aconseja, equivale a lo mismo.
  • Si algún hilo en una urdimbre ejecuta una instrucción de bar PTX (por ejemplo, desde _syncthreads ), es como si todos los hilos en la urdimbre tuvieran.
  • Cuando se bar.sync un bar.sync (como generado por el intrínseco __syncthreads() ), el recuento de llegada para ese bloque y barrera se incrementa por el tamaño de la urdimbre. Así es como se logran los puntos anteriores.
  • La divergencia de subprocesos (múltiples rutas) se maneja serializando la ejecución de las ramas. El orden de serialización es un factor que puede causar problemas.
  • Los hilos dentro de un warp no están sincronizados por __syncthreads() . La instrucción no causará que la urdimbre se bloquee y espere los hilos en caminos divergentes. La ejecución de la rama se serializa, por lo que solo cuando las ramas vuelvan a unirse o el código finalice, los hilos en la urdimbre se resincronizarán. Hasta eso, las ramas se ejecutan en secuencia e independientemente. De nuevo, solo un hilo en cada urdimbre del bloque necesita presionar __syncthreads() para que la ejecución continúe.

Estas declaraciones son respaldadas por documentación oficial y otras fuentes.

Interpretación y documentación

Dado que __syncthreads() actúa como una barrera para warps en un bloque en lugar de todos los hilos en un bloque, como se describe en la Guía de programación, parece que una simple salida anticipada estaría bien si al menos un hilo en cada warp golpea la barrera . (¡Pero eso no quiere decir que no se puedan generar bloqueos con lo intrínseco!). Esto también supone que __syncthreads() siempre generará una bar.sync a; Instrucción PTX y que la semántica de eso tampoco cambiará, así que no lo hagas en producción.

Un estudio interesante que encontré en realidad investiga qué sucede cuando vas en contra de las recomendaciones de la Guía de programación de CUDA, y descubrieron que si bien es posible causar un punto muerto al abusar de __syncthreads() en bloques condicionales, no todo el uso del intrínseco en el código condicional lo hará. De la Sección D.1 en el documento:

La Guía de programación recomienda que syncthreads () se use en código condicional solo si la condición se evalúa de forma idéntica en todo el bloque de subprocesos. El resto de esta sección investiga el comportamiento de syncthreads () cuando se infringe esta recomendación. Demostramos que syncthreads () funciona como una barrera para urdimbres, no hilos. Mostramos que cuando los hilos de una urdimbre se serializan debido a la divergencia de ramificación, cualquier syncthreads () en una ruta no espera a los subprocesos de la otra ruta, sino que solo espera otras deformaciones que se ejecutan dentro del mismo bloque de subprocesos.

Esta afirmación es concordante con el bit de la documentación de PTX citado por talonmies. Específicamente:

Las barreras se ejecutan sobre una base de urdimbre como si todos los hilos en una urdimbre estuvieran activos. Por lo tanto, si un hilo en un warp ejecuta una instrucción de barra, es como si todos los hilos en el warp hubieran ejecutado la instrucción de barra. Todos los hilos en la urdimbre se detienen hasta que se completa la barrera, y el recuento de llegada de la barrera se incrementa por el tamaño de la urdimbre (no por el número de hilos activos en la urdimbre).

De esto resulta claro por qué el recuento opcional de subprocesos b en la bar.sync a{, b}; la instrucción debe ser un múltiplo del tamaño de la urdimbre: cada vez que un hilo en una urdimbre ejecuta una instrucción de bar , la cuenta de arribo se incrementa por el tamaño de la urdimbre, no por el número de hilos en la urdimbre que realmente golpea la barrera . Los subprocesos que terminan temprano (seguidos de una ruta diferente) se contabilizaron como llegados de todos modos. Ahora, la siguiente oración en el pasaje citado dice que no se use __syncthreads() en el código condicional a menos que "se sepa que todos los hilos evalúan la condición de manera idéntica (el warp no diverge)". Esta parece ser una recomendación demasiado estricta (para la arquitectura actual), destinada a garantizar que el recuento de llegadas realmente refleje la cantidad real de hilos que golpean la barrera. Si al menos un hilo que golpea la barrera incrementa la cuenta de llegada para toda la deformación, es posible que realmente tenga un poco más de flexibilidad.

No hay ambigüedad en la documentación de PTX que bar.sync a; La instrucción generada por __syncthreads() espera a que todos los hilos en la matriz de hilos cooperativos (bloque) actual lleguen a la barrera a . Sin embargo, el punto es que la forma en que se determina actualmente "todos los hilos" incrementa el conteo de llegada en múltiplos de tamaño de urdimbre cuando se golpea la barrera (por defecto cuando no se especifica b ). Esta parte no es un comportamiento indefinido, al menos no con Parallel Thread Execution ISA Versión 4.2.

Tenga en cuenta que puede haber subprocesos inactivos en una urdimbre, incluso sin un condicional - "los últimos hilos de un bloque cuyo número de hilos no es un múltiplo del tamaño de la urdimbre". ( Notas de arquitectura SIMT ). Sin embargo, __syncthreads() no está prohibido en dichos bloques.

Ejemplos

Salida temprana versión 1:

__global__ void kernel(...)

    if (tidx >= N)
        return;      // OK for <32 threads to hit this, but if ALL
                     // threads in a warp hit this, THEN you are deadlocked
                     // (assuming there are other warps that sync)

    __syncthreads(); // If at least one thread on this path reaches this, the 
                     // arrival count for this barrier is incremented by 
                     // the number of threads in a warp, NOT the number of 
                     // threads that reach this in the current warp.
}

Esto no se estancará si al menos un hilo por warp golpea la sincronización, pero un posible problema es el orden de la serialización de la ejecución de rutas de código divergentes. Puede cambiar alrededor del kernel anterior para intercambiar las ramas de manera efectiva.

Salida temprana versión 2:

__global__ void kernel(...)

    if (tidx < N) {
        // do stuff

        __syncthreads();
    }
    // else return;
}

Todavía no hay un punto muerto si tiene al menos un hilo en la urdimbre que golpea la barrera, pero ¿es importante el orden de ejecución de la rama en este caso? No lo creo, pero probablemente sea una mala idea requerir un orden de ejecución en particular.

El documento demuestra esto en un ejemplo más complicado en comparación con una salida temprana trivial que también nos recuerda tener cuidado con la divergencia de la distorsión. Aquí la primera mitad del warp (thread id tid on [0,15]) escribe en alguna memoria compartida y ejecuta __syncthreads() , mientras que la otra mitad (id de thread tid on [16,31]) también ejecuta __syncthreads() pero ahora lee desde las ubicaciones de memoria compartida escritas por la primera mitad de la urdimbre. Ignorando la prueba de memoria compartida al principio, puede esperar un punto muerto en cualquiera de las barreras.

// incorrect code to demonstrate behavior of __syncthreads
if (tid < 16 ) {
  shared_array[tid] = tid;
  __syncthreads();
}
else {
  __syncthreads();
  output[tid] =
    shared_array[tid%16];
}

No hay punto muerto, lo que indica que __syncthreads() no sincroniza hilos divergentes dentro de un warp. Las rutas de código divergentes se serializan en un warp y solo se necesita un subproceso en una ruta de código para hacer que la llamada a __syncthreads() funcione al nivel de per-warp.

Sin embargo, el bit de memoria compartida muestra dónde puede entrar algún comportamiento impredecible. La segunda mitad del warp no obtiene los valores actualizados de la primera mitad porque la divergencia de ramas serializó la ejecución del warp y el bloque else se ejecutó primero . Entonces, la función no funciona bien, pero también muestra que __syncthreads() no sincroniza hilos divergentes en un warp.

Resumen

__syncthreads() no espera a todos los hilos en una urdimbre, y la llegada de una sola hebra en una urdimbre efectivamente cuenta que la urdimbre completa ha alcanzado la barrera. (Arquitectura actual).

Puede ser peligroso usar __syncthreads() en el código condicional debido a la serialización de la ejecución de subprocesos.

Use el código condicional intrínseco solo si comprende cómo funciona y cómo se maneja la divergencia de rama (que ocurre dentro de un warp).

Tenga en cuenta que no dije que siga adelante y use __syncthreads() de una manera incompatible con la forma en que está documentado.

Question

¿Es seguro usar __syncthreads() en un bloque donde he descartado hilos utilizando return ?

La documentación indica que __syncthreads() debe ser llamado por cada hilo en el bloque o de lo contrario llevará a un punto muerto, pero en la práctica nunca he experimentado dicho comportamiento.

Código de muestra:

__global__ void kernel(float* data, size_t size) {
    // Drop excess threads if user put too many in kernel call.
    // After the return, there are `size` active threads.
    if (threadIdx.x >= size) {
        return;
    }

    // ... do some work ...

    __syncthreads(); // Is this safe?

    // For the rest of the kernel, we need to drop one excess thread
    // After the return, there are `size - 1` active threads
    if (threadIdx.x + 1 == size) {
        return;
    }

     // ... do more work ...

    __syncthreads(); // Is this safe?
}