CUDA 7.5実験的__host__ __device__ lambdas

Question

私はexperimental device lambdasで少し演奏しましたが、CUDA 7.5で導入され、このblog post by Mark Harrisで宣伝されました。

次の例では、私の問題を示すのに必要でない多くのものを削除しました（私の実際の実装は少し良く見えます）。

テンプレートパラメータに応じて、デバイス上のベクトル（要素ごとに1つのスレッド）またはホスト（シリアル）のいずれかで動作するforeach関数を記述しようとしました。このforeach関数を使用すると、BLAS関数を簡単に実装できます。一例として、Iはベクトル（私は最終的に完全なコードを添付）の各成分にスカラーを割り当てる使用：

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

しかし、このコードは理由__host__ __device__ラムダのコンパイラエラーを与える：

ラムダのため

閉鎖タイプ - ラムダは私が手__device__か__global__機能

内で定義されていない限り（「ラムダ>ボイド」）、__global__関数テンプレートのインスタンス化のテンプレート引数の型で使用することはできませんその私は、ラムダ式から__device__を削除し、私は__host__（のみ__device__ラムダ）を削除した場合、私は何のコンパイルエラーを取得していない場合は、同じエラーが、私のようにラムダを定義した場合、この場合のホスト部分は...

を実行されません__host__または__device__を別々に指定すると、コードはコンパイルされ、期待どおりに動作します。

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

しかし、これはコードの重複を紹介し、実際にこの例の無用ラムダを使用しての全体的なアイデアを作ります。

私は何をしたいのか、これは実験的な機能のバグですか？実際には、__host__ __device__ラムダの定義は、programming guideの最初の例で明示的に言及されています。その単純な例（ラムダから一定の値を返す）であっても、ホストとデバイスの両方でラムダ式を使用する方法を見つけることができませんでした。私はCUDA 7.5の製品リリースを使用

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

：ここ

はオプション-std=c++11 --expt-extended-lambdaでコンパイルし、完全なコードです。

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

それがコンパイルされ、エラーなしで実行されますが、デバイスのバージョン（assignScalar3<true>）が実行されない：

更新私はassignScalar機能のために、この第三のバージョンを試してみました。実際には、__CUDA_ARCH__は（関数は__device__ではないので）常に定義されていないと考えましたが、定義されているコンパイルパスがあることを明示的にチェックしました。