update

ShengchenJ · ShengchenJ · commit dc8dc57e2cf0 · 2025-04-25T07:15:49.000-07:00
Signed-off-by: Chen, Sheng S &lt;sheng.s.chen@intel.com&gt;
diff --git a/clang/include/clang/DPCT/DPCTOptions.inc b/clang/include/clang/DPCT/DPCTOptions.inc
@@ -827,10 +827,10 @@ DPCT_ENUM_OPTION(
             "Experimental extension that allows use of SYCL prefetch APIs.\n",
             false),
         DPCT_OPTION_ENUM_VALUE(
-            "l0-ipc", int(ExperimentalFeatures::Exp_IPC),
+            "level_zero", int(ExperimentalFeatures::Exp_IPC),
             "Experimental extension that allows the use of Level Zero APIs "
-            "to migrate IPC code, with the DPCT_EXT_ONEAPI_BACKEND_LEVEL_ZERO "
-            "macro defined to enable these APIs.\n",
+            "to migrate the target code. For example, migrate the CUDA Inter "
+            "Process Communication(IPC) APIs. \n",
             false),
         DPCT_OPTION_ENUM_VALUE(
             "all", int(ExperimentalFeatures::Exp_All),
diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/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::useExtIPC())
+      OS << "#define ONEAPI_BACKEND_LEVEL_ZERO_EXT" << getNL();
+
     OS << "{{NEEDREPLACEP0}}";
 
     if (DpctGlobalInfo::getUsmLevel() == UsmLevel::UL_None)
diff --git a/clang/lib/DPCT/CMakeLists.txt b/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
diff --git a/clang/lib/DPCT/DPCT.cpp b/clang/lib/DPCT/DPCT.cpp
@@ -998,7 +998,7 @@ 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("l0-ipc"))
+          else if (Option.ends_with("level_zero"))
             Experimentals.addValue(ExperimentalFeatures::Exp_IPC);
         } else if (Option == "--no-dry-pattern") {
           NoDRYPattern.setValue(true);
diff --git a/clang/lib/DPCT/FileGenerator/GenHelperFunction.cpp b/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"
     ;
diff --git a/clang/lib/DPCT/FileGenerator/GenHelperFunction.h b/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);
diff --git a/clang/lib/DPCT/RuleInfra/MapNames.cpp b/clang/lib/DPCT/RuleInfra/MapNames.cpp
@@ -365,7 +365,7 @@ void MapNames::setExplicitNamespaceMap(
                                       HelperFeatureEnum::device_ext)},
       {"cudaIpcMemHandle_t",
        std::make_shared<TypeNameRule>(getDpctNamespace() +
-                                      "experimental::dpct_ipc_mem_handle_t")},
+                                      "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")},
diff --git a/clang/lib/DPCT/RulesLang/APINamesMisc.inc b/clang/lib/DPCT/RulesLang/APINamesMisc.inc
@@ -151,7 +151,7 @@ ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
         UNSUPPORT_FACTORY_ENTRY("cudaIpcGetMemHandle",
                                 Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                                 ARG("cudaIpcGetMemHandle"),
-                                ARG("--use-experimental-features=l0-ipc")))))
+                                ARG("--use-experimental-features=level_zero")))))
 
 ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
     HelperFeatureEnum::device_ext,
@@ -164,7 +164,7 @@ ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
         UNSUPPORT_FACTORY_ENTRY("cudaIpcCloseMemHandle",
                                 Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                                 ARG("cudaIpcCloseMemHandle"),
-                                ARG("--use-experimental-features=l0-ipc")))))
+                                ARG("--use-experimental-features=level_zero")))))
 
 ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
     HelperFeatureEnum::device_ext,
@@ -177,4 +177,4 @@ ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
         UNSUPPORT_FACTORY_ENTRY("cudaIpcOpenMemHandle",
                                 Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                                 ARG("cudaIpcOpenMemHandle"),
-                                ARG("--use-experimental-features=l0-ipc")))))
+                                ARG("--use-experimental-features=level_zero")))))
diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp
@@ -931,7 +931,7 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {
     if (CanonicalTypeStr == "cudaIpcMemHandle_st") {
       if (!DpctGlobalInfo::useExtIPC()) {
         report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
-               "cudaIpcMemHandle_t", "--use-experimental-features=l0-ipc");
+               "cudaIpcMemHandle_t", "--use-experimental-features=level_zero");
         return;
       }
     }
diff --git a/clang/runtime/dpct-rt/CMakeLists.txt b/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
diff --git a/clang/runtime/dpct-rt/include/dpct/dpct.hpp b/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)
diff --git a/clang/runtime/dpct-rt/include/dpct/util.hpp b/clang/runtime/dpct-rt/include/dpct/util.hpp
@@ -14,10 +14,7 @@
 #include <type_traits>
 #include <cassert>
 #include <cstdint>
-#if DPCT_EXT_ONEAPI_BACKEND_LEVEL_ZERO
-#include "level_zero/ze_api.h"
-#include "sycl/ext/oneapi/backend/level_zero.hpp"
-#endif
+
 // TODO: Remove these function definitions once they exist in the DPC++ compiler
 #if defined(__SYCL_DEVICE_ONLY__) && defined(__INTEL_LLVM_COMPILER)
 template <typename T>
@@ -1350,90 +1347,6 @@ inline uint32_t ternary_logic_op(uint32_t a, uint32_t b, uint32_t c,
   return result;
 }
 
-#ifdef DPCT_EXT_ONEAPI_BACKEND_LEVEL_ZERO
-#if defined(__linux__)
-namespace experimental {
-
-///  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.
-#ifndef SYS_pidfd_open
-#define SYS_pidfd_open 434
-#endif
-
-#ifndef SYS_pidfd_getfd
-#define SYS_pidfd_getfd 438
-#endif
-
-/// Process id and IPC memory handle structure for cross-process sharing.
-struct dpct_ipc_mem_handle_t {
-  pid_t pid;
-  ze_ipc_mem_handle_t handle;
-};
-
-/// Extracts native Level Zero context handle from SYCL context.
-/// \param [in] context SYCL context object
-/// \returns Native Level Zero context handle
-ze_context_handle_t get_ze_context(sycl::context context) {
-  return sycl::get_native<sycl::backend::ext_oneapi_level_zero>(context);
-}
-
-/// Retrieves native Level Zero device handle from SYCL device.
-/// \param [in] device SYCL device object
-/// \returns Native Level Zero device handle
-ze_device_handle_t get_ze_device(sycl::device device) {
-  return sycl::get_native<sycl::backend::ext_oneapi_level_zero>(device);
-}
-/// Acquires native Level Zero event handle from SYCL event.
-/// \param [in] event SYCL event object
-/// \returns Native Level Zero event handle
-ze_event_handle_t get_ze_event(sycl::event event) {
-  return sycl::get_native<sycl::backend::ext_oneapi_level_zero>(event);
-}
-
-/// Acquires IPC handle for shared memory region.
-/// \param [in] ptr Pointer to shared memory region
-/// \param [out] phipc Output IPC handle structure
-/// \returns Level Zero operation status code
-ze_result_t get_mem_ipc_handle(const void *ptr, dpct_ipc_mem_handle_t *phipc) {
-  phipc->pid = getpid();
-  return zeMemGetIpcHandle(
-      get_ze_context(dpct::get_current_device().get_context()), ptr,
-      &phipc->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) {
-  return zeMemCloseIpcHandle(
-      get_ze_context(dpct::get_current_device().get_context()), (char *)ptr);
-}
-/// Cross-process file descriptor translator.
-/// \param [in] phipc Source IPC handle structure
-/// \returns Local process file descriptor
-template <class T> int get_cur_pid(T phipc) {
-  int pidfd = syscall(SYS_pidfd_open, phipc.pid, 0);
-  int fd;
-  memcpy(&fd, (void *)&phipc.handle.data, sizeof(int));
-  int newfd = syscall(SYS_pidfd_getfd, pidfd, fd, 0);
-  return newfd;
-}
-/// Maps remote IPC memory to local address space.
-/// \param [in] hipc Source IPC handle structure
-/// \param [out] ptr Mapped memory pointer in local process
-/// \returns Level Zero operation status code
-ze_result_t open_mem_ipc_handle(dpct_ipc_mem_handle_t hipc, void **ptr) {
-  int newfd = get_cur_pid(hipc);
-  memcpy(&hipc.handle.data, &newfd, sizeof(newfd));
-  return zeMemOpenIpcHandle(
-      get_ze_context(dpct::get_current_device().get_context()),
-      get_ze_device(dpct::get_current_device()), hipc.handle, 0u, ptr);
-}
-
-} // namespace experimental
-#endif // __linux__
-#endif // DPCT_EXT_ONEAPI_BACKEND_LEVEL_ZERO
-
 #ifdef _WIN32
 #define DPCT_EXPORT __declspec(dllexport)
 #else
diff --git a/clang/runtime/dpct-rt/include/dpct/ze_utils.hpp b/clang/runtime/dpct-rt/include/dpct/ze_utils.hpp
@@ -0,0 +1,94 @@
+//==---- ze_util.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 __ZE_UTILS_HPP__
+#define __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 {
+
+///  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.
+#ifndef SYS_pidfd_open
+#define SYS_pidfd_open 434
+#endif
+
+#ifndef SYS_pidfd_getfd
+#define SYS_pidfd_getfd 438
+#endif
+
+/// Process id and IPC memory handle structure for cross-process sharing.
+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
+/// \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();
+  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);
+}
+
+/// 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));
+  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()),
+      ipc_ext_handle.handle, 0u, ptr);
+}
+
+} // namespace experimental
+#endif // __linux__
+#endif // ONEAPI_BACKEND_LEVEL_ZERO_EXT
+
+} // namespace dpct
+#endif // ! __ZE_UTILS_HPP__
diff --git a/clang/test/dpct/IPC/share_mem.cu b/clang/test/dpct/IPC/share_mem.cu
@@ -59,7 +59,7 @@ int sharedMemoryOpen(const char *name, size_t sz, sharedMemoryInfo *info) {
 
 typedef struct shmStruct_st {
   // CHECK: /*
-  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaIpcMemHandle_t is not supported, please try to remigrate with option: --use-experimental-features=l0-ipc.
+  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaIpcMemHandle_t is not supported, please try to remigrate with option: --use-experimental-features=level_zero.
   // CHECK-NEXT: */
   cudaIpcMemHandle_t memHandle;
 } shmStruct;
@@ -93,7 +93,7 @@ int childProcess(int id) {
   shm = (shmStruct *)info.addr;
   int *ptr;
   // CHECK: /*
-  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaIpcOpenMemHandle is not supported, please try to remigrate with option: --use-experimental-features=l0-ipc.
+  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaIpcOpenMemHandle is not supported, please try to remigrate with option: --use-experimental-features=level_zero.
   // CHECK-NEXT: */
   cudaIpcOpenMemHandle((void **)&ptr, shm->memHandle,
                        cudaIpcMemLazyEnablePeerAccess);
@@ -118,7 +118,7 @@ int parentProcess(char *app) {
   cudaMalloc(&ptr, DATA_SIZE);
   int *hostptr = (int *)malloc(DATA_SIZE);
   // CHECK: /*
-  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaIpcGetMemHandle is not supported, please try to remigrate with option: --use-experimental-features=l0-ipc.
+  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaIpcGetMemHandle is not supported, please try to remigrate with option: --use-experimental-features=level_zero.
   // CHECK-NEXT: */
   cudaIpcGetMemHandle(&shm->memHandle, ptr);
 
@@ -136,7 +136,7 @@ int parentProcess(char *app) {
   }
   std::cout << "verified Pass.\n";
   // CHECK: /*
-  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaIpcCloseMemHandle is not supported, please try to remigrate with option: --use-experimental-features=l0-ipc.
+  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaIpcCloseMemHandle is not supported, please try to remigrate with option: --use-experimental-features=level_zero.
   // CHECK-NEXT: */
   cudaIpcCloseMemHandle(ptr);
   return 0;
diff --git a/clang/test/dpct/IPC/share_mem_exp_option.cu b/clang/test/dpct/IPC/share_mem_exp_option.cu
@@ -1,7 +1,7 @@
 // UNSUPPORTED: system-windows
-// RUN: dpct --use-experimental-features=l0-ipc --format-range=none  -out-root %T/share_mem_exp_option %s --cuda-include-path="%cuda-path/include" --sycl-named-lambda -- -x cuda --cuda-host-only
+// RUN: dpct --use-experimental-features=level_zero --format-range=none  -out-root %T/share_mem_exp_option %s --cuda-include-path="%cuda-path/include" --sycl-named-lambda -- -x cuda --cuda-host-only
 // RUN: FileCheck --input-file %T/share_mem_exp_option/share_mem_exp_option.dp.cpp --match-full-lines %s
-// RUN: %if build_lit %{icpx -c -fsycl -DDPCT_EXT_ONEAPI_BACKEND_LEVEL_ZERO -DNO_BUILD_TEST  %T/share_mem_exp_option/share_mem_exp_option.dp.cpp -o %T/share_mem_exp_option/share_mem_exp_option.dp.o %}
+// RUN: %if build_lit %{icpx -c -fsycl  -DNO_BUILD_TEST  %T/share_mem_exp_option/share_mem_exp_option.dp.cpp -o %T/share_mem_exp_option/share_mem_exp_option.dp.o %}
 
 
 #include <cuda.h>
@@ -69,7 +69,7 @@ int sharedMemoryOpen(const char *name, size_t sz, sharedMemoryInfo *info)
 
 typedef struct shmStruct_st
 {
-  // CHECK: dpct::experimental::dpct_ipc_mem_handle_t memHandle;
+  // CHECK: dpct::experimental::ipc_mem_handle_ext_t memHandle;
   cudaIpcMemHandle_t memHandle;
 } shmStruct;
 
@@ -109,7 +109,7 @@ int childProcess(int id)
   }
   shm = (shmStruct *)info.addr;
   int *ptr;
-  // CHECK: dpct::experimental::open_mem_ipc_handle(*(dpct::experimental::dpct_ipc_mem_handle_t *)&shm->memHandle, (void **)&ptr);
+  // CHECK: dpct::experimental::open_mem_ipc_handle(*(dpct::experimental::ipc_mem_handle_ext_t *)&shm->memHandle, (void **)&ptr);
   cudaIpcOpenMemHandle((void **)&ptr, *(cudaIpcMemHandle_t *)&shm->memHandle,
                        cudaIpcMemLazyEnablePeerAccess);
 
@@ -132,7 +132,7 @@ int parentProcess(char *app)
   memset((void *)shm, 0, sizeof(*shm));
   cudaMalloc(&ptr, DATA_SIZE);
   int *hostptr = (int *)malloc(DATA_SIZE);
-  // CHECK: dpct::experimental::get_mem_ipc_handle(ptr, (dpct::experimental::dpct_ipc_mem_handle_t *)&shm->memHandle);
+  // CHECK: dpct::experimental::get_mem_ipc_handle(ptr, (dpct::experimental::ipc_mem_handle_ext_t *)&shm->memHandle);
   cudaIpcGetMemHandle((cudaIpcMemHandle_t *)&shm->memHandle, ptr);
 
   char *const args[] = {app, "0", NULL};
diff --git a/clang/test/dpct/check_header_files.cpp b/clang/test/dpct/check_header_files.cpp
diff --git a/clang/test/dpct/cuda-device-api.cu b/clang/test/dpct/cuda-device-api.cu
diff --git a/clang/test/dpct/help_option_check/lin/help_advanced.txt b/clang/test/dpct/help_option_check/lin/help_advanced.txt
diff --git a/clang/test/dpct/help_option_check/lin/help_all.txt b/clang/test/dpct/help_option_check/lin/help_all.txt
diff --git a/clang/test/dpct/help_option_check/win/help_advanced.txt b/clang/test/dpct/help_option_check/win/help_advanced.txt
diff --git a/clang/test/dpct/help_option_check/win/help_all.txt b/clang/test/dpct/help_option_check/win/help_all.txt

Original file line number	Diff line number	Diff line change
`@@ -150,6 +150,10 @@ const std::string CodePinSerializationBasicAllContentStr =`
`150`	`150`	`#include "clang/DPCT/codepin/serialization/basic.hpp.inc"`
`151`	`151`	`;`
`152`	`152`
	`153`	`+const std::string ZEUtilsAllContentStr =`
	`154`	`+#include "clang/DPCT/ze_utils.hpp.inc"`
	`155`	`+ ;`
	`156`	`+`
`153`	`157`	`const std::string CmakeAllContentStr =`
`154`	`158`	`#include "clang/DPCT/dpct.cmake.inc"`
`155`	`159`	`;`
Original file line number	Diff line number	Diff line change
`@@ -931,7 +931,7 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {`
`931`	`931`	`if (CanonicalTypeStr == "cudaIpcMemHandle_st") {`
`932`	`932`	`if (!DpctGlobalInfo::useExtIPC()) {`
`933`	`933`	`report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,`
`934`		`- "cudaIpcMemHandle_t", "--use-experimental-features=l0-ipc");`
	`934`	`+ "cudaIpcMemHandle_t", "--use-experimental-features=level_zero");`
`935`	`935`	`return;`
`936`	`936`	`}`
`937`	`937`	`}`