The error occurred when the source file is specified by a relative path.

[How to reproduce]
hipify-clang ../1/benchmark_curand_generate.cpp -- -v --cuda-path="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"
[Result]
[HIPIFY] error: no such file or directory: while copying ../1/benchmark_curand_generate.cpp to C:\Users\TT\AppData\Local\Temp../1/benchmark_curand_generate.cpp-3cbe3d.hip

Example from CUDA 8.0.44 sample (CUDASamples\0_Simple\matrixMulDrv\matrixMulDrv.cpp):

    CUjit_option *jitOptions = new CUjit_option[jitNumOptions];

where CUjit_option is enum, should be:

    hipJitOption *jitOptions = new hipJitOption[jitNumOptions];

Example from CUDA 8.0.44 sample (CUDASamples\0_Simple\matrixMulDrv\matrixMulDrv.cpp):

CUjit_option *jitOptions = new CUjit_option[jitNumOptions];

where CUjit_option is enum, should be:

hipJitOption *jitOptions = new hipJitOption[jitNumOptions];

A common situation is for people to have header files and .cu files with the same name: such as axpy.cu and axpy.h/axpy.cuh. There might even be axpy.cpp, too.

When you feed all of those files to hipify-clang at the same time, it will decide it wants to output all of them to axpy.hip, and output an arbitrary one of them to that file, losing the other.

The obvious solution might be to decide on a file extension for "hip header". We've got .cuh for CUDA, so let's have something like .hiph or similar for hip?

The practical effect of this bug is that the only way to convert any meaningfully-sized project is to call hipify-clang -o <something> <somefile> -- ... in a loop for every file. Makes multiple-input mode a bit unuseful, really.

Unrelatedly, it would also be nice to be able to specify an output directory when using multiple-file input mode...

One other solution might be to simply drop multiple-input-mode entirely, so it's clear to users that they have to write a loop script or suchlike.

• Passing -o-stats without also passing -print-stats silently does nothing. No stats anywhere, no warning, and no mention of this behaviour in the help text.
• Passing both options together does produce statistics output. And a segfault.
• Using -o-stats=wat.csv or -o-stats=wat.csv -print-stats with a nonexistent input source file causes hipify to silently do nothing and yield a 0 exit code.

What's the intended use-case of the stats output? I'm struggling to see why someone would want to know how many bytes of source hipify changed, or the total number of calls replaced. If they want to know what hipify did to their code, diff -u is extremely good at that.

Warnings about things it failed to convert make sense - but those should surely go through the clang diagnostics API, not this nonstandard output channel? Indeed, these are printed regardless of the stats option.

The stats system accounts for about a third of the code (excluding the lookup tables), and it's tightly coupled into all sorts of random places to keep the counters in sync. Maybe just get rid of it?

After translating any symbol in the source file forcibly include the hip header file in case it wasn’t.

Also add cuda.h translation.

[Example 1]
cudaMemcpyKind cuda_memcpy_kind(...);
[Example 2]
template
cudaMemcpyKind cuda_memcpy_kind(...);

I'm not 100% sure what's causing this, but the cudaMalloc in this short file doesn't get hipified:

#pragma once

#include <cuda_runtime.h>

namespace {

/**
 * Allocate GPU memory for `count` elements of type `T`.
 */
template<typename T>
static T* gpuMalloc(size_t count) {
    T* ret = nullptr;
    cudaMalloc(&ret, count * sizeof(T));
    return ret;
}

}

Assert occurs on ~llvm::Error() for llvm::Error e defined in LLVMCompat.cpp:28:

    llvm::Error e = replacements.add(rep);

Found on Windows 10 | clang 5.0.x (debug) | VS 2017 (15.5).
Hasn't yet verified on Linux.

Didn't work for me. I called it as follows:
First, cd to directory that contains cuda files. Then
../../../GPUOpen-ProfessionalCompute-Tools/HIP/bin/hipify "inplace" ./*.cu

OR

../../../GPUOpen-ProfessionalCompute-Tools/HIP/bin/hipify inplace ./*.cu

produce error:

error: could not open inplace at ../../../GPUOpen-ProfessionalCompute-Tools/HIP/bin/hipify line 189.

I am porting an application that exclusively uses hipMemcpyToSymbol for its kernels. When I run hipify-perl on this, it will happily convert it (except for the threadfence issue I noted in #32). However, when I try to compile the kernel I get the following for each kernel:

hipLaunchKernelGGL((Kernel1), dim3(numTBs), dim3(numThreadsPerBlock), 0, 0, );

This is related to the issue I raised in ROCm/HIP#257 -- currently hipcc requires at least one input arg to correctly compile a kernel, when this is not necessary if MemcpyToSymbol is used. But regardless, it doesn't seem like hipify should add a comma with nothing after it.

Matt

Duplicate of reopened #7

Examples:
// 1.
cudaEvent_t *event = NULL;
// typedef device_builtin struct CUevent_st *cudaEvent_t;
// 2.
CUevent *event = NULL;
// typedef struct CUevent_st *CUevent;

I have tried to compile hipify-compiler , but compilation has failed .
Compilation error message was as follows:

‘‘‘
main.cpp:134:19: error: no member
named 'dbgs' in namespace 'llvm'
DEBUG(llvm::dbgs() << "Skipped some replacements.\n");

‘‘‘

and here is the hipconfig message:

HIP version  : 1.3.17385

== hipconfig
HIP_PATH     : /opt/rocm
HIP_PLATFORM : hcc
CPP_CONFIG   :  -D__HIP_PLATFORM_HCC__=   -I/opt/rocm/include -I/opt/rocm/hcc/include

== hcc
HSA_PATH     : /opt/rocm/hsa
HCC_HOME     : /opt/rocm/hcc
HCC clang version 6.0.0  (based on HCC 1.0.17412-f590a25-821e6d8-64e7fc7 )
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm/hcc/bin
Can't exec "/opt/rocm/hcc/compiler/bin/llc": そのようなファイルやディレクトリはありません at /opt/rocm/hip/bin/hipconfig line 132.
HCC-cxxflags :  -hc -std=c++amp -I/opt/rocm/hcc/includeHCC-ldflags  :  -hc -std=c++amp -L/opt/rocm/hcc/lib -Wl,--rpath=/opt/rocm/hcc/lib -ldl -lm -lpthread -lunwind -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive

=== Environment Variables
PATH=/home/n-kazu/bin:/home/n-kazu/.local/bin:/opt/rocm/hcc/bin:/opt/rocm/hip/bin:/usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin:/usr/games:/usr/local/games:/opt/rocm/bin
HIP_PATH=/opt/rocm
HCC_HOME=/opt/rocm/hcc

== Linux Kernel
Hostname     : nkazu-kourindou
Linux nkazu-kourindou 4.11.0-kfd-compute-rocm-rel-1.6-180 ROCm-Developer-Tools/HIP#1 SMP Tue Oct 10 08:15:38 CDT 2017 x86_64 x86_64 x86_64 GNU/Linux
No LSB modules are available.
Distributor ID:	Ubuntu
Description:	Ubuntu 16.04.3 LTS
Release:	16.04
Codename:	xenial

Would you like to give some advice?

Thanks in advance.

Similar to #7 with the only difference that nested MACRO is an identifier, not a defined numeric literal.

Example:

checkCudaErrors(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, matrix_size.uiWB, matrix_size.uiHA, matrix_size.uiWA, &alpha, d_B, matrix_size.uiWB, d_A, matrix_size.uiWA, &beta, d_C, matrix_size.uiWB));

where both marked identifiers are MACRO:
#define checkCudaErrors(val) check ( (val), #val, _FILE_, _LINE_ )
#define cublasSgemm cublasSgemm_v2

After hipifying now:
checkCudaErrors(cublasSgemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, matrix_size.uiWB, matrix_size.uiHA, matrix_size.uiWA, &alpha, d_B, matrix_size.uiWB, d_A, matrix_size.uiWA, &beta, d_C, matrix_size.uiWB));

Should be:
checkCudaErrors(hipblasSgemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, matrix_size.uiWB, matrix_size.uiHA, matrix_size.uiWA, &alpha, d_B, matrix_size.uiWB, d_A, matrix_size.uiWA, &beta, d_C, matrix_size.uiWB));

Recently I try to use the Rocm and HIP on Vega20. I want to port my old cuda project to hip, but when I use $HIP_DIR/bin/hipexamine.sh or hipify --inplace , there is the same error :
/opt/rocm/hip/bin/hipify-clang: No such file or directory

I check the content in the /opt/rocm/hip/bin/:
ca finduncodep.sh hipcc_cmake_linker_helper hipconvertinplace-perl.sh hipdemangleatp hipexamine.sh hipify-perl
findcode.sh hipcc hipconfig hipconvertinplace.sh hipexamine-perl.sh hipify-cmakefile lpl

I can not find the hipify-clang and in the github project ,I also can not find the hipify-clang.
How can I get it ? Thanks !

1. Do not use JIT options on HCC path, call hipModuleLoadData instead.
2. NVCC path is unchanged, to call cuModuleLoadDataEx with all options.
3. Get rid of manual hipification, based on #ifdef #else for NVCC/HIP.
4. Update documentation accordingly.

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:

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

Example (Shoc's NeuralNet.cu):

CUDA_SAFE_CALL(cudaMemcpy(&D__TRAINING_DATA_X_2D[i_IMAGE_SIZE], TRAINING_DATA_X[i], sizeof(float)_IMAGE_SIZE, cudaMemcpyHostToDevice));

where:

define CUDA_SAFE_CALL(call) /.../

After hipifying now:
CUDA_SAFE_CALL(hipMemcpy(&D__TEST_DATA_X_2D[i_IMAGE_SIZE], TEST_DATA_X[i], sizeof(float)_IMAGE_SIZE, cudaMemcpyHostToDevice));

Should be:
CUDA_SAFE_CALL(hipMemcpy(&D__TEST_DATA_X_2D[i_IMAGE_SIZE], TEST_DATA_X[i], sizeof(float)_IMAGE_SIZE, hipMemcpyHostToDevice));

Currently hipify tries to create temporary files in the source folder. It goes ugly if someone tries to hipify an example from read-only folder. It is more reasonable to create temporary files in whether system temporary folder or have a way to set the folder for them.

error: error reading '/usr/local/cuda/NVIDIA_CUDA-9.2_Samples/0_Simple/matrixMul/matrixMul.cu.hipify-tmp'
1 error generated when compiling for host.

Looks like the mapping table converts to hipBLAS. Shouldn't this be rocBLAS instead, or did I miss the point?

[Synopsis]
An error occurred with hipify-clang built with clang version < 5.

[How to repro]

1. Copy the following files from: \tests\hipify-clang\unit_tests\libraries\cuRAND:
   benchmark_curand_generate.cpp
   cmdparser.hpp
   to the folder with hipify-clang bibary.

2. Run hipify-clang benchmark_curand_generate.cpp -- --cuda-path="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0"

3. Result:
   benchmark_curand_generate.cpp.hip:32:10: fatal error: 'cmdparser.hpp' file not found
   #include "cmdparser.hpp"
   ^~~~~~~~~~~~~~~

hipify-clang processes each of its input files as a "main" file using the clang API. This changes the behaviour of clang's warning system, yielding output like this when processing header files:

example.hip.h:1:9: warning: #pragma once in main file [-Wpragma-once-outside-header]
#pragma once
        ^
1 warning generated.

To reproduce, create any random header file with #pragma once and feed it to hipify-clang.

I think the solution here is to tell the API that no files are the main file - since the notion of a main file makes no sense for this application anyway. I haven't tested if this failure exists when using compilation-database mode.

I got the following error when I translated Deepbench cuda codes that contain thurst APIs.
/usr/local/cuda-8.0/include/thrust/detail/swap.h:25:1: error: expected unqualified-id
thrust_exec_check_disable

When I convert below code from CUDA to HIP by hipify (hipconvertinplace.sh)

# kernel.cu
__global__
void kernel(double z[], const double x[], const double y[])
{
    const auto i = (blockIdx.x * blockDim.x + threadIdx.x);

    z[i] = x[i] + y[i];
}

and compile them, I got an error.

kernel.cu:8:27: error: hc_grid_launch function must have grid_launch_parm type as first parameter, 'double *' type is not compatible
void kernel(double z[], const double x[], const double y[])
                          ^
kernel.cu:19:34: error: no matching function for call to 'kernel'
        hipLaunchKernel(HIP_KERNEL_NAME(kernel), dim3(block), dim3(thread), 0, 0, z, x, y);
                                        ^~~~~~
/opt/rocm/include/hip/hcc_detail/hip_runtime.h:529:30: note: expanded from macro 'HIP_KERNEL_NAME'
#define HIP_KERNEL_NAME(...) __VA_ARGS__
                             ^
/opt/rocm/include/hip/hcc_detail/hip_runtime.h:552:3: note: expanded from macro 'hipLaunchKernel'
  _kernelName (lp, ##__VA_ARGS__);\
  ^
kernel.cu:8:13: note: candidate function not viable: requires 3 arguments, but 4 were provided
void kernel(double z[], const double x[], const double y[])

"hipLaunchParm" seems missing so I can solve this error by manually adding "hipLaunchParm lp" at kernel's first argument.

Is it hard to automatically add "hipLaunchParm lp" by hipify converter?

(Part 2) Transformations are needed in cases of further possible assignments of already transfomed "extern" variable. Example:

extern __shared__ uint sRadix1[]; =>  HIP_DYNAMIC_SHARED(unsigned int, sRadix1);  // Part 1, done
uint * sRadix2 = sRadix1; => __attribute__((address_space(3))) uint* sRadix2 = sRadix1; // Part 2

First part of transformation (transformation of declarations with external linkage and shared attribute for IncompleteArrayType) has already done. Second part - all assignments of such "extern" var should be transformed as well.

sRadix2 can't be of type uint* after hipifying. It should be of type __attribute__((address_space(3))) uint* due to the changed type of sRadix1:

#define HIP_DYNAMIC_SHARED(type, var) \
__attribute__((address_space(3))) type* var = \
(__attribute__((address_space(3))) type*)__get_dynamicgroupbaseptr(); \

There might be more complicated situations when this "extern" var is used as a function argument (It is a 3rd part - to be investigated, example is needed).

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.y*gridDim.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

For example, sometimes people have functions that use the word cuda, such as initCuda(). Or even in comments etc...

Seems I can't get hipify-clang work by referencing the doc here:
https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md

opt/rocm/hip/bin/hipify-clang memory.cu -x cuda --cuda-path=/usr/local/cuda
LLVM ERROR: CommonOptionsParser: failed to parse command-line arguments. [CommonOptionsParser]: hipify-clang: Unknown command line argument '-x'. Try: '/opt/rocm/hip/bin/hipify-clang -help'
hipify-clang: Did you mean '-o'?
hipify-clang: Unknown command line argument '--cuda-path=/usr/local/cuda'. Try: '/opt/rocm/hip/bin/hipify-clang -help'
hipify-clang: Did you mean '-extra-arg=/usr/local/cuda'?

/opt/rocm/hip/bin/hipify-clang memory.cu
Error while trying to load a compilation database:
Could not auto-detect compilation database for file "memory.cu"
No compilation database found in /rocm/nwchem/src/tce/ccsd_t or any parent directory
fixed-compilation-database: Error while opening fixed database: No such file or directory
json-compilation-database: Error while opening JSON database: No such file or directory
Running without flags.
warning: /rocm/nwchem/src/tce/ccsd_t/memory.cu.hipify-tmp: 'linker' input unused [-Wunused-command-line-argument]
warning: argument unused during compilation: '-std=c++11' [-Wunused-command-line-argument]
error: unable to handle compilation, expected exactly one compiler job in ''
Error while processing /rocm/nwchem/src/tce/ccsd_t/memory.cu.hipify-tmp.

Do we use hipify-perl or hipify-clang nowadays?

Implement map of unsupported CUDA API refs analogously to supported ones, count them while hipifying and add the numbers to statistics.

STR:

1. Create an empty file called test.h
2. Create a file called test.cpp containing only #include "test.h"
3. Run: hipify-clang test.cpp -- -x cuda --cuda-path=<whatever> - works correctly
4. Run: hipify-clang -o test.hip test.cpp -- -x cuda --cuda-path=<whatever> - works correctly
5. Run: hipify-clang -o /any/other/directory/test.hip test.cpp -- -x cuda --cuda-path=<whatever> - Fails:

/home/chris/spectral/xcmake/src/../wat.hip.hip:1:10: fatal error: 'test.h' file not found
#include "test.h"
         ^~~~~~~~
1 error generated.
Error while processing /home/chris/spectral/xcmake/src/../wat.hip.hip.

It seems that the output path is being used to resolve includes, and hence relative includes fail if you set an output path in a different directory.

1. CONVERSION %: -9223372036854775708 in case of source file absence or zero length file.
2. TIME ELAPSED s: - XXXXX.XX in overall statistics in case of multiple input source files.
3. Other errors in case of previous compilation errors - to check.

Suggestions:

1. Print a message in case of hipification errors, that statistics is invalid.
2. Print the same message in total/globals section of statistics in case of multiple input files and at least one erroneous file.

Error occurs when '-o-dir' or '-temp-dir' or both are specified with a relative path to nonexisting folders to create.
Error occurs when hipify-clang is built against LLVM < 5.0. If LLVM >= 5.0 error doesn't occur.

[How to Repro]
hipify-clang test1.cu test2.cu -temp-dir=../1 -o-dir=../2 -save-temps

Some people presumably don't care about hipify, and it has several dependenceis (llvm, FileCheck, etc.). There's probably value in having a cmake option to turn the whole thing off.

With introducing temporary files in system temps by createTemporaryFile() option -inplace stopped working due to system error: "system cannot move the file to another disk".

• OS Ubuntu 18.04
• clang version 6.0.1-svn334776-1~exp1~20181018152737.116 (branches/release_60) from official llvm repos
• CUDA 9.0 Patch 4
• HIP: hip_nvcc from repos
• since HIP does not come with hipify clang, grabed and compiled
  hipify-clang built from commit 25aae88fd5029a5952524d617934e71194702b9c

reproduce by compiling square sample:
hipify-clang square.cu -- -x cuda --cuda-gpu-arch=sm_50 -L/usr/local/cuda/lib64 -lcudart_static -ldl -lrt -pthread -I/opt/rocm/include
will terminate without problems. Looking inside, I get
#include <hip/hipn^D^@^@^P^@^@^@¸p¥
square.cu.hip.txt

Using hipconvertinplace.sh for a multifile project a couple of files were transformed OK while some
resulted in

/home/milettri/Projects/gemmCUDA/gemmHIP/matmul/../helpers/cudaHelpers.hpp:5:21: warning: null character ignored [-Wnull-character]
#include <hip/hip+<U+0000><U+0000><U+0000><U+0008><U+0000><U+0000><U+0000><A8>U<U+000B>
                                  ^

reproduced on two different ubuntu 18.04 machines, one with ubuntu packaged clang, the other with official llvm packages. Building hipify-clang from release tag 1.9.1 resulted in the same issues.

Using hipify to convert a ".cu" file to a HIP ".cpp" file I noticed that the special variable warpSize within a kernel's code is transformed to hipWarpSize. I think it should stay as is (warpSize).

[How To Repro]
./hipify-clang -print-stats /srv/git/thrust/examples/cuda/fallback_allocator.cu -- -x cuda -I/srv/git/thrust -I/srv/git/HIP-TOT/include -D__CUDA_ARCH__=750 -v

Similar crash is on custom_temporary_allocation.cu.

Thrust sources:
https://github.com/thrust/thrust

[Sinopsys]
Looks like the crash in one of the hipify's matchers, as the crash occurs in clang::TagType::getDecl().

[Stack trace]

#0 0x0000000001bbf09a llvm::sys::PrintStackTrace(llvm::raw_ostream&) /srv/emankov/LLVM/3.8/lib/Support/Unix/Signals.inc:322:0
ROCm-Developer-Tools/HIP#1 0x0000000001bbf3af PrintStackTraceSignalHandler(void*) /srv/emankov/LLVM/3.8/lib/Support/Unix/Signals.inc:380:0
ROCm-Developer-Tools/HIP#2 0x0000000001bbdac7 llvm::sys::RunSignalHandlers() /srv/emankov/LLVM/3.8/lib/Support/Signals.cpp:44:0
ROCm-Developer-Tools/HIP#3 0x0000000001bbebc1 SignalHandler(int) /srv/emankov/LLVM/3.8/lib/Support/Unix/Signals.inc:210:0
ROCm-Developer-Tools/HIP#4 0x00007f59500a2330 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x10330)
ROCm-Developer-Tools/HIP#5 0x00000000019f91ca clang::TagType::getDecl() const /srv/emankov/LLVM/3.8/tools/clang/lib/AST/Type.cpp:2962:0
ROCm-Developer-Tools/HIP#6 0x000000000096d9f0 (anonymous namespace)::Cuda2HipCallback::run(clang::ast_matchers::MatchFinder::MatchResult const&) (/srv/emankov/git/HIP-TOT/dist/bin/hipify-clang+0x96d9f0)
ROCm-Developer-Tools/HIPIFY#3 0x000000000099354c clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::MatchVisitor::visitMatch(clang::ast_matchers::BoundNodes const&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:694:0
ROCm-Developer-Tools/HIPIFY#4 0x0000000000a484e9 clang::ast_matchers::internal::BoundNodesTreeBuilder::visitMatches(clang::ast_matchers::internal::BoundNodesTreeBuilder::Visitor*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchersInternal.cpp:47:0
ROCm-Developer-Tools/HIP#9 0x0000000000992b8d clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::matchWithFilter(clang::ast_type_traits::DynTypedNode const&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:554:0
ROCm-Developer-Tools/HIP#10 0x0000000000992d1d clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::matchDispatch(clang::Decl const*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:576:0
ROCm-Developer-Tools/HIP#11 0x00000000009973bd void clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::match<clang::Decl>(clang::Decl const&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:480:0
ROCm-Developer-Tools/HIP#12 0x0000000000993b7f clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:827:0
ROCm-Developer-Tools/HIP#13 0x00000000009d50fc clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseDeclStmt(clang::DeclStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1914:0
ROCm-Developer-Tools/HIP#14 0x00000000009ab89d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:103:0
ROCm-Developer-Tools/HIP#15 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#16 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#17 0x00000000009d4c21 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseCompoundStmt(clang::CompoundStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1931:0
ROCm-Developer-Tools/HIP#18 0x00000000009ab82d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:79:0
ROCm-Developer-Tools/HIP#19 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#20 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#21 0x00000000009dd469 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseForStmt(clang::ForStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1935:0
ROCm-Developer-Tools/HIP#22 0x00000000009ac463 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:807:0
ROCm-Developer-Tools/HIP#23 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#24 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#25 0x00000000009d4c21 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseCompoundStmt(clang::CompoundStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1931:0
ROCm-Developer-Tools/HIP#26 0x00000000009ab82d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:79:0
ROCm-Developer-Tools/HIPIFY#6 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#28 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#29 0x00000000009d49c9 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseCXXTryStmt(clang::CXXTryStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1929:0
ROCm-Developer-Tools/HIP#30 0x00000000009ab7f5 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:67:0
ROCm-Developer-Tools/HIP#31 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#32 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIPIFY#7 0x00000000009d4c21 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseCompoundStmt(clang::CompoundStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1931:0
ROCm-Developer-Tools/HIPIFY#8 0x00000000009ab82d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:79:0
ROCm-Developer-Tools/HIPIFY#9 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIPIFY#10 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#37 0x00000000009cc8c2 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseFunctionHelper(clang::FunctionDecl*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1795:0
ROCm-Developer-Tools/HIP#38 0x00000000009a8542 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseFunctionDecl(clang::FunctionDecl*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1800:0
ROCm-Developer-Tools/HIP#39 0x00000000009985d1 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/DeclNodes.inc:363:0
ROCm-Developer-Tools/HIP#40 0x0000000000993b92 clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:828:0
ROCm-Developer-Tools/HIP#41 0x00000000009ca6ca clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseDeclContextHelper(clang::DeclContext*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1231:0
ROCm-Developer-Tools/HIP#42 0x00000000009a92b4 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseTranslationUnitDecl(clang::TranslationUnitDecl*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1321:0
ROCm-Developer-Tools/HIP#43 0x000000000099893d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/DeclNodes.inc:491:0
ROCm-Developer-Tools/HIP#44 0x0000000000993b92 clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:828:0
ROCm-Developer-Tools/HIP#45 0x0000000000994492 clang::ast_matchers::MatchFinder::matchAST(clang::ASTContext&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:984:0
ROCm-Developer-Tools/HIP#46 0x0000000000993de9 clang::ast_matchers::internal::(anonymous namespace)::MatchASTConsumer::HandleTranslationUnit(clang::ASTContext&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:886:0
ROCm-Developer-Tools/HIP#47 0x0000000000bcfedd clang::ParseAST(clang::Sema&, bool, bool) /srv/emankov/LLVM/3.8/tools/clang/lib/Parse/ParseAST.cpp:170:0
ROCm-Developer-Tools/HIP#48 0x0000000000a4fed0 clang::ASTFrontendAction::ExecuteAction() /srv/emankov/LLVM/3.8/tools/clang/lib/Frontend/FrontendAction.cpp:540:0
ROCm-Developer-Tools/HIP#49 0x0000000000a4f995 clang::FrontendAction::Execute() /srv/emankov/LLVM/3.8/tools/clang/lib/Frontend/FrontendAction.cpp:443:0
ROCm-Developer-Tools/HIP#50 0x0000000000aa49c8 clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) /srv/emankov/LLVM/3.8/tools/clang/lib/Frontend/CompilerInstance.cpp:841:0
ROCm-Developer-Tools/HIP#51 0x0000000000bc8f07 clang::tooling::FrontendActionFactory::runInvocation(clang::CompilerInvocation*, clang::FileManager*, std::shared_ptr<clang::PCHContainerOperations>, clang::DiagnosticConsumer*) /srv/emankov/LLVM/3.8/tools/clang/lib/Tooling/Tooling.cpp:302:0
ROCm-Developer-Tools/HIP#52 0x0000000000bc8d9c clang::tooling::ToolInvocation::runInvocation(char const*, clang::driver::Compilation*, clang::CompilerInvocation*, std::shared_ptr<clang::PCHContainerOperations>) /srv/emankov/LLVM/3.8/tools/clang/lib/Tooling/Tooling.cpp:278:0
ROCm-Developer-Tools/HIPIFY#11 0x0000000000bc8c17 clang::tooling::ToolInvocation::run() /srv/emankov/LLVM/3.8/tools/clang/lib/Tooling/Tooling.cpp:263:0
ROCm-Developer-Tools/HIP#54 0x0000000000bc9d48 clang::tooling::ClangTool::run(clang::tooling::ToolAction*) /srv/emankov/LLVM/3.8/tools/clang/lib/Tooling/Tooling.cpp:422:0
ROCm-Developer-Tools/HIPIFY#12 0x0000000000965a2b main (/srv/emankov/git/HIP-TOT/dist/bin/hipify-clang+0x965a2b)
ROCm-Developer-Tools/HIP#56 0x00007f594f2a6f45 __libc_start_main /build/eglibc-oGUzwX/eglibc-2.19/csu/libc-start.c:321:0
ROCm-Developer-Tools/HIP#57 0x0000000000968258 _start (/srv/emankov/git/HIP-TOT/dist/bin/hipify-clang+0x968258)

Example:

include "cuda_runtime.h"

define MY_MACRO(func, flags) (func, flags)

...
cudaEvent_t *event = NULL;
MY_MACRO(cudaEventCreateWithFlags(event, cudaEventDisableTiming), NULL);

where cudaEventDisableTiming is a defined numeric literal and thus a nested MACRO:

define cudaEventDisableTiming 0x02 /**< Event will not record timing data */

After hipifying now:
MY_MACRO(hipEventCreateWithFlags(event, cudaEventDisableTiming), NULL);

Should be:
MY_MACRO(hipEventCreateWithFlags(event, hipEventDisableTiming), NULL);

Consider the following CUDA program:

__global__ void axpy_kernel(float a, float* x, float* y) {
    y[threadIdx.x] = a * x[threadIdx.x];
}

void axpy(float a, float* x, float* y) {
    axpy_kernel<<<1, 4>>> (a, x, y);
#ifdef SOME_MACRO
    axpy_kernel<<<1, 4>>> (a, x, y);
#endif
}

The result of hipifying this, if you don't pass -DSOME_MACRO is:

#include <hip/hip_runtime.h>
__global__ void axpy_kernel(float a, float* x, float* y) {
    y[hipThreadIdx_x] = a * x[hipThreadIdx_x];
}

void axpy(float a, float* x, float* y) {
    hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
#ifdef SOME_MACRO
    axpy_kernel<<<1, 4>>> (a, x, y);
#endif

}

Respecting conditional macros isn't the right thing to do with this sort of mechanised refactoring - what you really want to do is walk the entire tree applying your refactor, regardless of preprocessor conditionals.

This is going to present a relatively nasty obstacle to people with complicated CUDA programs they want to translate...

Consider the following highly uninteresting CUDA file:

__global__ void axpy_kernel(float a, float* x, float* y) {
    y[threadIdx.x] = a * x[threadIdx.x];
}

void axpy(float a, float* x, float* y) {
    axpy_kernel<<<1, 4>>> (a, x, y);
}

When processed with hipify-clang, the following warnings are produced:

[HIPIFY] warning: axpy.cu:2:7: the following reference is not handled: '__fetch_builtin_x' [function call].
[HIPIFY] warning: axpy.cu:2:28: the following reference is not handled: '__fetch_builtin_x' [function call].

The correct hip output is nevertheless produced.

In case unfamiliar, __fetch_builtin_x is the name clang uses in its representation of CUDA builtins like threadIdx or blockIdx. Clang defines a struct for each of those, with functions __fetch_builtin_x/y/z on them which delegate to the appropriate intrinsics. The definition of those can be found here.

So yes. Presumably Clang is desugaring the kernel launch into some more vanilla function calls that refer to those structs, and it's failing to convert properly. Shouldn't be too difficult, but the fact it's missing these builtin conversions might lead to more worrying problems in more complex cases...

missing 't' in line 328

- $ft{'stream'} += s/\bcudaStreamWaitEvent\b/hipStreamWaitEven/g;
+ $ft{'stream'} += s/\bcudaStreamWaitEvent\b/hipStreamWaitEvent/g;

https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/blob/master/bin/hipify#L328

This snippet of axpy.cu code is from Compiling CUDA with clang This code is verified that it can be successfully compiled by clang using:

clang++ axpy.cu -o axpy --cuda-gpu-arch=<GPU arch> -L<CUDA install path>/<lib64 or lib> -lcudart_static -ldl -lrt -pthread

I have tried to do:

hipify-clang square.cu -- -x cuda --cuda-path=/opt/cuda --cuda-gpu-arch=sm_30 -isystem /opt/cuda/samples/common/inc -I /opt/cuda/cuDNN and get a segfault from there

Then I tried to do the conversion by using the clang command I used in clang standalone compilation but stil got an error of:

Error while processing /home/kneternal/Work/test/clang_test/axpy.cu.hipify-tmp.

My version of LLVM+Clang is 7.0.0 and they are built and installed from source. Could I get some help about this?

When hipify inserts the #include <hip/hip_runtime.h> line into a header, it would be best to do it after any existing #pragma once or traditional include guard. Saves the user a bit of tedious cleanup.

One elegant way to implement this might be to replace the current insertion behaviour with logic to replace appropriate CUDA #include lines with hip ones. Unilaterally inserting this include into every file which you've made changes to is probably overkill - isn't it enough to replace #include <cuda> with the hip include?

I am currently porting a CUDA application that uses threadfence and threadfence_block to enforce ordering. According to the HIP documentation (http://rocm-documentation.readthedocs.io/en/latest/Programming_Guides/Kernel_language.html#memory-fence-instructions), these instructions are supported in HIP. However, while I run hipify-perl on this application, I get warnings like the following:

warning: unsupported device function : _threadfence();
…
warning: unsupported device function : _threadfence_block();

If these instructions are supported, then shouldn't these warnings not be printed? On the flip side, in hip/include/hip/hcc_detail/hip_runtime.h, there are __hip_hc_threadfence and __hip_hc_threadfence_block(). The comments suggest that these are currently empty stubs, but then there are implementations (in a single asm instructions) nevertheless. So I'm not sure if applications with threadfences are actually supported or not? If not, then perhaps that's why these warnings are appearing?

Matt

Hipify ought to be able to cope with whatever version of C++ is supported by the clang it was compiled against. Unfortunately, if you do something like:

hipify-clang concurentKernels.cu -- -x cuda --cuda-path=/opt/cuda --cuda-gpu-arch=sm_30 -isystem /opt/cuda/samples/common/inc -std=c++17

... you get compile errors if you use any C++17 features, since hipify is explicitly adding -std=c++11 internally. We probably want to make this conditional (only add -std=c++11 if the user hasn't provided their own -std option)

Hopefully a fairly easy fix once we find the right API for querying existing args...

socat doesn't propagate error code of TMP_SCRIPT, thus all lit tests pass, despite hipifying (1st) or FileCheck's (2nd) errors. Without TMP_SCRIPT and socat in run_test.sh lit testing works correctly with reporting errors if occured during 1st or 2nd.

Found on Ubuntu 16.04 | llvm 4.0 rel.

For instance:

// hipified to #include<hip/hip_runtime.h>
#include<cuda.h> // 1st cuda main include (Driver API)
// leaved unchanged
#include<cuda_runtime.h> // 2nd main include (Runtime API)

HIP has one header for both APIs hip_runtime.h, thus second cuda main include have to be eliminated.

Hello!
I can't build hipify-clang. I have this error "The dependency target "lit" of target "test-hipify" does not exist." Instruction: https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md#build-and-install

Can you help me? Can i install hipify-clang from deb-package?

hipconfig:

HIP version : 1.2.17263

== hipconfig
HIP_PATH : /opt/rocm
HIP_PLATFORM : hcc
CPP_CONFIG : -D__HIP_PLATFORM_HCC__= -I/opt/rocm/include -I/opt/rocm/hcc/include

== hcc
HSA_PATH : /opt/rocm/hsa
HCC_HOME : /opt/rocm/hcc
HCC clang version 5.0.0 (based on HCC 1.0.17262-7871d33-3953233-56b5abe )
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm/hcc/bin
LLVM (http://llvm.org/):
LLVM version 5.0.0svn
Optimized build.
Default target: x86_64-unknown-linux-gnu
Host CPU: haswell

Registered Targets:
amdgcn - AMD GCN GPUs
r600 - AMD GPUs HD2XXX-HD6XXX
x86 - 32-bit X86: Pentium-Pro and above
x86-64 - 64-bit X86: EM64T and AMD64
HCC-cxxflags : -hc -std=c++amp -I/opt/rocm/hcc-1.0/include -I/opt/rocm/includeHCC-ldflags : -hc -std=c++amp -L/opt/rocm/hcc-1.0/lib -Wl,--rpath=/opt/rocm/hcc-1.0/lib -ldl -lm -lpthread -lunwind -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive

=== Environment Variables
...

== Linux Kernel
Hostname : beamoflight-All-Series
Linux beamoflight-All-Series 4.9.0-kfd-compute-rocm-rel-1.6-77 ROCm/HIP#1 SMP Wed Jun 28 07:30:27 CDT 2017 x86_64 x86_64 x86_64 GNU/Linux
No LSB modules are available.
Distributor ID: Ubuntu
Description: Ubuntu 16.04.2 LTS
Release: 16.04
Codename: xenial