diff --git a/projects/clr/hipamd/src/hip_event_ipc.cpp b/projects/clr/hipamd/src/hip_event_ipc.cpp index 7ac8d63bee..38f259e488 100644 --- a/projects/clr/hipamd/src/hip_event_ipc.cpp +++ b/projects/clr/hipamd/src/hip_event_ipc.cpp @@ -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(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(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;