+ Perl part of [#1458]
+ Affected functions: hipFuncSetCacheConfig, hipFuncGetAttributes
+ Implement function generateHostFunctions() in hipify-clang for that purposes
+ Update hipify-perl accordingly
+ 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
+ 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
+ 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
* [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
+ 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
* [dtests] refactor windows specific changes
* Refactor hipMemoryAllocateCoherentDriver - PR- 1309
* Fix missing z in _putenv_s
* Revert "Fix missing z in _putenv_s"
This reverts commit 099a1b20a5c75c5f122d57c0ad2bca01745cdc9c.
* Refactor changes from PR 1299
* Update hipEnvVarDriver.cpp
* 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
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
* Enabled gcc for hip host code
* Adding tests for hip code + (gcc & g++), without kernels
* Excluding nvcc platforms for gcc and g++ tests + Addressing review comments
* minor code clean-up
* Add rocm include path
* Added relative path for library
* Hiding non supported functions for gcc
* Incorporating review comments
...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.
* Added support of hipOccupancyMaxActiveBlocksPerMultiprocessor & hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags APIs
* Taking into account of SGPR usage to determine the max active blocks in hipOccupancyMaxActiveBlocksPerMultiprocessor()
* Fix hipMemcpy-size test running out of Host Mem
The hipMemcpy-size uses a maxElem calculated from the total GPU mem /8. Then it will allocate 4 times that amount of host memory. This tests begins failing when there is not enough host memory, such as on systems with 32GB GPU mem, and 16GB RAM. This fixes the test if not enough host memory is available on the system.
* Add windows support to hipMemcpy-size fix
* avoid linking extra libs for windows
* HIPMemcpy-size Remove freeCPU including swap
* [hit] Workaround for %cc and %cxx mappings.
HIP CMakeLists.txt modifies CMAKE_C_COMPILER and CMAKE_CXX_COMPILER.
This messes up any dtests that want to test against cc/c++.
So hardcode %cc to /usr/bin/cc and %cxx to /usr/bin/c++ for now till
we come up with a better solution.
Change-Id: I7dce93ce8360191e612a94e3a735e5612ac27ab5
* [hit] Add auto-variable %hip-path to syntax for BUILD_CMD
Change-Id: Id097a183fbce2b2c9691d0180d3304dd17a4e016
* [HIP][tests] New testcases for module api
* [HIP][Tests]Support for CUDA devices
* Updated tests as per latest master & test GetGlobal to work on all platforms
* Add Max Texture 1D,2D,3D device properties
* Corrected testcase to use enums defined in hipDeviceAttribute_t
* Added texture 1D,2D and 3D support for NVIDIA path
[Reason] To be compatible with CUDA [#1133]
Update HIP code, hipify-clang, tests and docs
[TODO] Add support of the corresponding functions on nvcc fallback path
* Put 3-wide vector types on a ketogenic diet.
* Remove needless include.
* Do not be narrow-minded.
* Do not be narrow-minded.
* Put the C people on a diet too.
* Implement the hipOccupancyMaxPotentialBlockSize function
* Replaced hipGetDeviceProperties() call by ihipGetDeviceProperties() in ihipOccupancyMaxPotentialBlockSize()
* Add test for hipOccupancyMaxPotentialBlockSize in Module API
* Added extern declaration for ihipGetDeviceProperties() to be accessed inside ihipOccupancyMaxPotentialBlockSize()
* fixed hipOccupancyMaxPotentialBlockSize test build issue
* Fix hipOccupancyMaxPotentialBlockSize dtest
* Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest
* Revert "Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest"
This reverts commit 0480ff56f1441fc515d2c26ce33783e303423938.
* Disable hipOccupancyMaxPotentialBlockSize dtest on NVCC
* move extern declaration of ihipGetDeviceProperties to hip_module.cpp
* Update the limiation of 32 wavefronts per CU and 800/512 SGPRs for VI/pre-VI chips to calculate the occupancy
Evgeny Mankov
AlexBinXie
ansurya
Aryan Salmanpour
Sameer Sahasrabuddhe
Sarbojit2019
Rahul Garg
Maneesh Gupta
amd-lthakur
lthakur
wkwchau
Vladislav Sytchenko
Aaron Enye Shi
kpyzhov
Yaxun (Sam) Liu
Alex Voicu