Failure to link against a HIP-library using CMake (HIP_PLATFORM : nvidia)

[MihkuU]

Hi! I have been trying to create a shared-library, which uses HIP to enable fast computations with both AMD and NVidia GPUs. I think that it would be preferable to put the HIP-kernels into header files and the definitions into source-files as usual. This however requires using an option for enabling the 'relocatable device code' (-rdc=true for NVidia). The problem is that I am not able to get the build with CMake to work.

The library files: kernel.hpp:

#pragma once 

#include <hip/hip_runtime.h>

namespace kernel
{
    __device__ void hello();
}

kernel.cpp:

#include "kernel.hpp"
#include <stdio.h>

__device__ void kernel::hello()
{
    printf("Hello HIP\n");
}

libb.hpp:

#pragma once

namespace libb
{
    void printGPUInfo();
}

libb.cpp:

#include "libb.hpp"
#include "kernel.hpp"

#include <hip/hip_runtime.h>

__global__ void mainKernel()
{
    kernel::hello();
}

void libb::printGPUInfo()
{
    mainKernel<<<1,1>>>();
    hipDeviceSynchronize();
}

The main source file main.cpp:

#include "libb.hpp"

int main()
{
    libb::printGPUInfo();
}

By running the commands

hipcc --shared -rdc=true --compiler-options '-fPIC' -o libdyn.so libb.cpp kernel.cpp
hipcc main.cpp -o run -L$(pwd) -ldyn

everything seems to work fine, however I am not able to get it to run with CMake. I built the library inside a directory below (lib/) with the given CMakeLists.txt-file:

cmake_minimum_required(VERSION 3.25 FATAL_ERROR)

project(hiplib)

enable_language(HIP)

set(libb_srcs libb.cpp kernel.cpp)
set_source_files_properties(${libb_srcs} PROPERTIES LANGUAGE HIP)
set_source_files_properties(${libb_srcs} PROPERTIES COMPILE_FLAGS -rdc=true)

add_library(libb SHARED ${libb_srcs})

and then when I tried to link the main-executable against the library libdyn.so,I got the error:

/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.so: undefined reference to `__cudaRegisterLinkedBinary_ef17c27b_8_libb_cpp_de10737b'
/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.so: undefined reference to `__cudaRegisterLinkedBinary_eacc4d77_10_kernel_cpp_7d3f2376'

An analogous problem can be solved in CUDA (https://gist.github.com/gavinb/c993f71cf33d2354515c4452a3f8ef30), by setting specific CMake target properties. The solution didn't work for me with HIP and I am not able to overcome this linker error. Does any one of you have an idea of what might be the fix? Help would be much appreciated!

[cjatin]

Can you add these line to your cmake (right after the project declaration) and try again.

set(CMAKE_C_COMPILER hipcc) set(CMAKE_CXX_COMPILER hipcc)

[MihkuU]

Thank you for the reply. I tried but got the same error.

[cjatin]

I had a chat with the compiler folks, turns out -fgpu-rdc option is supported only with static libs and not dynamic lib.

so it needs to be: add_library(libb STATIC ${libb_srcs})

[MihkuU]

This gives me the following errors:

/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.a(libb.cpp.o): in function `hipDeviceSynchronize':
tmpxft_000022dd_00000000-6_libb.cudafe1.cpp:(.text+0x3ca): undefined reference to `cudaDeviceSynchronize'
/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.a(libb.cpp.o): in function `libb::printGPUInfo()':
tmpxft_000022dd_00000000-6_libb.cudafe1.cpp:(.text+0x44c): undefined reference to `__cudaPushCallConfiguration'
/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.a(libb.cpp.o): in function `__cudaUnregisterBinaryUtil()':
tmpxft_000022dd_00000000-6_libb.cudafe1.cpp:(.text+0x4b0): undefined reference to `__cudaUnregisterFatBinary'
/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.a(libb.cpp.o): in function `__nv_init_managed_rt_with_module(void**)':
tmpxft_000022dd_00000000-6_libb.cudafe1.cpp:(.text+0x4cf): undefined reference to `__cudaInitModule'
/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.a(libb.cpp.o): in function `__device_stub__Z10mainKernelv()':
tmpxft_000022dd_00000000-6_libb.cudafe1.cpp:(.text+0x543): undefined reference to `__cudaPopCallConfiguration'
/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.a(libb.cpp.o): in function `__nv_cudaEntityRegisterCallback(void**)':
tmpxft_000022dd_00000000-6_libb.cudafe1.cpp:(.text+0x650): undefined reference to `__cudaRegisterFunction'
/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.a(libb.cpp.o): in function `__sti____cudaRegisterAll()':
tmpxft_000022dd_00000000-6_libb.cudafe1.cpp:(.text+0x68c): undefined reference to `__cudaRegisterLinkedBinary_ef17c27b_8_libb_cpp_de10737b'
/usr/bin/ld: /home/alastor/tmp/hip_dynlink/lib/build/libdyn.a(libb.cpp.o): in function `cudaError cudaLaunchKernel<char>(char const*, dim3, dim3, void**, unsigned long, CUstream_st*)':
tmpxft_000022dd_00000000-6_libb.cudafe1.cpp:(.text+0x6dd): undefined reference to `cudaLaunchKernel'

Seems like something is still missing from the library?

[MihkuU]

I had a chat with the compiler folks, turns out -fgpu-rdc option is supported only with static libs and not dynamic lib.

so it needs to be: add_library(libb STATIC ${libb_srcs})

I thought the -fgpu-rdc is an option for the AMD platform? Im on nvidia right now and so the hipcc calls the nvcc CUDA compiler driver as I understand it. That's why I used the CUDA option -rdc=true. In a few days I will get to a computer with AMD GPU and will try there.

For now I got it to work by making the kernels inline and defining them entirely in header files.

[ppanchad-amd]

Hi @MihkuU, do you still need assistance with this ticket? If not, please close the ticket. Thanks!

[MihkuU]

Thanks for the reply, for now I have a workaround, so I'll close it.