Multiple symbol definition error for sqrt()

[catcondo]

I am creating a pipeline from 2 modules, one module holds the __raygen and _miss programs while the other holds the __closesthit and __intersection programs. However, if I use sqrt() in the 2nd module, I get the following run time error:

COMPILE ERROR: 
Info: Pipeline has 2 module(s), 4 entry function(s), 6 trace call(s), 0 continuation callable call(s), 0 direct callable call(s), 318 basic block(s) in entry functions, 5677 instruction(s) in entry functions, 136 non-entry function(s), 708 basic block(s) in non-entry functions, 11170 instruction(s) in non-entry functions
Error: Symbol '_ZSt4sqrtf' was defined multiple times. First seen in: '__raygen__blort_and_1_more_ID1'

The only #includes in that 2nd module are < cuda_runtime.h > and < optix.h >.

If I change that sqrt() call to sqrtf(), then the error no longer occurs, and everything runs fine. I will also mention that when I include additional files like <thrust/complex.h> in the 2nd module, even more pipeline link errors appear, despite the fact that my code does not even instantiate a complex value:

COMPILE ERROR:
Info: Pipeline has 2 module(s), 4 entry function(s), 6 trace call(s), 0 continuation callable call(s), 0 direct callable call(s), 318 basic block(s) in entry functions, 5677 instruction(s) in entry functions, 138 non-entry function(s), 710 basic block(s) in non-entry functions, 11172 instruction(s) in non-entry functions

Error: Symbol '_ZN6thrust6system6detail10sequential3tagC1Ev' was defined multiple times. First seen in: '__intersection__sphere_and_1_more_ID2'

Error: Symbol '_ZN6thrust6system6detail10sequential3tagC2Ev' was defined multiple times. First seen in: '__intersection__sphere_and_1_more_ID2'

Error: Symbol '_ZSt4sqrtf' was defined multiple times. First seen in: '__intersection__sphere_and_1_more_ID2'

Is this a limitation of the pipeline linker? Is the linker not really intended to link amongst multiple modules even though its interface appears to support that? Or are there certain files, eg cuda_runtime.h which should not be included in Optix programs? Thanks.

System Details
CENTOS 7 – 3.10.0-957.el7.x86_64 #1 SMP Thu Nov 8 23:39:32 UTC 2018 x86_64 x86_64 x86_64 GNU/Linux
Cuda 10.2
Optix 7.4
NVIDIA GeForce RTX 2080 Ti
GPU memory: 11264 MiB
Driver: 510.54
gcc (GCC) 4.8.5 20150623 (Red Hat 4.8.5-36)

[droettger]

The cuda_runtime.h header defines the CUDA Runtime API host functions, means there is not a single function defined with __device__ in it. There shouldn’t be any need to include that inside OptiX device code.
Just remove it and see if your device code still compiles.
If you’re referencing any CUDA device code function inside your OptiX device code which happens to be included by that header indirectly, include the resp. headers defining the required functions themselves instead.

That the multiple definitions happen with sqrt() and not sqrtf() would indicate that the former is not correctly declared as inline when it’s defined in a header. Similarly for the thrust headers then.

Note that OptiX imposes some restrictions on what CUDA functionality can be used inside device code (no shared memory, no barriers, no synchronization, and some other constructs): https://raytracing-docs.nvidia.com/optix7/guide/index.html#introduction#overview
Means if any of the thrust functions you’re planning to use is using any of these CUDA features, that will not work.
There are obviously no such restrictions when using thrust in native CUDA kernels running between optixLaunch calls, like you did in the past with OptiX Prime.

I’m using separate *.cu files and OptiX modules for all OptiX program domains inside my examples and have not experienced that issue, but I’m also only ever including optix.h and my own headers inside OptiX device code.

As workaround, it should work when combining the modules which result in the duplicate definition errors to one module.

Please read the related threads below. There are some cases where this is expected and one where it was not, but latter should be solved in a future driver release.

https://forums.developer.nvidia.com/t/geforce-driver-511-17-made-my-program-not-work/200327
https://forums.developer.nvidia.com/t/latest-driver-causing-multiple-symbol-definition-error/188884
https://forums.developer.nvidia.com/t/multiple-symbol-definition-error/191065

[catcondo]

Thank you for the detailed and helpful reply. I will give your suggestions a try.

[catcondo]

I’m using separate *.cu files and OptiX modules for all OptiX program domains inside my examples and have not experienced that issue, but I’m also only ever including optix.h and my own headers inside OptiX device code.

So I tried as suggested above and removed the < cuda_runtime.h > include, so that the only include is < optix.h >, and I am still experiencing the multiple symbol definition error if I call sqrt() from the __ closesthit __ / __ intersection __ .cu file.

The cuda_runtime.h header defines the CUDA Runtime API host functions, means there is not a single function defined with __device__ in it. There shouldn’t be any need to include that inside OptiX device code.

I needed to use cuda_runtime.h in some class headers that mentioned builtin types such as float3 in order that .cpp files using them would correctly compile. In other words, my ``make’’ uses g++ for .cpp files and nvcc for .cu files.

If you’re referencing any CUDA device code function inside your OptiX device code which happens to be included by that header indirectly, include the resp. headers defining the required functions themselves instead.

I will work on cleaning up my includes. However, as my original question mentioned, even something seemingly as basic as thrust/complex.h causes a PTX link error. That error refers to thrust::system::detail::sequential::tag::tag() being defined multiple times. Whatever in the world that has to do with complex numbers is beyond me. :)

Thanks.

[droettger]

If you have a minimal reproducer for the thrust issue, we can file a bug report to have that investigated.

[droettger]

Do you have --relocatable-device-code=true (-rdc) set in your NVCC compilation options for the CUDA to PTX translation?
If not, does the error still happen when you add it?

[catcondo]

I added that option to the .cu → .ptx translation step, and those run time link errors went away. Thank you very much for remembering my original post and identifying the fix!