OpenCL - Perché utilizzare READ_ONLY o WRITE_ONLY Buffer

Question

In OpenCL, ci sono prestazioni vantaggiose per contrassegnare i buffer come READ_ONLY o WRITE_ONLY?

Questo kernel è quello che vedo spesso (una è READ_ONLY e B è WRITE_ONLY):

__kernel void two_buffer_double(__global float* a, __global float* b) 
{ 
    int i = get_global_id(0); 
    b[i] = a[i] * 2; 
} 

Questo kernel sembra meglio, perché usa meno memoria globale (una è READ_WRITE):

__kernel void one_buffer_double(__global float* a) 
{ 
    int i = get_global_id(0); 
    a[i] = a[i] * 2; 
} 

I flag READ_ONLY e WRITE_ONLY esistono per facilitare il debug e gli errori di acquisizione?

Answer 1

Per rispondere dritto in avanti alla sua domanda direi: No, queste bandiere non solo esistono per aiutare con il debug e gli errori di cattura. Tuttavia è difficile dare un riferimento su come questi flag sono usati da qualsiasi implementazione e su come influenzano le performance.

mia comprensione (purtroppo non sostenuta da alcuna documentazione) è che quando si utilizzano questi flag si mette più vincoli su come verranno utilizzati i buffer e quindi si può aiutare il runtime/driver/compilatore per fare alcune ipotesi che potrebbe migliorare le prestazioni. Ad esempio, , immagino che sia che non ci dovrebbero essere preoccupazioni per la consistenza della memoria con un buffer di sola lettura mentre un kernel lo sta usando poiché i workitems non dovrebbero scrivere in esso. Pertanto alcuni controlli potrebbero essere saltati ... anche se in Opencl si suppone che tu ti prenda cura di te stesso usando le barriere e così via.

Si noti inoltre che dal momento che Opencl 1.2 sono stati introdotti altri flag correlati questa volta a come l'host deve accedere ai buffer. Ci sono:

CL_MEM_HOST_NO_ACCESS, 
CL_MEM_HOST_{READ, WRITE}_ONLY, 
CL_MEM_{USE, ALLOC, COPY}_HOST_PTR 

Sto indovinando che ancora una volta si deve aiutare il popolo di attuazione OpenCL per migliorare le prestazioni, ma credo che avremmo bisogno del contributo di alcuni esperti di AMD o NVIDIA.

Si prega di notare che tutto ciò che ho detto finora sono solo i miei pensieri e non sono basati su alcuna documentazione seria (non sono riuscito a trovarne).

D'altra parte vi posso dire per certo che la norma non costretti una sola lettura buffer da nello spazio costante come detto @Quonux. Potrebbe essere che alcune implementazioni lo facciano per piccoli buffer. Non dimentichiamo che la memoria dello spazio costante è piccola, quindi è possibile avere solo il buffer di lettura troppo grande per adattarsi.L'unico modo per assicurarsi che un buffer si trovi nella memoria dello spazio costante è usare la parola chiave costante nel codice del kernel come spiegato here. Ovviamente nel lato host, se si desidera utilizzare il buffer costante, è necessario utilizzare il flag di sola lettura.

Answer 2

Dipende,

una posizione READ_ONLY __global memoria è contenuto in "/ Constant cache di memoria dati globali" che è molto più veloce rispetto alla cache di normale o RAM su una GPU (vedi here), su una CPU esso non importa.

Non conosco alcun vantaggio di WRITE_ONLY, forse aiuta anche perché la GPU sa che è in grado di eseguire lo streaming dei dati senza necessità di memorizzazione nella cache.

Basta andare a misurare se il vostro dubbi ...

Answer 3

Nota che in realtà ce ne sono due tipi. Avete CL_MEM_READ_ONLY, CL_MEM_WRITE_ONLY e CL_MEM_READ_WRITE quando assegnate i vostri buffer ma poi avete anche __read_only, __write_only e __read_write per decorare i vostri puntatori nel codice del kernel con.

Questi potrebbero essere utilizzati sia per l'ottimizzazione che per il controllo degli errori. Diamo prima un'occhiata alle prestazioni. Se viene rilevato un buffer di sola scrittura, le scritture non devono essere memorizzate nella cache (come nella cache di scrittura), risparmiando più cache per le letture. Questo dipende molto dall'hardware della GPU e almeno l'hardware NVIDIA ha le istruzioni necessarie per implementarlo (i modificatori .cs e .lu). È possibile fare riferimento al loro PTX ISA. Non ho visto alcuna prova del compilatore che opera questa ottimizzazione, ad esempio:

__kernel void Memset4(__global __write_only unsigned int *p_dest, 
    const unsigned int n_dword_num) 
{ 
    unsigned int i = get_global_id(0); 
    if(i < n_dword_num) 
     p_dest[i] = 0; // this 
} 

viene compilato come:

st.global.u32 [%r10], %r11; // no cache operation specified 

questo ha un senso come CUDA non ha equivalenti per quelle qualificazioni in modo che il compilatore molto probabilmente ignora silenziosamente quelli. Ma non fa male metterli lì, potremmo diventare più fortunati in futuro. In CUDA, alcune di queste funzionalità sono esposte utilizzando la funzione __ldg e utilizzando i flag del compilatore per attivare/disattivare la memorizzazione nella cache dei trasferimenti globali di memoria in L1 (-Xptxas -dlcm=cg). È inoltre possibile utilizzare sempre asm se si scopre che il bypass della cache produce un vantaggio importante.

Per quanto riguarda il controllo degli errori, la scrittura su un buffer di sola lettura viene prontamente evitata utilizzando lo specificatore const nella dichiarazione del kernel. La disabilitazione della lettura da un buffer di sola scrittura non è possibile nella "C" pura.

Un altro possibile ottimizzazione si verifica quando si mappano i buffer nella memoria host. Quando si esegue il mapping di un buffer CL_MEM_READ_ONLY, la regione mappata potrebbe non essere inizializzata poiché l'host scriverà solo in quella memoria, in modo che il dispositivo possa leggerlo solo. Allo stesso modo, quando si annulla la mappatura di un buffer CL_MEM_WRITE_ONLY, il driver non deve copiare il contenuto (potenzialmente modificato dall'host) dalla memoria dell'host alla memoria del dispositivo. Non ho misurato questo.

Come nota a margine, ho provato ad utilizzare:

inline unsigned int n_StreamingLoad(__global __read_only const unsigned int *p_src) 
{ 
#ifdef NVIDIA 
    unsigned int n_result; 
    asm("ld.global.cs.u32 %r0, [%r1];" : "=r" (n_result) : "r" (p_src)); 
    return n_result; 
#else // NVIDIA 
    return *p_src; // generic 
#endif // NVIDIA 
} 

inline void StreamingWrite(__global __write_only unsigned int *p_dest, const unsigned int n_value) 
{ 
#ifdef NVIDIA 
    asm("st.global.cs.u32 [%r0], %r1;" : : "r" (p_dest), "r" (n_value) : "memory"); 
#else // NVIDIA 
    *p_dest = n_value; // generic 
#endif // NVIDIA 
} 

che vi dà circa 15 extra GB/sec anche su un semplice kernel memcpy con sm_35 dispositivi (testato su GTX 780 e K40). Non ho visto un'accelerazione notevole su sm_30 (non sono sicuro che sia stato pensato anche per essere supportato lì - anche se le istruzioni non sono state rimosse da ptx). Nota che devi definire tu stesso lo NVIDIA (o vedere Detect OpenCL device vendor in kernel code).