从 __global__ 函数中调用 __global__ 函数的替代方法？答案

【问题标题】：Alternatives to calling a __global__ function from within a __global__ function?从 __global__ 函数中调用 __global__ 函数的替代方法？
【发布时间】：2016-10-15 12:12:30
【问题描述】：

我有一些代码非常有用，值得作为自己的内核并独立运行。另一方面，我有另一段代码将第一个代码用作大型算法中的一个阶段。因此，我希望两者都通过内核启动调用第一个函数，都作为一个简单的devsion。

现在，您不能只从另一个 __global__ 函数调用 __global__ 函数；并且您不允许将函数标记为 __global__ __device__ （这真的太糟糕了，我原以为这是可能的）。那么你如何调用这个函数呢？

【问题讨论】：

@talonmies：完成。
没有完成，因为假设不正确。您可以从 CC 3.5 (Kepler) 开始从 __global__ 调用 __global__。甚至递归;)。该功能称为dynamic parallelism
@Drop：但我并不是说启动子网格，只是一个常规的函数调用。
这句话对我来说没有意义，我建议重新措辞以澄清问题：“因此，我希望两者都通过内核启动调用第一个函数，都作为一个简单的开发。”

标签： cuda

【解决方案1】：

我的解决方法是：

namespace detail {
    __device__ void foo(int) { /* ... */ }
} // namespace detail

__global__ void foo(int x) { detail::foo(x); }
__global__ void bar(int x) { detail::foo(baz(x));  }

可行，但有一个虚拟包装器，您需要复制函数头，这有时很乏味，所以我不太喜欢这种方法。

【讨论】：