cuda - tutorial - ¿__Syncthreads() sincroniza todos los hilos en la cuadrícula?




cuda wikipedia (2)

El __syncthreads() es una barrera de sincronización a nivel de bloque . Eso significa que es seguro usarlo cuando todos los hilos en un bloque llegan a la barrera. También es posible usar __syncthreads() en el código condicional, pero solo cuando todos los subprocesos evalúan de forma idéntica dicho código, de lo contrario, es probable que la ejecución se cuelgue o produzca efectos secundarios no deseados [4] .

Ejemplo de uso de __syncthreads() : ( source )

__global__ void globFunction(int *arr, int N) 
{
    __shared__ int local_array[THREADS_PER_BLOCK];  //local block memory cache           
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;

    //...calculate results
    local_array[threadIdx.x] = results;

    //synchronize the local threads writing to the local memory cache
    __syncthreads();

    // read the results of another thread in the current thread
    int val = local_array[(threadIdx.x + 1) % THREADS_PER_BLOCK];

    //write back the value to global memory
    arr[idx] = val;        
}

Para sincronizar todos los subprocesos en una cuadrícula actualmente no hay una llamada de API nativa. Una forma de sincronizar subprocesos en un nivel de cuadrícula es utilizando llamadas de kernel consecutivas, ya que en ese punto todos los subprocesos terminan y comienzan de nuevo desde el mismo punto. También comúnmente se llama sincronización de CPU o sincronización implícita. Así todos están sincronizados.

Ejemplo de uso de esta técnica ( source ):

Respecto a la segunda cuestión. , declara la cantidad de memoria compartida especificada por bloque. Tenga en cuenta que la cantidad de memoria compartida disponible se mide por SM . Por lo tanto, se debe tener mucho cuidado de cómo se utiliza la memoria compartida junto con la configuración de inicio .

... o solo los hilos en la urdimbre o bloque actual?

Además, cuando los hilos en un bloque particular se encuentran (en el núcleo) la siguiente línea

__shared__  float srdMem[128];

¿Sólo declararán este espacio una vez (por bloque)?

Obviamente, todos funcionan de forma asíncrona, de modo que si el subproceso 23 en el bloque 22 es el primer subproceso que llega a esta línea, y luego el subproceso 69 en el bloque 22 es el último que llega a esta línea, el subproceso 69 sabrá que ya se ha declarado.


Estoy de acuerdo con todas las respuestas aquí, pero creo que nos falta un punto importante aquí con la primera pregunta. No estoy respondiendo la segunda respuesta, ya que se respondió perfectamente en las respuestas anteriores.

La ejecución en GPU ocurre en unidades de urdimbre. Una deformación es un grupo de 32 hilos y, en un momento dado, cada hilo de una deformación particular ejecuta la misma instrucción. Si asigna 128 subprocesos en un bloque, su (128/32 =) 4 deformaciones para una GPU.

Ahora la pregunta es "Si todos los subprocesos ejecutan la misma instrucción, ¿por qué se necesita la sincronización?". La respuesta es que necesitamos sincronizar las deformaciones que pertenecen al bloque SAME . __syncthreads no sincroniza los hilos en un warp, ya están sincronizados. Sincroniza las deformaciones que pertenecen al mismo bloque.

Es por eso que la respuesta a su pregunta es: __syncthreads no sincroniza todos los subprocesos en una cuadrícula, pero los subprocesos que pertenecen a un bloque como cada bloque se ejecuta de forma independiente.

Si desea sincronizar una cuadrícula, divida su núcleo (K) en dos núcleos (K1 y K2) y llame a ambos. Se sincronizarán (K2 se ejecutará después de que K1 termine).