CUDA disabilita la cache L1 solo per una variabile

Question

Esiste un modo per i dispositivi CUDA 2.0 di disabilitare la cache L1 solo per una variabile specifica? So che è possibile disabilitare la cache L1 in fase di compilazione aggiungendo il flag -Xptxas -dlcm=cg a nvcc per tutte le operazioni di memoria. Tuttavia, voglio disabilitare la cache solo per le letture della memoria su una specifica variabile globale in modo che tutto il resto della memoria sia letto per passare attraverso la cache L1.

Sulla base di una ricerca eseguita nel web, una possibile soluzione è tramite il codice dell'assieme PTX.

Answer 1

come detto sopra è possibile utilizzare in linea PTX, ecco un esempio:

__device__ __inline__ double ld_gbl_cg(const double *addr) { 
    double return_value; 
    asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr)); 
    return return_value; 
} 

Si può facilmente variare a questo scambiando .f64 per .f32 (float) o .s32 (int), ecc, il vincolo di return_value "= d" per "= f" (float) o "= r" (int) etc. Si noti che l'ultimo vincolo prima (addr) - "l" - indica l'indirizzamento a 64 bit, se si utilizza l'indirizzamento a 32 bit, dovrebbe essere "r".

Answer 2

PTX in linea può essere utilizzato per caricare e memorizzare la variabile. Le istruzioni ld.cg e st.cg memorizzano solo i dati nella cache in L2. Gli operatori cache sono descritti nella sezione 8.7.8.1 Cache Operators del documento PTX ISA 2.3. Le istruzioni o gli interessi sono ld e st. Inline PTX è descritto in Using Inline PTX Assembly in CUDA.

Answer 3

Se si dichiara la variabile come volatile, verrà memorizzata nella cache L2 solo sulle GPU Fermi. Si noti che alcune ottimizzazioni del compilatore, come la rimozione di carichi ripetuti, non vengono eseguite su variabili volatili perché il compilatore presuppone che possano essere scritte da un altro thread.