CUDA：如何从主机函数返回设备 lambda

Question

我有一个虚函数，它根据派生类返回不同的 lambda：

class Base
{
public:
    virtual std::function<float()> foo(void) = 0;
};

class Derived : public Base
{
public:
    std::function<float()> foo(void) {
        return [] __device__ (void) {
            return 1.0f;
        };
    }
};

然后我想将此 lambda 传递给 CUDA 内核并从设备调用它。换句话说，我想这样做：

template<typename Func>
__global__ void kernel(Func f) {
    f();
}

int main(int argc, char** argv)
{
    Base* obj = new Derived;
    kernel<<<1, 1>>>(obj->foo());
    cudaDeviceSynchronize();
    return 0;
}

上面给出了这样的错误：calling a __host__ function("std::function<float ()> ::operator ()") from a __global__ function("kernel< ::std::function<float ()> > ") is not allowed

如您所见，我将 lambda 声明为__device__，但该foo()方法将其存储在 a 中std::function以便返回它。结果传给的kernel()是主机地址当然不行。我想这就是我的问题吧？所以我的问题是：

• 是否可以以某种方式创建__device__ std::function并从方法中返回它foo()？

• 如果这是不可能的，是否有其他方法可以动态选择 lambda 并将其传递给 CUDA 内核？使用所有可能的 lambda来硬编码多个调用kernel()并不是一种选择。

到目前为止，根据我所做的快速研究，CUDA 不具备/支持使函数返回设备 lambda 所需的必要语法。我只是希望我错了。：） 有任何想法吗？

提前致谢

Answer 1

在真正回答之前，我必须想知道你的问题是否不是一个XY问题。也就是说，我默认怀疑人们是否有充分的理由通过设备上的 lambda/函数指针执行代码。

但我不会这样回避你的问题...

是否可以以某种方式创建 a__device__ std::function并从 foo() 方法返回它？

简短的回答：不，尝试其他方法。

更长的答案：如果您想在设备端实现大部分标准库，那么也许您可以拥有一个类似设备端的std::function类。但我不确定这是否可能（很可能不可能），而且无论如何 - 除了经验丰富的库开发人员之外，这超出了每个人的能力。所以，做点别的事吧。

如果这是不可能的，是否有其他方法可以动态选择 lambda 并将其传递给 CUDA 内核？使用所有可能的 lambda 硬编码对 kernel() 的多次调用并不是一种选择。

首先，请记住 lambda 本质上是匿名类 - 因此，如果它们不捕获任何内容，它们就可以简化为函数指针，因为匿名类没有数据，只有operator().

因此，如果 lambda 具有相同的签名且没有捕获，您可以将它们转换为（非成员）函数指针并将其传递给函数；这绝对有效，请参阅nVIDIA 论坛上的这个简单示例。

另一种可能性是使用从类型 id 或其他此类键到这些类型的实例（或者更确切地说，到构造函数）的运行时映射。即使用工厂。但我不想深入讨论这个细节，以免这个答案比现在更长；这可能不是一个好主意。