如何在 Cuda 设备函数中使用 C++11 中的闭包参数声明函数？

Question

我正在使用 C++11 开发 Cuda（我认为 Cuda 还不支持更高的 C++ 版本）。我有一个传递给函数 Process() 的闭包对象，它为每次迭代调用闭包。

我了解 std:: 功能在 Cuda 中通常不可用。例如，当我尝试使用 std::function< float(uint32_t) > 时，出现此错误：

error: calling a host function("std::function ::function< ::, void, void> ") from a global function("_NV_ANON_NAMESPACE::LargeKernel") is not allowed

我可以用什么替换 lookupFunc 以便在 std::function 不可用的情况下进行编译？我能够通过创建一个函数模板来推断 lambda 函数的类型来解决这个问题。

此代码有效并显示了我所采用的工作：

//using lookupFunc = std::function< float(uint32_t) >;

template< typename Lambda > // Work around with function template
__device__
void Process(float       * const outData,
             const  int32_t      locationX,
             const Lambda /* lookupFunc */ lambda)
{
    float answer = 0.f;

    for( int32_t offset = -1 ; ++offset < 1024 ; )
    {
        const float value = lambda( offset );

        answer += value;
    }

    outData[ locationX ] = answer;
}

__global__
void LargeKernel(const float * const inData,
                 float       * const outData)
{
    constexpr uint32_t cellStride = 1;
    const     int32_t  locationX  = threadIdx.x + blockDim.x * blockIdx.x;
    const auto lambda
        = [locationX, inData, cellStride](const int32_t offset)
          {
              return inData[ locationX + offset + cellStride ];
          };

    Process( outData, locationX, lambda );
}

我也试过：

using lookupFunc = float(* const)(uint32_t);

但这给出了错误：

error: no suitable conversion function from "const lambda ->float" to "float (*)(uint32_t)" exists

如何在不使用模板的情况下声明 Process() 的第三个参数的类型？

Answer 1

有必要使用 lambda 吗？

否则，你可以模拟它声明一个struct

struct noLambda
 {
   std::int32_t const   locationX;
   float const * const  inData;
   std::uint32_t const  cellStride;

   noLambda (std::int32_t l0, float const * const i0, std::uint32_t c0)
      : locationX{l0}, inData{i0}, cellStride{c0}
    { }

   float operator() (std::int32_t const offset) const
    { return inData[ locationX + offset + cellStride ]; }
 };

所以进程的签名变成

void Process(float       * const outData,
             const  int32_t      locationX,
             const noLambda lambda)

并且可以如下调用

Process( outData, locationX, noLambda{locationX, inData, cellStride} );

（注意：代码未经测试）

Answer 2

相当于 std::function 的 CUDA 是 nvstd::function。从 CUDA 8.0 开始，nvstd::function 可用于主机和设备代码 - 但 "instances cannot be passed from host code to device code (and vice versa) at run time"。 CUDA programming guide.

中解释了它的使用

如何在 Cuda 设备函数中使用 C++11 中的闭包参数声明函数？

How to declare function with argument that is a closure in C++11 in a Cuda device function?

c++

closures

cuda

function-templates

c++11