CUDA的Lambda表达式

Lambda expressions with CUDA

本文关键字：表达式 Lambda CUDA      更新时间：2023-10-16

如果我在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__函数模板实例化的模板实参类型中使用，除非在__device__或__global__函数中定义了lambda

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

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

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

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

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

然而，CUDA 7.5引入了一个实验性设备lambda特性。此特性描述如下:

CUDA 7.5引入了一个实验性的特性:GPU lambda。GPU lambdas是匿名设备函数对象，您可以在主机代码中定义，通过使用__device__说明符注释它们。

为了启用此功能的编译，(目前，在CUDA 7.5中)有必要在nvcc编译命令行中指定--expt-extended-lambda。

使用设备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
为编译。

相关文章：