¿Puedo usar syncthreads () después de haber caído hilos?

Capacidad de cómputo 7.x (Volta) actualización:

Con la introducción de la Programación de subprocesos Independientes entre subprocesos en una deformación, CUDA es finalmente más estricta en la práctica, ahora coincidiendo 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 a nivel warp. En algunos casos, esto permitió que una barrera tuviera éxito sin ser ejecutada por cada hilo siempre y cuando al menos algún hilo en cada urdimbre alcanzara la barrera. Comenzando con Volta, el built-in de CUDA __syncthreads() y la barra de instrucciones PTX.la sincronización (y sus derivadas) se aplican por hilo y, por lo tanto, no tendrá éxito hasta que lleguen todos los hilos no salidos del bloque. El código que explote el comportamiento anterior probablemente se bloqueará y debe modificarse para garantizar que todos los subprocesos no salidos alcancen el barrera.

A continuación está la respuesta anterior, que divagaba sobre el comportamiento pre-Volta.

Actualización : Esta respuesta puede no añadir nada encima de talonmies' (dependiendo de su 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 podrían funcionar las cosas "bajo el capó" o lo que podría ser posible más allá de la documentación oficial, no hay nada que ver aquí. Dicho esto, todavía no recomiendo hacer suposiciones más allá de lo que está oficialmente documentado, especialmente en un entorno que espera soportar múltiples arquitecturas o futuras. Principalmente quería señalar que si bien esto se llama explícitamente como mala práctica por la Guía de Programación de CUDA, el comportamiento real de __syncthreads() puede ser algo diferente de cómo se describe y para mí eso es interesante. Lo último que quiero es esparcir desinformación, así que estoy abierto a la discusión y la revisión de mi respuesta!

Algunos hechos importantes

No hay TL; DR para esta respuesta ya que hay demasiado potencial para una mala interpretación, pero aquí hay algunos hechos relevantes para comenzar:{[35]]}

• __syncthreads() se comporta como una barrera para deformaciones en un bloque en lugar de todos los hilos en un bloque, aunque cuando se usa como se aconseja equivale a la misma cosa.
• Si cualquier hilo en un warp ejecuta una instrucción PTX bar (por ejemplo, de _syncthreads), es como si todos los hilos en la urdimbre tienen.
• Cuando se llama a un bar.sync (generado por el instrinsic __syncthreads()), el conteo de llegada para ese bloque y barrera se incrementa por el tamaño de urdimbre. Así es como se logran los puntos anteriores.
• La divergencia del hilo (rutas múltiples) se maneja mediante la serialización de la ejecución de las ramas. El orden de la serialización es un factor que puede causar problemas.
• Los hilos dentro de una urdimbre no están sincronizados por __syncthreads(). La instrucción no hará que la urdimbre se detenga y espere los hilos en caminos divergentes. La ejecución de la rama se serializa, por lo que solo cuando las ramas se vuelven a unir o el código termina, los hilos en la deformación se resincronan. Hasta eso, las ramas se ejecutan en secuencia e independientemente. Una vez más, solo un hilo en cada urdimbre del bloque necesita golpear __syncthreads() para que la ejecución continúe.

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

Interpretación y documentación

Dado que __syncthreads() actúa como una barrera para deformaciones 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 temprana estaría bien si al menos un hilo en cada urdimbre golpea la barrera. (¡Pero eso no quiere decir que no puedes causar estancamientos con lo intrínseco!) Esto también supone que __syncthreads() siempre generará una instrucción PTX simple bar.sync a; y que la semántica de eso tampoco cambiará, así que no hagas esto en producción.

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

La Programación Guide 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 viola esta recomendación. Demostramos que syncthreads () funciona como una barrera para warps, no para threads. Mostramos que cuando los hilos de una urdimbre son serializados debido a la divergencia de ramas, cualquier syncthreads () en un camino no espera hilos del otro camino, sino que solo espera otras deformaciones que se ejecutan dentro del mismo bloque de subprocesos.

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

Las barreras se ejecutan por urdimbre como si todos los hilos de una urdimbre estuvieran activos. Por lo tanto, si cualquier hilo en una urdimbre ejecuta una instrucción de barra, es como si todos los hilos en la urdimbre hubieran ejecutado la instrucción de barra. Todos los hilos en la urdimbre se estancan hasta que la barrera se completa, y el el conteo de llegadas para la barrera se incrementa por el tamaño de la urdimbre (no el número de hilos activos en la urdimbre).

Está claro de esto por qué el conteo de hilos opcional b en la instrucción bar.sync a{, b}; debe ser un múltiplo de tamaño de urdimbre whenever cada vez que un solo hilo en una urdimbre ejecuta una instrucción bar el conteo de llegada se incrementa por el tamaño de la urdimbre, no el número de hilos en la urdimbre que realmente golpean la barrera. Hilos que terminan temprano (seguido de un diferente path) fueron efectivamente contados como llegados de todos modos. Ahora, la siguiente oración en el pasaje citado dice entonces no usar __syncthreads() en código condicional a menos que "se sepa que todos los hilos evalúan la condición idénticamente (la urdimbre no diverge)."Esta parece ser una recomendación demasiado estricta (para la arquitectura actual), destinada a garantizar que el recuento de llegadas realmente refleje el número real de hilos que golpean la barrera. Si al menos un hilo golpea la barrera aumenta la cuenta de llegada para toda la urdimbre, usted realmente podría tener un poco más de flexibilidad.

No hay ambigüedad en la documentación de PTX de que la instrucción bar.sync a; generada por __syncthreads() espera a que todos los subprocesos en la matriz de subprocesos cooperativa actual (bloque) alcancen la barrera a. Sin embargo, el punto es que cómo "todos los hilos" se determina actualmente aumentando el conteo de llegadas en múltiplos de tamaño de urdimbre cada vez que se golpea la barrera (por defecto cuando no se especifica b). Esta parte no es un comportamiento indefinido, al menos no con Ejecución de Subprocesos Paralelos ISA Versión 4.2.

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

Ejemplos

Versión de salida temprana 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 bloqueará si al menos un hilo por urdimbre golpea la sincronización, pero un problema posible es el orden de serialización de la ejecución de rutas de código divergentes. Puede cambiar alrededor del núcleo anterior para intercambiar efectivamente las ramas.

Salida temprana versión 2:

__global__ void kernel(...)

    if (tidx < N) {
        // do stuff

        __syncthreads();
    }
    // else return;
}

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

El documento demuestra esto es un ejemplo más involucrado en comparación con una salida temprana trivial que también nos recuerda ser cautelosos en torno a la divergencia warp. Aquí la primera mitad de la urdimbre (thread id tid en [0,15]) escribe en alguna memoria compartida y ejecuta __syncthreads(), mientras que la otra mitad (thread id tid en [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, es posible que espere un punto muerto en cualquiera de los dos barrera.

// 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 un punto muerto, lo que indica que __syncthreads() no sincroniza hilos divergentes dentro de una urdimbre. Las rutas de código divergentes se serializan en una urdimbre y solo se necesita un hilo en una ruta de código para hacer que la llamada a __syncthreads() funcione en el nivel por urdimbre.

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

Resumen

__syncthreads() no espera todos los hilos en una urdimbre, y la llegada de un solo hilo en una urdimbre cuenta efectivamente toda la urdimbre como haber alcanzado la barrera. (Arquitectura actual).

Puede ser peligroso usar __syncthreads() en código condicional debido a cómo la ejecución de subprocesos divergentes se serializa.

Use el código intrínseco en condicional solo si comprende cómo funciona y cómo se maneja la divergencia de ramas (que ocurre dentro de una deformación).

Tenga en cuenta que no dije que siguiera adelante y usara __syncthreads() de una manera inconsistente con la forma en que está documentado.