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




toolkit gpu (4)

... 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.


Answers

__syncthreads() espera hasta que todos los subprocesos dentro del mismo bloque hayan alcanzado el comando y todos los subprocesos dentro de una deformación, lo que significa que todas las deformaciones que pertenecen a un bloqueo de subprocesos deben llegar a la declaración.

Si declara la memoria compartida en un kernel, la matriz solo será visible para un bloque de subproceso. Así que cada bloque tendrá su propio bloque de memoria compartida.


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 .


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).


En mi lugar de trabajo, estoy trabajando con una GTX 590, que contiene 512 núcleos CUDA, 16 multiprocesadores y que tiene un tamaño de deformación de 32. Así que esto significa que hay 32 núcleos CUDA en cada multiprocesador que funciona exactamente en el mismo código en la misma deformación . Y, finalmente, el máximo de hilos por tamaño de bloque es 1024.

Un GTX590 contiene 2 veces los números que mencionó, ya que hay 2 GPU en la tarjeta. A continuación, me enfoco en un solo chip.

Permítame decirle que entiendo la situación: por ejemplo, asigno N bloques con el tamaño máximo de threadPerBlock de 1024 en la GTX 590. Por lo que entiendo de la guía de programación de CUDA y de otras fuentes, los bloques primero son enumerados por el hardware . En este caso 16 de los N bloques están asignados a diferentes multiprocesadores.

Los bloques no están necesariamente distribuidos uniformemente entre los multiprocesadores (SM). Si programa exactamente 16 bloques, algunos de los SM pueden obtener 2 o 3 bloques, mientras que algunos de ellos quedan inactivos. No se por que

Cada bloque contiene 1024 subprocesos y el programador de hardware asigna 32 de estos subprocesos a los 32 núcleos en un solo multiprocesador.

La relación entre hilos y núcleos no es tan directa. Hay 32 ALU "básicas" en cada SM. Los que manejan cosas como el punto flotante de precisión simple y la mayoría de las instrucciones lógicas y de enteros de 32 bits. Pero solo hay 16 unidades de carga / almacenamiento, por lo que si la instrucción de deformación que se está procesando actualmente es una carga / almacén, debe programarse dos veces. Y solo hay 4 unidades de funciones especiales, que hacen cosas como la trigonometría. Así que estas instrucciones deben ser programadas 32/4 = 8 veces.

Los hilos en el mismo multiprocesador (warp) procesan la misma línea del código y utilizan la memoria compartida del multiprocesador actual.

No, puede haber muchos más de 32 subprocesos "en vuelo" al mismo tiempo en un solo SM.

Si los 32 subprocesos actuales encuentran una operación fuera del chip como las lecturas y escrituras en la memoria, se reemplazan con otro grupo de 32 subprocesos del bloque actual. Por lo tanto, en realidad hay 32 subprocesos en un solo bloque que se ejecutan exactamente en paralelo en un multiprocesador en un momento dado, no la totalidad de los 1024.

No, no son solo las operaciones de memoria las que hacen que se reemplacen las deformaciones. Las ALU también están profundamente canalizadas, por lo que se intercambiarán nuevas deformaciones a medida que se produzcan dependencias de datos para los valores que aún se encuentran en la tubería. Entonces, si el código contiene dos instrucciones donde la segunda usa la salida de la primera, la deformación se pondrá en espera mientras que el valor de la primera instrucción se abrirá paso a través de la tubería.

Finalmente, si un multiprocesador procesa completamente un bloque, un nuevo bloque de subprocesos de la lista de N bloques de subprocesos se conecta al multiprocesador actual.

Un multiprocesador puede procesar más de un bloque a la vez, pero un bloque no puede moverse a otro MP una vez que se haya iniciado el procesamiento. El número de subprocesos en un bloque que están actualmente en vuelo depende de cuántos recursos utiliza el bloque. La Calculadora de ocupación de CUDA le indicará cuántos bloques estarán en vuelo al mismo tiempo, en función del uso de recursos de su núcleo específico.

Y, finalmente, hay un total de 512 subprocesos que se ejecutan en paralelo en la GPU durante la ejecución del kernel CUDA. (Sé que si un bloque usa más registros que los disponibles en un solo multiprocesador, entonces se divide para funcionar en dos multiprocesadores, pero supongamos que cada bloque puede caber en un solo multiprocesador en nuestro caso).

No, un bloque no se puede dividir para trabajar en dos multiprocesadores. Un bloque completo siempre es procesado por un único multiprocesador. Si el multiprocesador dado no tiene suficientes recursos para procesar al menos un bloque con su kernel, obtendrá un error de inicio del kernel y su programa no se ejecutará en absoluto.

Depende de cómo se define un subproceso como "en ejecución". La GPU normalmente tendrá muchos más de 512 subprocesos que consumen varios recursos en el chip al mismo tiempo.

Vea la respuesta de @ harrism en esta pregunta: CUDA: ¿Cuántos hilos concurrentes en total?