c++ - 如何正确地将 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(具有内核、主机主机功能并包含具有设备功能的 header ):

#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 中的原型(prototype)。这应该会处理您看到的错误消息。

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

你的问题有点困惑:

Question: why is it that the DecoupledCallGpu is called from host function and not a kernel as it was supposed to?

我不知道你是被英语绊倒了还是这里有误解。实际的错误消息指出:

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

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

void DecoupledCallGpu(int *var);

这个原型(prototype)表示 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__ 函数，这是非法的。