From 9bb6fb0b79872e12f9945cf5351ae49496f02243 Mon Sep 17 00:00:00 2001
From: "Kudchadker, Saleel" <Saleel.Kudchadker@amd.com>
Date: Fri, 15 Aug 2025 15:32:12 -0700
Subject: [PATCH] SWDEV-537480 - Fix file handle cleanup (#523)

[ROCm/clr commit: b90f3c250719f53cd87427347a4c52dd0529b492]
---
 projects/clr/hipamd/src/hip_event.hpp     | 10 ++++++++++
 projects/clr/hipamd/src/hip_event_ipc.cpp | 17 ++++++++---------
 2 files changed, 18 insertions(+), 9 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_event.hpp b/projects/clr/hipamd/src/hip_event.hpp
index eed64a8701..75b870b605 100644
--- a/projects/clr/hipamd/src/hip_event.hpp
+++ b/projects/clr/hipamd/src/hip_event.hpp
@@ -24,6 +24,10 @@
 #include "hip_internal.hpp"
 #include "thread/monitor.hpp"
 
+#if !defined(_MSC_VER)
+#include <sys/mman.h>
+#endif
+
 // Internal structure for stream callback handler
 namespace hip {
 class StreamCallback {
@@ -206,6 +210,12 @@ class IPCEvent : public Event {
         amd::Os::shm_unlink(ipc_evt_.ipc_name_);
       }
     }
+#if !defined(_MSC_VER)
+    // Clean up the POSIX shared memory object
+    if (!ipc_evt_.ipc_name_.empty() && ipc_evt_.ipc_name_ != "dummy") {
+      shm_unlink(ipc_evt_.ipc_name_.c_str());
+    }
+#endif
   }
   IPCEvent() : Event(hipEventInterprocess) {}
   bool createIpcEventShmemIfNeeded();
diff --git a/projects/clr/hipamd/src/hip_event_ipc.cpp b/projects/clr/hipamd/src/hip_event_ipc.cpp
index 38f259e488..9114f17c57 100644
--- a/projects/clr/hipamd/src/hip_event_ipc.cpp
+++ b/projects/clr/hipamd/src/hip_event_ipc.cpp
@@ -38,15 +38,16 @@ bool IPCEvent::createIpcEventShmemIfNeeded() {
     return true;
   }
 
-  char name_template[] = "/tmp/eventXXXXXX";
 #if !defined(_MSC_VER)
-  int temp_fd = mkstemp(name_template);
+  static std::atomic<int> counter{0};
+  ipc_evt_.ipc_name_ = "/hip_" + std::to_string(getpid()) + "_" + std::to_string(counter++);
 #else
+  char name_template[] = "/hip_XXXXXX";
   _mktemp_s(name_template, sizeof(name_template));
-#endif
-
   ipc_evt_.ipc_name_ = name_template;
   ipc_evt_.ipc_name_.replace(0, 5, "/hip_");
+#endif
+
   if (!amd::Os::MemoryMapFileTruncated(
           ipc_evt_.ipc_name_.c_str(),
           const_cast<const void**>(reinterpret_cast<void**>(&(ipc_evt_.ipc_shmem_))),
@@ -54,10 +55,6 @@ bool IPCEvent::createIpcEventShmemIfNeeded() {
     return false;
   }
 
-#if !defined(_MSC_VER)
-  close(temp_fd);
-#endif
-
   ipc_evt_.ipc_shmem_->owners = 1;
   ipc_evt_.ipc_shmem_->read_index = -1;
   ipc_evt_.ipc_shmem_->write_index = 0;
@@ -145,6 +142,7 @@ hipError_t IPCEvent::enqueueRecordCommand(hip::Stream* stream, amd::Command* com
   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;
   }
@@ -154,7 +152,7 @@ hipError_t IPCEvent::enqueueRecordCommand(hip::Stream* stream, amd::Command* com
   while (!ipc_evt_.ipc_shmem_->read_index.compare_exchange_weak(expected, write_index)) {
     amd::Os::sleep(1);
   }
-  
+
   return hipSuccess;
 }
 
@@ -219,6 +217,7 @@ hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle)
 
   hip::Event* e = reinterpret_cast<hip::Event*>(*event);
   ihipIpcEventHandle_t* iHandle = reinterpret_cast<ihipIpcEventHandle_t*>(&handle);
+
   status = e->OpenHandle(iHandle);
   // Free the event in case of failure
   if (status != hipSuccess) {