Come utilizzare la memoria costante CUDA in un programmatore in modo piacevole?

Question

Sto lavorando a un'app di crunch numerico utilizzando il framework CUDA. Ho alcuni dati statici che devono essere accessibili a tutte le discussioni, così ho messo in memoria costante simili:

__device__ __constant__ CaseParams deviceCaseParams; 

uso il cudaMemcpyToSymbol chiamata per trasferire questi params dall'host al dispositivo:

void copyMetaData(CaseParams* caseParams) 
{ 
    cudaMemcpyToSymbol("deviceCaseParams", caseParams, sizeof(CaseParams)); 
} 

che funziona.

Ad ogni modo, sembra (per tentativi ed errori, e anche dalla lettura di messaggi in rete) che per qualche motivo insoddisfacente, la dichiarazione di deviceCaseParams e l'operazione di copia di esso (la chiamata a cudaMemcpyToSymbol) devono essere nello stesso file. Al momento ho questi due in un file .cu, ma voglio davvero avere il parametro struct in un file .cuh in modo che qualsiasi implementazione possa vederlo se lo desidera. Ciò significa che devo anche avere la funzione copyMetaData nel file di intestazione, ma questo mette a bada il collegamento (simbolo già definito) poiché entrambi i file .cpp e .cu includono questa intestazione (e quindi sia il compilatore MS C++ che nvcc lo compila).

Qualcuno ha qualche consiglio sul design qui?

Aggiornamento: Vedere i commenti

Answer 1

Con una CUDA up-to-date (ad esempio 3.2) si dovrebbe essere in grado di fare la memcpy dall'interno di un'unità di traduzione diversa, se stai cercando il simbolo a runtime (cioè passando una stringa come primo argomento a cudaMemcpyToSymbol come nel tuo esempio).

Inoltre, con i dispositivi di classe Fermi è possibile solo mallocare la memoria (cudaMalloc), copiare nella memoria del dispositivo e quindi passare l'argomento come puntatore const. Il compilatore riconoscerà se si accede ai dati in modo uniforme attraverso gli orditi e, in caso affermativo, utilizzerà la cache costante. Vedi la Guida alla Programmazione CUDA per maggiori informazioni. Nota: è necessario compilare con -arch=sm_20.

Answer 2

Se si utilizza CUDA pre-Fermi, si sarà scoperto che questo problema non si applica solo alla memoria costante, si applica a tutto ciò che si desidera sul lato CUDA delle cose. Gli unici due modi che ho trovato intorno a questo sono a uno:

1. Scrivere tutto CUDA in un unico file (br), o
2. Se avete bisogno di uscire di codice in file separati, limitarvi a intestazioni quale include il tuo singolo file .cu.

Se è necessario condividere codice tra CUDA e C/C++ o disporre di codice comune condiviso tra progetti, l'opzione 2 è l'unica scelta. Sembra molto innaturale, ma risolve il problema. È ancora possibile strutturare il codice, ma non in modo tipico C. L'overhead principale è che ogni volta che fai una compilazione compili il di tutto. Il lato positivo di questo (che penso sia probabilmente il motivo per cui funziona in questo modo) è che il compilatore CUDA ha accesso a tutto il codice sorgente in un colpo che è buono per l'ottimizzazione.