Help with CMake build of CUDA project, resolving `ptxas fatal : Unresolved extern function '_Z8XFluxRhoiiiPKfiii'`

Hello, I have a magnetohydrodynamics simulation that I wrote, and which I am struggling to build with CMake for the reason specified in the title.

After googling, I think the crux of the issue is that I have a family of __device__ functions defined in a file kernels_od_fluxes.cu which are called in __global__ kernels that are defined in separate files. Here is the CMakeLists.txt

cmake_minimum_required(VERSION 3.12)
project(imhd-CUDA LANGUAGES CUDA CXX)

find_package(CUDAToolkit REQUIRED)

set(CMAKE_VERBOSE_MAKEFILE on)

# set(CMAKE_CXX_KEEP_INTERMEDIATE_OBJECTS 1)
# set(CMAKE_CUDA_KEEP_INTERMEDIATE_OBJECTS 1)

add_library(helper_lib STATIC ../include/helper_functions.cu)
add_library(diffusion_lib STATIC ../include/diffusion.cu)
add_library(fluxes_lib STATIC ../include/kernels_od_fluxes.cu)
add_library(kernels_od_lib STATIC ../include/kernels_od.cu)
add_library(initialize_od_lib STATIC ../include/initialize_od.cu)
add_library(intvar_lib STATIC ../include/kernels_od_intvar.cu)
add_library(utils_lib STATIC ../include/utils.cpp)

add_executable(imhd-cuda main.cu)
add_executable(write-grid write_grid.cu)

target_link_libraries(fluxes_lib PRIVATE helper_lib)
target_link_libraries(intvar_lib PRIVATE fluxes_lib)
target_link_libraries(kernels_od_lib PRIVATE fluxes_lib diffusion_lib helper_lib)
target_link_libraries(imhd-cuda PRIVATE kernels_od_lib initialize_od_lib intvar_lib utils_lib)

target_link_libraries(write-grid PRIVATE kernels_od_lib)
target_link_libraries(write-grid PRIVATE initialize_od_lib)
target_link_libraries(write-grid PRIVATE utils_lib)

set_target_properties(kernels_od_lib PROPERTIES CUDA_ARCHITECTURES 75)
set_target_properties(initialize_od_lib PROPERTIES CUDA_ARCHITECTURES 75)
set_target_properties(fluxes_lib PROPERTIES CUDA_ARCHITECTURES 75)
set_target_properties(intvar_lib PROPERTIES CUDA_ARCHITECTURES 75)
set_target_properties(diffusion_lib PROPERTIES CUDA_ARCHITECTURES 75)
set_target_properties(helper_lib PROPERTIES CUDA_ARCHITECTURES 75)
set_target_properties(imhd-cuda PROPERTIES CUDA_ARCHITECTURES 75)
set_target_properties(write-grid PROPERTIES CUDA_ARCHITECTURES 75)

The device functions in question are defined in kernels_od_fluxes.cu, and the kernels which call them are defined in kernels_od.cu and kernels_od_intvar.cu. The order in which I’ve added the libraries, and linked them together, accurately represents the dependencies amongst the code, but when I run

cmake ../src
make clean
make 

I get the following:

~/Desktop/imhd-CUDA/build$ make
[  5%] Building CUDA object CMakeFiles/helper_lib.dir/home/matt/Desktop/imhd-CUDA/include/helper_functions.cu.o
[ 11%] Linking CUDA static library libhelper_lib.a
[ 11%] Built target helper_lib
[ 16%] Building CUDA object CMakeFiles/diffusion_lib.dir/home/matt/Desktop/imhd-CUDA/include/diffusion.cu.o
[ 22%] Linking CUDA static library libdiffusion_lib.a
[ 22%] Built target diffusion_lib
[ 27%] Building CUDA object CMakeFiles/fluxes_lib.dir/home/matt/Desktop/imhd-CUDA/include/kernels_od_fluxes.cu.o
[ 33%] Linking CUDA static library libfluxes_lib.a
[ 33%] Built target fluxes_lib
[ 38%] Building CUDA object CMakeFiles/kernels_od_lib.dir/home/matt/Desktop/imhd-CUDA/include/kernels_od.cu.o
ptxas fatal   : Unresolved extern function '_Z8XFluxRhoiiiPKfiii'
make[2]: *** [CMakeFiles/kernels_od_lib.dir/build.make:76: CMakeFiles/kernels_od_lib.dir/home/matt/Desktop/imhd-CUDA/include/kernels_od.cu.o] Error 255
make[1]: *** [CMakeFiles/Makefile2:179: CMakeFiles/kernels_od_lib.dir/all] Error 2
make: *** [Makefile:91: all] Error 2

What do I need to do in order to link these __device__ functions properly so that they can be resolved by the kernels? Any help is appreciated.

I’m not a CMake expert, and CMake is not a NVIDIA product. However the general thing you need is relocatable device code with device linking, to enable a kernel in one module to call a __device__ function defined in another module (compilation unit, file).

With a bit of searching you can find articles that discuss how to modify your CMakeLists.txt file in order to specify this kind of compilation. You will need something like:

  set_target_properties(${TARGETNAME} PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
1 Like