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;
}

以上代码出现了如下错误：从 __global__ 函数("kernel< ::std::function<float ()> > ") 调用 __host__ 函数("std::function<float ()> ::operator ()") 是不允许的。

可以看到，lambda被声明为__device__，但是foo()方法将其存储在std::function中以便返回它。因此，传递给kernel()的是主机地址，当然无法工作。这就是我的问题所在，对吧？所以我有以下问题：

• 是否有可能创建一个__device__ std::function并从foo()方法中返回它？

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

到目前为止，从我所做的快速研究来看，CUDA没有/支持需要的语法来使函数返回设备lambda。希望我是错的 :) 有什么想法吗？

提前感谢。

- AstrOne

---

我认为这不是语法问题。从我所看到的情况来看，设备不支持std::function，这就是编译错误的来源。 - talonmies

你可能想阅读这篇文章。我认为使用__device__ lambda作为内核参数/参数肯定是可行的，但你可能正在使用std::function，因为你想要“泛型化”它 - 你不喜欢每个lambda都有一个唯一的类型。我不认为你能用std::function来解决这个问题。使用functor可能会更容易些。 - Robert Crovella

Answer 1

在实际回答之前，我必须怀疑你的问题是否是一个XY问题。也就是说，我默认认为人们通过lambda/函数指针在设备上执行代码没有什么好的借口。

但我不会逃避你的问题...

引用：

“是否有可能创建一个__device__ std :: function并从foo()方法返回它？”

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

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

引用：

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

首先，记住Lambda本质上是匿名类 - 因此，如果它们不捕获任何内容，则可以将其简化为函数指针，因为匿名类没有数据，只有一个 operator()。

因此，如果这些Lambda具有相同的签名且没有捕获，您可以将它们转换为（非成员）函数指针，并将其传递给函数；这肯定有效，可以在nVIDIA论坛上看到这个简单的例子。

另一种可能性是使用类型ID或其他键的运行时映射到这些类型的实例，或者更确切地说，到构造函数。也就是说，使用一个工厂。但我不想详细介绍，以免答案变得比现在更长；而且这可能不是一个好主意。

- einpoklum

Answer 2

虽然我认为你无法使用返回设备lambda的虚函数来实现所需的功能，但是你可以通过将静态设备成员函数作为模板参数传递给内核来实现类似的功能。下面提供了一个示例。请注意，如果您喜欢，此示例中的类也可以是结构体。

#include <iostream>

// Operation: Element-wise logarithm
class OpLog {
    public:
    __device__ static void foo(int tid, float * x) {
        x[tid] = logf(x[tid]);
    };
};

// Operation: Element-wise exponential
class OpExp {
    public:
    __device__ static void foo(int tid, float * x) {
        x[tid] = expf(x[tid]);
    }
};

// Generic kernel
template < class Op >
__global__ void my_kernel(float * x) {
    int tid = threadIdx.x;
    Op::foo(tid,x);
}

// Driver
int main() {

    using namespace std;

    // length of vector
    int len = 10;

    // generate data
    float * h_x = new float[len];
    for(int i = 0; i < len; i++) {
        h_x[i] = rand()/float(RAND_MAX);
    }

    // inspect data
    cout << "h_x = [";
    for(int j = 0; j < len; j++) {
        cout << h_x[j] << " ";
    }
    cout << "]" << endl;

    // copy onto GPU
    float * d_x;
    cudaMalloc(&d_x, len*sizeof(float));
    cudaMemcpy(d_x, h_x, len*sizeof(float), cudaMemcpyHostToDevice);

    // Take the element-wise logarithm
    my_kernel<OpLog><<<1,len>>>(d_x);

    // get result
    cudaMemcpy(h_x, d_x, len*sizeof(float), cudaMemcpyDeviceToHost);
    cout << "h_x = [";
    for(int j = 0; j < len; j++) {
        cout << h_x[j] << " ";
    }
    cout << "]" << endl;

    // Take the element-wise exponential
    my_kernel<OpExp><<<1,len>>>(d_x);

    // get result
    cudaMemcpy(h_x, d_x, len*sizeof(float), cudaMemcpyDeviceToHost);
    cout << "h_x = [";
    for(int j = 0; j < len; j++) {
        cout << h_x[j] << " ";
    }
    cout << "]" << endl;


}

- tdoublep