如何编写一个可从主机和设备调用的内核函数？

Question

以下内容是我未能成功编译的意图：

__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".

希望我的陈述清晰明了，感谢您的建议。

- Hailiang Zhang

---

你肯定是指 __device__ __host__ void f() {} 吧？ - talonmies

我尝试了 "device host" 和 "host device" 的组合，但两者都失败了。 - Hailiang Zhang

以上代码是基于教程"http://www.uni-graz.at/~liebma/CUDA/NVISION08-Getting_Started_with_CUDA.pdf"编写的。 - Hailiang Zhang

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

- Eugene

---

CUDA_ARCH 将在两个调用中定义。预编译器代码在此上下文中毫无意义... - Henrique Mendonça

@HenriqueMendonça：我认为你错了。 - einpoklum

Answer 2

你正在查看的教程非常古老，2008年？它可能与您正在使用的CUDA版本不兼容。

您可以使用__global__，这意味着__host__ __device__，这样就可以正常运行：

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

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

- Adam

---

__global__ 指定了内核入口点，即一个函数，当使用启动参数调用时，它将自动并行化为 GPU 代码。 __host__ 和 __device__ 不用于修饰内核函数。你唯一能说 __global__ 意味着 __host__ __device__ 的情况是在 cuda 动态并行性 中，该功能仅适用于 cc 3.5 设备。即使在这种情况下，我认为说 __global__ 意味着 __host__ __device__ 是不严谨的。 - Robert Crovella

@RobertCrovella 我同意，我只是指在他的情境下它们是等效的，因为我的代码由于有内核变量，无法从主机调用。 - Adam