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
[ROCm/clr commit: 21c223f8df]
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
[ROCm/clr commit: 272ef9a7bf]
thread_rank() gives thread index in a block. Limit the range to the
current warp size.
Change-Id: Ib5c9831236096485cf99ba7ab0b911a3b10de31c
[ROCm/clr commit: bd7d40a4d8]
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 12dc02b4f
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
[ROCm/clr commit: dd4378611a]
- 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
[ROCm/clr commit: db8527f655]
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
[ROCm/clr commit: c46f843b99]
This is to avoid calling the HIP_INIT macro during the shutdown process.
Change-Id: I2e65f6e10491918a17445ee1e8ddd08286070358
[ROCm/clr commit: 5e3a29078d]
- Resolve stream once for event record. We should avoid calling
getStream again in addMarker
Change-Id: I78448c4f151ae10a5c8e8c248b2f4078b84191cb
[ROCm/clr commit: a22c45d635]
The vector with all kernels is preallocated on the executable init.
Thus, reduce the scope of global lock to the binary creation only.
Change-Id: I73035013a6562175069137e895bba815f466ee35
[ROCm/clr commit: 0640d36019]
Support gfx9-4-generic target to cover mi3XX.
Support features sramecc and xnack in generic target.
Improve some code formats.
Add more log on compiler.
Change-Id: I6b3c6af55c60cffd43ce6f17b75998f751b75713
[ROCm/clr commit: 3ad8f1b811]
Some libs use __HIP_NO_HALF_OPERATORS__ and __HIP_NO_HALF_CONVERSIONS__
which results in operators being hidden and can cause errors.
Change-Id: I83c194d7d727cba30b46d7c296f7d396549f5fca
[ROCm/clr commit: 98b33886cd]
On Windows, hipHostRegister may add a single object in the MemObjMap
that maps to memory that is allocated on different devices.
This change ensures that the offset that is returned from
getMemoryObject() is computed relative to the memory that is allocated
on the current device.
Change-Id: I5fd3af200bf6f4926fdeaea12dcb9d0154d3a843
[ROCm/clr commit: e80442fdbf]
- Header files inside rocclr/utils when included from hipamd or opencl should be included as #include "rocclr/utils/xxx.h" instead of "utils/xxx.h"
Change-Id: Ic0760c33b9d091f5620dec67e5482c9698d22093
[ROCm/clr commit: 78f62d3230]
Gerardo Hernandez
Marko Arandjelovic
Ioannis Assiouras
Sourabh Betigeri
zichguan-amd
Jaydeep Patel
Jatin Chaudhary
Konstantin Zhuravlyov
Vladana Stojiljkovic
Marko Arandjelovic
Rahul Manocha
Jacob Lambert
Anusha GodavarthySurya
Saleel Kudchadker
German Andryeyev
taosang2
Istvan Kiss
Maneesh Gupta
Todd tiantuo Li
Branislav Brzak