b9c8dd8ac6
Fixed CI build failure
2019-10-25 12:21:41 +05:30
b2238bacd4
Refactored the file as suggested
2019-10-25 10:44:38 +05:30
84fb936dfe
Update matmul.cpp
2019-10-25 09:22:07 +05:30
9a860e6766
Update hipExtModuleLaunchKernel.cpp
2019-10-25 09:19:49 +05:30
ff8d3fa446
Update profiling doc ( #1576 )
2019-10-24 17:51:55 +05:30
f53b1a1755
Adding New Analyze Target Merging with cppcheck ( #1583 )
2019-10-24 17:46:06 +05:30
170c4f0270
Add HIP checks in texture driver sample ( #1581 )
2019-10-24 17:45:51 +05:30
f25692b399
Hip vdi profiling header ( #1577 )
...
Add HIP-VDI profiling interface for GPU timing collection.
2019-10-24 17:45:42 +05:30
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
6c5fbf9b4a
Fix occupany APIs ( #1560 )
...
Addresses SWDEV-205006
2019-10-24 17:44:47 +05:30
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
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
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
0adcc890b0
Adding a directed test case for hipExtModuleLaunchKernel() api.
2019-10-24 15:06:28 +05:30
465581612e
Merge pull request #1559 from vsytch/win10_aligned_alloc
...
Fixes for hipMemcpy_simple on Windows
2019-10-23 13:10:59 -07:00
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
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
e2191e23e6
[HIPIFY][ #1569 ] Fix
2019-10-22 11:08:37 +03:00
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
9633cdbd8a
[HIPIFY][perl] Support of 'using namespace cub'
2019-10-21 17:15:05 +03:00
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
ff6057d1ff
[HIPIFY][tests] Set -I for CUDA path instead of --cuda-path for LLVM < 4
2019-10-20 20:08:56 +03:00
5bf1ff19ff
[HIPIFY][tests] Exclude all CUB tests if CUDA_CUB_ROOT_DIR is not set
2019-10-20 20:03:18 +03:00
432380aa5d
Remove extra #endif.
2019-10-18 16:40:29 -04:00
44a897a146
[HIPIFY][tests] Test clean-up
2019-10-18 18:55:52 +03:00
86f6756b02
[HIPIFY][CUB][ #1460 ] Add "using namespace cub" translation support
...
+ Add cub_03.cu
2019-10-18 18:51:40 +03:00
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
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
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
b3ea58abe7
[HIPIFY][cmake] Make CMakeLists use default 3.5.1 for Ubuntu 16.04
2019-10-17 21:21:24 +00:00
1165e6bd71
[HIPIFY][doc] Update README.md
...
+ Versions, testing
2019-10-17 22:26:48 +03:00
446718f990
Revert "hipcc defaults to code object v3 ( #1298 )"
...
This reverts commit c8f38a139a
2019-10-17 13:27:28 -04:00
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
a21fe1443b
Merge pull request #1544 from vsytch/master
...
QoL changes to the hipMemset family
2019-10-16 18:54:20 -07:00
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
86d0c5fa5a
hipMemset2D and hipMemset3D tests should be passing by default.
2019-10-16 11:02:38 -04:00
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
5555d46e66
[HIPIFY][CUB][ #1460 ] Implement cubFunctionTemplateDecl matcher
...
+ Add cub_02.cu test
+ Partial fixes #1460
2019-10-16 13:08:11 +03:00
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
c2aadd4d12
Update hipMathFunctions, hipTestHalf and hipTestNativeHalf tests to support Navi10 and Navi14. ( #1545 )
2019-10-16 10:51:48 +05:30
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
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
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
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
739530d53b
Add -fhip-new-launch-api to hipcc for HIP/VDI
2019-10-15 21:47:33 -04:00
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
f402b6d01a
Update indentation in the hipMemset3D test. Replace all tabs with four spaces.
2019-10-15 15:29:14 -04:00
c83b6adb33
Add async subtest to hipMemSet3D
2019-10-15 14:24:04 -04:00
39e42d4056
hipMemset2D test should pass only if both async and sync subtests pass.
2019-10-15 14:20:14 -04:00
0200aa3a21
Update the declarations of hipMemsetD8, hipMemsetD8Async, hipMemsetD16, hipMemsetD16Async. These functions are type aware and take in as their third argument the number of elements in the buffer, not the buffer size. Change the name of this argument from sizeBytes to count to align with the above description.
2019-10-15 14:18:42 -04:00
Anusha Godavarthy Surya
amd-lthakur
amd-lthakur
Rahul Garg
Jatin Chaudhary
gandryey
Alex Voicu
satyanveshd
searlmc1
Aryan Salmanpour
Evgeny Mankov
Vladislav Sytchenko
Aaron Enye Shi
kjayapra-amd
vsytch
kpyzhov
Nick Curtis
Xiaozhu Meng
Yaxun (Sam) Liu