Accesso alla memoria globale non coalescente causato dall'accesso indiretto in CUDA

Question

Il mio programma CUDA è affetto da un accesso di memoria globale non coalescente. Sebbene il thread idx-th si occupi solo della cella [idx] -th di un array, ci sono molti accessi indiretti alla memoria, come mostrato di seguito.

int idx=blockDim.x*blockIdx.x+threadIdx.x; 

.... = FF[m_front[m_fside[idx]]]; 

Per m_fisde [idx], abbiamo coalizzato accessi, ma ciò che realmente bisogno è FF [m_front [m_fside [idx]]]. C'è un accesso indiretto a due livelli.

Ho cercato di trovare alcuni pattern dei dati in m_front o m_fsied per fare in modo che questo fosse un accesso sequenziale diretto, ma ho scoperto che sono quasi "casuali".

C'è un modo possibile per affrontare questo?

Answer 1

Accelerazione memoria ad accesso casuale globale: Invalidare la linea di cache L1

Fermi e Keplero architetture supportano due tipi di carichi da memoria globale. La memorizzazione nella cache completa è la modalità predefinita , tenta di colpire in L1, quindi in L2, quindi in GMEM e la granularità del carico è di 128 byte. L2-only tenta di colpire in L2, quindi GMEM e la granularità del carico è 32-byte. Per alcuni modelli di accesso casuale, l'efficienza della memoria può essere aumentata invalidando L1 e sfruttando la granularità inferiore di L2. Questo può essere fatto compilando con l'opzione –Xptxas –dlcm=cg su nvcc.

Linee guida generali per accelerare l'accesso di memoria globale: disabilitare il supporto ECC

Fermi e GPU Kepler supporto Error Correcting Code (ECC), ed ECC è abilitato di default. ECC riduce la larghezza di banda della memoria di picco ed è richiesto per migliorare l'integrità dei dati in applicazioni come l'imaging medico e il cluster computing su larga scala. Se non è necessario, è possibile disabilitare per migliorare le prestazioni utilizzando l'utilità nvidia-smi su Linux (vedere link) o tramite il Pannello di controllo sui sistemi Microsoft Windows. Si noti che attivare o disattivare ECC richiede un riavvio per avere effetto.

Linee guida generali per accelerare l'accesso memoria globale su Kepler: l'utilizzo di cache di sola lettura dei dati

Keplero dispone di una cache di 48KB per i dati che è noto per essere di sola lettura per la durata della funzione. L'utilizzo del percorso di sola lettura è utile perché consente di scaricare il percorso della cache Shared/L1 e supporta l'accesso alla memoria non allineata a velocità massima pari a . L'utilizzo del percorso di sola lettura può essere gestito automaticamente dal compilatore (utilizzare la parola chiave const __restrict) o esplicitamente (utilizzare lo standard __ldg()) dal programmatore .