CUDA host e dispositivo che utilizza la stessa memoria __constant__

Question

Ho una funzione dispositivo/host che utilizza la memoria costante. Funziona correttamente sul dispositivo, ma su host sembra che questa memoria non sia inizializzata.

#include <iostream> 
#include <stdio.h> 


const __constant__ double vals[2] = { 0.0, 1000.0 }; 

__device__ __host__ double f(size_t i) 
{ 
    return vals[i]; 
} 

__global__ void kern() 
{ 
    printf("vals[%d] = %lf\n", threadIdx.x, vals[threadIdx.x]); 
} 

int main() { 
    std::cerr << f(0) << " " << f(1) << std::endl; 
    kern<<<1, 2>>>(); 
    cudaThreadSynchronize(); 
} 

Questo stampa (richiede CC 2.0 o superiore)

0 0 
vals[0] = 0.000000 
vals[1] = 1000.000000 

Qual è il problema e come posso ottenere entrambe le costanti di memoria del dispositivo e di accoglienza inizializzato allo stesso tempo?

Answer 1

Poiché CygnusX1 ha frainteso ciò che intendevo nel mio commento sulla risposta di MurphEngineer, forse dovrei pubblicare la mia risposta. Quello che ho dire era questo:

__constant__ double dc_vals[2] = { 0.0, 1000.0 }; 
     const double hc_vals[2] = { 0.0, 1000.0 }; 

__device__ __host__ double f(size_t i) 
{ 
#ifdef __CUDA_ARCH__ 
    return dc_vals[i]; 
#else 
    return hc_vals[i]; 
#endif 
} 

Questo ha lo stesso risultato di Cigno, ma è più flessibile di fronte a codice reale: esso consente di avere valori di runtime-definito nelle matrici costanti, per esempio e consente di utilizzare le funzioni API CUDA come cudaMemcpyToSymbol/cudsaMemcpyFromSymbol sull'array __constant__.

Una più realistica esempio completo:

#include <iostream> 
#include <stdio.h> 

__constant__ double dc_vals[2]; 
     const double hc_vals[2]; 

__device__ __host__ double f(size_t i) 
{ 
#ifdef __CUDA_ARCH__ 
    return dc_vals[i]; 
#else 
    return hc_vals[i]; 
#endif 
} 

__global__ void kern() 
{ 
    printf("vals[%d] = %lf\n", threadIdx.x, vals[threadIdx.x]); 
} 

int main() { 
    hc_vals[0] = 0.0; 
    hc_vals[1] = 1000.0; 

    cudaMemcpyToSymbol(dc_vals, hc_vals, 2 * sizeof(double), 0, cudaMemcpyHostToDevice); 

    std::cerr << f(0) << " " << f(1) << std::endl; 
    kern<<<1, 2>>>(); 
    cudaThreadSynchronize(); 
} 

Answer 2

L'utilizzo del qualificatore __constant__ alloca esplicitamente tale memoria sul dispositivo. Non c'è modo di accedere a quella memoria dall'host - nemmeno con la nuova roba di indirizzamento unificato di CUDA (che funziona solo per la memoria allocata con cudaMalloc() e i suoi amici). Qualificare la variabile con const dice semplicemente "questo è un puntatore costante a (...)".

Il modo corretto per farlo è, in effetti, disporre di due array: uno sull'host e uno sul dispositivo. Inizializza l'array host, quindi utilizza cudaMemcpyToSymbol() per copiare i dati sull'array di dispositivi in ​​fase di runtime. Per ulteriori informazioni su come eseguire questa operazione, vedere questo thread: http://forums.nvidia.com/index.php?showtopic=69724

Answer 3

Penso che MurphEngineer abbia spiegato bene perché non funziona.

Per risolvere rapidamente questo problema, è possibile seguire l'idea di harrism, qualcosa di simile:

#ifdef __CUDA_ARCH__ 
#define CONSTANT __constant__ 
#else 
#define CONSTANT 
#endif 

const CONSTANT double vals[2] = { 0.0, 1000.0 }; 

In questo modo la compilazione host creare un normale allineamento const ospite, mentre la compilazione dispositivo creerà un dispositivo __constant__ compilazione.

Si noti che con questo trucco potrebbe essere più difficile utilizzare l'API CUDA per accedere a quell'array di dispositivi con funzioni come cudaMemcpyToSymbol() se si decide di farlo.