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定义在__device__或__global__函数内部。

因此，问题是如何在设备lambda与__device__或__global__之间建立连接。

- spiritsaway

---

你尝试过任何东西吗？ - Kerrek SB

我尝试在lambda定义之前添加__device__，然后编译器会报错error: expected an expression。有什么想法吗？ - spiritsaway

我猜测如果这可以实现（我不确定是不是这样），你需要在全局作用域定义lambda，并用__device__ __host__装饰它，可能像这样：__device__ __host__ auto cudalambda = [](int a, int b)->int { return a + b;}; 然后在Thrust闭包中使用它。但这只是一个猜测。 - talonmies

我已经尝试过了，在我的帖子中出现了相同的错误信息。 - spiritsaway

好的，答案可能是你不能这样做。相反，编写一个函数对象。 - talonmies

Answer 1

在CUDA 7中不可能实现。引用自Mark Harris：（链接）：“这在CUDA中今天并不受支持，因为lambda是主机代码。将lambda从主机传递到设备是一个具有挑战性的问题，但这是我们将来会研究的东西。” 在CUDA 7中可以从设备代码调用thrust算法，而在这种情况下您可以将lambda传递给它们。使用（设备）lambda与thrust的示例在此处的parallelforall博客文章中给出。 然而，CUDA 7.5引入了一个实验性的“device lambda”功能。（链接）：GPU lambdas是匿名设备函数对象，您可以在主机代码中使用 __device__注释定义它们。 为了启用此功能的编译（目前仅支持CUDA 7.5），需要在nvcc编译命令行上指定--expt-extended-lambda。

- Robert Crovella

---

谢谢，我整天都在尝试弄清楚它，但是徒劳无功。所以我使用了一个函数对象包装器。 - spiritsaway

使用CUDA 7.5，可以在主机代码中使用[] __device__ (params...) { code ... }声明设备lambda函数，并将--expt-extended-lambda标志传递给nvcc编译器。 - Hopobcn

Answer 2

这段简单的代码使用设备lambda在CUDA 8.0 RC下工作，尽管该版本的设备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

用于编译。

- Vitality

---

看起来即使在CUDA 10.1（2019年），它们仍然需要--expt-extended-lambda。 - Jakub Narębski

是的，我希望他们能够摒弃这个标志。顺便说一下，它现在的拼写是：--extended-lambda。 - user2023370