Commit Graph

3652 Commits

Author SHA1 Message Date
Evgeny Mankov f68bee02f5 [HIPIFY][tests] Rename the ambiguous call as well 2019-10-25 16:07:31 +03:00
Evgeny Mankov 9529e1d91d [HIPIFY][tests] Fix ambiguous call to cusparseGetErrorString declared in cusparse.h 2019-10-25 16:04:20 +03:00
Alex Voicu 40522e2b6a Add missing operators, fix GCC compilation. (#1589) 2019-10-25 15:44:24 +05:30
Alex Voicu f909a393ff Fix deadlock, remove old __sync_* use. (#1584)
This fixes a deadlock introduced by the switch to TTAS loops, and is therefore mildly urgent (to prevent the CI from hoovering in the broken code).
2019-10-25 15:44:17 +05:30
Rahul Garg 66a3c874c8 [dtest] Fix hipMemset2D test (#1579)
Reverts changes made in #1399. This is a RT api test. For testing hipMemAllocPitch , a new test should be written and that should use correct memset API.
2019-10-25 15:44:05 +05:30
Rahul Garg 14b870d1ce Add hipMemcpy2DfromArray (#1510)
Adds hipMemcpy2DFromArray and hipMemcpy2DFromArrayAsync equivalent to cudaMemcpy2DFromArray and cudaMemcpy2DFromArrayAsync.
2019-10-25 15:43:33 +05:30
Rahul Garg ff8d3fa446 Update profiling doc (#1576) 2019-10-24 17:51:55 +05:30
Jatin Chaudhary f53b1a1755 Adding New Analyze Target Merging with cppcheck (#1583) 2019-10-24 17:46:06 +05:30
Rahul Garg 170c4f0270 Add HIP checks in texture driver sample (#1581) 2019-10-24 17:45:51 +05:30
gandryey f25692b399 Hip vdi profiling header (#1577)
Add HIP-VDI profiling interface for GPU timing collection.
2019-10-24 17:45:42 +05:30
Alex Voicu 26914ec76e Make CAS loops use the TTAS idiom. (#1573)
* Make CAS loops use the TTAS idiom.

* More efficient re-formulation of TTAS.

* Fix typo.

* The typo was not quite a typo
2019-10-24 17:45:20 +05:30
satyanveshd 6c5fbf9b4a Fix occupany APIs (#1560)
Addresses SWDEV-205006
2019-10-24 17:44:47 +05:30
searlmc1 15a699688e Improve performance of v2 arg handling (#1539)
* Improve performance of v2 arg handling

* Missing change to `std::string`
2019-10-24 17:44:05 +05:30
Alex Voicu 84d5b399f6 Improve scalar access into vector types. (#1531)
The improvement is based on the ideas here: https://t0rakka.silvrback.com/simd-scalar-accessor. It yields significantly better ISA when the base's .xyzw members are used.
2019-10-24 17:43:49 +05:30
Aryan Salmanpour 93c688a0c9 [hip] add support for implicit kernel argument for multi-grid sync (#1456)
* [hip] add support for implicit kernel argument for multi-grid sync

* modified code for calculating the prev_sum

* change the impCoopArg type to size_t

* add memory clean up

* launch init_gws and main kernels into two separate loops
2019-10-24 17:43:30 +05:30
Rahul Garg 465581612e Merge pull request #1559 from vsytch/win10_aligned_alloc
Fixes for hipMemcpy_simple on Windows
2019-10-23 13:10:59 -07:00
Evgeny Mankov 2435567e70 [HIPIFY][cmake][#1571] Take into account building hipify-clang as a part of building HIP while installing
[Algorithm]
  [Release]
    If CMAKE_INSTALL_PREFIX is set by the user:
       If BIN_INSTALL_DIR is set by HIP, use it as CMAKE_INSTALL_PREFIX, otherwise CMAKE_INSTALL_PREFIX is used unchanged.
    If the user does not set CMAKE_INSTALL_PREFIX (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT):
       If BIN_INSTALL_DIR is set by HIP, use it as CMAKE_INSTALL_PREFIX, otherwise use PROJECT_BINARY_DIR/bin for installation.
  [Debug]
    If CMAKE_INSTALL_PREFIX is set by the user:
       CMAKE_INSTALL_PREFIX is used unchanged.
    If the user does not set CMAKE_INSTALL_PREFIX (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT):
       use CMAKE_CURRENT_SOURCE_DIR/bin for installation.

Standalone build left unchanged: CMAKE_INSTALL_PREFIX is used if set.
2019-10-23 18:54:45 +03:00
Evgeny Mankov 7ab06b3892 [HIPIFY] Disable delayed template parsing
By implicit unconditional passing -fno-delayed-template-parsing option (which appeared in LLVM 3.8.0, thus doesn't need compatibility wrapping) to hipify-clang.

[Reason] To parse uncalled template functions otherwise they are not parsed without calling, thus not hipified.

Affects cub_03.cu test, which has uncalled global template function.
2019-10-22 19:07:37 +03:00
Evgeny Mankov e2191e23e6 [HIPIFY][#1569] Fix 2019-10-22 11:08:37 +03:00
Evgeny Mankov 3233a845f6 [HIPIFY][tests] Set max clang's CudaArch for corresponding CUDA major.minor version
[Reason] To support maximum CUDA features in offline tests

+ Add defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 restriction for atomicAdd on doubles in atomics.cu.
  So if LLVM < 7 and --cuda-gpu-arch doesn't work, __CUDA_ARCH__ is unset too (350 by default in clang);
  if LLVM >= 7 --cuda-gpu-arch is used and __CUDA_ARCH__ is set based on it.
2019-10-21 17:50:00 +03:00
Evgeny Mankov 9633cdbd8a [HIPIFY][perl] Support of 'using namespace cub' 2019-10-21 17:15:05 +03:00
Evgeny Mankov 9fc7afa738 [HIPIFY][tests] Set max clang's CudaArch for corresponding CUDA version
[Reason] To support maximum CUDA features in offline tests

+ Add CUDA_VERSION >= 800 restriction for atomics.cu

[TODO] Find a way to use or exclude atomicAdd for doubles if LLVM < 7, because
LLVM 6.0.1 and older do not use --cuda-gpu-arch in clang's Driver code at all (option is only declared)
2019-10-21 15:51:25 +03:00
Evgeny Mankov ff6057d1ff [HIPIFY][tests] Set -I for CUDA path instead of --cuda-path for LLVM < 4 2019-10-20 20:08:56 +03:00
Evgeny Mankov 5bf1ff19ff [HIPIFY][tests] Exclude all CUB tests if CUDA_CUB_ROOT_DIR is not set 2019-10-20 20:03:18 +03:00
Vladislav Sytchenko 432380aa5d Remove extra #endif. 2019-10-18 16:40:29 -04:00
Evgeny Mankov 44a897a146 [HIPIFY][tests] Test clean-up 2019-10-18 18:55:52 +03:00
Evgeny Mankov 86f6756b02 [HIPIFY][CUB][#1460] Add "using namespace cub" translation support
+ Add cub_03.cu
2019-10-18 18:51:40 +03:00
Evgeny Mankov eb6690bbba Merge pull request #1558 from aaronenyeshi/fix-hipify-cmake-version
[HIPIFY][cmake] Make CMakeLists use default 3.5.1 for Ubuntu 16.04
2019-10-18 06:39:35 +03:00
Rahul Garg 07eed1e5bf Merge pull request #1550 from yxsamliu/new-launch
Add -fhip-new-launch-api to hipcc for HIP/VDI
2019-10-17 19:07:32 -07:00
Vladislav Sytchenko f4440817cb _aligned_malloc() on Windows first takes size, then alignment, which is the opposite of how the similar function behaves on Linux. Memory allocated by it also has to be freed using _aligned_free(), unlike Linux where we can use regular free().
Edit aligned_alloc() macro and add a aligned_free() one to align with the above behaviour.
2019-10-17 18:58:32 -04:00
Aaron Enye Shi b3ea58abe7 [HIPIFY][cmake] Make CMakeLists use default 3.5.1 for Ubuntu 16.04 2019-10-17 21:21:24 +00:00
Evgeny Mankov 1165e6bd71 [HIPIFY][doc] Update README.md
+ Versions, testing
2019-10-17 22:26:48 +03:00
Rahul Garg 446718f990 Revert "hipcc defaults to code object v3 (#1298)"
This reverts commit c8f38a139a.
2019-10-17 13:27:28 -04:00
Evgeny Mankov 8c3dff7ab9 [HIPIFY][cmake] Add install rule for clang-resource-headers
+ Fix: set destination for all installing files to ${CMAKE_INSTALL_PREFIX}
2019-10-17 15:05:55 +03:00
Rahul Garg a21fe1443b Merge pull request #1544 from vsytch/master
QoL changes to the hipMemset family
2019-10-16 18:54:20 -07:00
Evgeny Mankov e557563947 [HIPIFY][CUB][#1460] Add cub:: namespace support in TemplateInstantiation of cudaLaunchKernel
+ Update cub_02.cu test accordingly
2019-10-16 19:02:13 +03:00
Vladislav Sytchenko 86d0c5fa5a hipMemset2D and hipMemset3D tests should be passing by default. 2019-10-16 11:02:38 -04:00
Evgeny Mankov 0a20048759 [HIPIFY] Refactor a couple of matcher functions
+ Separate out GetSubstrLocation function for finding substr SourceLocation in a given SourceRange
2019-10-16 13:43:56 +03:00
Evgeny Mankov 5555d46e66 [HIPIFY][CUB][#1460] Implement cubFunctionTemplateDecl matcher
+ Add cub_02.cu test
+ Partial fixes #1460
2019-10-16 13:08:11 +03:00
kjayapra-amd 9d571e3c9e Use the correct return type in runTest in 11_texture_driver sample. (#1546)
Fixes SWDEV-203394.
Currently in runTest() returns true, even if the texture reference copy does not happen. Using the existing testResult Flag to return from runTest().
2019-10-16 10:52:15 +05:30
vsytch c2aadd4d12 Update hipMathFunctions, hipTestHalf and hipTestNativeHalf tests to support Navi10 and Navi14. (#1545) 2019-10-16 10:51:48 +05:30
kpyzhov 9773f94c71 [hipcc] Temporary add -D_OPENMP to clang options to workaround cmake issue (#1540)
* Temporary add -D_OPENMP to clang options in hipcc to allow using CMake OpenMP detection with hip-clang (until updated CMake version is available).
2019-10-16 10:51:28 +05:30
Nick Curtis d16963c9d5 Guard against division by zero for no VGPR usage (e.g., in an empty kernel) (#1528)
* guard against division by zero for no VGPR usage (e.g., in an empty kernel)

* fix bracket format

* clean up parenthesis
2019-10-16 10:49:56 +05:30
Jatin Chaudhary b3351561c5 Adding code object manager to rtc (#1526)
Adding Code Object Manager file to rtc to resolve address of Bundled_code_object in libhiprtc.so
2019-10-16 10:49:16 +05:30
Xiaozhu Meng f9b8a01c77 Fix struct declaration for C (#1524)
This change is necessary for HPCToolkit to use Roctracer to produce code centric profiling view.
2019-10-16 10:48:55 +05:30
Yaxun (Sam) Liu 739530d53b Add -fhip-new-launch-api to hipcc for HIP/VDI 2019-10-15 21:47:33 -04:00
Vladislav Sytchenko cc5abec092 In the hipMemset2D and hipMemset3D tests synchronize with the default stream after performing an async memset. 2019-10-15 17:15:49 -04:00
Vladislav Sytchenko f402b6d01a Update indentation in the hipMemset3D test. Replace all tabs with four spaces. 2019-10-15 15:29:14 -04:00
Vladislav Sytchenko c83b6adb33 Add async subtest to hipMemSet3D 2019-10-15 14:24:04 -04:00
Vladislav Sytchenko 39e42d4056 hipMemset2D test should pass only if both async and sync subtests pass. 2019-10-15 14:20:14 -04:00