CUDA. Come srotolare i primi 32 thread in modo che vengano eseguiti in parallelo?

Question

So che "ogni warp contiene thread di ID thread incrementali consecutivi con il primo ordito contenente il thread 0", quindi i primi 32 thread dovrebbero essere nel primo warp. Inoltre, so che tutti i thread in un warp vengono eseguiti simultaneamente su qualsiasi multiprocessore streaming disponibile.

Come ho capito, per questo motivo non è necessario sincronizzare i thread se viene eseguito un solo ordito. Ma il codice qui sotto produce una risposta errata se rimuovo uno qualsiasi dei __syncthreads() nel penultimo blocco if. Ho provato a trovare la causa ma non ho trovato nulla. Spero davvero il tuo aiuto, quindi potresti dirmi cosa c'è di sbagliato in questo codice? Perché non posso lasciare solo l'ultima __syncthreads() e ottenere la risposta giusta?

#define BLOCK_SIZE 128 

__global__ void reduce (int * inData, int * outData) 
{ 
__shared__ int data [BLOCK_SIZE]; 
int tid = threadIdx.x; 
int i = blockIdx.x * blockDim.x + threadIdx.x; 

data [tid] = inData [i] + inData [i + blockDim.x/2 ]; 
__syncthreads(); 

for (int s = blockDim.x/4; s > 32; s >>= 1) 
{ 
    if (tid < s) 
    data [tid] += data [tid + s]; 
    __syncthreads(); 
} 

if (tid < 32) 
{ 
    data [tid] += data [tid + 32]; 
    __syncthreads(); 
    data [tid] += data [tid + 16]; 
    __syncthreads(); 
    data [tid] += data [tid + 8]; 
    __syncthreads(); 
    data [tid] += data [tid + 4]; 
    __syncthreads(); 
    data [tid] += data [tid + 2]; 
    __syncthreads(); 
    data [tid] += data [tid + 1]; 
    __syncthreads(); 
} 
if (tid == 0) 
    outData [blockIdx.x] = data [0]; 
} 

void main() 
{ 
... 
reduce<<<dim3(128), dim3(128)>>>(dev_data, dev_res); 
... 
} 

P.S. Sto usando GT560Ti

Answer 1

Si dovrebbe dichiarare la variabile di memoria condivisa come volatile:

__shared__ volatile int data [BLOCK_SIZE]; 

Il problema che state vedendo è un artefatto dell'architettura Fermi e ottimizzazione del compilatore. L'architettura Fermi manca di istruzioni per operare direttamente sulla memoria condivisa (erano presenti nella serie G80/90/GT200). Quindi tutto viene caricato per registrare, manipolare e archiviare nella memoria condivisa. Ma il compilatore è libero di dedurre che il codice potrebbe essere reso più veloce se una serie di operazioni sono state organizzate nel registro, senza carichi intermedi e negozi da/verso la memoria condivisa. Questo è perfettamente soddisfacente eccetto quando ci si basa sulla sincronizzazione implicita di thread all'interno dello stesso warp che manipola la memoria condivisa, come in questo tipo di codice di riduzione.

Dichiarando il buffer di memoria condivisa come volatile, si costringe il compilatore a imporre la scrittura della memoria condivisa dopo ogni fase della riduzione e viene ripristinata la sincronizzazione implicita dei dati tra i thread all'interno dell'ordito.

Questo problema è discusso nelle note di programmazione per Fermi che vengono spedite (o forse spedite) con il toolkit CUDA.