[HIPIFY] hipify-clang fails for application hipify-perl can convert about hipify HOT 27 CLOSED

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.

from hipify.

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?

from hipify.

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.

from hipify.

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 [email protected]
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

from hipify.

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?

from hipify.

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

from hipify.

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.

from hipify.

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

from hipify.

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

from hipify.

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.

from hipify.

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?

from hipify.

Sorry I attached the wrong zip file. Reattaching.

Matt
Interac-v2.zip

from hipify.

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.

from hipify.

ping

from hipify.

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

from hipify.

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

from hipify.

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

from hipify.

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

from hipify.

Nonrepro on ToT.

from hipify.

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

from hipify.

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.

from hipify.

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

from hipify.

Fixed with ROCm/HIP#357

from hipify.

Will give it another try, thanks!

Matt

from hipify.

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

from hipify.

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.

from hipify.

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

Matt

from hipify.