CUDA设备功能的多重定义

Multiple definition of CUDA device functions

本文关键字：定义 功能 CUDA      更新时间：2023-10-16

我正在尝试编译某些功能以在主机代码和设备CUDA代码中使用它们，但是我得到了多个定义链接错误。我试图实现的目标是：

我有一个带有以下内容的cudaconfig.h文件

cudaconfig.h 

#ifdef __CUDACC__
#define CUDA_CALLABLE_DEVICE __device__
#define CUDA_CALLABLE_HOST __host__
#define CUDA_CALLABLE __host__ __device__
#else
#define CUDA_CALLABLE_DEVICE
#define CUDA_CALLABLE_HOST
#define CUDA_CALLABLE
#endif

在我的foo.h文件中，我有一些功能，具有以下签名

#include "CudaConfig.h"
struct Bar {Eigen::Vector3d v;};
CUDA_CALLABLE_DEVICE Eigen::Vector3d &foo(Bar &aBar);

我在foo.cpp和foo.cu文件中实现它们。

foo.cpp 

#include "foo.h"
Eigen::Vector3d &foo(Bar &aBar) {aBar.v += {1,1,1}; return aBar.v;}

foo.cu 

#include "foo.h"
Eigen::Vector3d &foo(Bar &aBar) {aBar.v += {1,1,1}; return aBar.v;}

我需要将不同文件中的两个实现分开，因为当您从__device__函数中使用它时，eigen会禁用某些SIMD操作，因此出于性能原因，我不想在foo.cu文件中实现这两个操作。

我应该在.h文件中直接实现该函数，将它们标记为内联，以便我没有多个定义链接错误？当特征委员会禁用__device__代码的SIMD时，这是否会使__host__和__device__的功能不同，而与Inline期望不同？

这是发生的事情：

rthoni@rthoni-lt1:~/projects/nvidia/test_device_host$ cat test.cu
extern "C" {
__device__ void test_device_fn()
{
}
}
rthoni@rthoni-lt1:~/projects/nvidia/test_device_host$ nvcc test.cu -c -o test_cu.o
rthoni@rthoni-lt1:~/projects/nvidia/test_device_host$ objdump -t test_cu.o 
test_cu.o:     file format elf64-x86-64
SYMBOL TABLE:
0000000000000000 l    df *ABS*  0000000000000000 tmpxft_000004d9_00000000-5_test.cudafe1.cpp
0000000000000000 l    d  .text  0000000000000000 .text
0000000000000000 l    d  .data  0000000000000000 .data
0000000000000000 l    d  .bss   0000000000000000 .bss
0000000000000000 l     O .bss   0000000000000001 _ZL22__nv_inited_managed_rt
0000000000000008 l     O .bss   0000000000000008 _ZL32__nv_fatbinhandle_for_managed_rt
0000000000000000 l     F .text  0000000000000016 _ZL37__nv_save_fatbinhandle_for_managed_rtPPv
0000000000000010 l     O .bss   0000000000000008 _ZZL22____nv_dummy_param_refPvE5__ref
000000000000002f l     F .text  0000000000000016 _ZL22____nv_dummy_param_refPv
0000000000000000 l    d  __nv_module_id 0000000000000000 __nv_module_id
0000000000000000 l     O __nv_module_id 000000000000000f _ZL15__module_id_str
0000000000000018 l     O .bss   0000000000000008 _ZL20__cudaFatCubinHandle
0000000000000045 l     F .text  0000000000000022 _ZL26__cudaUnregisterBinaryUtilv
0000000000000067 l     F .text  000000000000001a _ZL32__nv_init_managed_rt_with_modulePPv
0000000000000000 l    d  .nv_fatbin 0000000000000000 .nv_fatbin
0000000000000000 l       .nv_fatbin 0000000000000000 fatbinData
0000000000000000 l    d  .nvFatBinSegment   0000000000000000 .nvFatBinSegment
0000000000000000 l     O .nvFatBinSegment   0000000000000018 _ZL15__fatDeviceText
0000000000000020 l     O .bss   0000000000000008 _ZZL31__nv_cudaEntityRegisterCallbackPPvE5__ref
0000000000000081 l     F .text  0000000000000026 _ZL31__nv_cudaEntityRegisterCallbackPPv
00000000000000a7 l     F .text  0000000000000045 _ZL24__sti____cudaRegisterAllv
0000000000000000 l    d  .init_array    0000000000000000 .init_array
0000000000000000 l    d  .note.GNU-stack    0000000000000000 .note.GNU-stack
0000000000000000 l    d  .eh_frame  0000000000000000 .eh_frame
0000000000000000 l    d  .comment   0000000000000000 .comment
0000000000000016 g     F .text  0000000000000019 test_device_fn
0000000000000000         *UND*  0000000000000000 _GLOBAL_OFFSET_TABLE_
0000000000000000         *UND*  0000000000000000 exit
0000000000000000         *UND*  0000000000000000 __cudaUnregisterFatBinary
0000000000000000         *UND*  0000000000000000 __cudaInitModule
0000000000000000         *UND*  0000000000000000 __cudaRegisterFatBinary
0000000000000000         *UND*  0000000000000000 atexit

您可以看到，即使该函数仅标记为__device__，nvcc仍将在对象文件中生成符号。

此行为是nvcc的错误。（我们的错误跟踪器中的＃845649）

有三种摆脱此错误的方法：

  • 让nvcc同时生成设备和主机代码
  • 更改您编译cu文件的方式只需构建设备代码
  • 将您的__device__功能包装在空名称空间中

在您的特定情况下，您似乎只能使其成为constexpr未装修功能：

constexpr Eigen::Vector3d &foo(Bar &aBar) noexcept {aBar.v += {1,1,1}; return aBar.v;}

用--expt-relaxed-constexpr调用nvcc：

--expt-relaxed-constexpr                   (-expt-relaxed-constexpr)       
        Experimental flag: Allow host code to invoke __device__ constexpr functions,
        and device code to invoke __host__ constexpr functions.Note that the behavior
        of this flag may change in future compiler releases.

这应该适用于设备和主机代码。

相关文章：