ROCm / HIPIFY

HIPIFY: Convert CUDA to Portable C++ Code
https://rocm.docs.amd.com/projects/HIPIFY/en/latest/
MIT License
521 stars 74 forks source link

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

Closed mattsinc closed 4 years ago

mattsinc commented 7 years ago

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:

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) )

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:

  1. Are these known issues with hipify-clang (kernel launch change, MACRO issues, parsing text for various CUDA calls)?
  2. 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 commented 7 years ago

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 commented 7 years ago

I've just fixed the issue with launching kernel in hipify-clang: https://github.com/ROCm-Developer-Tools/HIP/pull/195/commits/7ee8e2d51b2167de7d24f0e64f21a063d874570d 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 commented 7 years ago

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 commented 7 years ago

hipify-perl contains hipLaunchKernelGGLafter 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?) hipLaunchKernelGGLis 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 commented 7 years ago

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 commented 7 years ago

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 commented 7 years ago

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 commented 7 years ago

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 commented 7 years ago

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 commented 7 years ago

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 commented 7 years ago

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 commented 7 years ago

Sorry I attached the wrong zip file. Reattaching.

Matt Interac-v2.zip

emankov commented 6 years ago

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 commented 6 years ago

ping

mattsinc commented 6 years ago

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 commented 6 years ago

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

mattsinc commented 6 years ago

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 commented 6 years ago

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

emankov commented 6 years ago

Nonrepro on ToT.

mattsinc commented 6 years ago

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 commented 6 years ago

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 commented 6 years ago

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 commented 6 years ago

Fixed with ROCm-Developer-Tools/HIP#357

mattsinc commented 6 years ago

Will give it another try, thanks!

Matt

mattsinc commented 6 years ago

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 commented 6 years ago

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 commented 6 years ago

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

Matt