All checks should not occur in input file for FileCheck. The issue found on CHECK-NOT.
Change removes all lit checks in the hipified file based on regexp, and the resulted stdout is fed as stdin for FileCheck.
If in source CUDA file main header (cuda_runtime.h or cuda.h) is not presented, corresponding HIP main header (hip_runtime.h) should be explicitly included in output hipified file.
[Algorithm]
1. If #pragma once is presented, HIP main header should be placed just after it;
2. Otherwise if any other (not CUDA main) header is presented, HIP main header should be placed just before it;
3. Otherwise HIP main header should be placed in the beginning of output file.
P.S.
There might be one more situation when #ifndef #define ... #endif guard for the entire file is presented (make sense for *.h, *.hpp, *.cuh files). In this case HIP main include should be placed just after such #ifdef, or after #pragma once, if it is also presented. This situation will be handled in a separate change.
Env. var NVCUDASAMPLES_ROOT is changed to NVCUDASAMPLESX_Y_ROOT where X - major ver, Y - minor ver.
Reason: NVCUDASAMPLES_ROOT contains path to CUDA SDK installed last, while NVCUDASAMPLESX_Y_ROOT contains samples of the same version as of CUDA_TOOLKIT_ROOT_DIR.
// hipified to #include<hip/hip_runtime.h>
#include<cuda.h> // 1st cuda main include (Driver API)
// to eliminate
#include<cuda_runtime.h> // 2nd cuda main include (Runtime API)
HIP has one header hip_runtime.h for both CUDA APIs, thus second cuda main include directive is eliminated entirely.
All but the most contrived use of macros is now properly handled -
have a look at the new testcases this commit adds. You can have
macros in kernel calls, macros spanning chunks of your arguments,
the call, call parameters, or callee can all be macros or
partially macros.
It seems the test is already broken, but look how awesome the
error message is now:
/home/chris/HIP/tests/hipify-clang/axpy.cu:31:12: error: expected string not found in input
// CHECK: hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
^
<stdin>:31:2: note: scanning from here
//
^
<stdin>:33:2: note: possible intended match here
hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
^
Evgeny Mankov
Evgeny Mankov
Chris Kitching
Maneesh Gupta