如何正确链接 cuda 头文件与设备功能？

Question

我正在尝试将我的代码解耦，但有些失败。编译错误：

error: calling a __host__ function("DecoupledCallGpu") from a __global__ function("kernel") is not allowed

代码摘录：

ma​​in.c（调用了 cuda 主机函数）：

#include "cuda_compuations.h"
...
ComputeSomething(&var1,&var2);
...

cuda_computations.cu（具有内核、主机主控功能并包括具有设备功能的标头）：

#include "cuda_computations.h"
#include "decoupled_functions.cuh"
...
__global__ void kernel(){
...
DecoupledCallGpu(&var_kernel);
}

void ComputeSomething(int *var1, int *var2){
//allocate memory and etc..
...
kernel<<<20,512>>>();
//cleanup
...
}

decoupled_functions.cuh：

#ifndef _DECOUPLEDFUNCTIONS_H_
#define _DECOUPLEDFUNCTIONS_H_

void DecoupledCallGpu(int *var);

#endif

decoupled_functions.cu：

#include "decoupled_functions.cuh"

__device__ void DecoupledCallGpu(int *var){
  *var=0;
}

#endif

编译：

nvcc -g --ptxas-options=-v -arch=sm_30 -c cuda_computations.cu -o cuda_computations.o -lcudart

问题：为什么DecoupledCallGpu 是从主机函数调用而不是内核调用的？

P.S.：如果您需要，我可以分享它背后的实际代码。

Answer 1

将__device__ 装饰器添加到decoupled_functions.cuh 中的原型。这应该可以解决您看到的错误消息。

然后你需要在你的模块中使用separate compilation and linking。因此，不要使用-c 编译，而是使用-dc 编译。并且您的链接命令将需要修改。一个基本的例子是here。

你的问题有点混乱：

问题：为什么 DecoupledCallGpu 是从主机函数而不是内核调用的？

我不知道你是在说英语，还是这里有误会。实际的错误消息指出：

错误：不允许从 __global__ 函数（“内核”）调用 __host__ 函数（“DecoupledCallGpu”）

这是由于在编译单元内（即在模块内，在正在编译的文件内，即cuda_computations.cu），函数的唯一描述DecoupledCallGpu() 是原型中标题中提供的：

void DecoupledCallGpu(int *var);

这个原型表示CUDA C中的一个未修饰函数，这些函数是equivalent to__host__（仅）修饰函数：

__host__ void DecoupledCallGpu(int *var);

该编译单元不知道 decoupled_functions.cu 中的实际内容。

因此，当你有这样的内核代码时：

__global__ void kernel(){       //<- __global__ function
...
DecoupledCallGpu(&var_kernel);  //<- appears as a __host__ function to compiler
}

编译器认为您试图从 __global__ 函数调用 __host__ 函数，这是非法的。