Skip to content

Commit 1cff757

Browse files
ShengchenJShengchenJ
committed
update
Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>
1 parent 48014aa commit 1cff757

9 files changed

Lines changed: 57 additions & 48 deletions

File tree

‎clang/include/clang/DPCT/DPCTOptions.inc‎

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -827,11 +827,10 @@ DPCT_ENUM_OPTION(
827827
"Experimental extension that allows use of SYCL prefetch APIs.\n",
828828
false),
829829
DPCT_OPTION_ENUM_VALUE(
830-
"level_zero", int(ExperimentalFeatures::Exp_IPC),
831-
"Experimental extension that allows the use of Level Zero APIs "
832-
"to migrate the target code. For example, migrate the CUDA Inter "
833-
"Process Communication(IPC) APIs. \n",
834-
false),
830+
"level_zero", int(ExperimentalFeatures::Exp_LevelZero),
831+
"Experimental migration feature that enables the use of Level Zero "
832+
"APIs to migrate target code, like CUDA Inter-Process "
833+
"Communication (IPC) APIs.\n", false),
835834
DPCT_OPTION_ENUM_VALUE(
836835
"all", int(ExperimentalFeatures::Exp_All),
837836
"Enable all experimental extensions listed in this option.\n",

‎clang/lib/DPCT/AnalysisInfo.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -905,7 +905,7 @@ void DpctFileInfo::insertHeader(HeaderType Type, unsigned Offset,
905905
// Add the label for profiling macro "DPCT_PROFILING_ENABLED", which will be
906906
// replaced by "#define DPCT_PROFILING_ENABLED" or not in the post
907907
// replacement.
908-
if (DpctGlobalInfo::useExtIPC())
908+
if (DpctGlobalInfo::useExtLevelZero())
909909
OS << "#define ONEAPI_BACKEND_LEVEL_ZERO_EXT" << getNL();
910910

911911
OS << "{{NEEDREPLACEP0}}";

‎clang/lib/DPCT/AnalysisInfo.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1347,8 +1347,8 @@ class DpctGlobalInfo {
13471347
return getUsingExperimental<
13481348
ExperimentalFeatures::Exp_NonStandardSYCLBuiltins>();
13491349
}
1350-
static bool useExtIPC() {
1351-
return getUsingExperimental<ExperimentalFeatures::Exp_IPC>();
1350+
static bool useExtLevelZero() {
1351+
return getUsingExperimental<ExperimentalFeatures::Exp_LevelZero>();
13521352
}
13531353
static bool useExtPrefetch() {
13541354
return getUsingExperimental<ExperimentalFeatures::Exp_Prefetch>();

‎clang/lib/DPCT/CommandOption/ValidateArguments.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -101,7 +101,7 @@ enum class ExperimentalFeatures : unsigned int {
101101
Exp_ExperimentalFeaturesEnumSize,
102102
Exp_NonStandardSYCLBuiltins,
103103
Exp_Prefetch,
104-
Exp_IPC,
104+
Exp_LevelZero,
105105
Exp_All
106106
};
107107
enum class HelperFuncPreference : unsigned int { NoQueueDevice = 0 };

‎clang/lib/DPCT/DPCT.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1061,7 +1061,7 @@ int runDPCT(int argc, const char **argv) {
10611061
else if (Option.ends_with("prefetch"))
10621062
Experimentals.addValue(ExperimentalFeatures::Exp_Prefetch);
10631063
else if (Option.ends_with("level_zero"))
1064-
Experimentals.addValue(ExperimentalFeatures::Exp_IPC);
1064+
Experimentals.addValue(ExperimentalFeatures::Exp_LevelZero);
10651065
} else if (Option == "--no-dry-pattern") {
10661066
NoDRYPattern.setValue(true);
10671067
}

‎clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1825,8 +1825,8 @@ inline auto UseExtBindlessImages = [](const CallExpr *C) -> bool {
18251825
return DpctGlobalInfo::useExtBindlessImages();
18261826
};
18271827

1828-
inline auto UseExtL0IPC = [](const CallExpr *C) -> bool {
1829-
return DpctGlobalInfo::useExtIPC();
1828+
inline auto UseExtLevelZero = [](const CallExpr *C) -> bool {
1829+
return DpctGlobalInfo::useExtLevelZero();
18301830
};
18311831

18321832
inline auto UseExtGraph = [](const CallExpr *C) -> bool {

‎clang/lib/DPCT/RulesLang/APINamesMisc.inc‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -143,7 +143,7 @@ REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuGetExportTable",
143143
ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
144144
HelperFeatureEnum::device_ext,
145145
CONDITIONAL_FACTORY_ENTRY(
146-
clang::dpct::UseExtL0IPC,
146+
clang::dpct::UseExtLevelZero,
147147
CALL_FACTORY_ENTRY("cudaIpcGetMemHandle",
148148
CALL(MapNames::getDpctNamespace() +
149149
"experimental::get_mem_ipc_handle",
@@ -156,7 +156,7 @@ ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
156156
ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
157157
HelperFeatureEnum::device_ext,
158158
CONDITIONAL_FACTORY_ENTRY(
159-
clang::dpct::UseExtL0IPC,
159+
clang::dpct::UseExtLevelZero,
160160
CALL_FACTORY_ENTRY("cudaIpcCloseMemHandle",
161161
CALL("zeMemCloseIpcHandle",
162162
CALL("sycl::get_native<sycl::backend::ext_oneapi_level_zero>", MEMBER_CALL(CALL(MapNames::getDpctNamespace() +
@@ -171,7 +171,7 @@ ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
171171
ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
172172
HelperFeatureEnum::device_ext,
173173
CONDITIONAL_FACTORY_ENTRY(
174-
clang::dpct::UseExtL0IPC,
174+
clang::dpct::UseExtLevelZero,
175175
CALL_FACTORY_ENTRY("cudaIpcOpenMemHandle",
176176
CALL(MapNames::getDpctNamespace() +
177177
"experimental::open_mem_ipc_handle",

‎clang/lib/DPCT/RulesLang/RulesLang.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -929,7 +929,7 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {
929929
}
930930
}
931931
if (CanonicalTypeStr == "cudaIpcMemHandle_st") {
932-
if (!DpctGlobalInfo::useExtIPC()) {
932+
if (!DpctGlobalInfo::useExtLevelZero()) {
933933
report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
934934
"cudaIpcMemHandle_t", "--use-experimental-features=level_zero");
935935
return;

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

Lines changed: 42 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -20,16 +20,24 @@ namespace experimental {
2020

2121
namespace detail {
2222

23-
/// Covert remote fd to the local fd through IPC handle extension.
24-
/// \param [in] ipc_ext_handle The extension of the IPC handle
25-
/// \returns Local process file descriptor
26-
template <class T> int convert_fd_from_handle(T ipc_ext_handle) {
27-
int pidfd = syscall(434, ipc_ext_handle.pid,
23+
#ifndef _SYS_pidfd_open
24+
#define _SYS_pidfd_open 434 // syscall number for pidfd_open
25+
#endif
26+
27+
#ifndef _SYS_pidfd_getfd
28+
#define _SYS_pidfd_getfd 438 // syscall number for pidfd_getfd
29+
#endif
30+
31+
/// Obtain a duplicate of another process's file descriptor.
32+
/// \param [in] ext_handle IPC memory handle extension
33+
/// \returns obtained file descriptor
34+
template <class T> int get_fd_of_peer_process(T ext_handle) {
35+
int pidfd = syscall(_SYS_pidfd_open, ext_handle.pid,
2836
0); // obtain a file descriptor that refers to a
2937
// process(requires kernel 5.6+).
3038
if (pidfd < 0)
3139
return -1;
32-
return syscall(438, pidfd, *(int *)ipc_ext_handle.handle.data,
40+
return syscall(_SYS_pidfd_getfd, pidfd, *(int *)ext_handle.handle.data,
3341
0); // obtain a duplicate of another process's file
3442
// descriptor(requires kernel 5.6+).
3543
}
@@ -41,35 +49,37 @@ struct ipc_mem_handle_ext_t {
4149
ze_ipc_mem_handle_t handle;
4250
};
4351

44-
/// Acquires IPC handle for shared memory region.
45-
/// \param [in] ptr Pointer to shared memory region
46-
/// \param [out] handle_ptr Output IPC handle
47-
/// \returns Level Zero operation status code
48-
inline ze_result_t get_mem_ipc_handle(const void *ptr,
49-
ipc_mem_handle_ext_t *handle_ptr) {
50-
handle_ptr->pid = getpid();
51-
return zeMemGetIpcHandle(
52-
sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
53-
dpct::get_current_device().get_context()),
54-
ptr, &handle_ptr->handle);
52+
/// Creates an IPC memory handle for the specified allocation.
53+
/// \param [in] ptr Pointer to the device memory allocation
54+
/// \param [out] ext_handle_ptr IPC memory handle extension
55+
inline void get_mem_ipc_handle(const void *ptr,
56+
ipc_mem_handle_ext_t *ext_handle_ptr) {
57+
ext_handle_ptr->pid = getpid();
58+
auto ret =
59+
zeMemGetIpcHandle(sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
60+
dpct::get_current_device().get_context()),
61+
ptr, &ext_handle_ptr->handle);
62+
if (ret != ZE_RESULT_SUCCESS)
63+
throw std::runtime_error("The zeMemGetIpcHandle execution failed.");
5564
}
5665

57-
/// Maps remote IPC memory to local address space.
58-
/// \param [in] ipc_ext_handle The extension of the IPC handle
59-
/// \param [out] ptr Mapped memory pointer in local process
66+
/// Opens an IPC memory handle to retrieve a device pointer.
67+
/// \param [in] ext_handle IPC memory handle extension
68+
/// \param [out] pptr Pointer to device allocation in this process
6069
/// \returns Level Zero operation status code
61-
inline ze_result_t open_mem_ipc_handle(ipc_mem_handle_ext_t ipc_ext_handle,
62-
void **ptr) {
63-
int newfd = detail::convert_fd_from_handle(ipc_ext_handle);
64-
if (newfd < 0)
65-
throw std::runtime_error("Cannot convert fd from handle");
66-
*((int *)ipc_ext_handle.handle.data) = newfd;
67-
return zeMemOpenIpcHandle(
68-
sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
69-
dpct::get_current_device().get_context()),
70-
sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
71-
(sycl::device)dpct::get_current_device()),
72-
ipc_ext_handle.handle, 0u, ptr);
70+
inline void open_mem_ipc_handle(ipc_mem_handle_ext_t ext_handle, void **pptr) {
71+
int fd = detail::get_fd_of_peer_process(ext_handle);
72+
if (fd < 0)
73+
throw std::runtime_error("Cannot get file descriptor of peer process.");
74+
*((int *)ext_handle.handle.data) = fd;
75+
auto ret =
76+
zeMemOpenIpcHandle(sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
77+
dpct::get_current_device().get_context()),
78+
sycl::get_native<sycl::backend::ext_oneapi_level_zero>(
79+
(sycl::device)dpct::get_current_device()),
80+
ext_handle.handle, 0u, pptr);
81+
if (ret != ZE_RESULT_SUCCESS)
82+
throw std::runtime_error("The zeMemOpenIpcHandle execution failed.");
7383
}
7484

7585
} // namespace experimental

0 commit comments

Comments
 (0)