[HIPIFY] Incorrect HIP generated from CUDA for DNN based YOLO Algorithm

[rohanmahapatra]

Hi,

I have been trying to run Hipify-clang on a YOLO algorithm written using CUDA. However, hipify-clang does not correctly generate a HIP version of the CUDA file. Below are some of the errors I am facing

Some kernels are unable to convert keywords in CUDA: /u/p/r/preyesh/private/darknet_tiny_preyesh/src/activation_kernels.hip.cu:151:39: error: use of undeclared identifier 'gridDim' int id = (blockIdx.x + blockIdx.ygridDim.x) blockDim.x + threadIdx.x;

        But gridDim is a keyword in CUDA and error should not be flagged. 

   2. Does not recognize #define

       <command line>:1:11: note: expanded from here
       #define X 1
                         ^

   3. For kernels like batch_norm, we fixed the typecasting issues by replacing the below code
        l.scales = calloc(c, sizeof(float)); with 
        l.scales = (float *)calloc(c, sizeof(float));

However, even after these changes hipify-clang does not generate the correct code -- the .hip file which is generated is either empty or exactly the same as the .cu file. We are using roc-1.6.x branch of hipify-clang as that is the version compatible with the gem5 simulator.

Could you please tell us if we are missing something because of which were are not able to generate the .hip files correctly?

Thank you, Rohan Mahapatra

[emankov]

Could you please provide the source CUDA file and your command line?

[rohanmahapatra]

Hi,

I used the following source files (i used other .cu files present in the same directory as well) https://github.com/pjreddie/darknet/blob/master/src/activation_kernels.cu

And the command was as mentioned in https://github.com/ROCm-Developer-Tools/HIP/tree/roc-1.6.x/hipify-clang. I added my local paths where every it was required in the below command to include the source files and CUDA installation path.

./hipify-clang -print-stats sort_kernel.cu -- -x cuda -I/srv/git/HIP/include -I/usr/local/cuda-7.5/include -DX=1

Thanks, Rohan

[emankov]

Thanks!

First of all, please refer to the Readme.md from the master branch and clone sources from it as well. 2. hipify-clang is HIP independent tool now, so you may build it separately from HIP, and the include path to HIP (-I/srv/git/HIP/include) is not needed anymore.

Cuda 7.5, well ok... Could you please run hipify-clang -version and provide an output?

[emankov]

Hi @rohanmahapatra,

I've just checked your issue on my side, but only on Windows (with pthread.h port for Windows). I've tested it with CUDA 7.0, 7.5, 8.0, and 9.0. All is ok, except for one thing: #include "cuda.h", which is not CUDA's, but darknet's. So I changed the name to cuda_darknet.h. I'll think about avoiding such conflicts with CUDA headers.

hipify-clang was built from sources taken from the 'master' branch and based on LLVM 6.0.1. Nonetheless, I believe It should work on Linux as well. Tomorrow I'll let you know.

Here is my command line:

hipify-clang d:\GIT\darknet\src\activation_kernels.cu --cuda-path="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v7.5" -DCUDA_VERSION=7500 -Id:\GIT\darknet\include -Id:\GIT\darknet\src -DGPU -D_TIMESPEC_DEFINED -print-stats -o-dir=.

Here is a hipified result:

activation_kernels.cu.zip

[emankov]

Hi @rohanmahapatra,

Linux command line:

./hipify-clang ~/git/darknet/src/activation_kernels.cu -I/home/emankov/git/darknet/include -I/home/emankov/git/darknet/src -print-stats -DCUDA_VERSION=8000 -DGPU

With cuda.h to cuda_darknet.h workaround.

If you have multiple CUDA versions installed, --cuda-path should be specified.

So please check it is working on your side and also let me know of any other issues with hipification.

[rohanmahapatra]

Sorry for the delay - Running hipify-clang -version outputs

$ hipify-clang -version LLVM (http://llvm.org/): LLVM version 3.8.0 Optimized build. Default target: x86_64-unknown-linux-gnu Host CPU: (unknown)

I am using a simulator named GEM5 (https://github.com/gem5/gem5) which supports ROCM 1.6 and so will hipify-clang work with ROCM ROC-1.6 version or do I need the ROCM master branch as well?

Thanks, Rohan

[emankov]

Well, LLVM 3.8.0... Mmm, ok. Let's test with the oldest LLVM 3.8.0.

[emankov]

Ping

[rohanmahapatra]

Hello, I am there.

[emankov]

What about building hipify-clang from scratch (better based on llvm+clang 6.0.1 release)? hipify-clang as I said a HIP/Rocm independent tool.

[rohanmahapatra]

Hi, I tried building hipify-clang from scratch with the below command and used the versions you had asked me to use.

cmake -DHIPIFY_CLANG_TESTS=1 -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=../dist -DCMAKE_PREFIX_PATH=../../clang+llvm-6.0.1-x86_64-linux-gnu-ubuntu-16.04/ -DCUDA_TOOLKIT_ROOT_DIR=/p/hal/private/libs/cuda-8.0/ ..

Output is below. This was not as per what is present in the readme. -- HIP Platform: nvcc -- HIP Compiler: hcc -- HIP will be installed in: /nobackup/hipify-clang/HIP/dist -- Configuring done -- Generating done -- Build files have been written to: /nobackup/hipify-clang/HIP/hipify-clang

Thanks, Rohan

[emankov]

Output is below. This was not as per what is present in the readme.

You're right, cause you built HIP. Please refer to Readme to build hipify-clang only. As for the building HIP, by default, it doesn't build hipify-clang at all.

And if you want to build hipify-clang with testing support, setting -DHIPIFY_CLANG_TESTS=1 only is insufficient. -DCUDA_DNN_ROOT_DIR and -DLLVM_EXTERNAL_LIT should be specified as well.

[emankov]

ping

[rohanmahapatra]

I did use the README to build hipify-clang. Problem was that the output was not like what was expected.

Please give a time until this weekend and I will redo everything once again and then update you.

Thank you, Rohan

[emankov]

Sure, no problem, it's your ticket.

[rohanmahapatra]

Hi, Sorry for the delay. I installed LLVM + Clang 6.0.1 with CUDA 8 and followed the read me steps to build Hipify-clang in testing mode. However, when I use the below command as you had told earlier, I get one error as shown below and the output does not match to what is displayed in the README.

cmake -DHIPIFY_CLANG_TESTS=1 -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=../dist -DCMAKE_PREFIX_PATH=../../llvm-6.0.1.src/dist/ -DCUDA_TOOLKIT_ROOT_DIR=/p/hal/private/libs/cuda-8.0 -DCUDA_DNN_ROOT_DIR=/p/hal/private/libs/cuda-8.0/ -DLLVM_EXTERNAL_LIT=../../llvm-6.0.1.src/build/bin/llvm-lit .. -- Could NOT find LibXml2 (missing: LIBXML2_LIBRARIES LIBXML2_INCLUDE_DIR) -- Native target architecture is X86 -- Threads enabled. -- Doxygen disabled. -- Go bindings disabled. -- OCaml bindings disabled, need ctypes >=0.4. -- Could NOT find Python module pygments -- Could NOT find Python module pygments.lexers.c_cpp -- Found Python module yaml -- LLVM host triple: x86_64-unknown-linux-gnu -- LLVM default target triple: x86_64-unknown-linux-gnu -- Building with -fPIC -- Constructing LLVMBuild project information -- Linker detection: GNU ld -- Targeting AArch64 -- Targeting AMDGPU -- Targeting ARM -- Targeting BPF -- Targeting Hexagon -- Targeting Lanai -- Targeting Mips -- Targeting MSP430 -- Targeting NVPTX -- Targeting PowerPC -- Targeting Sparc -- Targeting SystemZ -- Targeting X86 -- Targeting XCore -- Configuring done -- Generating done -- Build files have been written to: /u/m/a/mahapatra/private/gpu_project/hipify/llvm-6.0.1.src/build [mahapatra@hylia] (98)$ make test-hipify make: *** No rule to make target 'test-hipify'. Stop. [mahapatra@hylia] (99)$

Also, I installed lit using the steps provided in the README. Should I point to LIT inside LLVM or in my /usr/local/bin/

[emankov]

It looks like you ran cmake for llvm...

[rohanmahapatra]

Yes, I followed the README from the master branch. I was supposed to follow the README right?

cmake \ -DCMAKE_INSTALL_PREFIX=../dist \ -DCMAKE_BUILD_TYPE=Release \ ..

[emankov]

So you ran cmake command intended for hipify-clang in llvm subdir, right? Well, what did you want to happen?

Building LLVM is building LLVM. Building hipify-clang is building hipify-clang. Hipify-clang depends on LLVM binaries, which might be used from packages or built from sources. Hipify-clang with testing support needs LLVM built from sources, distributive LLVM packages are insufficien in that case - it is mentioned in Readme.

Please, provide where is it written in Readme to run cmake in llvm tree for building hipify-clang (with or without testing support).

[rohanmahapatra]

Thank you, I got it working and could match your .hip which you generated. I will let you know in case I come across some issues.

[emankov]

Ok