SWDEV-545489, SWDEV-545082 - Free marker when ipc event is recorded (#784)

Co-authored-by: Rahul Manocha <rmanocha@amd.com>

[ROCm/clr commit: da068a2492]
Этот коммит содержится в:
Stojiljkovic, Vladana
2025-08-14 11:08:53 +02:00
коммит произвёл GitHub
родитель 76328ecfde
Коммит 5775e9202b
+14 -8
Просмотреть файл
@@ -134,14 +134,20 @@ hipError_t IPCEvent::enqueueRecordCommand(hip::Stream* stream, amd::Command* com
ipc_evt_.ipc_shmem_->owners_device_id = deviceId();
command->enqueue();
// device writes 0 to signal after the hipEventRecord command is completed
// the signal value is checked by WaitThenDecrementSignal cb
hipError_t status = ihipStreamOperation(
reinterpret_cast<hipStream_t>(stream), ROCCLR_COMMAND_STREAM_WRITE_VALUE,
&(ipc_evt_.ipc_shmem_->signal[offset]), 0, 0, 0, sizeof(uint32_t));
if (status != hipSuccess) {
return status;
}
// Set event_ in order to release marked command when event is destroyed
if (event_ != nullptr) {
event_->release();
}
event_ = &command->event();
// device writes 0 to signal after the hipEventRecord command is completed
// the signal value is checked by WaitThenDecrementSignal cb
hipError_t status =
ihipStreamOperation(reinterpret_cast<hipStream_t>(stream), ROCCLR_COMMAND_STREAM_WRITE_VALUE,
&(ipc_evt_.ipc_shmem_->signal[offset]), 0, 0, 0, sizeof(uint32_t));
if (status != hipSuccess) {
return status;
}
// Update read index to indicate new signal.
int expected = write_index - 1;