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);
Forse dovresti includere riferimenti specifici ai documenti che parlano di "thread persistenti". –
terminato: uno studio dei thread persistenti Programmazione GPU di stile per carichi di lavoro GPGPU – AmineMs
È possibile ottenere ulteriori utili informazioni se [si guarda la presentazione] (http://www.gputechconf.com/gtcnew/on-demand-gtc.php? sessionTopic = & searchByKeyword = S0157 & submit = & selezionare = + & sessionEvent = & sessionYear = 2012 & sessionFormat = # 1337). Ho letto che è più o meno come lo descrivete, cioè un kernel che non esce e continua a interrogare il lavoro da una coda nella memoria globale. –