如何制作可从主机和设备调用的内核函数？答案

【问题标题】：How to make a kernel function which callable from both the host and device?如何制作可从主机和设备调用的内核函数？
【发布时间】：2013-06-11 21:21:39
【问题描述】：

下面的试验提出了我的意图，编译失败：

__host__ __device__ void f(){}

int main()
{
    f<<<1,1>>>();
}

编译器投诉：

a.cu(5): error: a __device__ function call cannot be configured

1 error detected in the compilation of "/tmp/tmpxft_00001537_00000000-6_a.cpp1.ii".

希望我的陈述清楚，并感谢您的建议。

【问题讨论】：

你的意思是__device__ __host__ void f(){}？
我尝试了 "device host" 和 "host device" 的组合，都失败了
以上代码基于教程“uni-graz.at/~liebma/CUDA/…”

标签： cuda

【解决方案1】：

您需要创建一个 CUDA 内核入口点，例如__global__ 函数。比如：

#include <stdio.h>

__host__ __device__ void f() {
#ifdef __CUDA_ARCH__
    printf ("Device Thread %d\n", threadIdx.x);
#else
    printf ("Host code!\n");
#endif
}

__global__ void kernel() {
   f();
}

int main() {
   kernel<<<1,1>>>();
   if (cudaDeviceSynchronize() != cudaSuccess) {
       fprintf (stderr, "Cuda call failed\n");
   }
   f();
   return 0;
}

【讨论】：

CUDA_ARCH 将在两个调用中定义。在这种情况下，预编译器代码毫无意义......
@HenriqueMendonça：我believe you are mistaken。

【解决方案2】：

你看的教程太老了，2008 年？它可能与您使用的 CUDA 版本不兼容。

您可以使用__global__，这意味着__host__ __device__，这样可以：

__global__ void f()
{
    const int tid = threadIdx.x + blockIdx.x * blockDim.x;
}

int main()
{
    f<<<1,1>>>();
}

【讨论】：

__global__ 指定内核入口点，即在使用启动参数调用时将自动并行化为 GPU 代码的函数。 __host__ 和 __device__ 不用于装饰内核函数。你可以说__global__ 意味着__host__ __device__ 的唯一意义是cuda dynamic parallelism，它仅在cc 3.5 设备上可用。即使在那种情况下，我认为说__global__ 意味着__host__ __device__ 是草率的
@RobertCrovella 我同意，我只是说它们在他的上下文中是等价的，因为我的代码无论如何都不能从主机调用，因为它有内核变量。