带有 CUDA 的 Lambda 表达式

Question

如果我在 thrust::host 上使用 thrust::transform，则 lambda 用法很好

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 在 @987654327 中定义@或__global__函数

所以，问题在于如何使用__device__ 或__global__ 连接到设备lambda。

Answer 1

在 CUDA 7 中这是不可能的。引用Mark Harris:

现在 CUDA 不支持该功能，因为 lambda 是主机代码。将 lambda 从主机传递到设备是一个具有挑战性的问题，但我们将在未来的 CUDA 版本中对此进行研究。

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

使用 CUDA 7，可以从设备代码（例如 CUDA 内核或 __device__ 函子）调用推力算法。在这些情况下，您可以使用带推力的（设备）lambda。 parallelforall 博客文章here 中给出了一个示例。

但是，CUDA 7.5 引入了实验性设备 lambda 功能。此功能描述here：

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

为了启用此功能的编译，（当前使用 CUDA 7.5）需要在 nvcc 编译命令行上指定 --expt-extended-lambda。

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

用于编译。