Cuda atomics cambia bandiera

Ho un pezzo di codice seriale che fa qualcosa di simile

if( ! variable ) { do some initialization here variable = true; } 

Capisco che questo funziona perfettamente in seriale e verrà eseguito una sola volta. Quale operazione atomica sarebbe quella corretta qui in CUDA?

Mi sembra che quello che vuoi sia una “sezione critica” nel tuo codice. Una sezione critica consente a un thread di eseguire una sequenza di istruzioni impedendo a qualsiasi altro thread o threadblock di eseguire tali istruzioni.

Una sezione critica può essere utilizzata per controllare l’accesso a un’area di memoria, ad esempio, in modo da consentire un accesso senza conflitti a quell’area da un singolo thread.

Gli atomici da soli possono essere usati solo per un’operazione molto limitata, fondamentalmente singola, su una singola variabile. Ma l’atomica può essere usata per build una sezione critica.

Dovresti usare il seguente codice nel tuo kernel per controllare l’accesso al thread in una sezione critica:

 __syncthreads(); if (threadIdx.x == 0) acquire_semaphore(&sem); __syncthreads(); //begin critical section // ... your critical section code goes here //end critical section __syncthreads(); if (threadIdx.x == 0) release_semaphore(&sem); __syncthreads(); 

Prima del kernel definire queste funzioni di supporto e variabile del dispositivo:

 __device__ volatile int sem = 0; __device__ void acquire_semaphore(volatile int *lock){ while (atomicCAS((int *)lock, 0, 1) != 0); } __device__ void release_semaphore(volatile int *lock){ *lock = 0; __threadfence(); } 

Ho testato e utilizzato con successo il codice di cui sopra. Si noti che essenzialmente esegue l’ arbitraggio tra threadblock utilizzando thread 0 in ciascun threadblock come richiedente. È necessario condizionare ulteriormente (ad es. if (threadIdx.x < ...) ) il proprio codice di sezione critico se si desidera un solo thread nel threadblock vincente per eseguire il codice di sezione critico.

Avere più thread all'interno di un warp arbitrato per un semaforo presenta ulteriori complessità, quindi non consiglio questo approccio. Invece, fare in modo che ciascun threadblock sia arbitrario come ho mostrato qui e quindi controllare il comportamento all'interno del threadblock vincente utilizzando i comuni metodi di comunicazione / sincronizzazione a __syncthreads() ad esempio __syncthreads() , memoria condivisa, ecc.)

Si noti che questa metodologia sarà costosa per le prestazioni. Dovresti utilizzare solo le sezioni critiche quando non riesci a capire come parallelizzare altrimenti il ​​tuo algoritmo.

Infine, una parola di avvertimento. Come in qualsiasi architettura parallela filettata, l'uso improprio di sezioni critiche può portare a deadlock. In particolare, fare ipotesi sull'ordine di esecuzione di threadblock e / o warp all'interno di un threadblock è un approccio errato.