[SYCLomaitc] Add the IPC memory APIs migration to level zero APIs.

clang/include/clang/DPCT/DPCTOptions.inc

@@ -826,6 +826,11 @@ DPCT_ENUM_OPTION(
             "prefetch", int(ExperimentalFeatures::Exp_Prefetch),
             "Experimental extension that allows use of SYCL prefetch APIs.\n",
             false),
+        DPCT_OPTION_ENUM_VALUE(
+            "level_zero", int(ExperimentalFeatures::Exp_LevelZero),
+            "Experimental migration feature that enables the use of Level Zero "
+            "APIs to migrate target code, like CUDA Inter-Process "
+            "Communication (IPC) APIs.\n", false),
         DPCT_OPTION_ENUM_VALUE(
             "all", int(ExperimentalFeatures::Exp_All),
             "Enable all experimental extensions listed in this option.\n",

clang/lib/DPCT/AnalysisInfo.cpp

@@ -905,6 +905,9 @@ void DpctFileInfo::insertHeader(HeaderType Type, unsigned Offset,
     // Add the label for profiling macro "DPCT_PROFILING_ENABLED", which will be
     // replaced by "#define DPCT_PROFILING_ENABLED" or not in the post
     // replacement.
+    if (DpctGlobalInfo::useExtLevelZero())
+      OS << "#define ONEAPI_BACKEND_LEVEL_ZERO_EXT" << getNL();
+
     OS << "{{NEEDREPLACEP0}}";
 
     if (DpctGlobalInfo::getUsmLevel() == UsmLevel::UL_None)

clang/lib/DPCT/AnalysisInfo.h

@@ -1347,6 +1347,9 @@ class DpctGlobalInfo {
     return getUsingExperimental<
         ExperimentalFeatures::Exp_NonStandardSYCLBuiltins>();
   }
+  static bool useExtLevelZero() {
+    return getUsingExperimental<ExperimentalFeatures::Exp_LevelZero>();
+  }
   static bool useExtPrefetch() {
     return getUsingExperimental<ExperimentalFeatures::Exp_Prefetch>();
   }

clang/lib/DPCT/CMakeLists.txt

@@ -21,6 +21,7 @@ set(RUNTIME_HEADERS
   ${CMAKE_SOURCE_DIR}/../clang/runtime/dpct-rt/include/dpct/lib_common_utils.hpp
   ${CMAKE_SOURCE_DIR}/../clang/runtime/dpct-rt/include/dpct/ccl_utils.hpp
   ${CMAKE_SOURCE_DIR}/../clang/runtime/dpct-rt/include/dpct/sparse_utils.hpp
+  ${CMAKE_SOURCE_DIR}/../clang/runtime/dpct-rt/include/dpct/ze_utils.hpp
   ${CMAKE_SOURCE_DIR}/../clang/runtime/dpct-rt/include/dpct/fft_utils.hpp
   ${CMAKE_SOURCE_DIR}/../clang/runtime/dpct-rt/include/dpct/lapack_utils.hpp
   ${CMAKE_SOURCE_DIR}/../clang/runtime/dpct-rt/include/dpct/group_utils.hpp
@@ -71,6 +72,7 @@ set(PROCESS_FILES_OUTPUT
   ${CMAKE_BINARY_DIR}/tools/clang/include/clang/DPCT/rng_utils.hpp.inc
   ${CMAKE_BINARY_DIR}/tools/clang/include/clang/DPCT/lib_common_utils.hpp.inc
   ${CMAKE_BINARY_DIR}/tools/clang/include/clang/DPCT/ccl_utils.hpp.inc
+  ${CMAKE_BINARY_DIR}/tools/clang/include/clang/DPCT/ze_utils.hpp.inc
   ${CMAKE_BINARY_DIR}/tools/clang/include/clang/DPCT/sparse_utils.hpp.inc
   ${CMAKE_BINARY_DIR}/tools/clang/include/clang/DPCT/fft_utils.hpp.inc
   ${CMAKE_BINARY_DIR}/tools/clang/include/clang/DPCT/lapack_utils.hpp.inc

clang/lib/DPCT/CommandOption/ValidateArguments.h

@@ -101,6 +101,7 @@ enum class ExperimentalFeatures : unsigned int {
   Exp_ExperimentalFeaturesEnumSize,
   Exp_NonStandardSYCLBuiltins,
   Exp_Prefetch,
+  Exp_LevelZero,
   Exp_All
 };
 enum class HelperFuncPreference : unsigned int { NoQueueDevice = 0 };

clang/lib/DPCT/DPCT.cpp

@@ -1060,6 +1060,8 @@ int runDPCT(int argc, const char **argv) {
             Experimentals.addValue(ExperimentalFeatures::Exp_Graph);
           else if (Option.ends_with("prefetch"))
             Experimentals.addValue(ExperimentalFeatures::Exp_Prefetch);
+          else if (Option.ends_with("level_zero"))
+            Experimentals.addValue(ExperimentalFeatures::Exp_LevelZero);
         } else if (Option == "--no-dry-pattern") {
           NoDRYPattern.setValue(true);
         }

clang/lib/DPCT/FileGenerator/GenHelperFunction.cpp

@@ -150,6 +150,10 @@ const std::string CodePinSerializationBasicAllContentStr =
 #include "clang/DPCT/codepin/serialization/basic.hpp.inc"
     ;
 
+const std::string ZEUtilsAllContentStr =
+#include "clang/DPCT/ze_utils.hpp.inc"
+    ;
+
 const std::string CmakeAllContentStr =
 #include "clang/DPCT/dpct.cmake.inc"
     ;
@@ -206,6 +210,7 @@ void genHelperFunction(const clang::tooling::UnifiedPath &OutRoot) {
   GENERATE_ALL_FILE_CONTENT(Util, ".", util.hpp)
   GENERATE_ALL_FILE_CONTENT(RngUtils, ".", rng_utils.hpp)
   GENERATE_ALL_FILE_CONTENT(LibCommonUtils, ".", lib_common_utils.hpp)
+  GENERATE_ALL_FILE_CONTENT(ZEUtils, ".", ze_utils.hpp)
   GENERATE_ALL_FILE_CONTENT(CclUtils, ".", ccl_utils.hpp)
   GENERATE_ALL_FILE_CONTENT(SparseUtils, ".", sparse_utils.hpp)
   GENERATE_ALL_FILE_CONTENT(FftUtils, ".", fft_utils.hpp)

clang/lib/DPCT/FileGenerator/GenHelperFunction.h

@@ -45,6 +45,7 @@ extern const std::string MathDetailAllContentStr;
 extern const std::string MemoryDetailAllContentStr;
 extern const std::string CodePinAllContentStr;
 extern const std::string CodePinSerializationBasicAllContentStr;
+extern const std::string ZEUtilsAllContentStr;
 extern const std::string DpctCmakeHelperFileStr;
 void genHelperFunction(const clang::tooling::UnifiedPath &OutRoot);
 void genCmakeHelperFunction(const clang::tooling::UnifiedPath &OutRoot);

clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h

@@ -1825,6 +1825,10 @@ inline auto UseExtBindlessImages = [](const CallExpr *C) -> bool {
   return DpctGlobalInfo::useExtBindlessImages();
 };
 
+inline auto UseExtLevelZero = [](const CallExpr *C) -> bool {
+  return DpctGlobalInfo::useExtLevelZero();
+};
+
 inline auto UseExtGraph = [](const CallExpr *C) -> bool {
   return DpctGlobalInfo::useExtGraph();
 };

clang/lib/DPCT/RuleInfra/MapNames.cpp

@@ -363,6 +363,9 @@ void MapNames::setExplicitNamespaceMap(
       {"CUstreamCallback",
        std::make_shared<TypeNameRule>(getDpctNamespace() + "queue_callback",
                                       HelperFeatureEnum::device_ext)},
+      {"cudaIpcMemHandle_t",
+       std::make_shared<TypeNameRule>(getDpctNamespace() +
+                                      "experimental::ipc_mem_handle_ext_t")},
       {"char1", std::make_shared<TypeNameRule>("int8_t")},
       {"char2", std::make_shared<TypeNameRule>(getClNamespace() + "char2")},
       {"char3", std::make_shared<TypeNameRule>(getClNamespace() + "char3")},

clang/lib/DPCT/RulesLang/APINamesMisc.inc

@@ -139,3 +139,44 @@ REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuFuncSetAttribute",
 
 REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuGetExportTable",
                          getRemovedAPIWarningMessage("cuGetExportTable"))
+
+ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+    HelperFeatureEnum::device_ext,
+    CONDITIONAL_FACTORY_ENTRY(
+        clang::dpct::UseExtLevelZero,
+        CALL_FACTORY_ENTRY("cudaIpcGetMemHandle",
+                           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")))))
+
+ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+    HelperFeatureEnum::device_ext,
+    CONDITIONAL_FACTORY_ENTRY(
+        clang::dpct::UseExtLevelZero,
+        CALL_FACTORY_ENTRY("cudaIpcCloseMemHandle",
+                           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")))))
+
+ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
+    HelperFeatureEnum::device_ext,
+    CONDITIONAL_FACTORY_ENTRY(
+        clang::dpct::UseExtLevelZero,
+        CALL_FACTORY_ENTRY("cudaIpcOpenMemHandle",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::open_mem_ipc_handle",
+                                ARG(1), ARG(0))),
+        UNSUPPORT_FACTORY_ENTRY("cudaIpcOpenMemHandle",
+                                Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                                ARG("cudaIpcOpenMemHandle"),
+                                ARG("--use-experimental-features=level_zero")))))

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -348,7 +348,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType",
               "cudaExternalSemaphoreHandleType", "CUstreamCallback",
               "cudaHostFn_t", "__nv_half2", "__nv_half", "cudaGraphNodeType",
-              "CUsurfref", "CUdevice_P2PAttribute"))))))
+              "CUsurfref", "CUdevice_P2PAttribute", "cudaIpcMemHandle_t"))))))
           .bind("cudaTypeDef"),
       this);
 
@@ -928,6 +928,13 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {
                "cudaGraphNodeType", "--use-experimental-features=graph");
       }
     }
+    if (CanonicalTypeStr == "cudaIpcMemHandle_st") {
+      if (!DpctGlobalInfo::useExtLevelZero()) {
+        report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
+               "cudaIpcMemHandle_t", "--use-experimental-features=level_zero");
+        return;
+      }
+    }
 
     if (CanonicalTypeStr == "cudaGraphExecUpdateResult") {
       report(TL->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
@@ -2650,12 +2657,6 @@ void FunctionCallRule::runRule(const MatchFinder::MatchResult &Result) {
              MapNames::ITFName.at(FuncName), Msg);
       emplaceTransformation(new ReplaceStmt(CE, ""));
     }
-  } else if (FuncName == "cudaIpcGetEventHandle" ||
-             FuncName == "cudaIpcOpenEventHandle" ||
-             FuncName == "cudaIpcGetMemHandle" ||
-             FuncName == "cudaIpcOpenMemHandle" ||
-             FuncName == "cudaIpcCloseMemHandle") {
-    report(CE->getBeginLoc(), Diagnostics::IPC_NOT_SUPPORTED, false);
   } else if (FuncName == "__trap") {
     if (DpctGlobalInfo::useAssert()) {
       emplaceTransformation(new ReplaceStmt(CE, "assert(0)"));

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -65,11 +65,11 @@ ENTRY(cudaGetDeviceFlags, cudaGetDeviceFlags, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGetDeviceProperties, cudaGetDeviceProperties, true, NO_FLAG, P0, "Successful")
 ENTRY(cudaGetDeviceProperties_v2 , cudaGetDeviceProperties_v2, true, NO_FLAG, P0, "Successful")
 ENTRY(cudaInitDevice, cudaInitDevice, false, NO_FLAG, P4, "comment")
-ENTRY(cudaIpcCloseMemHandle, cudaIpcCloseMemHandle, true, NO_FLAG, P0, "DPCT1030")
-ENTRY(cudaIpcGetEventHandle, cudaIpcGetEventHandle, true, NO_FLAG, P0, "DPCT1030")
-ENTRY(cudaIpcGetMemHandle, cudaIpcGetMemHandle, true, NO_FLAG, P0, "DPCT1030")
-ENTRY(cudaIpcOpenEventHandle, cudaIpcOpenEventHandle, true, NO_FLAG, P0, "DPCT1030")
-ENTRY(cudaIpcOpenMemHandle, cudaIpcOpenMemHandle, true, NO_FLAG, P0, "DPCT1030")
+ENTRY(cudaIpcCloseMemHandle, cudaIpcCloseMemHandle, true, NO_FLAG, P0, "DPCT1119")
+ENTRY(cudaIpcGetEventHandle, cudaIpcGetEventHandle, false, NO_FLAG, P0, "DPCT1030")
+ENTRY(cudaIpcGetMemHandle, cudaIpcGetMemHandle, true, NO_FLAG, P0, "DPCT1119")
+ENTRY(cudaIpcOpenEventHandle, cudaIpcOpenEventHandle, false, NO_FLAG, P0, "DPCT1030")
+ENTRY(cudaIpcOpenMemHandle, cudaIpcOpenMemHandle, true, NO_FLAG, P0, "DPCT1119")
 ENTRY(cudaSetDevice, cudaSetDevice, true, NO_FLAG, P0, "DPCT1093")
 ENTRY(cudaSetDeviceFlags, cudaSetDeviceFlags, true, NO_FLAG, P0, "DPCT1026/DPCT1027")
 ENTRY(cudaSetValidDevices, cudaSetValidDevices, false, NO_FLAG, P4, "comment")

clang/runtime/dpct-rt/CMakeLists.txt

@@ -12,6 +12,7 @@ set(dpct_rt_files
   include/dpct/dnnl_utils.hpp
   include/dpct/dpl_utils.hpp
   include/dpct/rng_utils.hpp
+  include/dpct/ze_utils.hpp
   include/dpct/lib_common_utils.hpp
   include/dpct/ccl_utils.hpp
   include/dpct/sparse_utils.hpp

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

@@ -58,7 +58,7 @@ template <int Arg> class dpct_kernel_scalar;
 
 #include "bindless_images.hpp"
 #include "graph.hpp"
-
+#include "ze_utils.hpp"
 #define USE_DPCT_HELPER 1
 
 #if defined(_MSC_VER)

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

@@ -0,0 +1,85 @@
+//==---- ze_utils.hpp ---------------------------------*- C++
+//-*----------------==//
+//
+// Copyright (C) Intel Corporation
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// See https://llvm.org/LICENSE.txt for license information.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __DPCT_ZE_UTILS_HPP__
+#define __DPCT_ZE_UTILS_HPP__
+
+#ifdef ONEAPI_BACKEND_LEVEL_ZERO_EXT
+#if defined(__linux__)
+#include <level_zero/ze_api.h>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+#include <sycl/sycl.hpp>
+namespace dpct {
+namespace experimental {
+
+struct ipc_mem_handle_ext_t {
+  pid_t pid;
+  ze_ipc_mem_handle_t handle;
+};
+
+namespace detail {
+
+#ifndef _SYS_pidfd_open
+#define _SYS_pidfd_open 434 // syscall number for pidfd_open
+#endif
+
+#ifndef _SYS_pidfd_getfd
+#define _SYS_pidfd_getfd 438 // syscall number for pidfd_getfd
+#endif
+
+inline int get_fd_of_peer_process(ipc_mem_handle_ext_t ext_handle) {
+  int pidfd = syscall(_SYS_pidfd_open, ext_handle.pid,
+                      0); // obtain a file descriptor that refers to a
+                          // process(requires kernel 5.6+).
+  if (pidfd < 0)
+    return -1;
+  return syscall(_SYS_pidfd_getfd, pidfd, *(int *)ext_handle.handle.data,
+                 0); // obtain a duplicate of another process's file
+                     // descriptor(requires kernel 5.6+).
+}
+
+} // namespace detail
+
+/// Creates an IPC memory handle for the specified allocation.
+/// \param [in] ptr Pointer to the device memory allocation
+/// \param [out] ext_handle_ptr IPC memory handle extension
+inline ze_result_t get_mem_ipc_handle(const void *ptr,
+                                      ipc_mem_handle_ext_t *ext_handle_ptr) {
+  ext_handle_ptr->pid = getpid();
+  return zeMemGetIpcHandle(
+      sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
+          dpct::get_current_device().get_context()),
+      ptr, &ext_handle_ptr->handle);
+}
+
+/// Opens an IPC memory handle to retrieve a device pointer.
+/// \param [in] ext_handle IPC memory handle extension
+/// \param [out] pptr Pointer to device allocation in this process
+inline ze_result_t open_mem_ipc_handle(ipc_mem_handle_ext_t ext_handle,
+                                       void **pptr) {
+  int fd = detail::get_fd_of_peer_process(ext_handle);
+  if (fd < 0)
+    throw std::runtime_error("Cannot get file descriptor of peer process.");
+  *((int *)ext_handle.handle.data) = fd;
+
+  return zeMemOpenIpcHandle(
+      sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
+          dpct::get_current_device().get_context()),
+      sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
+          (sycl::device)dpct::get_current_device()),
+      ext_handle.handle, 0u, pptr);
+}
+
+} // namespace experimental
+} // namespace dpct
+
+#endif // __linux__
+#endif // ONEAPI_BACKEND_LEVEL_ZERO_EXT
+
+#endif // ! __DPCT_ZE_UTILS_HPP__