Verificare se il codice è in esecuzione sulla GPU o sulla CPU

Question

Qualcuno sa come verificare se il codice è in esecuzione sulla GPU o sulla CPU utilizzando Cuda?

__device__ __host__ double count_something(double variable) { 
    if (RUN_ON_GPU) { 
    use_cuda_variables(); 
    } else { 
    use_cpu_variables(); 
    } 
} 

Answer 1

Non v'è alcun modo per runtime controllo che l'architettura di un pezzo di codice è in esecuzione, ma c'è anche bisogno di sapere, perché può essere determinato al momento della compilazione e trattati di conseguenza. nvcc definisce diversi simboli del preprocessore che possono essere utilizzati per analizzare la traiettoria della compilazione mentre il codice è in fase di compilazione. Il simbolo della chiave è __CUDA_ARCH__ che non viene mai definito durante la compilazione del codice host e sempre definito durante la compilazione del codice dispositivo.

quindi è possibile scrivere una funzione come questa:

__device__ __host__ float function(float x) 
{ 
#ifdef __CUDA_ARCH__ 
    return 10.0f * __sinf(x); 
#else 
    return 10.0f * sin(x); 
#endif 
} 

che emetterà codice diverso a seconda che esso è stato compilato per la GPU o host. È possibile leggere una discussione più approfondita sulla sterzatura delle compilation in questo Stack Overflow question o nella sezione C language extensions della guida di programmazione CUDA.

Answer 2

Non riesco ad aggiungere il markdown del codice corretto nei commenti - ho deciso di aggiungere una risposta completa. L'utilizzo della sola definizione di controllo __CUDA_ARCH__ non è completamente corretto. In alcuni casi questo codice non funziona - Ho passato molto tempo nel debugging prima di trovare una soluzione (la documentazione CUDA non ne parlava ora).
__CUDA_ARCH__ può essere definito anche nel codice host, ma in tal caso è definito a 0. Quindi il controllo corretto è qualcosa del genere:

__device__ __host__ float function(float x) 
{ 
#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ > 0)) 
    // device code here 
    return 10.0f * __sinf(x); 
#else 
    // host code here 
    return 10.0f * sin(x); 
#endif 
}