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

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

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<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;