[SYCL] Introduce nd_range_view and integrate with handler-less path (#20531)

Introduced an nd_range_view type, which is a wrapper for nd_range and
range types. This allows for the ABI simplification. This PR changes the
handler-less related ABI to use the new type.

---------

Co-authored-by: Alexandr Konovalov <alexandr.konovalov@intel.com>

Author: slawekptak

sycl/include/sycl/detail/nd_range_view.hpp

@@ -0,0 +1,48 @@
+//==---- nd_range_view.hpp --- SYCL iteration with reference to ranges ---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/nd_range.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+class NDRDescT;
+
+inline namespace nd_range_view_v1 {
+
+// The structure to keep dimension and references to ranges unified for
+// all dimensions.
+class nd_range_view {
+
+public:
+  nd_range_view() = default;
+  nd_range_view(const nd_range_view &Desc) = default;
+  nd_range_view(nd_range_view &&Desc) = default;
+  nd_range_view &operator=(const nd_range_view &Desc) = default;
+  nd_range_view &operator=(nd_range_view &&Desc) = default;
+
+  template <int Dims_>
+  nd_range_view(sycl::nd_range<Dims_> &ExecutionRange)
+      : MGlobalSize(&(ExecutionRange.globalSize[0])),
+        MLocalSize(&(ExecutionRange.localSize[0])),
+        MOffset(&(ExecutionRange.offset[0])), MDims{size_t(Dims_)} {}
+
+  const size_t *MGlobalSize = nullptr;
+  const size_t *MLocalSize = nullptr;
+  const size_t *MOffset = nullptr;
+  size_t MDims = 1;
+};
+
+} // namespace nd_range_view_v1
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/nd_range.hpp

@@ -15,6 +15,12 @@
 namespace sycl {
 inline namespace _V1 {
 
+namespace detail {
+inline namespace nd_range_view_v1 {
+class nd_range_view;
+}
+} // namespace detail
+
 /// Defines the iteration domain of both the work-groups and the overall
 /// dispatch.
 ///
@@ -65,6 +71,8 @@ template <int Dimensions = 1> class nd_range {
   bool operator!=(const nd_range<Dimensions> &rhs) const {
     return !(*this == rhs);
   }
+
+  friend class sycl::_V1::detail::nd_range_view;
 };
 
 } // namespace _V1

sycl/include/sycl/queue.hpp

@@ -21,6 +21,7 @@
 #include <sycl/detail/id_queries_fit_in_int.hpp> // for checkValueRange
 #include <sycl/detail/info_desc_helpers.hpp>  // for is_queue_info_...
 #include <sycl/detail/kernel_desc.hpp>        // for KernelInfo
+#include <sycl/detail/nd_range_view.hpp>
 #include <sycl/detail/optional.hpp>
 #include <sycl/detail/owner_less_base.hpp> // for OwnerLessBase
 #include <sycl/device.hpp>                 // for device
@@ -64,18 +65,16 @@ template <backend BackendName, class SyclObjectT>
 auto get_native(const SyclObjectT &Obj)
     -> backend_return_t<BackendName, SyclObjectT>;
 
-template <int Dims>
 event __SYCL_EXPORT submit_kernel_direct_with_event_impl(
-    const queue &Queue, const nd_range<Dims> &Range,
+    const queue &Queue, const detail::nd_range_view &RangeView,
     detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     sycl::span<const event> DepEvents,
     const detail::KernelPropertyHolderStructTy &Props,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
-template <int Dims>
 void __SYCL_EXPORT submit_kernel_direct_without_event_impl(
-    const queue &Queue, const nd_range<Dims> &Range,
+    const queue &Queue, const detail::nd_range_view &RangeView,
     detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     sycl::span<const event> DepEvents,
@@ -167,7 +166,7 @@ template <detail::WrapAs WrapAs, typename LambdaArgType,
           typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
           typename KernelTypeUniversalRef, int Dims>
 auto submit_kernel_direct(
-    const queue &Queue, const nd_range<Dims> &Range,
+    const queue &Queue, const detail::nd_range_view &RangeView,
     KernelTypeUniversalRef &&KernelFunc, sycl::span<const event> DepEvents,
     const PropertiesT &ExtraProps =
         ext::oneapi::experimental::empty_properties_t{},
@@ -233,12 +232,12 @@ auto submit_kernel_direct(
 
   if constexpr (EventNeeded) {
     return submit_kernel_direct_with_event_impl(
-        Queue, Range, HostKernel, DeviceKernelInfoPtr, DepEvents,
+        Queue, RangeView, HostKernel, DeviceKernelInfoPtr, DepEvents,
         ParsedProperties, TlsCodeLocCapture.query(),
         TlsCodeLocCapture.isToplevel());
   } else {
     submit_kernel_direct_without_event_impl(
-        Queue, Range, HostKernel, DeviceKernelInfoPtr, DepEvents,
+        Queue, RangeView, HostKernel, DeviceKernelInfoPtr, DepEvents,
         ParsedProperties, TlsCodeLocCapture.query(),
         TlsCodeLocCapture.isToplevel());
   }
@@ -248,7 +247,7 @@ template <typename KernelName = detail::auto_name, bool EventNeeded = false,
           typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
           typename KernelTypeUniversalRef, int Dims>
 auto submit_kernel_direct_parallel_for(
-    const queue &Queue, const nd_range<Dims> &Range,
+    const queue &Queue, nd_range<Dims> Range,
     KernelTypeUniversalRef &&KernelFunc, sycl::span<const event> DepEvents = {},
     const PropertiesT &Props = ext::oneapi::experimental::empty_properties_t{},
     const detail::code_location &CodeLoc = detail::code_location::current()) {
@@ -271,8 +270,9 @@ auto submit_kernel_direct_parallel_for(
   return submit_kernel_direct<detail::WrapAs::parallel_for, TransformedArgType,
                               KernelName, EventNeeded, PropertiesT,
                               KernelTypeUniversalRef, Dims>(
-      Queue, Range, std::forward<KernelTypeUniversalRef>(KernelFunc), DepEvents,
-      Props, CodeLoc);
+      Queue, detail::nd_range_view(Range),
+      std::forward<KernelTypeUniversalRef>(KernelFunc), DepEvents, Props,
+      CodeLoc);
 }
 
 template <typename KernelName = detail::auto_name, bool EventNeeded = false,
@@ -287,7 +287,7 @@ auto submit_kernel_direct_single_task(
   return submit_kernel_direct<detail::WrapAs::single_task, void, KernelName,
                               EventNeeded, PropertiesT, KernelTypeUniversalRef,
                               1>(
-      Queue, nd_range<1>{1, 1},
+      Queue, detail::nd_range_view(),
       std::forward<KernelTypeUniversalRef>(KernelFunc), DepEvents, Props,
       CodeLoc);
 }

sycl/source/detail/ndrange_desc.hpp

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/detail/nd_range_view.hpp>
 #include <sycl/nd_range.hpp>
 #include <sycl/range.hpp>
 
@@ -32,6 +33,15 @@ class NDRDescT {
   NDRDescT(const NDRDescT &Desc) = default;
   NDRDescT(NDRDescT &&Desc) = default;
 
+  NDRDescT(const detail::nd_range_view &NDRangeView) : Dims{NDRangeView.MDims} {
+    if (!NDRangeView.MGlobalSize) {
+      init();
+    } else {
+      init(NDRangeView.MGlobalSize, NDRangeView.MLocalSize,
+           NDRangeView.MOffset);
+    }
+  }
+
   template <int Dims_>
   NDRDescT(sycl::range<Dims_> N, bool SetNumWorkGroups) : Dims{size_t(Dims_)} {
     if (SetNumWorkGroups) {
@@ -43,7 +53,7 @@ class NDRDescT {
         GlobalSize[I] = N[I];
       }
 
-      for (int I = Dims_; I < 3; ++I) {
+      for (size_t I = Dims_; I < 3; ++I) {
         GlobalSize[I] = 1;
       }
     }
@@ -53,19 +63,7 @@ class NDRDescT {
   NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::range<Dims_> LocalSizes,
            sycl::id<Dims_> Offset)
       : Dims{size_t(Dims_)} {
-    for (size_t I = 0; I < Dims_; ++I) {
-      GlobalSize[I] = NumWorkItems[I];
-      LocalSize[I] = LocalSizes[I];
-      GlobalOffset[I] = Offset[I];
-    }
-
-    for (int I = Dims_; I < 3; ++I) {
-      LocalSize[I] = LocalSizes[0] ? 1 : 0;
-    }
-
-    for (int I = Dims_; I < 3; ++I) {
-      GlobalSize[I] = 1;
-    }
+    init(&(NumWorkItems[0]), &(LocalSizes[0]), &(Offset[0]));
   }
 
   template <int Dims_>
@@ -109,6 +107,29 @@ class NDRDescT {
   std::array<size_t, 3> NumWorkGroups{0, 0, 0};
   std::array<size_t, 3> ClusterDimensions{1, 1, 1};
   size_t Dims = 0;
+
+private:
+  void init(const size_t *NumWorkItems, const size_t *LocalSizes,
+            const size_t *Offset) {
+    for (size_t I = 0; I < Dims; ++I) {
+      GlobalSize[I] = NumWorkItems[I];
+      LocalSize[I] = LocalSizes[I];
+      GlobalOffset[I] = Offset[I];
+    }
+
+    for (size_t I = Dims; I < 3; ++I) {
+      LocalSize[I] = LocalSizes[0] ? 1 : 0;
+    }
+
+    for (size_t I = Dims; I < 3; ++I) {
+      GlobalSize[I] = 1;
+    }
+  }
+
+  void init() {
+    size_t GlobalS = 1, LocalS = 1, Offset = 0;
+    init(&GlobalS, &LocalS, &Offset);
+  }
 };
 
 } // namespace detail

sycl/source/detail/queue_impl.hpp

@@ -360,27 +360,27 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
     return createSyclObjFromImpl<event>(ResEvent);
   }
 
-  template <int Dims>
   event submit_kernel_direct_with_event(
-      const nd_range<Dims> &Range, detail::HostKernelRefBase &HostKernel,
+      const detail::nd_range_view &RangeView,
+      detail::HostKernelRefBase &HostKernel,
       detail::DeviceKernelInfo *DeviceKernelInfo,
       sycl::span<const event> DepEvents,
       const detail::KernelPropertyHolderStructTy &Props,
       const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
     detail::EventImplPtr EventImpl = submit_kernel_direct_impl(
-        NDRDescT{Range}, HostKernel, DeviceKernelInfo,
+        NDRDescT(RangeView), HostKernel, DeviceKernelInfo,
         /*CallerNeedsEvent*/ true, DepEvents, Props, CodeLoc, IsTopCodeLoc);
     return createSyclObjFromImpl<event>(EventImpl);
   }
 
-  template <int Dims>
   void submit_kernel_direct_without_event(
-      const nd_range<Dims> &Range, detail::HostKernelRefBase &HostKernel,
+      const detail::nd_range_view &RangeView,
+      detail::HostKernelRefBase &HostKernel,
       detail::DeviceKernelInfo *DeviceKernelInfo,
       sycl::span<const event> DepEvents,
       const detail::KernelPropertyHolderStructTy &Props,
       const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
-    submit_kernel_direct_impl(NDRDescT{Range}, HostKernel, DeviceKernelInfo,
+    submit_kernel_direct_impl(NDRDescT(RangeView), HostKernel, DeviceKernelInfo,
                               /*CallerNeedsEvent*/ false, DepEvents, Props,
                               CodeLoc, IsTopCodeLoc);
   }

sycl/source/queue.cpp

@@ -463,80 +463,30 @@ void queue::ext_oneapi_set_external_event(const event &external_event) {
 
 const property_list &queue::getPropList() const { return impl->getPropList(); }
 
-template <int Dims>
 event submit_kernel_direct_with_event_impl(
-    const queue &Queue, const nd_range<Dims> &Range,
+    const queue &Queue, const detail::nd_range_view &RangeView,
     detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     sycl::span<const event> DepEvents,
     const detail::KernelPropertyHolderStructTy &Props,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
   return getSyclObjImpl(Queue)->submit_kernel_direct_with_event(
-      Range, HostKernel, DeviceKernelInfo, DepEvents, Props, CodeLoc,
+      RangeView, HostKernel, DeviceKernelInfo, DepEvents, Props, CodeLoc,
       IsTopCodeLoc);
 }
 
-template event __SYCL_EXPORT submit_kernel_direct_with_event_impl<1>(
-    const queue &Queue, const nd_range<1> &Range,
-    detail::HostKernelRefBase &HostKernel,
-    detail::DeviceKernelInfo *DeviceKernelInfo,
-    sycl::span<const event> DepEvents,
-    const detail::KernelPropertyHolderStructTy &Props,
-    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
-
-template event __SYCL_EXPORT submit_kernel_direct_with_event_impl<2>(
-    const queue &Queue, const nd_range<2> &Range,
-    detail::HostKernelRefBase &HostKernel,
-    detail::DeviceKernelInfo *DeviceKernelInfo,
-    sycl::span<const event> DepEvents,
-    const detail::KernelPropertyHolderStructTy &Props,
-    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
-
-template event __SYCL_EXPORT submit_kernel_direct_with_event_impl<3>(
-    const queue &Queue, const nd_range<3> &Range,
-    detail::HostKernelRefBase &HostKernel,
-    detail::DeviceKernelInfo *DeviceKernelInfo,
-    sycl::span<const event> DepEvents,
-    const detail::KernelPropertyHolderStructTy &Props,
-    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
-
-template <int Dims>
 void submit_kernel_direct_without_event_impl(
-    const queue &Queue, const nd_range<Dims> &Range,
+    const queue &Queue, const detail::nd_range_view &RangeView,
     detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     sycl::span<const event> DepEvents,
     const detail::KernelPropertyHolderStructTy &Props,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
   getSyclObjImpl(Queue)->submit_kernel_direct_without_event(
-      Range, HostKernel, DeviceKernelInfo, DepEvents, Props, CodeLoc,
+      RangeView, HostKernel, DeviceKernelInfo, DepEvents, Props, CodeLoc,
       IsTopCodeLoc);
 }
 
-template void __SYCL_EXPORT submit_kernel_direct_without_event_impl<1>(
-    const queue &Queue, const nd_range<1> &Range,
-    detail::HostKernelRefBase &HostKernel,
-    detail::DeviceKernelInfo *DeviceKernelInfo,
-    sycl::span<const event> DepEvents,
-    const detail::KernelPropertyHolderStructTy &Props,
-    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
-
-template void __SYCL_EXPORT submit_kernel_direct_without_event_impl<2>(
-    const queue &Queue, const nd_range<2> &Range,
-    detail::HostKernelRefBase &HostKernel,
-    detail::DeviceKernelInfo *DeviceKernelInfo,
-    sycl::span<const event> DepEvents,
-    const detail::KernelPropertyHolderStructTy &Props,
-    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
-
-template void __SYCL_EXPORT submit_kernel_direct_without_event_impl<3>(
-    const queue &Queue, const nd_range<3> &Range,
-    detail::HostKernelRefBase &HostKernel,
-    detail::DeviceKernelInfo *DeviceKernelInfo,
-    sycl::span<const event> DepEvents,
-    const detail::KernelPropertyHolderStructTy &Props,
-    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
-
 } // namespace _V1
 } // namespace sycl
 

sycl/test/abi/layout_nd_range_view.cpp

@@ -0,0 +1,18 @@
+// RUN: %clangxx -fsycl -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s
+// RUN: %clangxx -fsycl -fsycl-device-only -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s
+// REQUIRES: linux
+// UNSUPPORTED: libcxx
+
+// clang-format off
+
+#include <sycl/detail/nd_range_view.hpp>
+
+
+SYCL_EXTERNAL void nd_range_view(sycl::detail::nd_range_view_v1::nd_range_view) {}
+// CHECK: 0 | class sycl::detail::nd_range_view
+// CHECK-NEXT: 0 |   const size_t * MGlobalSize
+// CHECK-NEXT: 8 |   const size_t * MLocalSize
+// CHECK-NEXT: 16 |  const size_t * MOffset
+// CHECK-NEXT: 24 |  size_t MDims
+// CHECK-NEXT: | [sizeof=32, dsize=32, align=8,
+// CHECK-NEXT: |  nvsize=32, nvalign=8]

sycl/test/abi/sycl_symbols_linux.dump

@@ -2985,12 +2985,8 @@ _ZN4sycl3_V121__isgreaterequal_implEdd
 _ZN4sycl3_V121__isgreaterequal_implEff
 _ZN4sycl3_V122accelerator_selector_vERKNS0_6deviceE
 _ZN4sycl3_V128verifyUSMAllocatorPropertiesERKNS0_13property_listE
-_ZN4sycl3_V136submit_kernel_direct_with_event_implILi1EEENS0_5eventERKNS0_5queueERKNS0_8nd_rangeIXT_EEERNS0_6detail17HostKernelRefBaseEPNSA_16DeviceKernelInfoENS0_4spanIKS2_Lm18446744073709551615EEERKNSA_27kernel_launch_properties_v111PropsHolderIJNS0_3ext6oneapi12experimental23work_group_scratch_sizeENSK_5intel12experimental12cache_configENSM_17use_root_sync_keyENSM_23work_group_progress_keyENSM_22sub_group_progress_keyENSM_22work_item_progress_keyENSM_4cuda12cluster_sizeILi1EEENSW_ILi2EEENSW_ILi3EEEEEERKNSA_13code_locationEb
-_ZN4sycl3_V136submit_kernel_direct_with_event_implILi2EEENS0_5eventERKNS0_5queueERKNS0_8nd_rangeIXT_EEERNS0_6detail17HostKernelRefBaseEPNSA_16DeviceKernelInfoENS0_4spanIKS2_Lm18446744073709551615EEERKNSA_27kernel_launch_properties_v111PropsHolderIJNS0_3ext6oneapi12experimental23work_group_scratch_sizeENSK_5intel12experimental12cache_configENSM_17use_root_sync_keyENSM_23work_group_progress_keyENSM_22sub_group_progress_keyENSM_22work_item_progress_keyENSM_4cuda12cluster_sizeILi1EEENSW_ILi2EEENSW_ILi3EEEEEERKNSA_13code_locationEb
-_ZN4sycl3_V136submit_kernel_direct_with_event_implILi3EEENS0_5eventERKNS0_5queueERKNS0_8nd_rangeIXT_EEERNS0_6detail17HostKernelRefBaseEPNSA_16DeviceKernelInfoENS0_4spanIKS2_Lm18446744073709551615EEERKNSA_27kernel_launch_properties_v111PropsHolderIJNS0_3ext6oneapi12experimental23work_group_scratch_sizeENSK_5intel12experimental12cache_configENSM_17use_root_sync_keyENSM_23work_group_progress_keyENSM_22sub_group_progress_keyENSM_22work_item_progress_keyENSM_4cuda12cluster_sizeILi1EEENSW_ILi2EEENSW_ILi3EEEEEERKNSA_13code_locationEb
-_ZN4sycl3_V139submit_kernel_direct_without_event_implILi1EEEvRKNS0_5queueERKNS0_8nd_rangeIXT_EEERNS0_6detail17HostKernelRefBaseEPNS9_16DeviceKernelInfoENS0_4spanIKNS0_5eventELm18446744073709551615EEERKNS9_27kernel_launch_properties_v111PropsHolderIJNS0_3ext6oneapi12experimental23work_group_scratch_sizeENSK_5intel12experimental12cache_configENSM_17use_root_sync_keyENSM_23work_group_progress_keyENSM_22sub_group_progress_keyENSM_22work_item_progress_keyENSM_4cuda12cluster_sizeILi1EEENSW_ILi2EEENSW_ILi3EEEEEERKNS9_13code_locationEb
-_ZN4sycl3_V139submit_kernel_direct_without_event_implILi2EEEvRKNS0_5queueERKNS0_8nd_rangeIXT_EEERNS0_6detail17HostKernelRefBaseEPNS9_16DeviceKernelInfoENS0_4spanIKNS0_5eventELm18446744073709551615EEERKNS9_27kernel_launch_properties_v111PropsHolderIJNS0_3ext6oneapi12experimental23work_group_scratch_sizeENSK_5intel12experimental12cache_configENSM_17use_root_sync_keyENSM_23work_group_progress_keyENSM_22sub_group_progress_keyENSM_22work_item_progress_keyENSM_4cuda12cluster_sizeILi1EEENSW_ILi2EEENSW_ILi3EEEEEERKNS9_13code_locationEb
-_ZN4sycl3_V139submit_kernel_direct_without_event_implILi3EEEvRKNS0_5queueERKNS0_8nd_rangeIXT_EEERNS0_6detail17HostKernelRefBaseEPNS9_16DeviceKernelInfoENS0_4spanIKNS0_5eventELm18446744073709551615EEERKNS9_27kernel_launch_properties_v111PropsHolderIJNS0_3ext6oneapi12experimental23work_group_scratch_sizeENSK_5intel12experimental12cache_configENSM_17use_root_sync_keyENSM_23work_group_progress_keyENSM_22sub_group_progress_keyENSM_22work_item_progress_keyENSM_4cuda12cluster_sizeILi1EEENSW_ILi2EEENSW_ILi3EEEEEERKNS9_13code_locationEb
+_ZN4sycl3_V136submit_kernel_direct_with_event_implERKNS0_5queueERKNS0_6detail16nd_range_view_v113nd_range_viewERNS4_17HostKernelRefBaseEPNS4_16DeviceKernelInfoENS0_4spanIKNS0_5eventELm18446744073709551615EEERKNS4_27kernel_launch_properties_v111PropsHolderIJNS0_3ext6oneapi12experimental23work_group_scratch_sizeENSJ_5intel12experimental12cache_configENSL_17use_root_sync_keyENSL_23work_group_progress_keyENSL_22sub_group_progress_keyENSL_22work_item_progress_keyENSL_4cuda12cluster_sizeILi1EEENSV_ILi2EEENSV_ILi3EEEEEERKNS4_13code_locationEb
+_ZN4sycl3_V139submit_kernel_direct_without_event_implERKNS0_5queueERKNS0_6detail16nd_range_view_v113nd_range_viewERNS4_17HostKernelRefBaseEPNS4_16DeviceKernelInfoENS0_4spanIKNS0_5eventELm18446744073709551615EEERKNS4_27kernel_launch_properties_v111PropsHolderIJNS0_3ext6oneapi12experimental23work_group_scratch_sizeENSJ_5intel12experimental12cache_configENSL_17use_root_sync_keyENSL_23work_group_progress_keyENSL_22sub_group_progress_keyENSL_22work_item_progress_keyENSL_4cuda12cluster_sizeILi1EEENSV_ILi2EEENSV_ILi3EEEEEERKNS4_13code_locationEb
 _ZN4sycl3_V13ext5intel12experimental9pipe_base13get_pipe_nameB5cxx11EPKv
 _ZN4sycl3_V13ext5intel12experimental9pipe_base17wait_non_blockingERKNS0_5eventE
 _ZN4sycl3_V13ext5intel12experimental9pipe_base18get_pipe_name_implEPKv