+ 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.
[ROCm/hip commit: 277d3b8369]
directed_tests/runtimeApi/module/hipLaunchCooperativeKernel.tst - Disabling test temporarily until driver support is available.
directed_tests/runtimeApi/memory/hipArray.tst - Disabling test temporarily to reimplement it correctly.
[ROCm/hip commit: 87834500b9]
Added new memory API's hipMemAllocPitch, hipMemAllocHost, hipMemsetD16, hipMemsetD16Async, hipMemsetD8Async
Modified to support all scenarios hipMemcpyParam2DAsync, hipMemcpyParam2D.
[ROCm/hip commit: ba9c6e13e4]
[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.
[ROCm/hip commit: 6bb9913e8a]
+ Perl part of [#1458]
+ Affected functions: hipFuncSetCacheConfig, hipFuncGetAttributes
+ Implement function generateHostFunctions() in hipify-clang for that purposes
+ Update hipify-perl accordingly
[ROCm/hip commit: 9d1d4b78e3]
+ Affected functions: hipFuncSetCacheConfig, hipFuncGetAttributes
+ Add a corresponding Matcher cudaReinterpretCastArgFuncCall
+ Add reinterpret_cast.cu test
TODO: Do the same for hipify-perl
[ROCm/hip commit: dd6a92afb7]
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
[ROCm/hip commit: 3722d5b4b9]
+ Do not treat somenamespace::device_function_name as a device function
+ Fix generation of warnUnsupportedDeviceFunctions function in hipify-clang
+ Update hipify-perl based on hipify-clang -perl generation
+ Update device test math_functions.cu for hipify-perl
[Restrictions]
- hipify-perl is yet unable to handle function declarations in user namespaces
- hipify-perl is yet unable to handle using directive
[ROCm/hip commit: 4f59ec25fe]
* [hip][tests] add a unit test for testing hipLaunchCooperativeKernel
* use __ockl_grid_sync function
* remove already defined __ockl_grid_sync function
* use sync function for grid synchronization
[ROCm/hip commit: 48880a017e]
+ 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.
[ROCm/hip commit: 6602fadc16]
The implementation for OCKL AS was recently removed from the device
library since that feature is now superseded by hostcall.
[ROCm/hip commit: 70023c9075]
+ 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
[ROCm/hip commit: 24be21495d]
Compilation error being observed due to popen(), pclose() and setenv() linux calls on windows. Replaced with appropriate calls on windows.
[ROCm/hip commit: 9b31d26237]
Compilation error being observed on windows due to aligned_alloc() call. Mapped the call to _aligned_malloc() for windows.
[ROCm/hip commit: 9abae7114c]
* Removed unwanted #include sys/time.h , gettimeofday() and timeval variables and this also helps avavoid compilation error in windows due to gettimeofday() call equivalent of which is not available in windows
* Changed the Macro name from GPU_PRINT_TIME to MY_LAUNCH_MACRO
[ROCm/hip commit: e94c0592de]
Changed the third arg of the functions __hip_as_write_block and __ockl_as_write_block from ulong to uint64_t so as to fix the compilation error in windows
[ROCm/hip commit: d3ffad7c83]
...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.
[ROCm/hip commit: 25075729f9]
* Added query of hipDeviceAttributeHdpMemFlushCntl and hipDeviceAttributeHdpRegFlushCntl
* Added NVCC blocker for the hip*FlushCntl test cases
[ROCm/hip commit: e7447d5809]
* Added support of hipOccupancyMaxActiveBlocksPerMultiprocessor & hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags APIs
* Taking into account of SGPR usage to determine the max active blocks in hipOccupancyMaxActiveBlocksPerMultiprocessor()
[ROCm/hip commit: 4b18b321f7]
Evgeny Mankov
Jatin Chaudhary
ansurya
Sarbojit2019
Yaxun (Sam) Liu
AlexBinXie
Aryan Salmanpour
Sameer Sahasrabuddhe
Rahul Garg
Maneesh Gupta
amd-lthakur
lthakur
wkwchau
Vladislav Sytchenko
Aaron Enye Shi