CUDA 7.5 experimental __host__ __device__ lambdas

Question

Ho giocato un po 'con lo experimental device lambdas che è stato introdotto in CUDA 7.5 e promosso in questo blog post by Mark Harris.

Per il seguente esempio ho rimosso un sacco di cose che non sono necessarie per mostrare il mio problema (la mia implementazione effettiva sembra un po 'più bella ...).

Ho provato a scrivere una funzione foreach che opera su vettori sul dispositivo (1 thread per elemento) o host (seriale) in base a un parametro del modello. Con questa funzione foreach posso facilmente implementare le funzioni BLAS. Per fare un esempio io uso l'assegnazione di uno scalare ad ogni componente di un vettore (allego il codice completo alla fine):

template<bool onDevice> void assignScalar(size_t size, double* vector, double a) 
{ 
    auto assign = [=] __host__ __device__ (size_t index) { vector[index] = a; }; 
    if(onDevice) 
    { 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     foreachHost(size, assign); 
    } 
} 

Tuttavia, questo codice dà un errore del compilatore a causa del __host__ __device__ lambda:

il tipo di chiusura per una lambda ("lambda -> vuoto") non può essere utilizzato nel tipo di modello argomento di un modello di un'istanza funzione di __global__, a meno che il lambda è definito all'interno di un __device__ o __global__ funzione

ottengo il stesso errore se mi tolgo la __device__ dall'espressione lambda e ottengo nessun errore di compilazione se tolgo __host__ (solo __device__ lambda), ma in questo caso non viene eseguita la parte host ...

Se Mi definisco lambda come __host__ o __device__ separatamente, il codice viene compilato e funziona come previsto.

template<bool onDevice> void assignScalar2(size_t size, double* vector, double a) 
{ 
    if(onDevice) 
    { 
     auto assign = [=] __device__ (size_t index) { vector[index] = a; }; 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     auto assign = [=] __host__ (size_t index) { vector[index] = a; }; 
     foreachHost(size, assign); 
    } 
} 

Tuttavia, questo introduce codice duplicazione e rende effettivamente l'idea di utilizzare lambda inutile per questo esempio.

C'è un modo per realizzare ciò che voglio fare o si tratta di un bug nella funzionalità sperimentale? In realtà, la definizione di un lambda __host__ __device__ è esplicitamente menzionata nel primo esempio nello programming guide. Anche per quell'esempio più semplice (basta restituire un valore costante dal lambda) non sono riuscito a trovare un modo per usare l'espressione lambda su host e dispositivo.

Ecco il codice completo, compilare con le opzioni -std=c++11 --expt-extended-lambda:

#include <iostream> 
using namespace std; 

template<typename Operation> void foreachHost(size_t size, Operation o) 
{ 
    for(size_t i = 0; i < size; ++i) 
    { 
     o(i); 
    } 
} 

template<typename Operation> __global__ void kernel_foreach(Operation o) 
{ 
    size_t index = blockIdx.x * blockDim.x + threadIdx.x; 
    o(index); 
} 

template<typename Operation> void foreachDevice(size_t size, Operation o) 
{ 
    size_t blocksize = 32; 
    size_t gridsize = size/32; 
    kernel_foreach<<<gridsize,blocksize>>>(o); 
} 

__global__ void printFirstElementOnDevice(double* vector) 
{ 
    printf("dVector[0] = %f\n", vector[0]); 
} 

void assignScalarHost(size_t size, double* vector, double a) 
{ 
    auto assign = [=] (size_t index) { vector[index] = a; }; 
    foreachHost(size, assign); 
} 

void assignScalarDevice(size_t size, double* vector, double a) 
{ 
    auto assign = [=] __device__ (size_t index) { vector[index] = a; }; 
    foreachDevice(size, assign); 
} 

// compile error: 
template<bool onDevice> void assignScalar(size_t size, double* vector, double a) 
{ 
    auto assign = [=] __host__ __device__ (size_t index) { vector[index] = a; }; 
    if(onDevice) 
    { 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     foreachHost(size, assign); 
    } 
} 

// works: 
template<bool onDevice> void assignScalar2(size_t size, double* vector, double a) 
{ 
    if(onDevice) 
    { 
     auto assign = [=] __device__ (size_t index) { vector[index] = a; }; 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     auto assign = [=] __host__ (size_t index) { vector[index] = a; }; 
     foreachHost(size, assign); 
    } 
} 

int main() 
{ 
    size_t SIZE = 32; 

    double* hVector = new double[SIZE]; 
    double* dVector; 
    cudaMalloc(&dVector, SIZE*sizeof(double)); 

    // clear memory 
    for(size_t i = 0; i < SIZE; ++i) 
    { 
     hVector[i] = 0; 
    } 
    cudaMemcpy(dVector, hVector, SIZE*sizeof(double), cudaMemcpyHostToDevice); 

    assignScalarHost(SIZE, hVector, 1.0); 
    cout << "hVector[0] = " << hVector[0] << endl; 

    assignScalarDevice(SIZE, dVector, 2.0); 
    printFirstElementOnDevice<<<1,1>>>(dVector); 
    cudaDeviceSynchronize(); 

    assignScalar2<false>(SIZE, hVector, 3.0); 
    cout << "hVector[0] = " << hVector[0] << endl; 

    assignScalar2<true>(SIZE, dVector, 4.0); 
    printFirstElementOnDevice<<<1,1>>>(dVector); 
    cudaDeviceSynchronize(); 

// assignScalar<false>(SIZE, hVector, 5.0); 
// cout << "hVector[0] = " << hVector[0] << endl; 
// 
// assignScalar<true>(SIZE, dVector, 6.0); 
// printFirstElementOnDevice<<<1,1>>>(dVector); 
// cudaDeviceSynchronize(); 

    cudaError_t error = cudaGetLastError(); 
    if(error!=cudaSuccess) 
    { 
     cout << "ERROR: " << cudaGetErrorString(error); 
    } 
} 

ho usato la versione di produzione di CUDA 7.5.

Aggiornamento

Ho provato questa terza versione per la funzione assignScalar:

template<bool onDevice> void assignScalar3(size_t size, double* vector, double a) 
{ 
#ifdef __CUDA_ARCH__ 
#define LAMBDA_HOST_DEVICE __device__ 
#else 
#define LAMBDA_HOST_DEVICE __host__ 
#endif 

    auto assign = [=] LAMBDA_HOST_DEVICE (size_t index) { vector[index] = a; }; 
    if(onDevice) 
    { 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     foreachHost(size, assign); 
    } 
} 

Si compila e viene eseguito senza errori, ma non viene eseguita la versione del dispositivo (assignScalar3<true>). In realtà, ho pensato che __CUDA_ARCH__ sarà sempre indefinito (poiché la funzione non è __device__) ma ho verificato esplicitamente che esiste un percorso di compilazione in cui è definito.

Answer 1

Il compito che ho cercato di realizzare con gli esempi forniti nel la domanda è non è possibile con CUDA 7.5, anche se non è stato esplicitamente esclusa dai casi consentiti per il supporto lambda sperimentale.

NVIDIA ha annunciato che CUDA Toolkit 8.0 sosterrà __host__ __device__ lambda come una funzione sperimentale, secondo il post sul blog CUDA 8 Features Revealed.

Ho verificato che il mio esempio funziona con CIDA 8 Release Candidate (strumenti di compilazione Cuda, versione 8.0, V8.0.26).

Ecco il codice che ho finalmente usato, compilato con nvcc -std=c++11 --expt-extended-lambda:

#include <iostream> 
using namespace std; 

template<typename Operation> __global__ void kernel_foreach(Operation o) 
{ 
    size_t i = blockIdx.x * blockDim.x + threadIdx.x; 
    o(i); 
} 

template<bool onDevice, typename Operation> void foreach(size_t size, Operation o) 
{ 
    if(onDevice) 
    { 
     size_t blocksize = 32; 
     size_t gridsize = size/32; 
     kernel_foreach<<<gridsize,blocksize>>>(o); 
    } 
    else 
    { 
     for(size_t i = 0; i < size; ++i) 
     { 
      o(i); 
     } 
    } 
} 

__global__ void printFirstElementOnDevice(double* vector) 
{ 
    printf("dVector[0] = %f\n", vector[0]); 
} 

template<bool onDevice> void assignScalar(size_t size, double* vector, double a) 
{ 
    auto assign = [=] __host__ __device__ (size_t i) { vector[i] = a; }; 
    foreach<onDevice>(size, assign); 
} 

int main() 
{ 
    size_t SIZE = 32; 

    double* hVector = new double[SIZE]; 
    double* dVector; 
    cudaMalloc(&dVector, SIZE*sizeof(double)); 

    // clear memory 
    for(size_t i = 0; i < SIZE; ++i) 
    { 
     hVector[i] = 0; 
    } 
    cudaMemcpy(dVector, hVector, SIZE*sizeof(double), cudaMemcpyHostToDevice); 

    assignScalar<false>(SIZE, hVector, 3.0); 
    cout << "hVector[0] = " << hVector[0] << endl; 

    assignScalar<true>(SIZE, dVector, 4.0); 
    printFirstElementOnDevice<<<1,1>>>(dVector); 
    cudaDeviceSynchronize(); 

    cudaError_t error = cudaGetLastError(); 
    if(error!=cudaSuccess) 
    { 
     cout << "ERROR: " << cudaGetErrorString(error); 
    } 
}