Directly use the builtins. Use the elementwise versions since there's
no implied errno, regardless of -f[no]-math-errno.
I didn't change the cases unnecessarily casting. The bfloat and vector
cases should work directly.
[ROCm/clr commit: 1db9a7d48b]
* SWDEV-520352 - Remove HostThread and legacy monitor
Remove HostThread, semaphore and legacy monitor.
Make original logics of thread and command queue stricker.
Add more comments to make logics clearer.
Some other minor improvement.
Also part of SWDEV-458943.
[ROCm/clr commit: 96cadbc9e9]
Make sure that a newly created FatBinaryInfo is assigned to modules only after extractFatBinary has been called for the object.
[ROCm/clr commit: 1099e0a131]
Support programmatic query and change of scratch limit on
AMD devices.
Change-Id: Id5da355a77366f97868e462847f3916e87fd2af6
[ROCm/clr commit: 1113eff3f9]
* SWDEV-518831 - fix streams' sync issue in mthreads
1. Fix sync issue of null stream and non-null streams in
multithreads.
2. Remove assert(GetSubmissionBatch() == nullptr) as it
is invalid in multithreads.
3. Update getActiveQueues() to deal with the state of
being terminated.
[ROCm/clr commit: 27aad09bd4]
Do not use __ockl_activelane_u32() to calculate the index of the lane within the mask, as that would not work with divergent masks that have other bits on before the associated lane.
[ROCm/clr commit: 1a8d766836]
Also add flag: HIP_FORCE_SPIRV_CODEOBJECT to allow override to force use
SPIRV.
* use cache for already compiled code objects
* address review comments and use the two spirv isa names
[ROCm/clr commit: 07e57a1f0d]
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]
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]
Jayaprakash, Karthik
Dittakavi, Satyanvesh
Chaudhary, Jatin Jaikishan
Guan, Zichuan
Arsenault, Matthew
Searles, Mark
Godavarthy Surya, Anusha
Jiang, Julia
Sang, Tao
Betigeri, Sourabh
Assiouras, Ioannis
GunaShekar, Ajay
Hernandez, Gerardo
Brzak, Branislav
Arandjelovic, Marko
Xie, Pengda
Milanov, Aleksandar
Haehnle, Nicolai
Xie, Jiabao(Jimbo)
Patel, Jaydeepkumar
Stojiljkovic, Vladana
Andryeyev, German
Belton-Schure, Aidan