hcc equivalent of nvcc's "--device-link"

[atabish]

Hi, I want to compile and run cpu code that contains no hip/cuda code (but links against a library that does) with g++. I can do it if HIP_PLATFORM is nvcc as follows: hipcc -rdc=true -c -o temp.o somefile.cu hipcc -dlink -o somefile.o temp.o -lcudart ar cru libgpu.a somefile.o temp.o ranlib libgpu.a g++ main.cpp -L. -lgpu -o main -L/usr/local/cuda/lib64 -lcudart

But I can not do the same steps if HIP_PLATFORM is hcc. I am assuming it has to do with "-dlink" flag not being implemented in hcc. Please let me know if there is a way to repeat the above procedure for HIP_PLATFORM=hcc thanks

[gargrahul]

@atabish Could you try https://github.com/ROCm-Developer-Tools/HIP/blob/master/tests/src/g%2B%2B/hipMalloc.cpp or tests from https://github.com/ROCm-Developer-Tools/HIP/tree/master/tests/src/gcc ?

[atabish]

@gargrahul I was able to compile and run the test you referred to. But my question is more about separate compilation. My goal is to use HIP to create a library of my gpu code and then use g++ to compile and link my cpu code to that library. Let's say we have the following files (similar to the test you referred to):
hipMalloc.h

#include <hip/hip_runtime_api.h>
#include <iostream>

int foo();

hipMalloc.cpp (would also contain some device code)

int foo() {
    int* Ad;
    hipMalloc((void**)&Ad, 1024);
    std::cout<<"PASSED!"<<std::endl;
    return 0;
}

test.cpp (contains only host code)

#include "hipMalloc.h"
int main() {
    foo();
    return 0;
}

Now my goal is to pack hipMalloc.cpp into a library (lets say libgpu.a) and use that with my test.cpp. Something as follows:

1. /opt/rocm/bin/hipcc -fPIC -c -o temp.o hipMalloc.cpp
2. ar cru libgpu.a temp.o 
3. ranlib libgpu.a
4. g++ -D__HIP_PLATFORM_HCC__ -I/opt/rocm/hip/include -I/opt/rocm/include test.cpp -Wl,--rpath=/opt/rocm/hip/lib /opt/rocm/hip/lib/libhip_hcc.so -L. -lgpu  -o test -std=c++11

This gives the error

./libgpu.a(temp.o): In function `foo()':
(.text+0x15): undefined reference to `hipMalloc'
collect2: error: ld returned 1 exit status

My guess is that between steps 1 and 2 above, I need to somehow link temp.o to the necessary HIP libraries. Please let me know what is the right way to achieve this. Thanks

[SarbojitAMD]

@atabish, hipcc generates LLVM IR during compilation and final GPU codes get generates during linking hence hipcc should be used to create libgpu.a/lingpu.so Once library is generated using hipcc you can use it with host code and compile it g++.

For clear steps and some more details please refer link https://github.com/ROCm-Developer-Tools/HIP/issues/1055

Let me know if you have more questions.

[atabish]

@SarbojitAMD Thank you very much. How to modify your second step to create a static library? I tried

/opt/rocm/bin/hipcc -static -o libgpu.a temp.o /opt/rocm/hip/lib -lhip_hcc_static

and I got:

ld: attempted static link of dynamic object `/opt/rocm/lib/libhip_hcc.so'
clang-9: error: linker command failed with exit code 1 (use -v to see invocation)

[atabish]

@SarbojitAMD The shared library approach that you suggested in #1055 works well. My question now is, how to translate the first two steps to CMake:

1. hipcc -fPIC $(INC) -c gpu.cpp // will generate gpu.o object file
2. hipcc -shared -o libgpu.so gpu.o $(LIB) -i:libhip_hcc.so // will generate shared library name libgpu.so

If I just do:

hip_add_library(test_gpu SHARED ${gpu_src} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} HCC_OPTIONS ${MY_HCC_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS})

It seems like it compiles with hipcc but links with g++. The executable generated then crushes. I am looking for a generic CMake solution that I can build library which is runable on both AMD and NVidia cards.

Thanks for your help so far.

[Datta2324]

Hi, I am new to this I have ported CUDA code into HIP to run vector addition on Multiple GPUs using MPI and in makefile in CUDA the is a linker -lcudart . I am getting error with that in HIP . I have tried with hiprtc but it did not work, Can anyone help me

[fwyzard]

Now my goal is to pack hipMalloc.cpp into a library (lets say libgpu.a) and use that with my test.cpp. Something as follows:

1. /opt/rocm/bin/hipcc -fPIC -c -o temp.o hipMalloc.cpp
2. ar cru libgpu.a temp.o 
3. ranlib libgpu.a
4. g++ -D__HIP_PLATFORM_HCC__ -I/opt/rocm/hip/include -I/opt/rocm/include test.cpp -Wl,--rpath=/opt/rocm/hip/lib /opt/rocm/hip/lib/libhip_hcc.so -L. -lgpu  -o test -std=c++11

I know this is pretty old, but since I stumbled upon it while looking for a similar solution, here is what seems to work for me (using ROCm 5.4):

1. /opt/rocm/bin/hipcc -fPIC -fgpu-rdc -c -o temp.o hipMalloc.cpp
2. /opt/rocm/bin/hipcc -fPIC -fgpu-rdc --hip-link --emit-static-lib -o libgpu.a temp.o 
3. # not needed
4. g++ -D__HIP_PLATFORM_HCC__ -I/opt/rocm/include test.cpp -Wl,--rpath=/opt/rocm/lib -L/opt/rocm/lib -lamdhip64 -L. -lgpu -o test -std=c++11

[ppanchad-amd]

@atabish Apologies for the lack of response. Can you please test with latest ROCm 6.1.0 (HIP 6.1)? If resolved, please close ticket. Thanks!