* Enables warp sync builtins by default
* Removes HIP_ENABLE_WARP_SYNC_BUILTINS; that macro will no longer have an effect. Instead, we will now be able to disable the builtins with the macro: HIP_DISABLE_WARP_SYNC_BUILTINS
Since the sub-buffer(virtual memory that is mapped to device memory) is associated with device memory, it should utilize the device context instead of the host context. The original implementation caused hipMemcpyPeer to not take the P2P path, as the memory object was treated as host memory.
* SWDEV-527299 - Support HIP_POINTER_ATTRIBUTE_CONTEXT
As HIP enables UVA by default, it seems we can simply expose the context to support this feature.
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.
* 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.
* 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.
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.
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
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
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.
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
Brzak, Branislav
Sang, Tao
Jayaprakash, Karthik
Belton-Schure, Aidan
Patel, Jaydeepkumar
Xie, Pengda
Assiouras, Ioannis
Hernandez, Gerardo
Arandjelovic, Marko
Huang, AnZhong
Chaudhary, Jatin Jaikishan
Dittakavi, Satyanvesh
Guan, Zichuan
Arsenault, Matthew
Searles, Mark
Godavarthy Surya, Anusha
Jiang, Julia
Betigeri, Sourabh
GunaShekar, Ajay
Milanov, Aleksandar
Haehnle, Nicolai