syncthreads condizionali e stallo (o non)

Question

un follow-up a Q: EarlyExit e DroppedThreads

Secondo i link di cui sopra, il codice qui sotto dovrebbe deadlock.
Si prega di spiegare perché questo NON si blocca. (Cuda 5 su un Fermi)

__device__ int add[144]; 
__device__ int result; 

add<<<1,96>>>(); // the calling 

__global__ void add() { 
for(idx=72>>1; idx>0; idx>>=1) { 
    if(thrdIdx < idx) 
    add[thrdIdx]+= add[thrdIdx+idx]; 
    else 
    return; 
    __syncthreads(); 
} 

if(thrdIdx == 0) 
    result= add[0]; 
} 

Answer 1

Questo è tecnicamente un programma mal definito.

La maggior parte, ma non tutte (ad esempio, G80 non lo è), le GPU NVIDIA supportano l'uscita anticipata in questo modo perché l'hardware mantiene un conteggio thread attivo per ogni blocco e questo conteggio viene utilizzato per la sincronizzazione della barriera anziché il thread iniziale conta per il blocco.

Pertanto, quando viene raggiunto il numero __syncthreads() nel codice, l'hardware non attende i thread che sono già stati restituiti e il programma viene eseguito senza deadlock.

Un uso più comune di questo stile è:

__global__ void foo(int n, ...) { 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 
    if (idx >= n) return; 
    ... // do some computation with remaining threads 
} 

Nota importante: conta barriera vengono aggiornate per-ordito (vedi here), non per-thread. Quindi potresti avere il caso in cui, per esempio, solo pochi (o zero) thread ritornano presto. Ciò significa che il conteggio delle barriere non viene decrementato. Tuttavia, finché almeno un filo di ciascun ordito raggiunge la barriera, non si bloccherà.

Quindi, in generale, è necessario utilizzare le barriere con attenzione. Ma in particolare, (semplici) i primi schemi di uscita come questo funzionano.

Modifica: per il caso specifico.

Iterazione Idx == 36: 2 orditi attivi in ​​modo che il conteggio delle uscite di barriera sia 64. Tutti i fili da ordito 0 raggiungono la barriera, conteggio incrementale da 0 a 32. 4 fili dalla barriera di curvatura 1, conteggio incrementale da 32 a 64, e gli orditi 0 e 1 vengono rilasciati dalla barriera. Leggi il link qui sopra per capire perché questo accade.

Iteration Idx == 18: 1 warp attivo in modo che il conteggio delle uscite di barriera sia 32. 18 thread da warp 0 raggiungere barriera, conteggio incrementale da 0 a 32. La barriera è soddisfatta e l'ordito 0 viene rilasciato.

Ecc ...