[HIPIFY] hipify-clang fails for application hipify-perl can convert #21

mattsinc · 2017-10-11T00:27:23Z

I tried to run hipify-clang on a simple program (attached for reference). Interac doesn't have any reference calls, and hipify-perl is able to convert it perfectly except that it uses hipLaunchKernel instead of hipLaunchKernelGGL. Unfortunately, hipify-clang is not able to convert it perfectly. Here are some of the errors I'm seeing with hipify-clang:

Doesn't convert kernel launches
The macros that take the block ID x/y/z and thread ID x/y/z fail with repeated error messages like:

note: expanded from macro ‘THREAD ID’
#define THREAD_ID ( (THREADS_PER_BLOCK * BLOCK_ID) + threadIdx.x + (THREADS_PER_BLOCK_X * threadIdx.x) + (THREADS_PER_BLOCK_X * THREADS_PER_BLOCK_Y * threadIdx.z) )

[HIPIFY] warning: interac.cu:205:63: the following reference is not handled: 'pyHostToDevice' [enum constant ref].
[HIPIFY] warning: interac.cu:206:75: the following reference is not handled: 'pyHostToDevice' [enum constant ref].
[HIPIFY] warning: interac.cu:216:5: the following reference is not handled: 'CUevent_st' [struct var ptr].
[HIPIFY] warning: interac.cu:216:5: the following reference is not handled: 'CUevent_st' [struct var ptr].
[HIPIFY] warning: interac.cu:231:5: the following reference is not handled: 'adSynchronize' [function call].
[HIPIFY] warning: interac.cu:234:63: the following reference is not handled: 'pyDeviceToHost' [enum constant ref].
[HIPIFY] warning: interac.cu:237:5: the following reference is not handled: 'tElapsedTime' [function call].
[HIPIFY] warning: interac.cu:274:5: the following reference is not handled: 'tDestroy' [function call].
[HIPIFY] warning: interac.cu:275:5: the following reference is not handled: 'tDestroy' [function call].

So it seems like hipify-clang is having some problems parsing the text. This is the command I ran: /opt/rocm/bin/hipify-clang -o interac.hip.cpp interac.cu -- -x cuda -l/opt/rocm/hip/include/ -l/usr/local/cuda-8.0/targets/x86_64-linux/include/

Here is the hipify-clang info: (hipify-clang --version)

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

So my questions are:

Are these known issues with hipify-clang (kernel launch change, MACRO issues, parsing text for various CUDA calls)?
Is it expected that we should have to run hipify-perl after running hipify-clang to get the remainder of the calls converted properly?

Thanks,
Matt
interac.zip

emankov · 2017-10-12T11:24:20Z

https://github.com/ROCm-Developer-Tools/HIP/blob/master/bin/hipify-perl already support hipLaunchKernelGGL. What HIP repository has been cloned? Could you please attach the resulted files also?

emankov · 2017-10-12T17:14:42Z

I've just fixed the issue with launching kernel in hipify-clang: ROCm/HIP@7ee8e2d
Awaiting the merge into master, thus you may update the changed source on your side and rebuild hipify-clang, and then let me know if it is eliminated. Second part of the issue (thread/block Idx) is in progress, looks like a macro expansion issue too.

mattsinc · 2017-10-12T20:19:11Z

Hi Evgeny,

I have attached a zip file with all 4 versions of Interac -- the original CUDA file, the resultant hipify-perl version, the resultant hipify-clang version, and the version that I ran hipify-perl on after hipify-clang.

I cloned this HIP repository. At the time the head was:

commit a08e30bb20b7f32d17d377eecc63da191243fda0
Merge: a3bc662 0ff34b5
Author: Maneesh Gupta maneesh.gupta@amd.com
Date: Fri Sep 22 12:01:40 2017 +0530

Merge branch 'master' into roc-1.6.x

Change-Id: I8c5861c83032c6006731595ec40e09fdc9102749

Perhaps the GGL support was added to hipify-perl after this?

Thanks,
Matt
interac-all.zip

emankov · 2017-10-13T18:24:54Z

hipify-perl contains hipLaunchKernelGGL after the merge you clonned: https://github.com/ROCm-Developer-Tools/HIP/blob/master/bin/hipify-perl. Moreover in the attached interac.hip.cpp.perl (after perl only, right?) hipLaunchKernelGGL is presented: hipLaunchKernelGGL(HIP_KERNEL_NAME(interac_tm), dim3(grid_size), dim3(block_size ), 0, 0, device_accounts, device_transactions, NUM_TRANSACTIONS);

Actually, HIP_KERNEL_NAME is not needed here (I'll remove it).

As for hipLaunchKernelGGL issue in hipify-clang, have you had a chance to test the fix?

mattsinc · 2017-10-13T18:34:23Z

Sorry, I made a mistake in my wording before -- for the kernel launch command, I updated it to hipLaunchKernelGGL manually. Everything else was automatically done by hipify. It was hipLaunchKernel before I manually changed it. I just re-ran hipify-perl on interac (the version of the HIP repo I linked above) and this is what it output for the kernel launch:

if (useTM) {
    hipLaunchKernel(HIP_KERNEL_NAME(interac_tm), dim3(grid_size), dim3(block_size ), 0, 0, device_accounts, device_transactions, NUM_TRANSACTIONS);
} else {
    hipLaunchKernel(HIP_KERNEL_NAME(interac_atomic), dim3(grid_size), dim3(block_size ), 0, 0, device_accounts, device_transactions, NUM_TRANSACTIONS);
}

I'm not sure why it's doing hipLaunchKernel instead of hipLaunchKernelGGL. I just checked the code in my version of hipify-perl, and it looks like HIP_KERNEL_NAME and hipLaunchKernel are both still there:

        # Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax:                                                                                                                                                         
        $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), $5, $6, /g;
        $kernelName = $1 if $k;

        # Handle the <<numBlocks, blockDim, sharedSize>>> syntax:                                                                                                                                                                 
        $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), $5, 0, /g;
        $kernelName = $1 if $k;

        # Handle the <<numBlocks, blockDim>>> syntax:                                                                                                                                                                             
        $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>([\s\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), 0, 0, /g;
        $kernelName = $1 if $k;

I realize this version of HIP does not include your fixes, but that's strange about hipify-perl -- according to you, I should have the fixes for hipLaunchKernelGGL ...

I need to update/reinstall HIP in order to test the fixes, and I'm not an admin on the system I'm running on, so I haven't had a chance to test the fixes yet. Hopefully this update will resolve the hipify-perl issue too.

Matt

emankov · 2017-10-13T18:51:34Z

Well, actually there is no need in reinstalling HIP, only sources sync are actually needed (git pull origin master). You may apply the above partial fix to Cuda2Hip.cpp and rebuild only hipify-clang. The change with the support of standalone hipify-clang's build has not been merged into master yet, thus here is the attached CMakeLists.txt file to be replaced in hipify-clang folder. Then just do the same build steps as in README.md but in hipify-clang folder.

mattsinc · 2017-10-13T19:06:01Z

I decided to install my own HIP locally on the machine, in case we need to iterate multiple times on a solution. hipify-perl seems to work now, although it seems to have some extra parentheses where the HIP_KERNEL_NAME used to be:

// execute the kernel
if (useTM) {
    hipLaunchKernelGGL(**(interac_tm)**, dim3(grid_size), dim3(block_size ), 0, 0, device_accounts, device_transactions, NUM_TRANSACTIONS);
} else {
    hipLaunchKernelGGL(**(interac_atomic)**, dim3(grid_size), dim3(block_size ), 0, 0, device_accounts, device_transactions, NUM_TRANSACTIONS);
}

I assume this is because the "$1$2" commands on lines 492, 496, and 500 of hipify-perl have parentheses around them (or because this is required)?

I will work on installing and trying your new hipify-clang next.

Matt

emankov · 2017-10-13T19:35:18Z

Parentheses around kernel name are ok and should work, as kernel name may contain characters which are unaccustomed for Identifiers in C/C++ languages, AFAIR, comma, for instance.

mattsinc · 2017-10-13T21:57:43Z

I made the local change to Cuda2Hip.cpp, then remade and reinstalled hipify-clang as you directed. After re-running hipify-clang on the application, I'm now only getting failures for the MACRO. Did you also fix the other issues and I missed it? Because I thought the commit you asked me to try locally only fixes the kernel launch issue.

Here is the error output:

[HIPIFY] warning: interac.cu:52:18: the following reference is not handled: '__fetch_builtin_x' [function call].
[HIPIFY] warning: interac.cu:52:18: the following reference is not handled: '__fetch_builtin_y' [function call].
[HIPIFY] warning: interac.cu:52:18: the following reference is not handled: '__fetch_builtin_z' [function call].
[HIPIFY] warning: interac.cu:52:18: the following reference is not handled: '__fetch_builtin_x' [function call].
[HIPIFY] warning: interac.cu:52:18: the following reference is not handled: '__fetch_builtin_y' [function call].
[HIPIFY] warning: interac.cu:52:18: the following reference is not handled: '__fetch_builtin_z' [function call].
[HIPIFY] warning: interac.cu:74:18: the following reference is not handled: '__fetch_builtin_x' [function call].
[HIPIFY] warning: interac.cu:74:18: the following reference is not handled: '__fetch_builtin_y' [function call].
[HIPIFY] warning: interac.cu:74:18: the following reference is not handled: '__fetch_builtin_z' [function call].
[HIPIFY] warning: interac.cu:74:18: the following reference is not handled: '__fetch_builtin_x' [function call].
[HIPIFY] warning: interac.cu:74:18: the following reference is not handled: '__fetch_builtin_y' [function call].
[HIPIFY] warning: interac.cu:74:18: the following reference is not handled: '__fetch_builtin_z' [function call].
[HIPIFY] warning: interac.cu:218:5: the following reference is not handled: 'CUevent_st' [struct var ptr].
[HIPIFY] warning: interac.cu:218:5: the following reference is not handled: 'CUevent_st' [struct var ptr].

Matt
interac-all.zip

emankov · 2017-10-14T09:02:35Z

These are warnings, let's leave them for now, I'll think about their unwanted triggering improvement. As for the issue with nontranslated thread/block Idx within macro definitions, the fix is not yet ready.

emankov · 2017-10-14T09:13:33Z

Matt,
In the attached interac.hip.cpp.clang kernel launchers are not hipified. Is this file the resulted file of hipify-clang with the above patch applied?

mattsinc · 2017-10-17T03:16:09Z

Sorry I attached the wrong zip file. Reattaching.

Matt
Interac-v2.zip

[HIPIFY] CUDA RT Textures and Arrays support update [HIPIFY] cmake changes [HIPIFY][#199][Partial fix] Fix for cudaLaunchKernel transformation

emankov · 2018-01-16T17:26:03Z

Hi Matt,

I've just checked interac.cu on ToT hipify-clang, all problems reported in this issue are finally eliminated/ I'm going to close the ticket. I want to add interac.cu to our lit testing, is it possible? There are no any legal notices in it.

emankov · 2018-01-30T12:22:58Z

ping

mattsinc · 2018-01-30T17:39:48Z

Hi Evgeny,

Sorry for the delayed response. The reason for my delay is that I'm checking with the original authors of that benchmark about the copyright on it -- they didn't release the benchmark(s) with a license, but they still own the copyright so I need to determine what their terms of use are before it could be included in the test suite. I have messaged them but haven't heard anything back. I will ping them again.

Matt

emankov · 2018-01-30T17:45:32Z

Ok, thanks! Anyway, I suppose we may close this issue as resolved. Could you double check it on your side?

mattsinc · 2018-01-30T17:51:23Z

Just to be clear: you are asking me to update my version of hipify-clang and then re-run it on this benchmark, to make sure the warnings are all gone, right?

If so, I have a few higher priority things but I will try to get that tonight.

Matt

emankov · 2018-02-01T11:11:01Z

It is correctly hipified by ToT hipify-clang now. So, closing the issue and just waiting the answer about licence.

emankov · 2018-02-01T11:11:41Z

Nonrepro on ToT.

mattsinc · 2018-02-06T22:42:13Z

Which branch should I use? I'm not sure what ToT means, sorry ...

In regards to the copyright, this is the response I got:

"... Tor and I have no problem with you using Interac and Hashtable for your work at AMD. You can assume that it is released under the same BSD license as GPGPU-Sim. "

Matt

emankov · 2018-02-06T23:46:59Z

Top of the tree, master branch. The issues reported are nonrepro now, but you may check it on your side.

As for the license, I'll look into GPGPU-Sim, thanks.

mattsinc · 2018-02-07T00:03:16Z

Does hipify-clang require a newer version of llvm now? The README says 3.8+, but I'm getting the following error:

/home/msinclai/downloads/HIP/hipify-clang/src/main.cpp:134:19: error: no member
named 'dbgs' in namespace 'llvm'
DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
~~~~~~^
1 error generated.
hipify-clang/CMakeFiles/hipify-clang.dir/build.make:206: recipe for target 'hipify-clang/CMakeFiles/hipify-clang.dir/src/main.cpp.o' failed

This is what I ran:

cmake -DCMAKE_INSTALL_PREFIX=../inst -DCMAKE_BUILD_TYPE=Release -DBUILD_HIPIFY_CLANG=ON -DCMAKE_PREFIX_PATH="/home/msinclai/downloads/HIP/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-16.04" ..
make -j5 install

emankov · 2018-02-07T21:12:00Z

Fixed with ROCm/HIP#357

mattsinc · 2018-02-07T21:17:10Z

Will give it another try, thanks!

Matt

mattsinc · 2018-02-07T21:34:51Z

Were there changes in how to get the output of hipify-clang? I don't see anything on the README about it, but when I run hipify-clang on the same benchmark now, all I see is a interac.cu.hip file created, but the file has no changes from the original CUDA file.

This is what I ran:

/home/msinclai/downloads/HIP/inst/bin/hipify-clang -print-stats interac.cu -- -x cuda -I/home/msinclai/downloads/HIP/include/ -I/usr/local/cuda-8.0/include/

I see the stats output saying that everything was converted:

[HIPIFY] info: file 'interac.cu' statistics:
CONVERTED refs count: 19
UNCONVERTED refs count: 0
CONVERSION %: 100
REPLACED bytes: 276
TOTAL bytes: 9370
CHANGED lines of code: 17
TOTAL lines of code: 281
CODE CHANGED (in bytes) %: 3
CODE CHANGED (in lines) %: 6
TIME ELAPSED s: 0.03
....

But no actual HIP code. I'm guessing something has changed slightly with the command for hipify-clang that I'm missing?

My guess is that the reason the code is not being generated is that I'm seeing errors like this:

/home/msinclai/benchmarks/gpu-tm-tests/interac/test/interac.cu.hipify-tmp:47:1: error:
unknown type name 'global'
global void interac_atomic( account * accounts, transaction * transa...

Is this a separate, known problem?

Matt

emankov · 2018-02-08T01:43:58Z

Yes, your guess is right, hipified file doesn't contain replacements due to errors. It was done intentionally, as we couldn't garantee correct hipification in case of any compilation errors.

As for the erros you observe, it is because of the changed compilation paradigm. Now instead of regular include path to CUDA headers, path to CUDA root should be specified by the option --cuda-path. It is reflected in readme: https://github.com/ROCm-Developer-Tools/HIP/tree/master/hipify-clang#running-and-using-hipify-clang.

Looking ahead, if you change -I/usr/local/cuda-8.0/include/ to --cuda-path=/usr/local/cuda-8.0 you'll get another, this time cmake error CUDA 8.0 is not supported by clang 3.8. Please install clang 4.0 or higher. That is the reason why you've got the compilation errors. There are 2 options here: install CUDA 7.5 or 7.0 and hipify (I hope :)) successfully by clang 3.8.x, or install clang 4.0 or higher and hipify based on CUDA 8.0. CUDA/clang interoperability (based on clang internals) is coded in https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/CMakeLists.txt:

        if (CUDA_VERSION VERSION_LESS "7.0")
            message(STATUS "Please install CUDA 7.0 or higher.")
        elseif ((CUDA_VERSION VERSION_EQUAL "7.0") OR (CUDA_VERSION VERSION_EQUAL "7.5"))
            message(STATUS "Please install clang 3.8 or higher.")
        elseif (CUDA_VERSION VERSION_EQUAL "8.0")
            message(STATUS "Please install clang 4.0 or higher.")
        elseif (CUDA_VERSION VERSION_EQUAL "9.0")
            message(STATUS "Please install clang 6.0 or higher.")
        elseif (CUDA_VERSION VERSION_EQUAL "9.1")
            message(STATUS "Please install clang 7.0 or higher.")
        endif()

And, btw, long ago there is no need in specifying path to HIP include files. hipify-clang is HIP independent standalone tool, which also might be built now, you know, separately. And we are going to move it in a separate git repo soon and supply it with stable releases.

P.S.
Sorry for so many changes in hipification process since your last using hipify-clang, and thank you for using it - it helps us to make the tool better.

mattsinc · 2018-02-08T18:39:59Z

Thanks Evgeny -- I must not have refreshed the README page, as once I did, the changes you highlighted above were there.

Matt

emankov self-assigned this Oct 14, 2017

emankov referenced this issue in ROCm/HIP Oct 17, 2017

Merge pull request #195 from emankov/master

e8de9dd

[HIPIFY] CUDA RT Textures and Arrays support update [HIPIFY] cmake changes [HIPIFY][#199][Partial fix] Fix for cudaLaunchKernel transformation

emankov closed this as completed Feb 1, 2018

emankov transferred this issue from ROCm/HIP Apr 3, 2020

emankov added perl perl related bug Something isn't working labels Apr 3, 2020

GowriShankarEAAS mentioned this issue Aug 9, 2023

[HIPIFY] clang_cuda_runtime_wrapper.h:41:10: fatal error: 'cmath' file not found #983

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[HIPIFY] hipify-clang fails for application hipify-perl can convert #21

[HIPIFY] hipify-clang fails for application hipify-perl can convert #21

mattsinc commented Oct 11, 2017

emankov commented Oct 12, 2017

emankov commented Oct 12, 2017

mattsinc commented Oct 12, 2017

emankov commented Oct 13, 2017

mattsinc commented Oct 13, 2017 •

edited

emankov commented Oct 13, 2017

mattsinc commented Oct 13, 2017

emankov commented Oct 13, 2017

mattsinc commented Oct 13, 2017

emankov commented Oct 14, 2017

emankov commented Oct 14, 2017

mattsinc commented Oct 17, 2017

emankov commented Jan 16, 2018

emankov commented Jan 30, 2018

mattsinc commented Jan 30, 2018 •

edited

emankov commented Jan 30, 2018

mattsinc commented Jan 30, 2018

emankov commented Feb 1, 2018

emankov commented Feb 1, 2018

mattsinc commented Feb 6, 2018 •

edited

emankov commented Feb 6, 2018

mattsinc commented Feb 7, 2018

emankov commented Feb 7, 2018

mattsinc commented Feb 7, 2018

mattsinc commented Feb 7, 2018 •

edited

emankov commented Feb 8, 2018

mattsinc commented Feb 8, 2018

[HIPIFY] hipify-clang fails for application hipify-perl can convert #21

[HIPIFY] hipify-clang fails for application hipify-perl can convert #21

Comments

mattsinc commented Oct 11, 2017

emankov commented Oct 12, 2017

emankov commented Oct 12, 2017

mattsinc commented Oct 12, 2017

emankov commented Oct 13, 2017

mattsinc commented Oct 13, 2017 • edited

emankov commented Oct 13, 2017

mattsinc commented Oct 13, 2017

emankov commented Oct 13, 2017

mattsinc commented Oct 13, 2017

emankov commented Oct 14, 2017

emankov commented Oct 14, 2017

mattsinc commented Oct 17, 2017

emankov commented Jan 16, 2018

emankov commented Jan 30, 2018

mattsinc commented Jan 30, 2018 • edited

emankov commented Jan 30, 2018

mattsinc commented Jan 30, 2018

emankov commented Feb 1, 2018

emankov commented Feb 1, 2018

mattsinc commented Feb 6, 2018 • edited

emankov commented Feb 6, 2018

mattsinc commented Feb 7, 2018

emankov commented Feb 7, 2018

mattsinc commented Feb 7, 2018

mattsinc commented Feb 7, 2018 • edited

emankov commented Feb 8, 2018

mattsinc commented Feb 8, 2018

mattsinc commented Oct 13, 2017 •

edited

mattsinc commented Jan 30, 2018 •

edited

mattsinc commented Feb 6, 2018 •

edited

mattsinc commented Feb 7, 2018 •

edited