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
- Use correct header in device_library_decl
- use std:: instead of __hip_internal:: for host compilation
- hide device specific stuff behind __clang__ and __HIP__ check
Change-Id: I2f3647e00555ed0e79f9954a459c41394c3cd49b
- 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
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
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
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
This change fixes random segfaults in graph tests that
are seen after the change make internal callbacks non-blocking.
The callback thread that decreases the GraphExec ref count
may now run after the runtime shutdown. This can cause a segfault
because the hip::device that is accessed in GraphExec destructor
is already destroyed during runtime shutdown. This patch ensures
that the hip::device object stays alive until after the
callback thread completes.
Change-Id: I75a6ac01f27a0b2250bbd10ed389ebfb322927af
Fixes#123. find_program doesn't follow CMP0074 and thus ignores LLVM_ROOT and Clang_ROOT. This change adds LLVM_ROOT and Clang_ROOT to the search path of find_program for llvm-mc and clang in hiprtc to mimics previous add_package behaviour.
Caveat: cmake-specific variables like CMAKE_PREFIX_PATH will take precedence over paths specified with HINTS for find_program, there's no way to change the ordering unless we skip cmake-specific variables all together using NO_CMAKE_PATH and NO_CMAKE_ENVIRONMENT_PATH.
Change-Id: I1fedb60cda09744416e19b3c6e3e0c5c9045f8e7
This removes almost all uses of the deprecated
__AMDGCN_WAVEFRONT_SIZE macro, which is unavailable
when targeting SPIR-V, and adds a SPIR-V compatible
formulation of warpSize (which should end up as the
sole definition of warpSize once we remove support
for treating it as a compile time constant). It
is incomplete in that the cooperative_groups
implementation will need additional surgery.
Squashed commit of the following:
commit 6840826c3fec8516857dc4f2092d84358550f588
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri Dec 6 23:36:32 2024 +0000
Add deprecation warning for constexpr uses of `warpSize`.
commit a72307a7353034c2de53fd164e016967945fd0d1
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri Dec 6 23:12:14 2024 +0000
Prepare HIP RT for SPIR-V.
commit 5e40dd746ac4f8c93b521ef048ff9d494905ba95
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri Dec 6 22:46:05 2024 +0000
Revert stale change.
commit 231fe91c53dba4cabd832fc84eaa6ddb402271a0
Merge: a48905ec9 78f62d323
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri Dec 6 22:37:24 2024 +0000
Merge branch 'amd-staging' of https://github.com/ROCm/clr into amd-staging
commit a48905ec9cfe0e017cc64943195be82b530117d7
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Tue Sep 17 03:14:56 2024 +0100
Add scaffolding for SPIR-V support.
Change-Id: I2e84bbe90df58a5f9a8709b619905f04fa5b96dc
- Added missing validation as graph node should not be created
if parameters are invalid
- Fix conversion of input params to graphNode params
Change-Id: I37ab04942b5fb2eb07386850cb7dbbf26f9ca967
If a module is loaded on one device, hipModuleGetFunction and other similar APIs should be able to run successfully from another device.
Change-Id: I96084cbd6c6dcf2a81019779a6ab1842ef2f35d1
German Andryeyev
Jimbo Xie
Jatin Chaudhary
Anusha GodavarthySurya
David Salinas
harkgill-amd
Aidan Belton-Schure
Ioannis Assiouras
Satyanvesh Dittakavi
taosang2
Steven Chung
Ajay
Rahul Manocha
Jacob Lambert
Saleel Kudchadker
Gerardo Hernandez
Marko Arandjelovic
Sourabh Betigeri
zichguan-amd
Jaydeep Patel
Konstantin Zhuravlyov
Vladana Stojiljkovic
Marko Arandjelovic