Barriera interblocco su CUDA

Question

Desidero implementare una barriera Inter-block su CUDA, ma riscontrare un problema serio.

Non riesco a capire perché non funziona.

#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; 
} 

Infatti, anche se riscrivo l'attesa() come il seguente

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

il programma esce normalmente. Ma mi aspetto di ottenere un ciclo infinito in questo caso.

Answer 1

Sfortunatamente, ciò che si desidera ottenere (comunicazione/sincronizzazione tra blocchi) non è strettamente possibile in CUDA. La guida alla programmazione CUDA afferma che "i blocchi di thread sono necessari per l'esecuzione indipendente: deve essere possibile eseguirli in qualsiasi ordine, in parallelo o in serie". Il motivo di questa limitazione è di consentire flessibilità nel programmatore del blocco di thread e di consentire al codice di scalare agnosticamente il numero di core. L'unico metodo di sincronizzazione tra blocchi supportato è quello di avviare un altro kernel: i kernel launches (all'interno dello stesso flusso) sono punti di sincronizzazione impliciti.

Il codice viola la regola di indipendenza del blocco perché presuppone implicitamente che i blocchi di thread del kernel vengano eseguiti contemporaneamente (vedere in parallelo). Ma non c'è alcuna garanzia che lo facciano. Per capire perché questo è importante per il tuo codice, consideriamo un'ipotetica GPU con un solo core. Supponiamo anche che tu voglia solo lanciare due blocchi di thread. Il tuo kernel spinloop in realtà si bloccherà in questa situazione. Se il blocco zero del blocco è pianificato per primo sul core, verrà interrotto per sempre quando arriva alla barriera, perché il blocco thread non ha mai la possibilità di aggiornare il contatore. Poiché il blocco del blocco zero non viene mai sostituito (i blocchi di thread vengono eseguiti al termine del loro completamento), il thread di thread viene bloccato da uno dei componenti principali durante la rotazione.

Alcune persone hanno provato schemi come il tuo e hanno avuto successo perché lo scheduler è riuscito a programmare in modo sereno i blocchi in modo tale che le ipotesi funzionassero. Ad esempio, c'è stato un momento in cui l'avvio di tutti i blocchi di thread in cui una GPU ha SM significava che i blocchi venivano effettivamente eseguiti contemporaneamente. Ma sono rimasti delusi quando una modifica al driver o al runtime o GPU CUDA ha invalidato tale ipotesi, infrangendo il loro codice.

Per l'applicazione, provare a trovare una soluzione che non dipenda dalla sincronizzazione tra blocchi, poiché (salvo una modifica di significazione al modello di programmazione CUDA) non è possibile.

Answer 2

Sembra un problema di ottimizzazioni del compilatore. Io non sono bravo con la lettura PTX-code, ma sembra che il compilatore ha omesso il while -loop a tutti (anche quando si 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; 

In caso di codice della CPU, viene impedito tale comportamento dichiarando la variabile con il prefisso volatile. Ma anche se si dichiara count come int __device__ count (e opportunamente modificare il codice), aggiungendo volatile identificatore solo rompe la compilazione (con errori di Loke argument of type "volatile int *" is incompatible with parameter of type "void *")

Suggerisco guardando threadFenceReduction esempio da CUDA SDK. Loro stanno facendo più o meno lo stesso di te, ma il blocco per fare la sommatoria finale viene scelto in runtime, piuttosto che predefinito, e lo while -loop viene eliminato, perché spin-lock sulla variabile globale dovrebbe essere molto lento.

Answer 3

È possibile bloccare la sincronizzazione. Vedi questo paper.
Il documento non è molto dettagliato su come funziona, ma si basa sul funzionamento di __syncthreads(); per creare la barriera di pausa per il blocco corrente, ... mentre si attende che gli altri blocchi raggiungano il punto di sincronizzazione.

Un elemento che non è indicato nella carta è che la sincronizzazione è possibile solo se il numero di blocchi è abbastanza piccolo o il numero di SM è abbastanza grande per l'attività a portata di mano. Ad esempio, se hai 4 SM e stai provando a sincronizzare 5 blocchi, il kernel si bloccherà.

Con il loro approccio, sono stato in grado di distribuire un lungo compito seriale tra molti blocchi, risparmiando facilmente il 30% di tempo su un approccio a blocco singolo. Ad esempio, la sincronizzazione dei blocchi ha funzionato per me.