【发布时间】：2015-12-10 14:42:31
【问题描述】：
我在experimental device lambdas 上玩了一下，它是在 CUDA 7.5 中引入并在此 blog post by Mark Harris 中推广的。
对于以下示例，我删除了很多不需要的东西来显示我的问题（我的实际实现看起来更好......）。
我尝试编写一个 foreach 函数，它可以根据模板参数对设备上的向量（每个元素 1 个线程）或主机（串行）上的向量进行操作。使用这个 foreach 函数，我可以轻松实现 BLAS 函数。作为一个例子，我使用为向量的每个分量分配一个标量（我在最后附上了完整的代码）：
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__ lambda：此代码给出了编译器错误：
lambda 的闭包类型（“lambda ->void”）不能用于 __global__ 函数模板实例化的模板参数类型，除非 lambda 在 __device__ 或 __global__ 函数中定义
如果我从 lambda 表达式中删除 __device__，我会得到相同的错误，如果我删除 __host__（只有 __device__ lambda），我不会收到编译错误，但在这种情况下，主机部分不会被执行。 ..
如果我将 lambda 分别定义为 __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 );
}
}
但是，这会引入代码重复，实际上使使用 lambdas 的整个想法在本示例中毫无用处。
有没有办法完成我想做的事情，或者这是实验功能中的错误？实际上，在programming guide 的第一个示例中明确提到了定义__host__ __device__ lambda。即使对于那个更简单的示例（仅从 lambda 返回一个常量值），我也找不到在主机和设备上使用 lambda 表达式的方法。
这里是完整代码，编译选项-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);
}
}
我使用的是 CUDA 7.5 的生产版本。
更新
我为 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 );
}
}
编译运行没有错误，但是设备版本（assignScalar3<true>）没有执行。实际上，我认为__CUDA_ARCH__ 将始终未定义（因为函数不是__device__），但我明确检查了定义它的编译路径。
【问题讨论】：
-
我认为该错误具有指导意义，它可能是文档中未明确说明的进一步实施限制。如果您按照报告错误的建议，并将
assignScalar模板函数标记为__host__ __device__，我认为您可以解决这个特定问题。这将引发编译器警告，可以安全地忽略这些警告，或者可以使用__CUDA_ARCH__宏来解决，以进行干净的编译。到那时，我认为您可能会偶然发现某种实现错误。我目前没有其他信息。 -
我会说该错误具有误导性，因为如果您检查示例
assignScalar2，它是不正确的。那里的 lambda 以相同的方式使用，并且未在__device__或__global__函数中定义。 -
@RobertCrovella 正如你所说，制作
assignScalar函数可以解决错误，但不是问题，因为该函数仅从主机调用（实际上，当我调用主机和设备 foreach 时，我听从建议）。但是你的评论让我想到了第三个版本，我将添加到问题中。 -
“该函数仅从主机调用”是我提到的实现错误。实际上设备版本开始被调用，但内核启动过程在运行时由于某种原因失败。如果您使用 gdb 和分析器四处寻找，您会发现这一点。
-
第二，编程指南中的限制 6：“解析函数时，CUDA 编译器为该函数中的每个
__device__lambda 分配一个计数器值。此计数器值用于替换命名类型传递给宿主编译器。因此，是否在函数中定义__device__lambda 不应取决于__CUDA_ARCH__的特定值，或__CUDA_ARCH__未定义。"