Skip to content

Commit bf827a2

Browse files
author
Maxime France-Pillois
authored
[SYCL][Graph] Throw exception when using bindless images extension in a graph (#294)
* [SYCL][Graph] Throw exception when using bindless images extension in a graph Throws an invalid exception when using bindless images extension in a graph. Adds Unitests to test the exception throwing.
1 parent 04941fd commit bf827a2

File tree

3 files changed

+222
-3
lines changed

3 files changed

+222
-3
lines changed

‎sycl/include/sycl/ext/oneapi/experimental/graph.hpp‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -37,7 +37,8 @@ enum class UnsupportedGraphFeatures {
3737
sycl_ext_oneapi_kernel_properties,
3838
sycl_ext_oneapi_enqueue_barrier,
3939
sycl_ext_oneapi_memcpy2d,
40-
sycl_ext_oneapi_device_global
40+
sycl_ext_oneapi_device_global,
41+
sycl_ext_oneapi_bindless_images
4142
};
4243

4344
constexpr const char *
@@ -58,6 +59,8 @@ UnsupportedFeatureToString(UnsupportedGraphFeatures Feature) {
5859
return "sycl_ext_oneapi_memcpy2d";
5960
case UGF::sycl_ext_oneapi_device_global:
6061
return "sycl_ext_oneapi_device_global";
62+
case UGF::sycl_ext_oneapi_bindless_images:
63+
return "sycl_ext_oneapi_bindless_images";
6164
default:
6265
return {};
6366
}

‎sycl/source/handler.cpp‎

Lines changed: 24 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -903,6 +903,9 @@ void handler::ext_oneapi_memset2d_impl(void *Dest, size_t DestPitch, int Value,
903903
void handler::ext_oneapi_copy(
904904
void *Src, ext::oneapi::experimental::image_mem_handle Dest,
905905
const ext::oneapi::experimental::image_descriptor &Desc) {
906+
throwIfGraphAssociated<
907+
ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
908+
sycl_ext_oneapi_bindless_images>();
906909
MSrcPtr = Src;
907910
MDstPtr = Dest.raw_handle;
908911

@@ -936,7 +939,9 @@ void handler::ext_oneapi_copy(
936939
ext::oneapi::experimental::image_mem_handle Dest, sycl::range<3> DestOffset,
937940
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
938941
sycl::range<3> CopyExtent) {
939-
942+
throwIfGraphAssociated<
943+
ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
944+
sycl_ext_oneapi_bindless_images>();
940945
MSrcPtr = Src;
941946
MDstPtr = Dest.raw_handle;
942947

@@ -969,6 +974,9 @@ void handler::ext_oneapi_copy(
969974
void handler::ext_oneapi_copy(
970975
ext::oneapi::experimental::image_mem_handle Src, void *Dest,
971976
const ext::oneapi::experimental::image_descriptor &Desc) {
977+
throwIfGraphAssociated<
978+
ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
979+
sycl_ext_oneapi_bindless_images>();
972980
MSrcPtr = Src.raw_handle;
973981
MDstPtr = Dest;
974982

@@ -1002,6 +1010,9 @@ void handler::ext_oneapi_copy(
10021010
const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest,
10031011
sycl::range<3> DestOffset, sycl::range<3> DestExtent,
10041012
sycl::range<3> CopyExtent) {
1013+
throwIfGraphAssociated<
1014+
ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
1015+
sycl_ext_oneapi_bindless_images>();
10051016
MSrcPtr = Src.raw_handle;
10061017
MDstPtr = Dest;
10071018

@@ -1034,6 +1045,9 @@ void handler::ext_oneapi_copy(
10341045
void handler::ext_oneapi_copy(
10351046
void *Src, void *Dest,
10361047
const ext::oneapi::experimental::image_descriptor &Desc, size_t Pitch) {
1048+
throwIfGraphAssociated<
1049+
ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
1050+
sycl_ext_oneapi_bindless_images>();
10371051
MSrcPtr = Src;
10381052
MDstPtr = Dest;
10391053

@@ -1069,6 +1083,9 @@ void handler::ext_oneapi_copy(
10691083
const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
10701084
size_t DeviceRowPitch, sycl::range<3> HostExtent,
10711085
sycl::range<3> CopyExtent) {
1086+
throwIfGraphAssociated<
1087+
ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
1088+
sycl_ext_oneapi_bindless_images>();
10721089
MSrcPtr = Src;
10731090
MDstPtr = Dest;
10741091

@@ -1102,13 +1119,19 @@ void handler::ext_oneapi_copy(
11021119

11031120
void handler::ext_oneapi_wait_external_semaphore(
11041121
sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle) {
1122+
throwIfGraphAssociated<
1123+
ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
1124+
sycl_ext_oneapi_bindless_images>();
11051125
MImpl->MInteropSemaphoreHandle =
11061126
(sycl::detail::pi::PiInteropSemaphoreHandle)SemaphoreHandle.raw_handle;
11071127
setType(detail::CG::SemaphoreWait);
11081128
}
11091129

11101130
void handler::ext_oneapi_signal_external_semaphore(
11111131
sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle) {
1132+
throwIfGraphAssociated<
1133+
ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
1134+
sycl_ext_oneapi_bindless_images>();
11121135
MImpl->MInteropSemaphoreHandle =
11131136
(sycl::detail::pi::PiInteropSemaphoreHandle)SemaphoreHandle.raw_handle;
11141137
setType(detail::CG::SemaphoreSignal);

‎sycl/unittests/Extensions/CommandGraph.cpp‎

Lines changed: 194 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -311,6 +311,160 @@ void addMemcpy2D(experimental::detail::modifiable_command_graph &G, queue &Q,
311311
ASSERT_EQ(ExceptionCode, sycl::errc::invalid);
312312
}
313313

314+
/// Tries to add nodes including images bindless copy instructions
315+
/// to the graph G. It tests that an invalid exception has been thrown
316+
/// Since sycl_ext_oneapi_bindless_images extension can not be used
317+
/// along with SYCL Graph.
318+
///
319+
/// @param G Modifiable graph to add commands to.
320+
/// @param Q Queue to submit nodes to.
321+
/// @param Img Image memory
322+
/// @param HostData Host Pointer to the memory
323+
/// @param ImgUSM USM Pointer to Image memory
324+
/// @param Pitch image pitch
325+
/// @param Desc Image descriptor
326+
template <OperationPath PathKind>
327+
void addImagesCopies(experimental::detail::modifiable_command_graph &G,
328+
queue &Q, sycl::ext::oneapi::experimental::image_mem Img,
329+
std::vector<sycl::float4> HostData, void *ImgUSM,
330+
size_t Pitch,
331+
sycl::ext::oneapi::experimental::image_descriptor Desc) {
332+
// simple copy Host to Device
333+
std::error_code ExceptionCode = make_error_code(sycl::errc::success);
334+
try {
335+
if constexpr (PathKind == OperationPath::RecordReplay) {
336+
Q.submit([&](handler &CGH) {
337+
CGH.ext_oneapi_copy(HostData.data(), Img.get_handle(), Desc);
338+
});
339+
}
340+
if constexpr (PathKind == OperationPath::Shortcut) {
341+
Q.ext_oneapi_copy(HostData.data(), Img.get_handle(), Desc);
342+
}
343+
if constexpr (PathKind == OperationPath::Explicit) {
344+
G.add([&](handler &CGH) {
345+
CGH.ext_oneapi_copy(HostData.data(), Img.get_handle(), Desc);
346+
});
347+
}
348+
} catch (exception &Exception) {
349+
ExceptionCode = Exception.code();
350+
}
351+
ASSERT_EQ(ExceptionCode, sycl::errc::invalid);
352+
353+
// simple copy Device to Host
354+
ExceptionCode = make_error_code(sycl::errc::success);
355+
try {
356+
if constexpr (PathKind == OperationPath::RecordReplay) {
357+
Q.submit([&](handler &CGH) {
358+
CGH.ext_oneapi_copy(Img.get_handle(), HostData.data(), Desc);
359+
});
360+
}
361+
if constexpr (PathKind == OperationPath::Shortcut) {
362+
Q.ext_oneapi_copy(Img.get_handle(), HostData.data(), Desc);
363+
}
364+
if constexpr (PathKind == OperationPath::Explicit) {
365+
G.add([&](handler &CGH) {
366+
CGH.ext_oneapi_copy(Img.get_handle(), HostData.data(), Desc);
367+
});
368+
}
369+
} catch (exception &Exception) {
370+
ExceptionCode = Exception.code();
371+
}
372+
ASSERT_EQ(ExceptionCode, sycl::errc::invalid);
373+
374+
// simple copy Host to Device USM
375+
ExceptionCode = make_error_code(sycl::errc::success);
376+
try {
377+
if constexpr (PathKind == OperationPath::RecordReplay) {
378+
Q.submit([&](handler &CGH) {
379+
CGH.ext_oneapi_copy(HostData.data(), ImgUSM, Desc, Pitch);
380+
});
381+
}
382+
if constexpr (PathKind == OperationPath::Shortcut) {
383+
Q.ext_oneapi_copy(HostData.data(), ImgUSM, Desc, Pitch);
384+
}
385+
if constexpr (PathKind == OperationPath::Explicit) {
386+
G.add([&](handler &CGH) {
387+
CGH.ext_oneapi_copy(HostData.data(), ImgUSM, Desc, Pitch);
388+
});
389+
}
390+
} catch (exception &Exception) {
391+
ExceptionCode = Exception.code();
392+
}
393+
ASSERT_EQ(ExceptionCode, sycl::errc::invalid);
394+
395+
// subregion copy Host to Device
396+
ExceptionCode = make_error_code(sycl::errc::success);
397+
try {
398+
if constexpr (PathKind == OperationPath::RecordReplay) {
399+
Q.submit([&](handler &CGH) {
400+
CGH.ext_oneapi_copy(HostData.data(), {0, 0, 0}, {0, 0, 0},
401+
Img.get_handle(), {0, 0, 0}, Desc, {0, 0, 0});
402+
});
403+
}
404+
if constexpr (PathKind == OperationPath::Shortcut) {
405+
Q.ext_oneapi_copy(HostData.data(), {0, 0, 0}, {0, 0, 0}, Img.get_handle(),
406+
{0, 0, 0}, Desc, {0, 0, 0});
407+
}
408+
if constexpr (PathKind == OperationPath::Explicit) {
409+
G.add([&](handler &CGH) {
410+
CGH.ext_oneapi_copy(HostData.data(), {0, 0, 0}, {0, 0, 0},
411+
Img.get_handle(), {0, 0, 0}, Desc, {0, 0, 0});
412+
});
413+
}
414+
} catch (exception &Exception) {
415+
ExceptionCode = Exception.code();
416+
}
417+
ASSERT_EQ(ExceptionCode, sycl::errc::invalid);
418+
419+
// subregion copy Device to Host
420+
ExceptionCode = make_error_code(sycl::errc::success);
421+
try {
422+
if constexpr (PathKind == OperationPath::RecordReplay) {
423+
Q.submit([&](handler &CGH) {
424+
CGH.ext_oneapi_copy(Img.get_handle(), {0, 0, 0}, Desc, HostData.data(),
425+
{0, 0, 0}, {0, 0, 0}, {0, 0, 0});
426+
});
427+
}
428+
if constexpr (PathKind == OperationPath::Shortcut) {
429+
Q.ext_oneapi_copy(Img.get_handle(), {0, 0, 0}, Desc, HostData.data(),
430+
{0, 0, 0}, {0, 0, 0}, {0, 0, 0});
431+
}
432+
if constexpr (PathKind == OperationPath::Explicit) {
433+
G.add([&](handler &CGH) {
434+
CGH.ext_oneapi_copy(Img.get_handle(), {0, 0, 0}, Desc, HostData.data(),
435+
{0, 0, 0}, {0, 0, 0}, {0, 0, 0});
436+
});
437+
}
438+
} catch (exception &Exception) {
439+
ExceptionCode = Exception.code();
440+
}
441+
ASSERT_EQ(ExceptionCode, sycl::errc::invalid);
442+
443+
// subregion copy Host to Device USM
444+
ExceptionCode = make_error_code(sycl::errc::success);
445+
try {
446+
if constexpr (PathKind == OperationPath::RecordReplay) {
447+
Q.submit([&](handler &CGH) {
448+
CGH.ext_oneapi_copy(HostData.data(), {0, 0, 0}, ImgUSM, {0, 0, 0}, Desc,
449+
Pitch, {0, 0, 0}, {0, 0, 0});
450+
});
451+
}
452+
if constexpr (PathKind == OperationPath::Shortcut) {
453+
Q.ext_oneapi_copy(HostData.data(), {0, 0, 0}, ImgUSM, {0, 0, 0}, Desc,
454+
Pitch, {0, 0, 0}, {0, 0, 0});
455+
}
456+
if constexpr (PathKind == OperationPath::Explicit) {
457+
G.add([&](handler &CGH) {
458+
CGH.ext_oneapi_copy(HostData.data(), {0, 0, 0}, ImgUSM, {0, 0, 0}, Desc,
459+
Pitch, {0, 0, 0}, {0, 0, 0});
460+
});
461+
}
462+
} catch (exception &Exception) {
463+
ExceptionCode = Exception.code();
464+
}
465+
ASSERT_EQ(ExceptionCode, sycl::errc::invalid);
466+
}
467+
314468
bool depthSearchSuccessorCheck(
315469
std::shared_ptr<sycl::ext::oneapi::experimental::detail::node_impl> Node) {
316470
if (Node->MSuccessors.size() > 1)
@@ -1238,7 +1392,7 @@ TEST_F(CommandGraphTest, Reductions) {
12381392
{
12391393
try {
12401394
Graph.add([&](handler &CGH) {
1241-
CGH.parallel_for<class TestKernel>(
1395+
CGH.parallel_for<class CustomTestKernel>(
12421396
range<1>{1}, reduction(&ReduVar, int{0}, sycl::plus()),
12431397
[=](item<1> idx, auto &Sum) {});
12441398
});
@@ -1250,6 +1404,45 @@ TEST_F(CommandGraphTest, Reductions) {
12501404
sycl::exception);
12511405
}
12521406

1407+
TEST_F(CommandGraphTest, BindlessExceptionCheck) {
1408+
auto Ctxt = Queue.get_context();
1409+
1410+
// declare image data
1411+
size_t Height = 13;
1412+
size_t Width = 7;
1413+
size_t Depth = 11;
1414+
size_t N = Height * Width * Depth;
1415+
std::vector<sycl::float4> DataIn(N);
1416+
1417+
// Extension: image descriptor - can use the same for both images
1418+
sycl::ext::oneapi::experimental::image_descriptor Desc(
1419+
{Width, Height, Depth}, sycl::image_channel_order::rgba,
1420+
sycl::image_channel_type::fp32);
1421+
1422+
// Extension: allocate memory on device and create the handle
1423+
// Input images memory
1424+
sycl::ext::oneapi::experimental::image_mem ImgMem(Desc, Dev, Ctxt);
1425+
// Extension: returns the device pointer to USM allocated pitched memory
1426+
size_t Pitch = 0;
1427+
auto ImgMemUSM = sycl::ext::oneapi::experimental::pitched_alloc_device(
1428+
&Pitch, Desc, Queue);
1429+
1430+
Graph.begin_recording(Queue);
1431+
1432+
addImagesCopies<OperationPath::RecordReplay>(Graph, Queue, ImgMem, DataIn,
1433+
ImgMemUSM, Pitch, Desc);
1434+
1435+
addImagesCopies<OperationPath::Shortcut>(Graph, Queue, ImgMem, DataIn,
1436+
ImgMemUSM, Pitch, Desc);
1437+
1438+
Graph.end_recording();
1439+
1440+
addImagesCopies<OperationPath::Explicit>(Graph, Queue, ImgMem, DataIn,
1441+
ImgMemUSM, Pitch, Desc);
1442+
1443+
sycl::free(ImgMemUSM, Ctxt);
1444+
}
1445+
12531446
class MultiThreadGraphTest : public CommandGraphTest {
12541447
public:
12551448
MultiThreadGraphTest()

0 commit comments

Comments
 (0)