Skip to content

Commit 8922f41

Browse files
mfrancepilloismfrancepillois
committed
[SYCL][Graph] Implement exceptions for incompatible extensions (#276)
* [SYCL][Graph] Implement exceptions for incompatible extensions Throws an invalid exception when trying to use the following extensions along with Graph. - sycl_ext_oneapi_enqueue_barrier - sycl_ext_oneapi_memcpy2d - sycl_ext_codeplay_kernel_fusion - sycl_ext_oneapi_kernel_properties - sycl_ext_oneapi_device_global Closes Issue: #154 * [SYCL][Graph] Implement exceptions for incompatible extensions Adds info to exception message Moves tests from e2e to unitests when possible * [SYCL][Graph] Implement exceptions for incompatible extensions Corrects some typos and adds comments. * [SYCL][Graph] Implement exceptions for incompatible extensions Used a template function to throw exception instead of a parametrized function. * [SYCL][Graph] Implement exceptions for incompatible extensions Moves Sycl-extension enum definition. Limits graph recording to non-explicit path in the new tests. * [SYCL][Graph] Implement exceptions for incompatible extensions Updates Linux ABI dump file with the new handler function throwing exception.
1 parent 4b980d0 commit 8922f41

File tree

7 files changed

+620
-1
lines changed

7 files changed

+620
-1
lines changed

‎sycl/include/sycl/handler.hpp

Lines changed: 43 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -110,8 +110,19 @@ class pipe;
110110
}
111111

112112
namespace ext::oneapi::experimental::detail {
113+
// List of sycl experimental extensions
114+
// This enum is used to define the extension from which a function is called.
115+
// This is used in handler::throwIfGraphAssociated() to specify
116+
// the message of the thrown expection.
117+
enum SyclExtensions {
118+
sycl_ext_oneapi_kernel_properties,
119+
sycl_ext_oneapi_enqueue_barrier,
120+
sycl_ext_oneapi_memcpy2d,
121+
sycl_ext_oneapi_device_global
122+
};
123+
113124
class graph_impl;
114-
}
125+
} // namespace ext::oneapi::experimental::detail
115126
namespace detail {
116127

117128
class handler_impl;
@@ -2085,6 +2096,7 @@ class __SYCL_EXPORT handler {
20852096
std::enable_if_t<
20862097
ext::oneapi::experimental::is_property_list<PropertiesT>::value>
20872098
single_task(PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) {
2099+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
20882100
single_task_lambda_impl<KernelName, KernelType, PropertiesT>(Props,
20892101
KernelFunc);
20902102
}
@@ -2095,6 +2107,7 @@ class __SYCL_EXPORT handler {
20952107
ext::oneapi::experimental::is_property_list<PropertiesT>::value>
20962108
parallel_for(range<1> NumWorkItems, PropertiesT Props,
20972109
_KERNELFUNCPARAM(KernelFunc)) {
2110+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
20982111
parallel_for_lambda_impl<KernelName, KernelType, 1, PropertiesT>(
20992112
NumWorkItems, Props, std::move(KernelFunc));
21002113
}
@@ -2105,6 +2118,7 @@ class __SYCL_EXPORT handler {
21052118
ext::oneapi::experimental::is_property_list<PropertiesT>::value>
21062119
parallel_for(range<2> NumWorkItems, PropertiesT Props,
21072120
_KERNELFUNCPARAM(KernelFunc)) {
2121+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
21082122
parallel_for_lambda_impl<KernelName, KernelType, 2, PropertiesT>(
21092123
NumWorkItems, Props, std::move(KernelFunc));
21102124
}
@@ -2115,6 +2129,7 @@ class __SYCL_EXPORT handler {
21152129
ext::oneapi::experimental::is_property_list<PropertiesT>::value>
21162130
parallel_for(range<3> NumWorkItems, PropertiesT Props,
21172131
_KERNELFUNCPARAM(KernelFunc)) {
2132+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
21182133
parallel_for_lambda_impl<KernelName, KernelType, 3, PropertiesT>(
21192134
NumWorkItems, Props, std::move(KernelFunc));
21202135
}
@@ -2125,6 +2140,7 @@ class __SYCL_EXPORT handler {
21252140
ext::oneapi::experimental::is_property_list<PropertiesT>::value>
21262141
parallel_for(nd_range<Dims> Range, PropertiesT Properties,
21272142
_KERNELFUNCPARAM(KernelFunc)) {
2143+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
21282144
parallel_for_impl<KernelName>(Range, Properties, std::move(KernelFunc));
21292145
}
21302146

@@ -2137,6 +2153,7 @@ class __SYCL_EXPORT handler {
21372153
detail::AreAllButLastReductions<RestT...>::value &&
21382154
ext::oneapi::experimental::is_property_list<PropertiesT>::value>
21392155
parallel_for(range<1> Range, PropertiesT Properties, RestT &&...Rest) {
2156+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
21402157
detail::reduction_parallel_for<KernelName>(*this, Range, Properties,
21412158
std::forward<RestT>(Rest)...);
21422159
}
@@ -2148,6 +2165,7 @@ class __SYCL_EXPORT handler {
21482165
detail::AreAllButLastReductions<RestT...>::value &&
21492166
ext::oneapi::experimental::is_property_list<PropertiesT>::value>
21502167
parallel_for(range<2> Range, PropertiesT Properties, RestT &&...Rest) {
2168+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
21512169
detail::reduction_parallel_for<KernelName>(*this, Range, Properties,
21522170
std::forward<RestT>(Rest)...);
21532171
}
@@ -2159,6 +2177,7 @@ class __SYCL_EXPORT handler {
21592177
detail::AreAllButLastReductions<RestT...>::value &&
21602178
ext::oneapi::experimental::is_property_list<PropertiesT>::value>
21612179
parallel_for(range<3> Range, PropertiesT Properties, RestT &&...Rest) {
2180+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
21622181
detail::reduction_parallel_for<KernelName>(*this, Range, Properties,
21632182
std::forward<RestT>(Rest)...);
21642183
}
@@ -2213,6 +2232,7 @@ class __SYCL_EXPORT handler {
22132232
int Dims, typename PropertiesT>
22142233
void parallel_for_work_group(range<Dims> NumWorkGroups, PropertiesT Props,
22152234
_KERNELFUNCPARAM(KernelFunc)) {
2235+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
22162236
parallel_for_work_group_lambda_impl<KernelName, KernelType, Dims,
22172237
PropertiesT>(NumWorkGroups, Props,
22182238
KernelFunc);
@@ -2223,6 +2243,7 @@ class __SYCL_EXPORT handler {
22232243
void parallel_for_work_group(range<Dims> NumWorkGroups,
22242244
range<Dims> WorkGroupSize, PropertiesT Props,
22252245
_KERNELFUNCPARAM(KernelFunc)) {
2246+
throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
22262247
parallel_for_work_group_lambda_impl<KernelName, KernelType, Dims,
22272248
PropertiesT>(
22282249
NumWorkGroups, WorkGroupSize, Props, KernelFunc);
@@ -2530,6 +2551,8 @@ class __SYCL_EXPORT handler {
25302551
/// until all commands previously submitted to this queue have entered the
25312552
/// complete state.
25322553
void ext_oneapi_barrier() {
2554+
throwIfGraphAssociated<ext::oneapi::experimental::detail::SyclExtensions::
2555+
sycl_ext_oneapi_enqueue_barrier>();
25332556
throwIfActionIsCreated();
25342557
setType(detail::CG::Barrier);
25352558
}
@@ -2615,6 +2638,8 @@ class __SYCL_EXPORT handler {
26152638
typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
26162639
void ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src,
26172640
size_t SrcPitch, size_t Width, size_t Height) {
2641+
throwIfGraphAssociated<ext::oneapi::experimental::detail::SyclExtensions::
2642+
sycl_ext_oneapi_memcpy2d>();
26182643
throwIfActionIsCreated();
26192644
if (Width > DestPitch)
26202645
throw sycl::exception(sycl::make_error_code(errc::invalid),
@@ -2793,6 +2818,8 @@ class __SYCL_EXPORT handler {
27932818
void memcpy(ext::oneapi::experimental::device_global<T, PropertyListT> &Dest,
27942819
const void *Src, size_t NumBytes = sizeof(T),
27952820
size_t DestOffset = 0) {
2821+
throwIfGraphAssociated<ext::oneapi::experimental::detail::SyclExtensions::
2822+
sycl_ext_oneapi_device_global>();
27962823
if (sizeof(T) < DestOffset + NumBytes)
27972824
throw sycl::exception(make_error_code(errc::invalid),
27982825
"Copy to device_global is out of bounds.");
@@ -2825,6 +2852,8 @@ class __SYCL_EXPORT handler {
28252852
memcpy(void *Dest,
28262853
const ext::oneapi::experimental::device_global<T, PropertyListT> &Src,
28272854
size_t NumBytes = sizeof(T), size_t SrcOffset = 0) {
2855+
throwIfGraphAssociated<ext::oneapi::experimental::detail::SyclExtensions::
2856+
sycl_ext_oneapi_device_global>();
28282857
if (sizeof(T) < SrcOffset + NumBytes)
28292858
throw sycl::exception(make_error_code(errc::invalid),
28302859
"Copy from device_global is out of bounds.");
@@ -3346,8 +3375,21 @@ class __SYCL_EXPORT handler {
33463375
"handler::require() before it can be used.");
33473376
}
33483377

3378+
template <typename PropertiesT>
3379+
std::enable_if_t<
3380+
ext::oneapi::experimental::is_property_list<PropertiesT>::value>
3381+
throwIfGraphAssociatedAndKernelProperties() {
3382+
if (!std::is_same_v<PropertiesT,
3383+
ext::oneapi::experimental::detail::empty_properties_t>)
3384+
throwIfGraphAssociated<ext::oneapi::experimental::detail::SyclExtensions::
3385+
sycl_ext_oneapi_kernel_properties>();
3386+
}
3387+
33493388
// Set value of the gpu cache configuration for the kernel.
33503389
void setKernelCacheConfig(sycl::detail::pi::PiKernelCacheConfig);
3390+
3391+
template <ext::oneapi::experimental::detail::SyclExtensions ExtensionT>
3392+
void throwIfGraphAssociated();
33513393
};
33523394
} // namespace _V1
33533395
} // namespace sycl

‎sycl/source/detail/fusion/fusion_wrapper_impl.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,11 @@ bool fusion_wrapper_impl::is_in_fusion_mode() const {
2727
}
2828

2929
void fusion_wrapper_impl::start_fusion() {
30+
if (MQueue->getCommandGraph()) {
31+
throw sycl::exception(sycl::make_error_code(errc::invalid),
32+
"SYCL kernel fusion can NOT be started "
33+
"on a queue that is in a recording state.");
34+
}
3035
detail::Scheduler::getInstance().startFusion(MQueue);
3136
}
3237

‎sycl/source/detail/graph_impl.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -647,6 +647,12 @@ modifiable_command_graph::finalize(const sycl::property_list &) const {
647647
bool modifiable_command_graph::begin_recording(queue &RecordingQueue) {
648648
auto QueueImpl = sycl::detail::getSyclObjImpl(RecordingQueue);
649649

650+
if (QueueImpl->is_in_fusion_mode()) {
651+
throw sycl::exception(sycl::make_error_code(errc::invalid),
652+
"SYCL queue in kernel in fusion mode "
653+
"can NOT be recorded.");
654+
}
655+
650656
if (QueueImpl->get_context() != impl->getContext()) {
651657
throw sycl::exception(sycl::make_error_code(errc::invalid),
652658
"begin_recording called for a queue whose context "

‎sycl/source/handler.cpp

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -794,6 +794,8 @@ void handler::verifyUsedKernelBundle(const std::string &KernelName) {
794794
}
795795

796796
void handler::ext_oneapi_barrier(const std::vector<event> &WaitList) {
797+
throwIfGraphAssociated<ext::oneapi::experimental::detail::SyclExtensions::
798+
sycl_ext_oneapi_enqueue_barrier>();
797799
throwIfActionIsCreated();
798800
MCGType = detail::CG::BarrierWaitlist;
799801
MEventsWaitWithBarrier.resize(WaitList.size());
@@ -1338,5 +1340,49 @@ handler::getCommandGraph() const {
13381340
return MQueue->getCommandGraph();
13391341
}
13401342

1343+
template void handler::throwIfGraphAssociated<
1344+
ext::oneapi::experimental::detail::SyclExtensions::
1345+
sycl_ext_oneapi_kernel_properties>();
1346+
template void handler::throwIfGraphAssociated<
1347+
ext::oneapi::experimental::detail::SyclExtensions::
1348+
sycl_ext_oneapi_enqueue_barrier>();
1349+
template void
1350+
handler::throwIfGraphAssociated<ext::oneapi::experimental::detail::
1351+
SyclExtensions::sycl_ext_oneapi_memcpy2d>();
1352+
template void handler::throwIfGraphAssociated<
1353+
ext::oneapi::experimental::detail::SyclExtensions::
1354+
sycl_ext_oneapi_device_global>();
1355+
1356+
template <ext::oneapi::experimental::detail::SyclExtensions ExtensionT>
1357+
void handler::throwIfGraphAssociated() {
1358+
std::string ExceptionMsg = "";
1359+
1360+
if constexpr (ExtensionT ==
1361+
ext::oneapi::experimental::detail::SyclExtensions::
1362+
sycl_ext_oneapi_kernel_properties) {
1363+
ExceptionMsg = "sycl_ext_oneapi_kernel_properties";
1364+
}
1365+
if constexpr (ExtensionT ==
1366+
ext::oneapi::experimental::detail::SyclExtensions::
1367+
sycl_ext_oneapi_enqueue_barrier) {
1368+
ExceptionMsg = "sycl_ext_oneapi_enqueue_barrier";
1369+
}
1370+
if constexpr (ExtensionT == ext::oneapi::experimental::detail::
1371+
SyclExtensions::sycl_ext_oneapi_memcpy2d) {
1372+
ExceptionMsg = "sycl_ext_oneapi_memcpy2d";
1373+
}
1374+
if constexpr (ExtensionT ==
1375+
ext::oneapi::experimental::detail::SyclExtensions::
1376+
sycl_ext_oneapi_device_global) {
1377+
ExceptionMsg = "sycl_ext_oneapi_device_global";
1378+
}
1379+
1380+
if (MGraph || MQueue->getCommandGraph()) {
1381+
throw sycl::exception(sycl::make_error_code(errc::invalid),
1382+
"The feature " + ExceptionMsg +
1383+
" is not yet available "
1384+
"along with SYCL Graph extension.");
1385+
}
1386+
}
13411387
} // namespace _V1
13421388
} // namespace sycl

‎sycl/test-e2e/Graph/graph_exception_global_device_extension.cpp

Lines changed: 150 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,150 @@
1+
// REQUIRES: level_zero, gpu
2+
// RUN: %{build} -o %t.out
3+
// RUN: %{run} %t.out
4+
//
5+
// The test checks that invalid exception is thrown
6+
// when trying to use sycl_ext_oneapi_device_global
7+
// along with Graph.
8+
9+
#include "graph_common.hpp"
10+
11+
using TestProperties = decltype(sycl::ext::oneapi::experimental::properties{});
12+
13+
sycl::ext::oneapi::experimental::device_global<int, TestProperties>
14+
MemcpyDeviceGlobal;
15+
sycl::ext::oneapi::experimental::device_global<int, TestProperties>
16+
CopyDeviceGlobal;
17+
18+
enum OperationPath { Explicit, RecordReplay, Shortcut };
19+
20+
template <OperationPath PathKind> void test() {
21+
queue Q;
22+
int MemcpyWrite = 42, CopyWrite = 24, MemcpyRead = 1, CopyRead = 2;
23+
24+
exp_ext::command_graph Graph{Q.get_context(), Q.get_device()};
25+
26+
if constexpr (PathKind != OperationPath::Explicit) {
27+
Graph.begin_recording(Q);
28+
}
29+
30+
// Copy from device globals before having written anything.
31+
std::error_code ExceptionCode = make_error_code(sycl::errc::success);
32+
try {
33+
if constexpr (PathKind == OperationPath::Shortcut) {
34+
Q.memcpy(&MemcpyRead, MemcpyDeviceGlobal);
35+
}
36+
if constexpr (PathKind == OperationPath::RecordReplay) {
37+
Q.submit([&](handler &CGH) {
38+
return CGH.memcpy(&MemcpyRead, MemcpyDeviceGlobal);
39+
});
40+
}
41+
if constexpr (PathKind == OperationPath::Explicit) {
42+
Graph.add([&](handler &CGH) {
43+
return CGH.memcpy(&MemcpyRead, MemcpyDeviceGlobal);
44+
});
45+
}
46+
} catch (exception &Exception) {
47+
ExceptionCode = Exception.code();
48+
}
49+
assert(ExceptionCode == sycl::errc::invalid);
50+
51+
ExceptionCode = make_error_code(sycl::errc::success);
52+
try {
53+
if constexpr (PathKind == OperationPath::Shortcut) {
54+
Q.copy(CopyDeviceGlobal, &CopyRead);
55+
}
56+
if constexpr (PathKind == OperationPath::RecordReplay) {
57+
Q.submit(
58+
[&](handler &CGH) { return CGH.copy(CopyDeviceGlobal, &CopyRead); });
59+
}
60+
if constexpr (PathKind == OperationPath::Explicit) {
61+
Graph.add(
62+
[&](handler &CGH) { return CGH.copy(CopyDeviceGlobal, &CopyRead); });
63+
}
64+
} catch (exception &Exception) {
65+
ExceptionCode = Exception.code();
66+
}
67+
assert(ExceptionCode == sycl::errc::invalid);
68+
69+
// Write to device globals and then read their values.
70+
ExceptionCode = make_error_code(sycl::errc::success);
71+
try {
72+
if constexpr (PathKind == OperationPath::Shortcut) {
73+
Q.memcpy(MemcpyDeviceGlobal, &MemcpyWrite);
74+
}
75+
if constexpr (PathKind == OperationPath::RecordReplay) {
76+
Q.submit([&](handler &CGH) {
77+
return CGH.memcpy(MemcpyDeviceGlobal, &MemcpyWrite);
78+
});
79+
}
80+
if constexpr (PathKind == OperationPath::Explicit) {
81+
Graph.add([&](handler &CGH) {
82+
return CGH.memcpy(MemcpyDeviceGlobal, &MemcpyWrite);
83+
});
84+
}
85+
} catch (exception &Exception) {
86+
ExceptionCode = Exception.code();
87+
}
88+
assert(ExceptionCode == sycl::errc::invalid);
89+
90+
ExceptionCode = make_error_code(sycl::errc::success);
91+
try {
92+
if constexpr (PathKind == OperationPath::Shortcut) {
93+
Q.copy(&CopyWrite, CopyDeviceGlobal);
94+
} else if constexpr (PathKind == OperationPath::RecordReplay) {
95+
Q.submit(
96+
[&](handler &CGH) { return CGH.copy(&CopyWrite, CopyDeviceGlobal); });
97+
} else if constexpr (PathKind == OperationPath::Explicit) {
98+
Graph.add(
99+
[&](handler &CGH) { return CGH.copy(&CopyWrite, CopyDeviceGlobal); });
100+
}
101+
} catch (exception &Exception) {
102+
ExceptionCode = Exception.code();
103+
}
104+
assert(ExceptionCode == sycl::errc::invalid);
105+
106+
ExceptionCode = make_error_code(sycl::errc::success);
107+
try {
108+
if constexpr (PathKind == OperationPath::Shortcut) {
109+
Q.memcpy(&MemcpyRead, MemcpyDeviceGlobal);
110+
} else if constexpr (PathKind == OperationPath::RecordReplay) {
111+
Q.submit([&](handler &CGH) {
112+
return CGH.memcpy(&MemcpyRead, MemcpyDeviceGlobal);
113+
});
114+
} else if constexpr (PathKind == OperationPath::Explicit) {
115+
Graph.add([&](handler &CGH) {
116+
return CGH.memcpy(&MemcpyRead, MemcpyDeviceGlobal);
117+
});
118+
}
119+
} catch (exception &Exception) {
120+
ExceptionCode = Exception.code();
121+
}
122+
assert(ExceptionCode == sycl::errc::invalid);
123+
124+
ExceptionCode = make_error_code(sycl::errc::success);
125+
try {
126+
if constexpr (PathKind == OperationPath::Shortcut) {
127+
Q.copy(CopyDeviceGlobal, &CopyRead);
128+
} else if constexpr (PathKind == OperationPath::RecordReplay) {
129+
Q.submit(
130+
[&](handler &CGH) { return CGH.copy(CopyDeviceGlobal, &CopyRead); });
131+
} else if constexpr (PathKind == OperationPath::Explicit) {
132+
Graph.add(
133+
[&](handler &CGH) { return CGH.copy(CopyDeviceGlobal, &CopyRead); });
134+
}
135+
} catch (exception &Exception) {
136+
ExceptionCode = Exception.code();
137+
}
138+
assert(ExceptionCode == sycl::errc::invalid);
139+
140+
if constexpr (PathKind != OperationPath::Explicit) {
141+
Graph.end_recording();
142+
}
143+
}
144+
145+
int main() {
146+
test<OperationPath::Explicit>();
147+
test<OperationPath::RecordReplay>();
148+
test<OperationPath::Shortcut>();
149+
return 0;
150+
}

0 commit comments

Comments
 (0)