Utilizzo memoria costante nel codice CUDA

Question

Non riesco a capirlo da solo, qual è il modo migliore per garantire che la memoria utilizzata nel mio kernel sia costante. C'è una domanda simile a http://stackoverflow...r-pleasant-way. Sto lavorando con GTX580 e sto compilando solo per funzionalità 2.0. Il mio kernel assomiglia

__global__ Foo(const int *src, float *result) {...} 

eseguo il seguente codice sull'host:

cudaMalloc(src, size); 
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice); 
Foo<<<...>>>(src, result); 

il modo alternativo è quello di aggiungere

__constant__ src[size]; 

a .CU di file, rimuovere src puntatore da il kernel ed esegui

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice); 
Foo<<<...>>>(result); 

Sono questi due modi equivalenti o il primo non garantisce l'utilizzo della memoria costante anziché della memoria globale? dimensioni cambia dinamicamente, quindi il secondo modo non è utile nel mio caso.

Answer 1

Il secondo modo è l'unico modo per garantire che l'array sia compilato in memoria costante CUDA e accessibile correttamente tramite la cache di memoria costante. Ma dovresti chiederti come si accede ai contenuti di quell'array all'interno di un blocco di thread. Se ogni thread accederà all'array in modo uniforme, ci sarà un vantaggio in termini di prestazioni nell'utilizzo della memoria costante, poiché esiste un meccanismo di trasmissione dalla cache della memoria costante (salva anche la larghezza di banda della memoria globale perché la memoria costante viene archiviata nella DRAM preassegnata e nella cache riduce il conteggio delle transazioni DRAM). Ma se l'accesso è casuale, può esserci serializzazione dell'accesso alla memoria locale che influirà negativamente sulle prestazioni.

Le cose tipiche che potrebbero essere valide per la memoria __constant__ sono i coefficienti del modello, i pesi e altri valori costanti che devono essere impostati in fase di esecuzione. Sulle GPU Fermi, l'elenco degli argomenti del kernel è memorizzato nella memoria costante, per esempio. Ma se il contenuto è un accesso non uniforme e il tipo o la dimensione dei membri non è costante da chiamata a chiamata, è preferibile la memoria globale normale.

Inoltre, tenere presente che esiste un limite di 64 KB di memoria costante per contesto GPU, quindi non è pratico memorizzare grandi quantità di dati nella memoria costante. Se hai bisogno di molta memoria di sola lettura con cache, potrebbe valere la pena provare a legare i dati a una texture e vedere come è la performance. Sulle schede pre-Fermi, solitamente fornisce un utile guadagno in termini di prestazioni, su Fermi i risultati possono essere meno prevedibili rispetto alla memoria globale a causa del miglioramento del layout della cache in tale architettura.

Answer 2

Il primo metodo garantisce che la memoria sia costante all'interno della funzione Foo. I due non sono equivalenti, il secondo garantisce di contenerlo dopo l'inizializzazione. Se hai bisogno di dinamiche di te per usare qualcosa di simile alla prima.