+ Add a corresponding test kernel_launch_01.cu
+ Add isBefore() check to avoid crash on Replacement with negative length
TODO:
+ Compatibility with former LLVM versions
+ More complicated kernel launch tests
+ Add one matcher (will be more)
+ Update Maps and Statistics
+ Add cub_01.cu unit test
+ Update lit harness to support standalone CUB
+ Update README.md
+ Update hipify-perl (only CUB header is supported for now)
[IMPORTANT]
clang (and hipify-clang) works correctly only with official NVLabs version on GitHub.
Compilation of CUB from official CUDA release has conflicts with THRUST.
Thus, to compile CUB sources, option "-I" should be specified to the cloned CUB from NVLAB on GitHub.
[REASON]
1. hip-clang is fine with the templated kernel launch, brackets are unneeded: HIP_KERNEL_NAME(...) __VA_ARGS__
2. HCC is not, thus: HIP_KERNEL_NAME(...) (__VA_ARGS__)
[TODO] Clean-up entirely kernel name wrapping when HCC is finally obsolete.
+ Update perl generation, hipify-perl, and affected tests accordingly.
+ Perl part of [#1458]
+ Affected functions: hipFuncSetCacheConfig, hipFuncGetAttributes
+ Implement function generateHostFunctions() in hipify-clang for that purposes
+ Update hipify-perl accordingly
+ Merge of cudaSymbolFuncCall and cudaReinterpretCastArgFuncCall matchers into a single cudaHostFuncCall matcher
+ More const std::string declarations
+ Formatting
+ Affected functions: hipFuncSetCacheConfig, hipFuncGetAttributes
+ Add a corresponding Matcher cudaReinterpretCastArgFuncCall
+ Add reinterpret_cast.cu test
TODO: Do the same for hipify-perl
cudaMemcpyToSymbol, cudaMemcpyToSymbolAsync, cudaGetSymbolSize, cudaGetSymbolAddress, cudaMemcpyFromSymbol, cudaMemcpyFromSymbolAsync
+ Perl part of [#1441]
+ Implement function generateSymbolFunctions() in hipify-clang for that purposes
+ Update hipify-perl
TODO: Eliminate dim3() issue in hipify-perl as well
cudaMemcpyToSymbol, cudaMemcpyToSymbolAsync, cudaGetSymbolSize, cudaGetSymbolAddress, cudaMemcpyFromSymbol, cudaMemcpyFromSymbolAsync
+ Add a corresponding cudaSymbolFuncCall matcher.
+ Add device_symbols.cu test for the above 6 functions, update existed.
+ Fix dim3() type cast issue, update affected tests.
TODO: Do the same in hipify-perl
+ [perl] Add generation of sub countSupportedDeviceFunctions
+ [FIX] Do not perform any replacement for supported device functions
+ [perl] Sync hipify-perl accordingly
+ 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.
Evgeny Mankov
Evgeny Mankov
Maneesh Gupta