Barrera entre bloques en CUDA

Question

Quiero implementar una barrera entre bloques en CUDA, pero con un problema grave.

No puedo entender por qué no funciona.

#include <iostream> 
#include <cstdlib> 
#include <ctime> 

#define SIZE 10000000 
#define BLOCKS 100 

using namespace std; 

struct Barrier { 
    int *count; 

    __device__ void wait() { 
     atomicSub(count, 1); 
     while(*count) 
      ; 
    } 

    Barrier() { 
     int blocks = BLOCKS; 
     cudaMalloc((void**) &count, sizeof(int)); 
     cudaMemcpy(count, &blocks, sizeof(int), cudaMemcpyHostToDevice); 
    } 

    ~Barrier() { 
     cudaFree(count); 
    } 
}; 


__global__ void sum(int* vec, int* cache, int *sum, Barrier barrier) 
{ 
    int tid = blockIdx.x; 

    int temp = 0; 
    while(tid < SIZE) { 
     temp += vec[tid]; 
     tid += gridDim.x; 
    } 

    cache[blockIdx.x] = temp; 

    barrier.wait(); 

    if(blockIdx.x == 0) { 
     for(int i = 0 ; i < BLOCKS; ++i) 
      *sum += cache[i]; 
    } 
} 

int main() 
{ 
    int* vec_host = (int *) malloc(SIZE * sizeof(int));  
    for(int i = 0; i < SIZE; ++i) 
     vec_host[i] = 1; 

    int *vec_dev; 
    int *sum_dev; 
    int *cache; 
    int sum_gpu = 0; 

    cudaMalloc((void**) &vec_dev, SIZE * sizeof(int)); 
    cudaMemcpy(vec_dev, vec_host, SIZE * sizeof(int), cudaMemcpyHostToDevice); 
    cudaMalloc((void**) &sum_dev, sizeof(int)); 
    cudaMemcpy(sum_dev, &sum_gpu, sizeof(int), cudaMemcpyHostToDevice); 
    cudaMalloc((void**) &cache, BLOCKS * sizeof(int)); 
    cudaMemset(cache, 0, BLOCKS * sizeof(int)); 

    Barrier barrier; 
    sum<<<BLOCKS, 1>>>(vec_dev, cache, sum_dev, barrier); 

    cudaMemcpy(&sum_gpu, sum_dev, sizeof(int), cudaMemcpyDeviceToHost); 

    cudaFree(vec_dev); 
    cudaFree(sum_dev); 
    cudaFree(cache); 
    free(vec_host); 
    return 0; 
} 

De hecho, incluso si vuelvo a escribir la espera() de la siguiente manera

__device__ void wait() { 
     while(*count != 234124) 
      ; 
    } 

el programa sale normalmente. Pero espero obtener un ciclo infinito en este caso.

Answer 1

Desafortunadamente, lo que quiere lograr (comunicación/sincronización entre bloques) no es estrictamente posible en CUDA. La guía de programación de CUDA establece que "se requieren bloques de subprocesos para ejecutar de forma independiente: debe ser posible ejecutarlos en cualquier orden, en paralelo o en serie". El motivo de esta restricción es permitir flexibilidad en el programador de bloques de hilos y permitir que el código se escale agnósticamente con la cantidad de núcleos. El único método de sincronización interbloque admitido es ejecutar otro kernel: los lanzamientos del kernel (dentro de la misma secuencia) son puntos de sincronización implícitos.

Su código infringe la regla de independencia del bloque porque supone implícitamente que los bloques de hilos del kernel se ejecutan simultáneamente (cf. en paralelo). Pero no hay garantía de que lo hagan. Para ver por qué esto es importante para su código, consideremos una GPU hipotética con solo un núcleo. También asumiremos que solo desea ejecutar dos bloques de hilos. Su núcleo de spinloop se estancará en esta situación. Si el bloque de hilos cero está programado en el núcleo primero, se repetirá siempre cuando llegue a la barrera, porque el bloque de hilos uno nunca tiene la oportunidad de actualizar el contador. Debido a que el bloque de hilos cero nunca se intercambia (los bloques de subprocesos se ejecutan hasta su finalización) priva al bloque de hilos uno de los núcleos mientras gira.

Algunas personas han probado esquemas como el suyo y han tenido éxito porque el planificador ha programado aleatoriamente bloques de tal manera que las suposiciones funcionen. Por ejemplo, hubo un momento en el que se lanzaron tantos bloques de subprocesos como una GPU. SM significa que los bloques se ejecutaron realmente al mismo tiempo. Pero se decepcionaron cuando un cambio en el controlador o el tiempo de ejecución CUDA o GPU invalidaron esa suposición, rompiendo su código.

Para su aplicación, intente encontrar una solución que no dependa de la sincronización entre bloques, porque (salvo un cambio de significación en el modelo de programación CUDA) simplemente no es posible.

Answer 2

Parece el problema de optimización del compilador. No soy bueno con la lectura de PTX-código, pero parece que el compilador ha omitido el -loop while en absoluto (incluso cuando se compila con -O0):

.loc 3 41 0 
cvt.u64.u32  %rd7, %ctaid.x; // Save blockIdx.x to rd7 
ld.param.u64 %rd8, [__cudaparm__Z3sumPiS_S_7Barrier_cache]; 
mov.s32  %r8, %ctaid.x; // Now calculate ouput address 
mul.wide.u32 %rd9, %r8, 4; 
add.u64  %rd10, %rd8, %rd9; 
st.global.s32 [%rd10+0], %r5; // Store result to cache[blockIdx.x] 
.loc 17 128 0 
ld.param.u64 %rd11, [__cudaparm__Z3sumPiS_S_7Barrier_barrier+0]; // Get *count to rd11 
mov.s32  %r9, -1; // put -1 to r9 
atom.global.add.s32  %r10, [%rd11], %r9; // Do AtomicSub, storing the result to r10 (will be unused) 
cvt.u32.u64  %r11, %rd7; // Put blockIdx.x saved in rd7 to r11 
mov.u32  %r12, 0; // Put 0 to r12 
setp.ne.u32  %p3, %r11, %r12; // if(blockIdx.x == 0) 
@%p3 bra $Lt_0_5122; 
ld.param.u64 %rd12, [__cudaparm__Z3sumPiS_S_7Barrier_sum]; 
ld.global.s32 %r13, [%rd12+0]; 
mov.s64  %rd13, %rd8; 
mov.s32  %r14, 0; 

En el caso del código de la CPU, se evita que este tipo de comportamiento al declarar la variable con el prefijo volatile. Pero incluso si declaramos count como int __device__ count (y cambiar apropiadamente el código), añadiendo volatile especificador solo rompe la compilación (con errores Loke argument of type "volatile int *" is incompatible with parameter of type "void *")

Sugiero mirar threadFenceReduction ejemplo de CUDA SDK. Allí están haciendo prácticamente lo mismo que tú, pero el bloque para hacer la suma final se elige en tiempo de ejecución, en lugar de predefinirse, y el while -loop se elimina, porque spin-lock en la variable global debería ser muy lento.

Answer 3

La sincronización de bloque a bloque es posible. Vea esto paper.
El documento no entra en gran detalle sobre cómo funciona, pero se basa en el funcionamiento de __syncthreads(); para crear la barrera de pausa para el bloque actual, ... mientras espera que los otros bloques lleguen al punto de sincronización.

Un elemento que no se menciona en el documento es que la sincronización solo es posible si el número de bloques es lo suficientemente pequeño o si el número de SM es lo suficientemente grande para la tarea en cuestión. es decir, si tiene 4 SM e intenta sincronizar 5 bloques, ... el kernel se estancará.

Con su enfoque, he podido extender una larga tarea en serie entre muchos bloques, ahorrando fácilmente un 30% de tiempo en un enfoque de bloque único. es decir, la sincronización de bloques funcionó para mí.