c++ - 从全局内核调用宿主函数

Question

我试图从内核中调用一个函数，但我不断收到错误。我的函数是一个模板参数。 这是错误:

 calling a __host__ function("mul<int> ") from a __global__ function("multiply<int, &mul<int> > ") is not allowed. 

这是我的代码:

    template<typename T> T mul(T a, T b)
{
    return a * b;
}
template <typename T, T (*myop)(T, T)>
static __global__ void multiply(int *in, int *out)
{...
 a = myop(b, c);
 ...
}
int main()
{
multiply<T, mul<T> ><<<Blocks, 1024>>>(...);
}

有人知道这是怎么回事吗？

Answer 1

要在设备代码中使用的函数必须具有 __device__ qualifier :

$ cat t1018.cu
#include <stdio.h>
template<typename T>
__host__ __device__ T mul(T a, T b)
{
    return a * b;
}

template <typename T, T (*myop)(T, T) >
static __global__ void multiply()
{
 T b=2;
 T c=5;
 T a = myop(b, c);
 printf("a = %d\n", a);

}
int main()
{
  multiply<int, mul<int> ><<<1, 1>>>();
  cudaDeviceSynchronize();
}
$ nvcc -o t1018 t1018.cu
$ ./t1018
a = 10
$