使用CUDA的Lambda表达式

Question

如果我使用thrust::transform上thrust::host,拉姆达用法是罚款

thrust::transform(thrust::host, a, a+arraySize,b,d,[](int a, int b)->int
{
    return a + b;
});

但是,如果我thrust::host改为thrust::device,代码将不会通过编译器.这是VS2013上的错误:

lambda的闭包类型("lambda [](int,int) - > int")不能在__global__函数模板实例化的模板参数类型中使用,除非lambda是在一个__device__或__global__函数内定义的

所以,问题是如何使用__device__或__global__连接到设备lambdas.

Answer 1

在CUDA 7中,这是不可能的.引自马克哈里斯:

今天在CUDA中不支持,因为lambda是主机代码.将lambdas从主机传递到设备是一个具有挑战性的问题,但我们将对未来的CUDA版本进行调查.

您在CUDA 7中可以做的是从您的设备代码中调用推力算法,在这种情况下,您可以将lambda传递给它们......

使用CUDA 7,可以从设备代码(例如CUDA内核或__device__仿函数)调用推力算法.在这些情况下,你可以使用(设备)lambdas与推力.这里的parallelforall博客文章中给出了一个例子 .

但是,CUDA 7.5引入了实验设备lambda特性.此功能在此处描述:

CUDA 7.5引入了一个实验性功能:GPU lambda.GPU lambdas是匿名设备函数对象,您可以在主机代码中定义它们,方法是使用__device__说明符对它们进行注释.

为了启用此功能的编译(目前,使用CUDA 7.5),必须--expt-extended-lambda在nvcc编译命令行上指定.

Answer 2

这个使用设备lambda的简单代码可以在CUDA 8.0 RC下运行，尽管此版本CUDA的设备lambda仍处于试验阶段：

#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/transform.h>

using namespace thrust::placeholders;

int main(void)
{
    // --- Input data 
    float a = 2.0f;
    float x[4] = { 1, 2, 3, 4 };
    float y[4] = { 1, 1, 1, 1 };

    thrust::device_vector<float> X(x, x + 4);
    thrust::device_vector<float> Y(y, y + 4);

    thrust::transform(X.begin(), 
                      X.end(),  
                      Y.begin(), 
                      Y.begin(),
                      [=] __host__ __device__ (float x, float y) { return a * x + y; }      // --- Lambda expression 
                     );        

    for (size_t i = 0; i < 4; i++) std::cout << a << " * " << x[i] << " + " << y[i] << " = " << Y[i] << std::endl;

    return 0;
}

切记使用

--expt-extended-lambda

进行编译。