Skip to content

Commit f0df89a

Browse files
AidanBeltonSAidanBeltonS
authored
[SYCL][CUDA] Add experimental context and device interop (#6202)
This PR is adds part of the CUDA-backend spec interop proposed in KhronosGroup/SYCL-Docs#197. The changes work with the CUDA CTS interop checks KhronosGroup/SYCL-CTS#336. This PR just adds the context and device interop. Further PRs will be submitted to fill in other missing interop functionality. The changes are experimental so as to not break oneMKL, and oneDNN which use CUDA interops. The experimental interop is enabled by defining the SYCL_EXT_ONEAPI_BACKEND_CUDA_EXPERIMENTAL macro and including the header file sycl/ext/oneapi/experimental/backend/cuda.hpp. llvm-test-suite: intel/llvm-test-suite#1041
1 parent 58c9d3a commit f0df89a

File tree

11 files changed

+434
-26
lines changed

11 files changed

+434
-26
lines changed

‎sycl/include/CL/sycl/backend.hpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,12 @@
1818
#include <CL/sycl/detail/backend_traits_opencl.hpp>
1919
#endif
2020
#if SYCL_EXT_ONEAPI_BACKEND_CUDA
21+
#ifdef SYCL_EXT_ONEAPI_BACKEND_CUDA_EXPERIMENTAL
22+
#include <sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp>
23+
#else
2124
#include <CL/sycl/detail/backend_traits_cuda.hpp>
2225
#endif
26+
#endif
2327
#if SYCL_EXT_ONEAPI_BACKEND_HIP
2428
#include <CL/sycl/detail/backend_traits_hip.hpp>
2529
#endif

‎sycl/include/CL/sycl/context.hpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,8 @@ class platform;
2727
namespace detail {
2828
class context_impl;
2929
}
30+
template <backend Backend, class SyclT>
31+
auto get_native(const SyclT &Obj) -> backend_return_t<Backend, SyclT>;
3032

3133
/// The context class represents a SYCL context on which kernel functions may
3234
/// be executed.
@@ -230,6 +232,10 @@ class __SYCL_EXPORT context {
230232
pi_native_handle getNative() const;
231233

232234
std::shared_ptr<detail::context_impl> impl;
235+
236+
template <backend Backend, class SyclT>
237+
friend auto get_native(const SyclT &Obj) -> backend_return_t<Backend, SyclT>;
238+
233239
template <class Obj>
234240
friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
235241

‎sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp

Lines changed: 143 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,143 @@
1+
//===------- backend_traits_cuda.hpp - Backend traits for CUDA ---*-C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// This file defines the specializations of the sycl::detail::interop,
10+
// sycl::detail::BackendInput and sycl::detail::BackendReturn class templates
11+
// for the CUDA backend but there is no sycl::detail::InteropFeatureSupportMap
12+
// specialization for the CUDA backend.
13+
//===----------------------------------------------------------------------===//
14+
15+
#pragma once
16+
17+
#include <CL/sycl/accessor.hpp>
18+
#include <CL/sycl/context.hpp>
19+
#include <CL/sycl/detail/backend_traits.hpp>
20+
#include <CL/sycl/device.hpp>
21+
#include <CL/sycl/event.hpp>
22+
#include <CL/sycl/kernel_bundle.hpp>
23+
#include <CL/sycl/queue.hpp>
24+
25+
#include <vector>
26+
27+
typedef int CUdevice;
28+
typedef struct CUctx_st *CUcontext;
29+
typedef struct CUstream_st *CUstream;
30+
typedef struct CUevent_st *CUevent;
31+
typedef struct CUmod_st *CUmodule;
32+
33+
__SYCL_INLINE_NAMESPACE(cl) {
34+
namespace sycl {
35+
namespace detail {
36+
37+
// TODO the interops for context, device, event, platform and program
38+
// may be removed after removing the deprecated 'get_native()' methods
39+
// from the corresponding classes. The interop<backend, queue> specialization
40+
// is also used in the get_queue() method of the deprecated class
41+
// interop_handler and also can be removed after API cleanup.
42+
template <> struct interop<backend::ext_oneapi_cuda, context> {
43+
using type = CUcontext;
44+
};
45+
46+
template <> struct interop<backend::ext_oneapi_cuda, device> {
47+
using type = CUdevice;
48+
};
49+
50+
template <> struct interop<backend::ext_oneapi_cuda, event> {
51+
using type = CUevent;
52+
};
53+
54+
template <> struct interop<backend::ext_oneapi_cuda, queue> {
55+
using type = CUstream;
56+
};
57+
58+
template <> struct interop<backend::ext_oneapi_cuda, platform> {
59+
using type = std::vector<CUdevice>;
60+
};
61+
62+
#ifdef __SYCL_INTERNAL_API
63+
template <> struct interop<backend::ext_oneapi_cuda, program> {
64+
using type = CUmodule;
65+
};
66+
#endif
67+
68+
template <typename DataT, int Dimensions, typename AllocatorT>
69+
struct BackendInput<backend::ext_oneapi_cuda,
70+
buffer<DataT, Dimensions, AllocatorT>> {
71+
using type = DataT *;
72+
};
73+
74+
template <typename DataT, int Dimensions, typename AllocatorT>
75+
struct BackendReturn<backend::ext_oneapi_cuda,
76+
buffer<DataT, Dimensions, AllocatorT>> {
77+
using type = DataT *;
78+
};
79+
80+
template <> struct BackendInput<backend::ext_oneapi_cuda, context> {
81+
using type = CUcontext;
82+
};
83+
84+
template <> struct BackendReturn<backend::ext_oneapi_cuda, context> {
85+
using type = std::vector<CUcontext>;
86+
};
87+
88+
template <> struct BackendInput<backend::ext_oneapi_cuda, device> {
89+
using type = CUdevice;
90+
};
91+
92+
template <> struct BackendReturn<backend::ext_oneapi_cuda, device> {
93+
using type = CUdevice;
94+
};
95+
96+
template <> struct BackendInput<backend::ext_oneapi_cuda, event> {
97+
using type = CUevent;
98+
};
99+
100+
template <> struct BackendReturn<backend::ext_oneapi_cuda, event> {
101+
using type = CUevent;
102+
};
103+
104+
template <> struct BackendInput<backend::ext_oneapi_cuda, queue> {
105+
using type = CUstream;
106+
};
107+
108+
template <> struct BackendReturn<backend::ext_oneapi_cuda, queue> {
109+
using type = CUstream;
110+
};
111+
112+
template <> struct BackendInput<backend::ext_oneapi_cuda, platform> {
113+
using type = std::vector<CUdevice>;
114+
};
115+
116+
template <> struct BackendReturn<backend::ext_oneapi_cuda, platform> {
117+
using type = std::vector<CUdevice>;
118+
};
119+
120+
#ifdef __SYCL_INTERNAL_API
121+
template <> struct BackendInput<backend::ext_oneapi_cuda, program> {
122+
using type = CUmodule;
123+
};
124+
125+
template <> struct BackendReturn<backend::ext_oneapi_cuda, program> {
126+
using type = CUmodule;
127+
};
128+
#endif
129+
130+
template <> struct InteropFeatureSupportMap<backend::ext_oneapi_cuda> {
131+
static constexpr bool MakePlatform = false;
132+
static constexpr bool MakeDevice = true;
133+
static constexpr bool MakeContext = true;
134+
static constexpr bool MakeQueue = false;
135+
static constexpr bool MakeEvent = false;
136+
static constexpr bool MakeBuffer = false;
137+
static constexpr bool MakeKernel = false;
138+
static constexpr bool MakeKernelBundle = false;
139+
};
140+
141+
} // namespace detail
142+
} // namespace sycl
143+
} // __SYCL_INLINE_NAMESPACE(cl)

‎sycl/include/sycl/ext/oneapi/experimental/backend/cuda.hpp

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
//==--------- cuda.hpp - SYCL CUDA backend ---------------------------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
#include <CL/sycl/backend.hpp>
12+
#include <CL/sycl/context.hpp>
13+
14+
#include <vector>
15+
16+
__SYCL_INLINE_NAMESPACE(cl) {
17+
namespace sycl {
18+
namespace ext {
19+
namespace oneapi {
20+
namespace cuda {
21+
22+
// Implementation of ext_oneapi_cuda::make<device>
23+
inline __SYCL_EXPORT device make_device(pi_native_handle NativeHandle) {
24+
return sycl::detail::make_device(NativeHandle, backend::ext_oneapi_cuda);
25+
}
26+
27+
} // namespace cuda
28+
} // namespace oneapi
29+
} // namespace ext
30+
31+
// CUDA context specialization
32+
template <>
33+
inline auto get_native<backend::ext_oneapi_cuda, context>(const context &C)
34+
-> backend_return_t<backend::ext_oneapi_cuda, context> {
35+
// create a vector to be returned
36+
backend_return_t<backend::ext_oneapi_cuda, context> ret;
37+
38+
// get the native CUDA context from the SYCL object
39+
auto native = reinterpret_cast<
40+
backend_return_t<backend::ext_oneapi_cuda, context>::value_type>(
41+
C.getNative());
42+
ret.push_back(native);
43+
44+
return ret;
45+
}
46+
47+
// Specialisation of non-free context get_native
48+
template <>
49+
inline backend_return_t<backend::ext_oneapi_cuda, context>
50+
context::get_native<backend::ext_oneapi_cuda>() const {
51+
return sycl::get_native<backend::ext_oneapi_cuda, context>(*this);
52+
}
53+
54+
// Specialisation of interop_handles get_native_context
55+
template <>
56+
inline backend_return_t<backend::ext_oneapi_cuda, context>
57+
interop_handle::get_native_context<backend::ext_oneapi_cuda>() const {
58+
#ifndef __SYCL_DEVICE_ONLY__
59+
return std::vector{reinterpret_cast<CUcontext>(getNativeContext())};
60+
#else
61+
// we believe this won't be ever called on device side
62+
return {};
63+
#endif
64+
}
65+
66+
// CUDA device specialization
67+
template <>
68+
inline device make_device<backend::ext_oneapi_cuda>(
69+
const backend_input_t<backend::ext_oneapi_cuda, device> &BackendObject) {
70+
pi_native_handle NativeHandle = static_cast<pi_native_handle>(BackendObject);
71+
return ext::oneapi::cuda::make_device(NativeHandle);
72+
}
73+
74+
} // namespace sycl
75+
} // __SYCL_INLINE_NAMESPACE(cl)

0 commit comments

Comments
 (0)