Compiling programs that use dynamic parallelism (in Thrust) with device link time optimization

Accelerated Computing CUDA CUDA Programming and Performance
,
1

I am attempting to enable link time optimization for device code in an existing project that already uses separable compilation. My understanding is that I need to ensure -dlto is present in both the compilation and linking stage when calling nvcc. However, even with this flag the linking step still gives an error. Based on searching, I have tracked the problem down to Thrust, but I am unsure exactly what is causing the error. Any help would be appreciated. A minimal reproducible example is below, along with the erroneous output.

I am using the following docker container to run the code
nvcr.io/nvidia/cuda:12.1.0-devel-ubuntu20.04

compile.sh

#!/bin/bash

# using dockerfile nvcr.io/nvidia/cuda:12.1.0-devel-ubuntu20.04

FLAGS="-I."

echo "Compiling lib.cu..."
nvcc $FLAGS -dc lib.cu -o lib.o

echo "Linking..."
nvcc $FLAGS -dlink lib.o -o lib_dlink.o

lib.cu

#include <thrust/sort.h>

void sort() {
    int *a = nullptr;
    thrust::sort(a, a);
}

output

Compiling lib.cu...
Linking...
nvlink info    : Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust6system6detail10sequential3seqE
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust6system3cpp3parE
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust8cuda_cub3parE
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust8cuda_cub10par_nosyncE
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders2_1E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders2_2E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders2_3E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders2_4E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders2_5E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders2_6E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders2_7E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders2_8E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders2_9E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust12placeholders3_10E
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust3seqE
Info: no reference to variable __nv_static_34__08ad0cab_6_lib_cu_07cd89d6_592459__ZN43_INTERNAL_08ad0cab_6_lib_cu_07cd89d6_5924596thrust6deviceE
2

I don’t see any errors. The lines in the output that begin with Info: don’t represent errors, AFAIK.