Thread persistenti in OpenCL e CUDA

Question

Ho letto alcuni articoli che parlano di "thread persistenti" per GPGPU, ma non lo capisco davvero. Qualcuno può darmi un esempio o mostrarmi l'uso di questa moda di programmazione?

Quello che tengo nella mia mente dopo aver letto e googling "le discussioni persistenti":

Presistent Le discussioni si tratta di non più di un ciclo while che mantengono thread in esecuzione e calcolando un sacco di gruppo di opere.

È corretto? Grazie in anticipo

Riferimento: http://www.idav.ucdavis.edu/publications/print_pub?pub_id=1089 http://developer.download.nvidia.com/GTC/PDF/GTC2012/PresentationPDF/S0157-GTC2012-Persistent-Threads-Computing.pdf

Answer 1

CUDA sfrutta il modello di programmazione SIMD (Single Instruction Multiple Data). I thread di calcolo sono organizzati in blocchi e i blocchi di thread sono assegnati a un diverso Streaming Multiprocessor (SM). L'esecuzione di un blocco di thread su un SM viene eseguita disponendo i thread in warps di thread 32: ogni warp opera in lock-step ed esegue esattamente le stesse istruzioni su dati diversi. .

In genere, per riempire la GPU, il kernel viene avviato con molti più blocchi che possono essere effettivamente ospitati sugli SM. Poiché non tutti i blocchi possono essere ospitati su un SM, un pianificatore di lavoro esegue un cambio di contesto quando un blocco ha terminato il calcolo. Dovrebbe essere notato che la commutazione dei blocchi è gestita interamente in hardware dallo scheduler, e il programmatore non ha modo di influenzare il modo in cui i blocchi sono programmati sull'SM. Questo espone una limitazione per tutti quegli algoritmi che non si adattano perfettamente a un modello di programmazione SIMD e per i quali esiste uno squilibrio di lavoro. Infatti, un blocco A non verrà sostituito da un altro blocco B sullo stesso SM fino a quando l'ultimo thread del blocco A non avrà terminato l'esecuzione.

Anche se CUDA non espone l'utilità di pianificazione hardware al programmatore, lo stile di thread persistenti ignora lo scheduler hardware facendo affidamento su una coda di lavoro. Al termine di un blocco, controlla la coda per ulteriori lavori e continua a farlo fino a quando non viene lasciato alcun lavoro, a quel punto il blocco si ritira. In questo modo, il kernel viene lanciato con un numero di blocchi pari al numero di SM disponibili.

I fili persistenti tecnica è meglio illustrato dal seguente esempio, che è stato preso dalla presentazione

“GPGPU” computing and the CUDA/OpenCL Programming Model

Un altro esempio più dettagliato è disponibile nel documento

Understanding the efficiency of ray traversal on GPUs

// Persistent thread: Run until work is done, processing multiple work per thread 
// rather than just one. Terminates when no more work is available 

// count represents the number of data to be processed 

__global__ void persistent(int* ahead, int* bhead, int count, float* a, float* b) 
{ 
    int local_input_data_index, local_output_data_index; 
while ((local_input_data_index = read_and_increment(ahead)) < count) 
{         
     load_locally(a[local_input_data_index]); 

     do_work_with_locally_loaded_data(); 

     int out_index = read_and_increment(bhead); 

     write_result(b[out_index]); 
    } 
} 

// Launch exactly enough threads to fill up machine (to achieve sufficient parallelism 
// and latency hiding) 
persistent<<numBlocks,blockSize>>(ahead_addr, bhead_addr, total_count, A, B); 

Answer 2

abbastanza facile da capire. Di solito ogni articolo di lavoro ha elaborato una piccola quantità di lavoro. Se si desidera salvare il tempo di salvataggio del gruppo di lavoro, è possibile consentire a un elemento di lavoro di elaborare molto lavoro utilizzando un ciclo. Ad esempio, hai un'immagine, ed è 1920x1080, hai un workitem di 1920, e ogni oggetto di lavoro elabora una colonna di 1080 pixel usando il loop.