Множественное определение функций устройства CUDA

Question

Я пытаюсь скомпилировать некоторые функции, чтобы использовать их в коде хоста и коде устройства, но я получаю ошибку связывания с несколькими определениями. Я пытаюсь достичь следующего:

У меня есть файл 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;}

Мне нужно разделить обе реализации в разных файлах, так как Eigen отключает некоторые операции SIMD при использовании его из функции __device__, поэтому я не хочу реализовывать обе в файле foo.cu по соображениям производительности.

Должен ли я реализовать эту функцию непосредственно в файле .h, помечая их как встроенные, чтобы у меня не было ошибки связывания с несколькими определениями? Поскольку Eigen отключает SIMD для кода __device__, не приведет ли это к тому, что функции __host__ и __device__ будут отличаться от ожидаемых inline?

Answer 1

Вот что происходит:

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 в нашем трекере ошибок)

Есть 3 способа избавиться от этой ошибки:

• Пусть nvcc генерирует код устройства и хоста
• Измените способ компиляции cu файлов, чтобы просто создать код устройства
• Оберните вашу функцию __device__ в пустое пространство имен

Answer 2

В вашем конкретном случае похоже, что вы можете просто сделать ее constexpr недекорированной функцией:

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

и вызвать nvcc с помощью --expt-relaxed-constexpr:

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

это должно работать как для устройства, так и для кода хоста.