CUDA/openCL; rami riscrittura come senza ramificazioni espressione

Question

maggior parte del tempo un ramo è richiesto in un programma CUDA o OpenCL, come:

for (int i=0; i<width; i++) 
{ 
    if(i % threadIdx.x == 0) 
    quantity += i*i; 
} 

il codice può sempre (o almeno la maggior parte del tempo) essere riscritta in stile non ramificati:

for (int i=0; i<width; i++) 
{ 
    quantity += i*i* (i % threadIdx.x != 0); 
} 

il compromesso sembra essere sia in esecuzione in un singolo slot ordito contro facendo più calcoli su tutti i fili (nel secondo caso, la somma viene eseguita sempre, solo che a volte il valore è zero)

Supponendo che le operazioni di ramificazione occuperanno più intervalli di curvatura per ogni ramo possibile, ci si aspetterebbe che il secondo sia costantemente migliore del primo, Ora la mia domanda è; Posso fare affidamento sul compilatore per ottimizzare 1) in 2) ogni volta che ha senso, o non esiste un criterio ampiamente applicabile, il che implica che non si può decidere in generale quale sia il migliore senza provare e profilare?

Answer 1

Non ho molti ricordi su CUDA, ma perché non parallelizzi il tuo ciclo? Dovresti usare le operazioni atomiche [1] per aggiungere il tuo calcolo. Spero che questo ti possa aiutare! Scusa se non è il caso.

1. Atomic Operazioni: http://supercomputingblog.com/cuda/cuda-tutorial-4-atomic-operations/

Answer 2

Nella mia esperienza - è totalmente fino al compilatore-scrittori di ottimizzare questo tipo di casi limite.

Quindi, posso pensare a casi in cui 1) non può essere impostato su 2)? Eccone uno: ho scritto kernel dove è stato più efficiente eseguire certi pezzi dei calcoli ogni 10 thread o qualcosa del genere, nel qual caso tale ottimizzazione non può essere dedotta, anche se esiste un'operazione matematica (una divisione e una sottrazione) che può produrre lo stesso risultato indipendentemente dal condizionale rispetto a "correre su tutti, ma produrre risultati zero".

Tuttavia, anche se il controllo di threadId == 0 è uno scenario abbastanza comune, non ho idea se sia effettivamente ottimizzato per. Scommetto che dipende dall'implementazione E anche dal dispositivo stesso (CPU vs GPU).

Dovrete provarlo per scoprire veramente cosa funziona meglio, non solo per la ragione sopra, ma anche perché lo scheduler di lavoro potrebbe comportarsi in modo diverso in base a quanto è costoso pianificare/avviare/interrompere un set di thread al contrario di averli tutti eseguiti (e la maggior parte fornisce un risultato di zero/identità).

Spero che questo aiuti!

Answer 3

Le operazioni di modulo sono ragionevolmente costose: sono ragionevolmente sicuro che l'aggiunta del modulo richiederebbe più tempo di una semplice istruzione che esegue solo 1 thread. La tua singola dichiarazione di diramazione, una if senza else, bloccherà solo gli altri thread mentre quella se la statistica è in esecuzione. Poiché gpus è ottimizzato per il cambio di contesto molto veloce, ci dovrebbe essere un costo molto basso per questo.

Si sconsiglia tuttavia l'uso di istruzioni di ramificazione lunghe: un eccessivo calcolo seriale sulla GPU (ovvero un thread che esegue tutto il lavoro) annulla il vantaggio del parallelismo.