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]
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]
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]
- hipStreamWaitEvent may not resolve streams
- Correct usage of flag passed to streamWait function
Change-Id: I2ee163615d303b98937c1035d60da283cce6f677
[ROCm/clr commit: 940347ad42]
- This change tries to save extra synchronization packets we may insert
as we didnt track the completion signals for every command. We track
the current enqueued command until it exits the enqueue stage. We also
record the exit scope to know if we flushed the caches
- Handle correct release scopes and store completion signal as HW events
- Use a new finishCommand implementation to only wait for the command
passed as the argument
Change-Id: Ie4350c5dd24f5d48dfa6ccbabd892f0544caadcc
[ROCm/clr commit: e03e4f3b5d]
Since hipMemMap can be called for multiple device handles on the same virtual memory, the same is true for hipMemUnmap, meaning that virtual memory can be "partially unmapped".
This means that the unmap function can be called for a specific part of the reserved address, meaning that only the designated subbuffer should be released. If unmap is called on the entire reserved memory, then all subbuffers should be released.
The main point is that for every hsa_amd_vmem_map, there should be a corresponding hsa_amd_vmem_unmap. Otherwise, if entire memory is unmapped by a single unmap call, then HSA will report the memory as "in use" if an attempt is made to delete it.
Change-Id: I039308eafb820decfb1c09f60347f26cdad1a362
[ROCm/clr commit: 3ec1d2d2f1]
1) Add Linker APIs to runtime to support SPIRV linking
2) Migrate Internal implementations to runtime and share with rtc
3) Add Support to bundled and unbundled SPIRV Code object linking.
Change-Id: Ic1fd4431f842a208a2468e8aec54a65b5fa6b0e3
[ROCm/clr commit: 5930f047bb]
This change removes the stream callback from hipStreamWaitEvent and
uses a stream memory wait operation instead. This allows the
hipStreamWaitEvent to be non-blocking on the host.
Change-Id: Ie5530febda5a5bcb5daa0db8a01249d6b137fd43
[ROCm/clr commit: 721c5800ca]
Add initial implementation of virtual memory heap with
dynamic virtual memory mapping support for memory pools.
DEBUG_HIP_MEM_POOL_VMHEAP controls the new method.
Change-Id: I8dc5be2e0f34ab472f1800f43bb6243639a5e500
[ROCm/clr commit: 296dce5570]
- Also add a cache, which allows compiled code objects to be reused
instead of compiling again. This should improve performance on
multigpu systems.
Change-Id: Ib135d616c076b77f8aaf28de275d408b38021d89
[ROCm/clr commit: 0391aec14a]
There are 2 functional changes to this patch:
* Use GPU timing for internal markers for HIP.
* Measure CPU time closer to GPU timer, to reduce delta between GPU/CPU timestamp measurements.
There are some smaller non-functional updates:
* waifForFence -> waitForFence typo
* Remove unused drmProfiling
Change-Id: I4c5fa600a842ab60e454888779edcac8449a902a
[ROCm/clr commit: 179801a750]
Resolved an issue where hipEventSynchronize and hipStreamWaitEvent APIs
did not function correctly for events created with the hipEventInterprocess flag.
The bug caused the event to be incorrectly marked as "recorded,"
leading to these APIs failing to wait for the event as expected.
Change-Id: Ic9fdfaab2393beb93d6e0b83661545e902a63499
[ROCm/clr commit: 1cdfbfd270]
hipGetLastError should return the error by any of the previous APIs
in the same host thread to match the CUDA behavior, whereas
hipExtGetLastError will return the error by the immediate previous API.
This Ext API was added earlier to facilitate the existing HIP apps which
are following the current behavior of hipGetLastError
Change-Id: I61e95b1fc136cc761e2434e02187b7ed2598b733
[ROCm/clr commit: 4b443f8133]
Xie, Pengda
Chaudhary, Jatin Jaikishan
Hernandez, Gerardo
Haehnle, Nicolai
Xie, Jiabao(Jimbo)
Patel, Jaydeepkumar
Stojiljkovic, Vladana
Brzak, Branislav
Andryeyev, German
Betigeri, Sourabh
Sang, Tao
Belton-Schure, Aidan
Arandjelovic, Marko
Godavarthy Surya, Anusha
Rakesh Roy
Ioannis Assiouras
Jacob Lambert
Julia Jiang
Saleel Kudchadker
Rahul Manocha
Li, Todd tiantuo
Jimbo Xie
Satyanvesh Dittakavi