CUDA 7.5 deneysel __host__ __device__ lambdas

Question

CUDA 7.5'te tanıtılan ve bu blog post by Mark Harris'da tanıtılan experimental device lambdas ile biraz oynadım.

Aşağıdaki örnekte, sorunumu göstermek için gerekli olmayan birçok şeyi kaldırdım (gerçek uygulamam biraz daha güzel görünüyor ...).

Şablon parametresine bağlı olarak aygıttaki vektörlerde (öğe başına 1 iş parçacığı) veya ana makine (seri) üzerinde çalışan bir foreach işlevi yazmaya çalıştım. Bu foreach fonksiyonu ile BLAS fonksiyonlarını kolayca uygulayabilirim. Bir vektör, her bileşen için bir skaler atama kullanmak Bir örnek olarak (Sonunda tam kod takmak): Ancak

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); 
    } 
} 

, bu kod nedeniyle __host__ __device__ lambda bir derleyici hata verir:

bir lambda ("lambda -> geçersiz") için kapatma tipi lambda bir __device__ veya

alıyorum __global__ işlev içinde tanımlanan sürece bir __global__ işlevi şablon örnekleme şablon argümanı tipinde kullanılamaz

Ben lambda olarak tanımlarsanız ... aynı hata ben lambda ifadeden __device__ kaldırırsanız ve ben __host__ (sadece __device__ lambda) kaldırırsanız hiçbir derleme hatası alıyorum, ama bu durumda konak parçası yürütülmez __host__ veya __device__ ayrı olarak, kod derler ve beklendiği gibi çalışır.

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); 
    } 
} 

Ancak bu kod tekrarını tanıtır ve aslında bu örnek için yararsız lambdas kullanma fikrine yapar.

Yapmak istediklerimi gerçekleştirmenin bir yolu var mı, yoksa bu deneysel özellikteki bir hata mı? Aslında, __host__ __device__ lambda'nın tanımlanması, programming guide'daki ilk örnekte açıkça belirtilmiştir. Bu basit örnek için bile (sadece lambdadan sabit bir değer döndürün) lambda ifadesini hem ana bilgisayar hem de aygıtta kullanmanın bir yolunu bulamadım.

#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); 
    } 
} 

Ben CUDA 7.5 üretim salınımını kullandı:

İşte seçenekleri -std=c++11 --expt-extended-lambda ile derlemek, tam koddur.

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); 
    } 
} 

Bu derler ve hatasız çalışır ancak cihaz sürümü (assignScalar3<true>) yürütülmez:

Güncelleme ı assignScalar fonksiyonu için bu üçüncü versiyonu çalıştı. Aslında, __CUDA_ARCH__'un her zaman tanımlanamayacağını düşündüm (işlev __device__ değil) ancak açıkça tanımlandığı bir derleme yolu olduğunu kontrol ettim. açıkça deneysel lambda desteği için izin vakalardan bertaraf edilmemiş olsa