Utilizzo di costanti con CUDA

Question

Qual è il modo migliore di utilizzare le costanti in CUDA?

Un modo è quello di definire costanti nella memoria costante, come:

// CUDA global constants 
__constant__ int M; 

int main(void) 
{ 
    ... 
    cudaMemcpyToSymbol("M", &M, sizeof(M)); 
    ... 
} 

Un modo alterativa sarebbe usare il preprocessore C:

#define M = ... 

penserei definire costanti con il preprocessore C è molto più veloce. Quali sono quindi i vantaggi dell'utilizzo della memoria costante su un dispositivo CUDA?

Answer 1

1. costanti che sono noti al momento della compilazione devono essere definiti usando macro preprocessore (ad esempio #define) o tramite C/C++ const variabili in ambito globale/file.
2. L'utilizzo di __constant__ memory può essere utile per i programmi che utilizzano determinati valori che non cambiano per la durata del kernel e per i quali sono presenti determinati modelli di accesso (ad esempio, tutti i thread accedono allo stesso valore nello stesso momento). Questo non è migliore o più veloce delle costanti che soddisfano i requisiti del precedente articolo 1.
3. Se il numero di scelte da effettuare da un programma sono relativamente poco numerosi, e queste scelte influenzano l'esecuzione del kernel, un possibile approccio per l'ottimizzazione in fase di compilazione potrebbe essere quella di utilizzare templated code/kernels

Answer 2

regolare C/Costanti di stile C++: le costanti CUDA C (a sua volta una modifica di C99) sono entità assolute di tempo di compilazione. Ciò non sorprende, dato che la quantità di ottimizzazione che si verifica in NVCC è MOLTO coinvolta data la natura dell'elaborazione della GPU.

#define: i macro sono come sempre molto ineleganti ma utili in un pizzico.

Lo specifier variabile __constant__ è, tuttavia, un animale completamente nuovo e qualcosa di un termine improprio a mio parere. Io sedare quella Nvidia ha here nello spazio sottostante:

Il __constant__ qualificazione, eventualmente utilizzati insieme con __device__, dichiara una variabile che:

• risiede nello spazio di memoria costante,
• La durata di un'applicazione,
• È accessibile da tutti i thread all'interno della griglia e dall'host tramite la libreria di runtime (cudaGetSymbolAddress()/ cudaGetSymbolSize()/cudaMemcpyToSymbol()/cudaMemcpyFromSymbol()).

documentazione Nvidia specifica che __constant__ è disponibile a registro velocità livello (quasi zero latenza) a condizione che sia la stessa costante accedano tutte fili di ordito.

Sono dichiarati a livello globale nel codice CUDA. TUTTAVIA, in base all'esperienza personale (e attualmente in corso), bisogna fare attenzione con questo specificatore quando si tratta di compilare separatamente, come la separazione del codice CUDA (.cu e.file cuh) dal codice C/C++ inserendo funzioni wrapper in intestazioni in stile C.

A differenza delle tradizionali variabili "costanti" specificate, tuttavia, queste vengono inizializzate in fase di esecuzione dal codice host che alloca la memoria del dispositivo e infine avvia il kernel. Ripeto che sto attualmente lavorando codice che dimostra questi possono essere impostati in runtime utilizzando cudaMemcpyToSymbol() prima dell'esecuzione del kernel.

Sono abbastanza pratici per non dire altro data la velocità del livello di cache L1 che è garantita per l'accesso.