Quando utilizzare volatile con memoria CUDA condivisa

Question

In quali circostanze si deve utilizzare la parola chiave volatile con una memoria condivisa del kernel CUDA? Capisco che volatile dice al compilatore di non memorizzare nella cache tutti i valori, ma la mia domanda è circa il comportamento con una matrice comune:

__shared__ float products[THREADS_PER_ACTION]; 

// some computation 
products[threadIdx.x] = localSum; 

// wait for everyone to finish their computation 
__syncthreads(); 

// then a (basic, ugly) reduction: 
if (threadIdx.x == 0) { 
    float globalSum = 0.0f; 
    for (i = 0; i < THREADS_PER_ACTION; i++) 
     globalSum += products[i]; 
} 

Devo products ad essere volatili in questo caso? Ogni voce dell'array è accessibile solo da un singolo thread, tranne alla fine, dove tutto viene letto dal thread 0. È possibile che il compilatore possa memorizzare l'intero array, e quindi ho bisogno che sia il volatile, o lo memorizzerà solo nella cache elementi?

Grazie!

Answer 1

Se non si dichiara un array condiviso come volatile, il compilatore è libero di ottimizzare le posizioni nella memoria condivisa individuandole in registri (il cui ambito è specifico per un singolo thread), per qualsiasi thread, a sua scelta . Questo è vero se si accede a quel particolare elemento condiviso da un solo thread o meno. Pertanto, se si utilizza la memoria condivisa come veicolo di comunicazione tra i thread di un blocco, è meglio dichiararlo volatile.

Ovviamente se ogni thread accede solo ai propri elementi di memoria condivisa, e mai a quelli associati ad un altro thread, allora questo non ha importanza e l'ottimizzazione del compilatore non interromperà nulla.

Nel tuo caso, dove si dispone di una sezione di codice in cui ogni thread accede ai propri elementi di memoria condivisa e l'unico accesso tra thread si verifica in una posizione ben compresa, è possibile utilizzare un memory fence function per forzare il compilatore per rimuovere qualsiasi valore temporaneamente memorizzato nei registri, tornare alla matrice condivisa. Quindi potresti pensare che __threadfence_block() potrebbe essere utile, ma nel tuo caso, __syncthreads()already has memory-fencing functionality built in. Pertanto, la tua chiamata __syncthreads() è sufficiente per forzare la sincronizzazione dei thread e per forzare i valori memorizzati nella cache di registrazione nella memoria condivisa per essere reinseriti nella memoria condivisa.

A proposito, se quella riduzione alla fine del codice è di preoccupazione per le prestazioni, si potrebbe considerare l'utilizzo di un metodo di riduzione parallelo per accelerarlo.