*Lay foundation to batch packets efficiently for graphs
*Dynamically copy packets with max threshold set with
DEBUG_HIP_GRAPH_BATCH_SIZE, if not stagger packet copy with pow2
*Default threshold for DEBUG_HIP_GRAPH_BATCH_SIZE is 256
*If TS are not collected for a signal for reuse, create a new signal.
This can potentially increase signal footprint if the handler doesn't run
fast enough.
During a process tear-down we wait on all signals before releasing them:
VirtualGPU::HwQueueTracker::~HwQueueTracker() {
for (auto& signal : signal_list_) {
CpuWaitForSignal(signal);
signal->release();
}
[...]
}
In the case where we exit the process after a GPU error that did not
cause an abort (ulimit -c == 0), waiting for the signal can be skipped.
With the device on the error state, no progress is made, and the signal
is probably never going to be modified again:
inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool yield = false) {
[...]
if (HIP_SKIP_ABORT_ON_GPU_ERROR && amd::Device::IsGPUInError()) {
ClPrint(amd::LOG_ERROR, amd::LOG_SIG,
"Device not Stable, while waiting for Signal ="
"(0x%lx) for %d ns",
signal.handle, kTimeout4Secs);
return true;
}
[...]
}
However, after calling CpuWaitForSignal, when calling "release", we can
end-up on a signal dtor which also tries to wait on the signal. Because
the GPU is the error state, we never receive the signal, and hang the
process during tear down. This happens with the ProfilingSignal dtor:
ProfilingSignal::~ProfilingSignal() {
if (signal_.handle != 0) {
if (hsa_signal_load_relaxed(signal_) > 0) {
LogError("Runtime shouldn't destroy a signal that is still busy!");
if (hsa_signal_wait_scacquire(signal_, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
kUnlimitedWait, HSA_WAIT_STATE_BLOCKED) != 0) {
}
}
hsa_signal_destroy(signal_);
}
}
This dtor should check that the GPU is not in the error state before
trying to wait, which is what this patch implements.
Bug: SWDEV-555043
Bug: SWDEV-553435
Bug: SWDEV-553679
Bug: SWDEV-555119
* libhsakmt: fix UB due to signed integer literal in 1 << 31
Bit shift operations on signed numbers should not shift into or beyond
the signed bit as this results in Undefined Behaviour.
Signed-off-by: Sunday Clement <Sunday.Clement@amd.com>
* libhsakmt: Fix UB due to signed integer literal in 1 << x
Bit Shifting an unsigned integer is undefined behavior.
BUG: SWDEV-532853
Signed-off-by: Sunday Clement <Sunday.Clement@amd.com>
* rocr: Fix UB in various places due signed integer in bit shift
Bit shifting signed integers into or beyond the sign bit is undefined.
Signed-off-by: Sunday Clement <Sunday.Clement@amd.com>
* rocr: Change signed integer literals to unsigned
Changing the signed integers in the macro expressions throughout the file
to avoid overflow.
Signed-off-by: Sunday Clement <Sunday.Clement@amd.com>
---------
Signed-off-by: Sunday Clement <Sunday.Clement@amd.com>
Co-authored-by: Flora Cui <flora.cui@amd.com>
Increased the AMDSMI_MAX_DEVICES to 64 to accomodate all
devices in CPX mode. The link type has been modified in
amd-smi to match with rocm-smi types, updated the same
for drm tests.
---------
Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com>
Increased the AMDSMI_MAX_DEVICES to 64 to accomodate all
devices in CPX mode. The link type has been modified in
amd-smi to match with rocm-smi types, updated the same
for drm tests.
---------
Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com>
[ROCm/amdsmi commit: 6715c5aa92]
* SWDEV-541623 - cuda parity hipLaunchCooperativeKernelMultiDevice and hipExtLaunchMultiKernelMultiDevice
numDevices does not match the system devices
* SWDEV-541623 - enable Unit_hipExtLaunchMultiKernelMultiDevice_Negative_MultiKernelSameDevice
---------
Co-authored-by: agunashe <ajay.gunashekar@amd.com>
Minimum required cmake version of test/CMakeList.txt is bumped from 2.8
to 3.16. This alignes with the version used in CMakeList.txt and will
enable building with cmake 4.
[ROCm/rccl commit: 0f6fec1553]
Minimum required cmake version of test/CMakeList.txt is bumped from 2.8
to 3.16. This alignes with the version used in CMakeList.txt and will
enable building with cmake 4.
This is to prevent calling catch2 macros from outside catch2 TEST_CASE
that can lead to undefined bahavior. This change also disables
hipGetProcAddress tests that are not supported on static build.
Co-authored-by: Ioannis Assiouras <Ioannis.Assiouras@amd.com>
SaleelK
corey-derochie-amd
JC
Swati Rawat
Atul Kulkarni
xuchen-amd
Yiltan
Venkateshwar Reddy Kandula
Ioannis Assiouras
lancesix
David Yat Sin
Sunday Clement
Surya Periaswamy
spolifroni-amd
Nilesh M Negi
Kanangot Balakrishnan, Bindhiya
Mario Limonciello
Williams, Justin
Maisam Arif
amd-hsivasun
Julia Jiang
systems-assistant[bot]
cfreeamd
nawrinsu
Laura Promberger
Saeed, Oosman