Fix pytorch 2.5 issues, by defining reduce sync operations for type __half in amd_hip_fp16.h and not in
amd_warp_sync_functions.h which is problematic in case __half does not get included before that header.
Only define types not supported by cuda if HIP_ENABLE_EXTRA_WARP_SYNC_TYPES is defined, to avoid portability issues
[ROCm/clr commit: 66496258b4]
Explicitly nulling the pointer causes us to report the error below
instead of keeping a dangling pointer around that will most likely lead
to a subsequent segfault.
[ROCm/clr commit: 199b0f1086]
This now has host conversions too, which is directly from Christopher's
work on fcbx.
Signed-off-by: Christopher M. Riedl
* add const to func parameter
* do not depend on builtins, use gfx950 detection
[ROCm/clr commit: 628777b73d]
Currently, we check if there's enough system RAM even if we don't allocate on host device. This is incorrect logic.
We should not check for this size on windows because PAL checks for memory allocation. See SWDEV-467263.
Co-authored-by: Jimbo Xie <jiabaxie@amd.com>
[ROCm/clr commit: 0d6e554d92]
* SWDEV-505795 - Return the same ptr from hipIpcOpenMemHandle if it is called multiple times
* Move initialization outside of if statement
[ROCm/clr commit: e91cb4f320]
Fix data validation issue of rocFFT when dynamic queue on.
ReleaseHwQueue() can be called only when no command in HostQueue.
The checking condition need be protected by lock.
[ROCm/clr commit: 18d191fd1d]
Add VmHeapArray class to reduce the pressure on VA reservation, since
multiple memory pools can be active at the same time.
[ROCm/clr commit: e974f7fde1]
Also removes asserts in cooperative groups shfl functions since
__hip_bfloat16 shfl is present now
Change-Id: I57578b6e68dccc10c2ddcd194e9cc18bc7732ce1
[ROCm/clr commit: 376f23b86a]
Needs further debugging but for now can test the change
Need to verify if this fixes all the below issues-
SWDEV-512754, SWDEV-511675, SWDEV-511055, SWDEV-504085, SWDEV-499503
Also verify original issues
SWDEV-471863, SWDEV-490991
Change-Id: Ic845f851de1b98e8ed9aa0f07afddec3858119e9
[ROCm/clr commit: f1b8ee7b7f]
- For D2H cases avoid passing dependent signals to SDMA, the signals
take a while to resolve on SDMA engine
Change-Id: I569635228af977847f201c82ca897002f8f2f4a8
[ROCm/clr commit: 78d0ff2dbc]
This reverts commit f674ba58f0.
Reason for revert: 6.4 Preview changes need not be merged to amd-staging as of now
Change-Id: I86452adfed14655f72d90440a486089743cc6587
[ROCm/clr commit: 5da8ce45ab]
This reverts commit 82f78ce187.
Reason for revert: 6.4 Preview changes need not be merged to amd-staging as of now
Change-Id: Ifba0c8a248bc40deaa9c59b7f2901531300e5ea4
[ROCm/clr commit: 4206405514]
This reverts commit 04dc7ca51f.
Reason for revert: 6.4 Preview changes need not be merged to amd-staging as of now
Change-Id: I04af8603053338f08c396e78ff8a6715e641ca19
[ROCm/clr commit: 3fa6049c46]
Using target_link_libraries does not properly link the hipRTC-header.o
into libhiprtc for static build. Change to use target_sources instead.
This does not affect the linkage in the shared build.
Change-Id: I626f9eacc1637b792a50e7ddddb5db09e704ac4a
[ROCm/clr commit: 8f54aeb765]
Also part of SWDEV-510994.
1. Fix atomicMin/Max_system() for float and double.
2. Remove logics of gfx941 which isn't supported.
Change-Id: Iacfdc1bc13e8da2f5df8751bb315b37d33cea667
[ROCm/clr commit: d91e1f19d0]
Milanov, Aleksandar
Hernandez, Gerardo
Yao, Longlong
Haehnle, Nicolai
Jiang, Julia
Sang, Tao
Chaudhary, Jatin Jaikishan
Andryeyev, German
Xie, Jiabao(Jimbo)
Patel, Jaydeepkumar
Stojiljkovic, Vladana
Brzak, Branislav
Betigeri, Sourabh
Arandjelovic, Marko
Mallya, Ameya Keshava
MartinezFernandez, Juan
GunaShekar, Ajay
Belton-Schure, Aidan
Dittakavi, Satyanvesh
Gupta, Maneesh
Godavarthy Surya, Anusha
Saleel Kudchadker
Pengda Xie
Rakesh Roy
Ioannis Assiouras
Jacob Lambert