不支持外部调用 - CUDA答案

【问题标题】：External calls are not supported - CUDA不支持外部调用 - CUDA
【发布时间】：2014-01-19 23:54:34
【问题描述】：

目标是调用另一个文件中可用的设备函数，当我编译 global 内核时，它显示以下错误 *External calls are not supported (found non-inlined call to _Z6GoldenSectionCUDA )*.

有问题的代码（不是完整的代码，而是出现问题的地方），猫范数.h

# ifndef NORM_H_
# define NORM_H_
# include<stdio.h>

__device__ double invcdf(double prob, double mean, double stddev);

#endif

cat norm.cu

# include <norm.h>

__device__ double invcdf(double prob, double mean, double stddev) {
    return (mean + stddev*normcdfinv(prob));
       }

猫测试.cu

# include <norm.h>
# include <curand.h>
# include <curand_kernel.h>

__global__ void phase2Kernel(double* out_profit, struct strategyHolder* strategy) {
       curandState seedValue;
       curand_init(threadIdx.x, 0, 0, &seedValue);
       double randomD = invcdf(curand_uniform_double( &seedValue ), 300, 80);
    }

nvcc -c norm.cu -o norm.o -I"."
nvcc -c test.cu -o test.o -I"."

【问题讨论】：

请发布一个示例来重现问题（请参阅here 以获得指导），您的代码对我有用。
在实际代码中（与您在此处显示的相反），您是显式调用构造函数还是依赖于类的默认构造函数？
能否指定您使用的 nvcc 参数？
贴出真实代码。我正在使用Makefile进行编译，我使用的命令是nvcc -c file.cu -o file.o
@Bala 我不明白为什么将上述代码的一部分放在单独的.h 文件中，然后将其包含在.cu 文件中会导致编译错误。您能否使用重现错误的确切文件分区编辑您的问题？

标签： c++ cuda gpgpu nvidia

【解决方案1】：

您正在尝试进行单独编译，这需要一些特殊的命令行选项。有关详细信息，请参阅NVCC manual，但这里是如何让您的示例编译。我已针对 sm_20，但您可以针对 sm_20 或更高版本，具体取决于您拥有的 GPU。旧设备 (sm_1x) 上无法单独编译。

您不需要在头文件中将__device__ 函数声明为extern，但如果您有任何静态设备变量，则需要将它们声明为extern
如下所示，通过编译为设备生成可重定位代码（-dc 是-c 的设备等价物，有关详细信息，请参阅manual）
```
nvcc -arch=sm_20 -dc norm.cu -o norm.o -I.
nvcc -arch=sm_20 -dc test.cu -o test.o -I.
```
通过在最终主机链接之前调用 nvlink 来链接代码的设备部分
```
nvlink -arch=sm_20 norm.o test.o -o final.o
```

【讨论】：

是否需要将__device__函数声明为extern？我正在使用不带extern 的-rdc=true 选项编译上述代码。有什么我没看到的吗？
@Tom extern 我猜并不是真的需要。它可以按照您提供的链接中指出的说明进行编译。
extern 对于静态定义的符号绝对是必需的，以防止链接期间出现重复错误，但对于函数来说，前向声明应该足够了。
正如 talonmies 所说，函数声明中不需要extern，我太仓促了。我已经修改了答案以向未来的访问者反映这一点。