update

Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>

Author: ShengchenJ

clang/lib/DPCT/RulesLang/APINamesMisc.inc

@@ -148,23 +148,25 @@ ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
                            CALL(MapNames::getDpctNamespace() +
                                     "experimental::get_mem_ipc_handle",
                                 ARG(1), ARG(0))),
-        UNSUPPORT_FACTORY_ENTRY("cudaIpcGetMemHandle",
-                                Diagnostics::TRY_EXPERIMENTAL_FEATURE,
-                                ARG("cudaIpcGetMemHandle"),
-                                ARG("--use-experimental-features=level_zero")))))
+        UNSUPPORT_FACTORY_ENTRY(
+            "cudaIpcGetMemHandle", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+            ARG("cudaIpcGetMemHandle"),
+            ARG("--use-experimental-features=level_zero")))))
 
 ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
     HelperFeatureEnum::device_ext,
     CONDITIONAL_FACTORY_ENTRY(
         clang::dpct::UseExtL0IPC,
         CALL_FACTORY_ENTRY("cudaIpcCloseMemHandle",
-                           CALL(MapNames::getDpctNamespace() +
-                                    "experimental::close_mem_ipc_handle",
+                           CALL("zeMemCloseIpcHandle",
+                               CALL("sycl::get_native<sycl::backend::ext_oneapi_level_zero>",  MEMBER_CALL(CALL(MapNames::getDpctNamespace() +
+                                                 "get_current_device"),
+                                            false, "get_context")),
                                 ARG(0))),
-        UNSUPPORT_FACTORY_ENTRY("cudaIpcCloseMemHandle",
-                                Diagnostics::TRY_EXPERIMENTAL_FEATURE,
-                                ARG("cudaIpcCloseMemHandle"),
-                                ARG("--use-experimental-features=level_zero")))))
+        UNSUPPORT_FACTORY_ENTRY(
+            "cudaIpcCloseMemHandle", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+            ARG("cudaIpcCloseMemHandle"),
+            ARG("--use-experimental-features=level_zero")))))
 
 ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
     HelperFeatureEnum::device_ext,

clang/runtime/dpct-rt/include/dpct/ze_utils.hpp

@@ -1,4 +1,5 @@
-//==---- ze_util.hpp ---------------------------------*- C++ -*----------------==//
+//==---- ze_utils.hpp ---------------------------------*- C++
+//-*----------------==//
 //
 // Copyright (C) Intel Corporation
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
@@ -11,12 +12,13 @@
 
 #ifdef ONEAPI_BACKEND_LEVEL_ZERO_EXT
 #if defined(__linux__)
-#include "level_zero/ze_api.h"
-#include "sycl/ext/oneapi/backend/level_zero.hpp"
+#include <level_zero/ze_api.h>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
 #include <sycl/sycl.hpp>
 namespace dpct {
 namespace experimental {
 
+namespace detail {
 ///  System call number definitions for kernel compatibility.
 ///  SYS_pidfd_open: Process file descriptor opener (requires kernel 5.6+).
 ///  SYS_pidfd_getfd: Cross-process FD fetcher system call.
@@ -28,56 +30,46 @@ namespace experimental {
 #define SYS_pidfd_getfd 438
 #endif
 
-/// Process id and IPC memory handle structure for cross-process sharing.
+/// Covert remote fd to the local fd through IPC handle extension.
+/// \param [in] ipc_ext_handle The extension of the IPC handle
+/// \returns Local process file descriptor
+template <class T> int convert_fd_from_handle(T ipc_ext_handle) {
+  int pidfd = syscall(SYS_pidfd_open, ipc_ext_handle.pid, 0);
+  if (pidfd < 0)
+    return -1;
+  return syscall(SYS_pidfd_getfd, pidfd, *(int *)ipc_ext_handle.handle.data, 0);
+}
+
+} // namespace detail
+
 struct ipc_mem_handle_ext_t {
   pid_t pid;
   ze_ipc_mem_handle_t handle;
 };
 
 /// Acquires IPC handle for shared memory region.
 /// \param [in] ptr Pointer to shared memory region
-/// \param [out] phipc Output IPC handle
+/// \param [out] handle_ptr Output IPC handle
 /// \returns Level Zero operation status code
-ze_result_t get_mem_ipc_handle(const void *ptr,
-                               ipc_mem_handle_ext_t *ipc_ext_handle) {
-  ipc_ext_handle->pid = getpid();
+inline ze_result_t get_mem_ipc_handle(const void *ptr,
+                                      ipc_mem_handle_ext_t *handle_ptr) {
+  handle_ptr->pid = getpid();
   return zeMemGetIpcHandle(
       sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
           dpct::get_current_device().get_context()),
-      ptr, &ipc_ext_handle->handle);
-}
-
-/// Releases resources associated with IPC handle.
-/// \param [in] ptr Pointer to shared memory region
-/// \returns Level Zero operation status code
-ze_result_t close_mem_ipc_handle(const void *ptr) {
-  if (ptr == nullptr) {
-    return ZE_RESULT_SUCCESS;
-  }
-  return zeMemCloseIpcHandle(
-      sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
-          dpct::get_current_device().get_context()),
-      (char *)ptr);
-}
-
-/// Covert remote fd to the local fd through the IPC handle extension.
-/// \param [in] ipc_ext_handle The extension of the IPC handle
-/// \returns Local process file descriptor
-template <class T> int convert_fd_pidfd_from_handle(T ipc_ext_handle) {
-  int pidfd = syscall(SYS_pidfd_open, ipc_ext_handle.pid, 0);
-  int fd;
-  memcpy(&fd, (void *)&ipc_ext_handle.handle.data, sizeof(int));
-  return syscall(SYS_pidfd_getfd, pidfd, fd, 0);
+      ptr, &handle_ptr->handle);
 }
 
 /// Maps remote IPC memory to local address space.
 /// \param [in] ipc_ext_handle The extension of the IPC handle
 /// \param [out] ptr Mapped memory pointer in local process
 /// \returns Level Zero operation status code
-ze_result_t open_mem_ipc_handle(ipc_mem_handle_ext_t ipc_ext_handle,
-                                void **ptr) {
-  int newfd = convert_fd_pidfd_from_handle(ipc_ext_handle);
-  memcpy(&ipc_ext_handle.handle.data, &newfd, sizeof(newfd));
+inline ze_result_t open_mem_ipc_handle(ipc_mem_handle_ext_t ipc_ext_handle,
+                                       void **ptr) {
+  int newfd = detail::convert_fd_from_handle(ipc_ext_handle);
+  if (newfd < 0)
+    throw std::runtime_error("Cannot convert fd from handle");
+  *((int *)ipc_ext_handle.handle.data) = newfd;
   return zeMemOpenIpcHandle(
       sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
           dpct::get_current_device().get_context()),

clang/test/dpct/IPC/share_mem_exp_option.cu

@@ -150,7 +150,7 @@ int parentProcess(char *app)
     }
   }
   std::cout << "verified Pass.\n";
-  // CHECK: dpct::experimental::close_mem_ipc_handle(ptr);
+  // CHECK: zeMemCloseIpcHandle(sycl::get_native<sycl::backend::ext_oneapi_level_zero>(dpct::get_current_device().get_context()), ptr);
   cudaIpcCloseMemHandle(ptr);
   return 0;
 }