+ Add a corresponding matcher cudaDeviceFuncCall to match only (__device__ or __global__) and not __host__ functions.
+ Add a corresponding device functions mapping:
only unsupported are listed, cause supported are exactly the same as of CUDA and do not need transformation;
make FindAndReplace for device functions separated from host API calls.
+ Add a test to distinguish device functions and user-defined.
+ Start to translate preprocessor's false conditional blocks too:
based on clang's https://reviews.llvm.org/D66597;
available only starting from LLVM 10.0 or trunk.
+ Option -skip-excluded-preprocessor-conditional-blocks for skipping excluded conditional blocks:
the default behavior for hipify-clang built with LLVM < 10.0;
false by default for hipify-clang built with LLVM 10 or trunk.
+ Add 4 preprocessor unit tests, 2 of which are LLVM 10.0 only
+ Update couple of existing tests by setting -skip-excluded-preprocessor-conditional-blocks option:
update lit testing accordingly
...while including HIP main header file, which is inserted now after #indef controlling macro, or after #pragma once, if it's occurred earlier.
+ Add a couple of unit tests.
ToDo: Check backward compatibility on older clang versions.
+ Option '-roc' is added to translate to ROCm instead of HIP
+ Update hipification mechanism accordingly
+ Populate all translation maps with additional field for ROCm identifiers
+ There are 4 degrees of CUDA support now: full support, only by HIP, only by ROCm, unsupported
+ StringRef issue, which is finally moved to LLVM from Clang
+ Renamed getBeginLoc() and getEndLoc() resolution for Expr and TypeLoc classes
+ Support all the previous LLVM versions via LLVCompat
[Reasons]
+ Simplify the support
+ Optimize compilation and code generation time
[Additionally]
+ Update minimum required version of cmake
+ Fix MSVC compile options force update
Tested with CUDA 8.0, 9.0, 9.1 and 9.2.
Only 8.0 works with LLVM 7.0, due to the changes in LLVM trunc since released 6.0, which works fine with CUDA 8.0 and 9.0.
So, nothing to do in hipify-clang, hope that all the CUDA 9.x related issues will be fixed in 7.0 release.
Due to support of cuRAND headers.
+ compound test on all headers is added;
+ missing entities are added with updating the doc;
+ a couple cuRAND tests are added (https://github.com/ROCmSoftwarePlatform/rocRAND/tree/master/benchmark):
- the following CUDA entities are still unsupported by hipRAND:
curandMakeMTGP32Constants
curandMakeMTGP32KernelState
curandGetDirectionVectors32
curandDirectionVectorSet_t
CURAND_DIRECTION_VECTORS_32_JOEKUO6
curandStateSobol64_t
curandStateScrambledSobol64_t
curandGenerateLongLong
- and the following - by HIP:
cudaRuntimeGetVersion
- those entities are handled by CHECK-NOT directive for now.
File location have to be verified, otherwise location of the first found '#pragma once' in any included header even system will be erroneously handled, which might lead to attempt to including hip_runtime.h in it.
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.
// 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.
Much pretty. Very wow
This gives users all the usual power when it comes to manipulating
clang diagnostics. People can pass -Werror can have hipify fail if
it doesn't completely translate a file, for example. Much nicer
than reinventing the wheel.
Most of what hipify does is really just replacing CUDA idenitifers
with HIP ones. CUDA function calls, preprocessor macro calls,
enum references, types, etc.
This is problematic: calls/types/enum-refs require name resolution
for the AST matcher to work. This fails in the presence of code
deleted by the preprocessor, and in two-pass template compilation.
Instead, we can simply hook the lexer and have it rewrite the
identifiers for us.
This approach means identifier transformations will work correctly
regardless of where they appear (and we get to delete lots of code)
- Fixes#260
- Helps a bit with #207 - it will still fail to translate kernel
calls in preprocessor-ignored code, but everything except kerel
launches should translate correctly now, even in
preprocessor-deleted code.
Evgeny Mankov
Evgeny Mankov
Maneesh Gupta
Chris Kitching