queue.hpp

Go to the documentation of this file.

//==-------------------- queue.hpp - SYCL queue ----------------------------==//
//
// 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/access/access.hpp>             // for target, access...
#include <sycl/accessor.hpp>                  // for accessor
#include <sycl/aspects.hpp>                   // for aspect
#include <sycl/async_handler.hpp>             // for async_handler
#include <sycl/backend_types.hpp>             // for backend, backe...
#include <sycl/buffer.hpp>                    // for buffer
#include <sycl/context.hpp>                   // for context
#include <sycl/detail/assert_happened.hpp>    // for AssertHappened
#include <sycl/detail/cg_types.hpp>           // for check_fn_signa...
#include <sycl/detail/common.hpp>             // for code_location
#include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEP...
#include <sycl/detail/export.hpp>             // for __SYCL_EXPORT
#include <sycl/detail/info_desc_helpers.hpp>  // for is_queue_info_...
#include <sycl/detail/kernel_desc.hpp>        // for KernelInfo
#include <sycl/detail/owner_less_base.hpp>    // for OwnerLessBase
#include <sycl/detail/pi.h>                   // for pi_mem_advice
#include <sycl/device.hpp>                    // for device
#include <sycl/device_selector.hpp>           // for device_selector
#include <sycl/event.hpp>                     // for event
#include <sycl/exception.hpp>                 // for make_error_code
#include <sycl/exception_list.hpp>            // for defaultAsyncHa...
#include <sycl/ext/oneapi/bindless_images_descriptor.hpp> // for image_descriptor
#include <sycl/ext/oneapi/bindless_images_interop.hpp> // for interop_semaph...
#include <sycl/ext/oneapi/bindless_images_memory.hpp>  // for image_mem_handle
#include <sycl/ext/oneapi/device_global/device_global.hpp> // for device_global
#include <sycl/ext/oneapi/device_global/properties.hpp> // for device_image_s...
#include <sycl/ext/oneapi/experimental/graph.hpp>       // for command_graph...
#include <sycl/ext/oneapi/properties/properties.hpp>    // for empty_properti...
#include <sycl/handler.hpp>                             // for handler, isDev...
#include <sycl/id.hpp>                                  // for id
#include <sycl/kernel.hpp>                              // for auto_name
#include <sycl/kernel_handler.hpp>                      // for kernel_handler
#include <sycl/nd_range.hpp>                            // for nd_range
#include <sycl/property_list.hpp>                       // for property_list
#include <sycl/range.hpp>                               // for range
 
#include <cstddef>     // for size_t
#include <functional>  // for function
#include <memory>      // for shared_ptr, hash
#include <stdint.h>    // for int32_t
#include <tuple>       // for tuple
#include <type_traits> // for remove_all_ext...
#include <variant>     // for hash
#include <vector>      // for vector
 
// having _TWO_ mid-param #ifdefs makes the functions very difficult to read.
// Here we simplify the KernelFunc param is simplified to be
// _KERNELFUNCPARAM(KernelFunc) Once the queue kernel functions are defined,
// these macros are #undef immediately.
// replace _KERNELFUNCPARAM(KernelFunc) with   KernelType KernelFunc
//                                     or     const KernelType &KernelFunc
#ifdef __SYCL_NONCONST_FUNCTOR__
#define _KERNELFUNCPARAM(a) KernelType a
#else
#define _KERNELFUNCPARAM(a) const KernelType &a
#endif
 
namespace sycl {
inline namespace _V1 {
 
// Forward declaration
class context;
class device;
class event;
class queue;
 
template <backend BackendName, class SyclObjectT>
auto get_native(const SyclObjectT &Obj)
    -> backend_return_t<BackendName, SyclObjectT>;
 
namespace detail {
class queue_impl;
 
#if __SYCL_USE_FALLBACK_ASSERT
inline event submitAssertCapture(queue &, event &, queue *,
                                 const detail::code_location &);
#endif
} // namespace detail
 
namespace ext {
namespace oneapi {
namespace experimental {
// State of a queue with regards to graph recording,
// returned by info::queue::state
enum class queue_state { executing, recording };
} // namespace experimental
} // namespace oneapi
} // namespace ext
 
class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
public:
  explicit queue(const property_list &PropList = {})
      : queue(default_selector_v, detail::defaultAsyncHandler, PropList) {}
 
  queue(const async_handler &AsyncHandler, const property_list &PropList = {})
      : queue(default_selector_v, AsyncHandler, PropList) {}
 
  template <typename DeviceSelector,
            typename =
                detail::EnableIfSYCL2020DeviceSelectorInvocable<DeviceSelector>>
  explicit queue(const DeviceSelector &deviceSelector,
                 const async_handler &AsyncHandler,
                 const property_list &PropList = {})
      : queue(detail::select_device(deviceSelector), AsyncHandler, PropList) {}
 
  template <typename DeviceSelector,
            typename =
                detail::EnableIfSYCL2020DeviceSelectorInvocable<DeviceSelector>>
  explicit queue(const DeviceSelector &deviceSelector,
                 const property_list &PropList = {})
      : queue(detail::select_device(deviceSelector),
              detail::defaultAsyncHandler, PropList) {}
 
  template <typename DeviceSelector,
            typename =
                detail::EnableIfSYCL2020DeviceSelectorInvocable<DeviceSelector>>
  explicit queue(const context &syclContext,
                 const DeviceSelector &deviceSelector,
                 const property_list &propList = {})
      : queue(syclContext, detail::select_device(deviceSelector, syclContext),
              propList) {}
 
  template <typename DeviceSelector,
            typename =
                detail::EnableIfSYCL2020DeviceSelectorInvocable<DeviceSelector>>
  explicit queue(const context &syclContext,
                 const DeviceSelector &deviceSelector,
                 const async_handler &AsyncHandler,
                 const property_list &propList = {})
      : queue(syclContext, detail::select_device(deviceSelector, syclContext),
              AsyncHandler, propList) {}
 
  __SYCL2020_DEPRECATED("SYCL 1.2.1 device selectors are deprecated. Please "
                        "use SYCL 2020 device selectors instead.")
  queue(const device_selector &DeviceSelector,
        const property_list &PropList = {})
      : queue(DeviceSelector.select_device(), detail::defaultAsyncHandler,
              PropList) {}
 
  __SYCL2020_DEPRECATED("SYCL 1.2.1 device selectors are deprecated. Please "
                        "use SYCL 2020 device selectors instead.")
  queue(const device_selector &DeviceSelector,
        const async_handler &AsyncHandler, const property_list &PropList = {})
      : queue(DeviceSelector.select_device(), AsyncHandler, PropList) {}
 
  explicit queue(const device &SyclDevice, const property_list &PropList = {})
      : queue(SyclDevice, detail::defaultAsyncHandler, PropList) {}
 
  explicit queue(const device &SyclDevice, const async_handler &AsyncHandler,
                 const property_list &PropList = {});
 
  __SYCL2020_DEPRECATED("SYCL 1.2.1 device selectors are deprecated. Please "
                        "use SYCL 2020 device selectors instead.")
  queue(const context &SyclContext, const device_selector &DeviceSelector,
        const property_list &PropList = {});
 
  __SYCL2020_DEPRECATED("SYCL 1.2.1 device selectors are deprecated. Please "
                        "use SYCL 2020 device selectors instead.")
  queue(const context &SyclContext, const device_selector &DeviceSelector,
        const async_handler &AsyncHandler, const property_list &PropList = {});
 
  queue(const context &SyclContext, const device &SyclDevice,
        const property_list &PropList = {});
 
  queue(const context &SyclContext, const device &SyclDevice,
        const async_handler &AsyncHandler, const property_list &PropList = {});
 
#ifdef __SYCL_INTERNAL_API
  queue(cl_command_queue ClQueue, const context &SyclContext,
        const async_handler &AsyncHandler = {});
#endif
 
  queue(const queue &RHS) = default;
 
  queue(queue &&RHS) = default;
 
  queue &operator=(const queue &RHS) = default;
 
  queue &operator=(queue &&RHS) = default;
 
  bool operator==(const queue &RHS) const { return impl == RHS.impl; }
 
  bool operator!=(const queue &RHS) const { return !(*this == RHS); }
 
#ifdef __SYCL_INTERNAL_API
  cl_command_queue get() const;
#endif
 
  context get_context() const;
 
  device get_device() const;
 
  ext::oneapi::experimental::queue_state ext_oneapi_get_state() const;
 
  ext::oneapi::experimental::command_graph<
      ext::oneapi::experimental::graph_state::modifiable>
  ext_oneapi_get_graph() const;
 
  __SYCL2020_DEPRECATED(
      "is_host() is deprecated as the host device is no longer supported.")
  bool is_host() const;
 
  template <typename Param>
  typename detail::is_queue_info_desc<Param>::return_type get_info() const;
 
  template <typename Param>
  typename detail::is_backend_info_desc<Param>::return_type
  get_backend_info() const;
 
private:
  // A shorthand for `get_device().has()' which is expected to be a bit quicker
  // than the long version
  bool device_has(aspect Aspect) const;
 
public:
  template <typename T>
  std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event> submit(
      T CGF,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
#if __SYCL_USE_FALLBACK_ASSERT
    auto PostProcess = [this, &CodeLoc](bool IsKernel, bool KernelUsesAssert,
                                        event &E) {
      if (IsKernel && !device_has(aspect::ext_oneapi_native_assert) &&
          KernelUsesAssert && !device_has(aspect::accelerator)) {
        // __devicelib_assert_fail isn't supported by Device-side Runtime
        // Linking against fallback impl of __devicelib_assert_fail is
        // performed by program manager class
        // Fallback assert isn't supported for FPGA
        submitAssertCapture(*this, E, /* SecondaryQueue = */ nullptr, CodeLoc);
      }
    };
 
    return submit_impl_and_postprocess(CGF, CodeLoc, PostProcess);
#else
    return submit_impl(CGF, CodeLoc);
#endif // __SYCL_USE_FALLBACK_ASSERT
  }
 
  template <typename T>
  std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event> submit(
      T CGF, queue &SecondaryQueue,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
#if __SYCL_USE_FALLBACK_ASSERT
    auto PostProcess = [this, &SecondaryQueue, &CodeLoc](
                           bool IsKernel, bool KernelUsesAssert, event &E) {
      if (IsKernel && !device_has(aspect::ext_oneapi_native_assert) &&
          KernelUsesAssert && !device_has(aspect::accelerator)) {
        // Only secondary queues on devices need to be added to the assert
        // capture.
        // __devicelib_assert_fail isn't supported by Device-side Runtime
        // Linking against fallback impl of __devicelib_assert_fail is
        // performed by program manager class
        // Fallback assert isn't supported for FPGA
        submitAssertCapture(*this, E, &SecondaryQueue, CodeLoc);
      }
    };
 
    return submit_impl_and_postprocess(CGF, SecondaryQueue, CodeLoc,
                                       PostProcess);
#else
    return submit_impl(CGF, SecondaryQueue, CodeLoc);
#endif // __SYCL_USE_FALLBACK_ASSERT
  }
 
  event ext_oneapi_submit_barrier(
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event ext_oneapi_submit_barrier(
      const std::vector<event> &WaitList,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  void wait(
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    wait_proxy(CodeLoc);
  }
 
  void wait_and_throw(
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    wait_and_throw_proxy(CodeLoc);
  }
 
  void wait_proxy(const detail::code_location &CodeLoc);
  void wait_and_throw_proxy(const detail::code_location &CodeLoc);
 
  void throw_asynchronous();
 
  template <typename PropertyT> bool has_property() const noexcept;
 
  template <typename PropertyT> PropertyT get_property() const;
 
  template <typename T>
  event fill(
      void *Ptr, const T &Pattern, size_t Count,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit([&](handler &CGH) { CGH.fill<T>(Ptr, Pattern, Count); },
                  CodeLoc);
  }
 
  template <typename T>
  event fill(
      void *Ptr, const T &Pattern, size_t Count, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.fill<T>(Ptr, Pattern, Count);
        },
        CodeLoc);
  }
 
  template <typename T>
  event fill(
      void *Ptr, const T &Pattern, size_t Count,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.fill<T>(Ptr, Pattern, Count);
        },
        CodeLoc);
  }
 
  event memset(
      void *Ptr, int Value, size_t Count,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event memset(
      void *Ptr, int Value, size_t Count, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event memset(
      void *Ptr, int Value, size_t Count, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event memcpy(
      void *Dest, const void *Src, size_t Count,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event memcpy(
      void *Dest, const void *Src, size_t Count, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event memcpy(
      void *Dest, const void *Src, size_t Count,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T>
  event copy(
      const T *Src, T *Dest, size_t Count,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, Count * sizeof(T));
  }
 
  template <typename T>
  event copy(
      const T *Src, T *Dest, size_t Count, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, Count * sizeof(T), DepEvent);
  }
 
  template <typename T>
  event copy(
      const T *Src, T *Dest, size_t Count, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, Count * sizeof(T), DepEvents);
  }
 
  __SYCL2020_DEPRECATED("use the overload with int Advice instead")
  event mem_advise(
      const void *Ptr, size_t Length, pi_mem_advice Advice,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event mem_advise(
      const void *Ptr, size_t Length, int Advice,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event mem_advise(
      const void *Ptr, size_t Length, int Advice, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event mem_advise(
      const void *Ptr, size_t Length, int Advice,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  event prefetch(
      const void *Ptr, size_t Count,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit([=](handler &CGH) { CGH.prefetch(Ptr, Count); }, CodeLoc);
  }
 
  event prefetch(
      const void *Ptr, size_t Count, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [=](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.prefetch(Ptr, Count);
        },
        CodeLoc);
  }
 
  event prefetch(
      const void *Ptr, size_t Count, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [=](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.prefetch(Ptr, Count);
        },
        CodeLoc);
  }
 
  template <typename T = unsigned char,
            typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
  event ext_oneapi_memcpy2d(
      void *Dest, size_t DestPitch, const void *Src, size_t SrcPitch,
      size_t Width, size_t Height,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [=](handler &CGH) {
          CGH.ext_oneapi_memcpy2d<T>(Dest, DestPitch, Src, SrcPitch, Width,
                                     Height);
        },
        CodeLoc);
  }
 
  template <typename T = unsigned char,
            typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
  event ext_oneapi_memcpy2d(
      void *Dest, size_t DestPitch, const void *Src, size_t SrcPitch,
      size_t Width, size_t Height, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T = unsigned char,
            typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
  event ext_oneapi_memcpy2d(
      void *Dest, size_t DestPitch, const void *Src, size_t SrcPitch,
      size_t Width, size_t Height, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T>
  event ext_oneapi_copy2d(
      const T *Src, size_t SrcPitch, T *Dest, size_t DestPitch, size_t Width,
      size_t Height,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T>
  event ext_oneapi_copy2d(
      const T *Src, size_t SrcPitch, T *Dest, size_t DestPitch, size_t Width,
      size_t Height, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T>
  event ext_oneapi_copy2d(
      const T *Src, size_t SrcPitch, T *Dest, size_t DestPitch, size_t Width,
      size_t Height, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T = unsigned char,
            typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
  event ext_oneapi_memset2d(
      void *Dest, size_t DestPitch, int Value, size_t Width, size_t Height,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T = unsigned char,
            typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
  event ext_oneapi_memset2d(
      void *Dest, size_t DestPitch, int Value, size_t Width, size_t Height,
      event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T = unsigned char,
            typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
  event ext_oneapi_memset2d(
      void *Dest, size_t DestPitch, int Value, size_t Width, size_t Height,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T>
  event ext_oneapi_fill2d(
      void *Dest, size_t DestPitch, const T &Pattern, size_t Width,
      size_t Height,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T>
  event ext_oneapi_fill2d(
      void *Dest, size_t DestPitch, const T &Pattern, size_t Width,
      size_t Height, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T>
  event ext_oneapi_fill2d(
      void *Dest, size_t DestPitch, const T &Pattern, size_t Width,
      size_t Height, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current());
 
  template <typename T, typename PropertyListT>
  event memcpy(
      ext::oneapi::experimental::device_global<T, PropertyListT> &Dest,
      const void *Src, size_t NumBytes, size_t Offset,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    if (sizeof(T) < Offset + NumBytes)
      throw sycl::exception(make_error_code(errc::invalid),
                            "Copy to device_global is out of bounds.");
 
    if (!detail::isDeviceGlobalUsedInKernel(&Dest)) {
      // device_global is unregistered so we need a fallback. We let the handler
      // implement this fallback.
      return submit(
          [&](handler &CGH) {
            CGH.depends_on(DepEvents);
            return CGH.memcpy(Dest, Src, NumBytes, Offset);
          },
          CodeLoc);
    }
 
    constexpr bool IsDeviceImageScoped = PropertyListT::template has_property<
        ext::oneapi::experimental::device_image_scope_key>();
    return memcpyToDeviceGlobal(&Dest, Src, IsDeviceImageScoped, NumBytes,
                                Offset, DepEvents);
  }
 
  template <typename T, typename PropertyListT>
  event memcpy(
      ext::oneapi::experimental::device_global<T, PropertyListT> &Dest,
      const void *Src, size_t NumBytes, size_t Offset, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, NumBytes, Offset,
                        std::vector<event>{DepEvent});
  }
 
  template <typename T, typename PropertyListT>
  event memcpy(
      ext::oneapi::experimental::device_global<T, PropertyListT> &Dest,
      const void *Src, size_t NumBytes = sizeof(T), size_t Offset = 0,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, NumBytes, Offset, std::vector<event>{});
  }
 
  template <typename T, typename PropertyListT>
  event memcpy(
      void *Dest,
      const ext::oneapi::experimental::device_global<T, PropertyListT> &Src,
      size_t NumBytes, size_t Offset, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    if (sizeof(T) < Offset + NumBytes)
      throw sycl::exception(make_error_code(errc::invalid),
                            "Copy from device_global is out of bounds.");
 
    if (!detail::isDeviceGlobalUsedInKernel(&Src)) {
      // device_global is unregistered so we need a fallback. We let the handler
      // implement this fallback.
      return submit([&](handler &CGH) {
        CGH.depends_on(DepEvents);
        return CGH.memcpy(Dest, Src, NumBytes, Offset);
      });
    }
 
    constexpr bool IsDeviceImageScoped = PropertyListT::template has_property<
        ext::oneapi::experimental::device_image_scope_key>();
    return memcpyFromDeviceGlobal(Dest, &Src, IsDeviceImageScoped, NumBytes,
                                  Offset, DepEvents);
  }
 
  template <typename T, typename PropertyListT>
  event memcpy(
      void *Dest,
      const ext::oneapi::experimental::device_global<T, PropertyListT> &Src,
      size_t NumBytes, size_t Offset, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, NumBytes, Offset,
                        std::vector<event>{DepEvent});
  }
 
  template <typename T, typename PropertyListT>
  event memcpy(
      void *Dest,
      const ext::oneapi::experimental::device_global<T, PropertyListT> &Src,
      size_t NumBytes = sizeof(T), size_t Offset = 0,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, NumBytes, Offset, std::vector<event>{});
  }
 
  template <typename T, typename PropertyListT>
  event copy(
      const std::remove_all_extents_t<T> *Src,
      ext::oneapi::experimental::device_global<T, PropertyListT> &Dest,
      size_t Count, size_t StartIndex, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, Count * sizeof(std::remove_all_extents_t<T>),
                        StartIndex * sizeof(std::remove_all_extents_t<T>),
                        DepEvents);
  }
 
  template <typename T, typename PropertyListT>
  event copy(
      const std::remove_all_extents_t<T> *Src,
      ext::oneapi::experimental::device_global<T, PropertyListT> &Dest,
      size_t Count, size_t StartIndex, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, Count * sizeof(std::remove_all_extents_t<T>),
                        StartIndex * sizeof(std::remove_all_extents_t<T>),
                        DepEvent);
  }
 
  template <typename T, typename PropertyListT>
  event copy(
      const std::remove_all_extents_t<T> *Src,
      ext::oneapi::experimental::device_global<T, PropertyListT> &Dest,
      size_t Count = sizeof(T) / sizeof(std::remove_all_extents_t<T>),
      size_t StartIndex = 0,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, Count * sizeof(std::remove_all_extents_t<T>),
                        StartIndex * sizeof(std::remove_all_extents_t<T>));
  }
 
  template <typename T, typename PropertyListT>
  event copy(
      const ext::oneapi::experimental::device_global<T, PropertyListT> &Src,
      std::remove_all_extents_t<T> *Dest, size_t Count, size_t StartIndex,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, Count * sizeof(std::remove_all_extents_t<T>),
                        StartIndex * sizeof(std::remove_all_extents_t<T>),
                        DepEvents);
  }
 
  template <typename T, typename PropertyListT>
  event copy(
      const ext::oneapi::experimental::device_global<T, PropertyListT> &Src,
      std::remove_all_extents_t<T> *Dest, size_t Count, size_t StartIndex,
      event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, Count * sizeof(std::remove_all_extents_t<T>),
                        StartIndex * sizeof(std::remove_all_extents_t<T>),
                        DepEvent);
  }
 
  template <typename T, typename PropertyListT>
  event copy(
      const ext::oneapi::experimental::device_global<T, PropertyListT> &Src,
      std::remove_all_extents_t<T> *Dest,
      size_t Count = sizeof(T) / sizeof(std::remove_all_extents_t<T>),
      size_t StartIndex = 0,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return this->memcpy(Dest, Src, Count * sizeof(std::remove_all_extents_t<T>),
                        StartIndex * sizeof(std::remove_all_extents_t<T>));
  }
 
  event ext_oneapi_copy(
      void *Src, ext::oneapi::experimental::image_mem_handle Dest,
      const ext::oneapi::experimental::image_descriptor &DestImgDesc,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, DestImgDesc); },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent,
      ext::oneapi::experimental::image_mem_handle Dest,
      sycl::range<3> DestOffset,
      const ext::oneapi::experimental::image_descriptor &DestImgDesc,
      sycl::range<3> CopyExtent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset,
                              DestImgDesc, CopyExtent);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, ext::oneapi::experimental::image_mem_handle Dest,
      const ext::oneapi::experimental::image_descriptor &DestImgDesc,
      event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_copy(Src, Dest, DestImgDesc);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent,
      ext::oneapi::experimental::image_mem_handle Dest,
      sycl::range<3> DestOffset,
      const ext::oneapi::experimental::image_descriptor &DestImgDesc,
      sycl::range<3> CopyExtent, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset,
                              DestImgDesc, CopyExtent);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, ext::oneapi::experimental::image_mem_handle Dest,
      const ext::oneapi::experimental::image_descriptor &DestImgDesc,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_copy(Src, Dest, DestImgDesc);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent,
      ext::oneapi::experimental::image_mem_handle Dest,
      sycl::range<3> DestOffset,
      const ext::oneapi::experimental::image_descriptor &DestImgDesc,
      sycl::range<3> CopyExtent, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset,
                              DestImgDesc, CopyExtent);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      ext::oneapi::experimental::image_mem_handle Src, void *Dest,
      const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc); },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      ext::oneapi::experimental::image_mem_handle Src, sycl::range<3> SrcOffset,
      const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest,
      sycl::range<3> DestOffset, sycl::range<3> DestExtent,
      sycl::range<3> CopyExtent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
                              DestExtent, CopyExtent);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      ext::oneapi::experimental::image_mem_handle Src, void *Dest,
      const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
      event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      ext::oneapi::experimental::image_mem_handle Src, sycl::range<3> SrcOffset,
      const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest,
      sycl::range<3> DestOffset, sycl::range<3> DestExtent,
      sycl::range<3> CopyExtent, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
                              DestExtent, CopyExtent);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      ext::oneapi::experimental::image_mem_handle Src, void *Dest,
      const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      ext::oneapi::experimental::image_mem_handle Src, sycl::range<3> SrcOffset,
      const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest,
      sycl::range<3> DestOffset, sycl::range<3> DestExtent,
      sycl::range<3> CopyExtent, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
                              DestExtent, CopyExtent);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, void *Dest,
      const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
      size_t DeviceRowPitch,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, sycl::range<3> SrcOffset, void *Dest,
      sycl::range<3> DestOffset,
      const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
      size_t DeviceRowPitch, sycl::range<3> HostExtent,
      sycl::range<3> CopyExtent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc,
                              DeviceRowPitch, HostExtent, CopyExtent);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, void *Dest,
      const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
      size_t DeviceRowPitch, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      ext::oneapi::experimental::image_mem_handle Src,
      ext::oneapi::experimental::image_mem_handle Dest,
      const ext::oneapi::experimental::image_descriptor &ImageDesc,
      event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_copy(Src, Dest, ImageDesc);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      ext::oneapi::experimental::image_mem_handle Src,
      ext::oneapi::experimental::image_mem_handle Dest,
      const ext::oneapi::experimental::image_descriptor &ImageDesc,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_copy(Src, Dest, ImageDesc);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      ext::oneapi::experimental::image_mem_handle Src,
      ext::oneapi::experimental::image_mem_handle Dest,
      const ext::oneapi::experimental::image_descriptor &ImageDesc,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, ImageDesc); },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, sycl::range<3> SrcOffset, void *Dest,
      sycl::range<3> DestOffset,
      const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
      size_t DeviceRowPitch, sycl::range<3> HostExtent,
      sycl::range<3> CopyExtent, event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc,
                              DeviceRowPitch, HostExtent, CopyExtent);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, void *Dest,
      const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
      size_t DeviceRowPitch, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch);
        },
        CodeLoc);
  }
 
  event ext_oneapi_copy(
      void *Src, sycl::range<3> SrcOffset, void *Dest,
      sycl::range<3> DestOffset,
      const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
      size_t DeviceRowPitch, sycl::range<3> HostExtent,
      sycl::range<3> CopyExtent, const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc,
                              DeviceRowPitch, HostExtent, CopyExtent);
        },
        CodeLoc);
  }
 
  event ext_oneapi_wait_external_semaphore(
      sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle);
        },
        CodeLoc);
  }
 
  event ext_oneapi_wait_external_semaphore(
      sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle,
      event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle);
        },
        CodeLoc);
  }
 
  event ext_oneapi_wait_external_semaphore(
      sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle);
        },
        CodeLoc);
  }
 
  event ext_oneapi_signal_external_semaphore(
      sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle);
        },
        CodeLoc);
  }
 
  event ext_oneapi_signal_external_semaphore(
      sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle,
      event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle);
        },
        CodeLoc);
  }
 
  event ext_oneapi_signal_external_semaphore(
      sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            typename PropertiesT>
  std::enable_if_t<
      ext::oneapi::experimental::is_property_list<PropertiesT>::value, event>
  single_task(
      PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc),
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    static_assert(
        (detail::check_fn_signature<std::remove_reference_t<KernelType>,
                                    void()>::value ||
         detail::check_fn_signature<std::remove_reference_t<KernelType>,
                                    void(kernel_handler)>::value),
        "sycl::queue.single_task() requires a kernel instead of command group. "
        "Use queue.submit() instead");
 
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.template single_task<KernelName, KernelType, PropertiesT>(
              Properties, KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event single_task(
      _KERNELFUNCPARAM(KernelFunc),
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return single_task<KernelName, KernelType>(
        ext::oneapi::experimental::empty_properties_t{}, KernelFunc, CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            typename PropertiesT>
  std::enable_if_t<
      ext::oneapi::experimental::is_property_list<PropertiesT>::value, event>
  single_task(
      event DepEvent, PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc),
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    static_assert(
        (detail::check_fn_signature<std::remove_reference_t<KernelType>,
                                    void()>::value ||
         detail::check_fn_signature<std::remove_reference_t<KernelType>,
                                    void(kernel_handler)>::value),
        "sycl::queue.single_task() requires a kernel instead of command group. "
        "Use queue.submit() instead");
 
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.template single_task<KernelName, KernelType, PropertiesT>(
              Properties, KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event single_task(
      event DepEvent, _KERNELFUNCPARAM(KernelFunc),
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return single_task<KernelName, KernelType>(
        DepEvent, ext::oneapi::experimental::empty_properties_t{}, KernelFunc,
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            typename PropertiesT>
  std::enable_if_t<
      ext::oneapi::experimental::is_property_list<PropertiesT>::value, event>
  single_task(
      const std::vector<event> &DepEvents, PropertiesT Properties,
      _KERNELFUNCPARAM(KernelFunc),
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    static_assert(
        (detail::check_fn_signature<std::remove_reference_t<KernelType>,
                                    void()>::value ||
         detail::check_fn_signature<std::remove_reference_t<KernelType>,
                                    void(kernel_handler)>::value),
        "sycl::queue.single_task() requires a kernel instead of command group. "
        "Use queue.submit() instead");
 
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.template single_task<KernelName, KernelType, PropertiesT>(
              Properties, KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event single_task(
      const std::vector<event> &DepEvents, _KERNELFUNCPARAM(KernelFunc),
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return single_task<KernelName, KernelType>(
        DepEvents, ext::oneapi::experimental::empty_properties_t{}, KernelFunc,
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename... RestT>
  event parallel_for(range<1> Range, RestT &&...Rest) {
    return parallel_for_impl<KernelName>(Range, Rest...);
  }
 
  template <typename KernelName = detail::auto_name, typename... RestT>
  event parallel_for(range<2> Range, RestT &&...Rest) {
    return parallel_for_impl<KernelName>(Range, Rest...);
  }
 
  template <typename KernelName = detail::auto_name, typename... RestT>
  event parallel_for(range<3> Range, RestT &&...Rest) {
    return parallel_for_impl<KernelName>(Range, Rest...);
  }
 
  template <typename KernelName = detail::auto_name, typename... RestT>
  event parallel_for(range<1> Range, event DepEvent, RestT &&...Rest) {
    return parallel_for_impl<KernelName>(Range, DepEvent, Rest...);
  }
 
  template <typename KernelName = detail::auto_name, typename... RestT>
  event parallel_for(range<2> Range, event DepEvent, RestT &&...Rest) {
    return parallel_for_impl<KernelName>(Range, DepEvent, Rest...);
  }
 
  template <typename KernelName = detail::auto_name, typename... RestT>
  event parallel_for(range<3> Range, event DepEvent, RestT &&...Rest) {
    return parallel_for_impl<KernelName>(Range, DepEvent, Rest...);
  }
 
  template <typename KernelName = detail::auto_name, typename... RestT>
  event parallel_for(range<1> Range, const std::vector<event> &DepEvents,
                     RestT &&...Rest) {
    return parallel_for_impl<KernelName>(Range, DepEvents, Rest...);
  }
 
  template <typename KernelName = detail::auto_name, typename... RestT>
  event parallel_for(range<2> Range, const std::vector<event> &DepEvents,
                     RestT &&...Rest) {
    return parallel_for_impl<KernelName>(Range, DepEvents, Rest...);
  }
 
  template <typename KernelName = detail::auto_name, typename... RestT>
  event parallel_for(range<3> Range, const std::vector<event> &DepEvents,
                     RestT &&...Rest) {
    return parallel_for_impl<KernelName>(Range, DepEvents, Rest...);
  }
 
  // While other shortcuts with offsets are able to go through parallel_for(...,
  // RestT &&...Rest), those that accept dependency events vector have to be
  // overloaded to allow implicit construction from an init-list.
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dim>
  event parallel_for(range<Dim> Range, id<Dim> WorkItemOffset,
                     const std::vector<event> &DepEvents,
                     _KERNELFUNCPARAM(KernelFunc)) {
    static_assert(1 <= Dim && Dim <= 3, "Invalid number of dimensions");
    return parallel_for_impl<KernelName>(Range, WorkItemOffset, DepEvents,
                                         KernelFunc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
  event parallel_for_impl(range<Dims> Range, id<Dims> WorkItemOffset,
                          _KERNELFUNCPARAM(KernelFunc)) {
    // Actual code location needs to be captured from KernelInfo object.
    const detail::code_location CodeLoc = {};
    return submit(
        [&](handler &CGH) {
          CGH.template parallel_for<KernelName>(Range, WorkItemOffset,
                                                KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
  event parallel_for_impl(range<Dims> Range, id<Dims> WorkItemOffset,
                          event DepEvent, _KERNELFUNCPARAM(KernelFunc)) {
    // Actual code location needs to be captured from KernelInfo object.
    const detail::code_location CodeLoc = {};
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.template parallel_for<KernelName>(Range, WorkItemOffset,
                                                KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
  event parallel_for_impl(range<Dims> Range, id<Dims> WorkItemOffset,
                          const std::vector<event> &DepEvents,
                          _KERNELFUNCPARAM(KernelFunc)) {
    // Actual code location needs to be captured from KernelInfo object.
    const detail::code_location CodeLoc = {};
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.template parallel_for<KernelName>(Range, WorkItemOffset,
                                                KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, int Dims,
            typename PropertiesT, typename... RestT>
  std::enable_if_t<
      detail::AreAllButLastReductions<RestT...>::value &&
          ext::oneapi::experimental::is_property_list<PropertiesT>::value,
      event>
  parallel_for(nd_range<Dims> Range, PropertiesT Properties, RestT &&...Rest) {
    using KI = sycl::detail::KernelInfo<KernelName>;
    constexpr detail::code_location CodeLoc(
        KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
        KI::getColumnNumber());
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, int Dims,
            typename... RestT>
  std::enable_if_t<detail::AreAllButLastReductions<RestT...>::value, event>
  parallel_for(nd_range<Dims> Range, RestT &&...Rest) {
    return parallel_for<KernelName>(
        Range, ext::oneapi::experimental::empty_properties_t{}, Rest...);
  }
 
  template <typename KernelName = detail::auto_name, int Dims,
            typename... RestT>
  event parallel_for(nd_range<Dims> Range, event DepEvent, RestT &&...Rest) {
    using KI = sycl::detail::KernelInfo<KernelName>;
    constexpr detail::code_location CodeLoc(
        KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
        KI::getColumnNumber());
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.template parallel_for<KernelName>(Range, Rest...);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, int Dims,
            typename... RestT>
  event parallel_for(nd_range<Dims> Range, const std::vector<event> &DepEvents,
                     RestT &&...Rest) {
    using KI = sycl::detail::KernelInfo<KernelName>;
    constexpr detail::code_location CodeLoc(
        KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
        KI::getColumnNumber());
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.template parallel_for<KernelName>(Range, Rest...);
        },
        CodeLoc);
  }
 
  template <typename SrcT, int SrcDims, access_mode SrcMode, target SrcTgt,
            access::placeholder IsPlaceholder, typename DestT>
  event copy(
      accessor<SrcT, SrcDims, SrcMode, SrcTgt, IsPlaceholder> Src,
      std::shared_ptr<DestT> Dest,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.require(Src);
          CGH.copy(Src, Dest);
        },
        CodeLoc);
  }
 
  template <typename SrcT, typename DestT, int DestDims, access_mode DestMode,
            target DestTgt, access::placeholder IsPlaceholder>
  event copy(
      std::shared_ptr<SrcT> Src,
      accessor<DestT, DestDims, DestMode, DestTgt, IsPlaceholder> Dest,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.require(Dest);
          CGH.copy(Src, Dest);
        },
        CodeLoc);
  }
 
  template <typename SrcT, int SrcDims, access_mode SrcMode, target SrcTgt,
            access::placeholder IsPlaceholder, typename DestT>
  event copy(
      accessor<SrcT, SrcDims, SrcMode, SrcTgt, IsPlaceholder> Src, DestT *Dest,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.require(Src);
          CGH.copy(Src, Dest);
        },
        CodeLoc);
  }
 
  template <typename SrcT, typename DestT, int DestDims, access_mode DestMode,
            target DestTgt, access::placeholder IsPlaceholder>
  event copy(
      const SrcT *Src,
      accessor<DestT, DestDims, DestMode, DestTgt, IsPlaceholder> Dest,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.require(Dest);
          CGH.copy(Src, Dest);
        },
        CodeLoc);
  }
 
  template <typename SrcT, int SrcDims, access_mode SrcMode, target SrcTgt,
            access::placeholder IsSrcPlaceholder, typename DestT, int DestDims,
            access_mode DestMode, target DestTgt,
            access::placeholder IsDestPlaceholder>
  event copy(
      accessor<SrcT, SrcDims, SrcMode, SrcTgt, IsSrcPlaceholder> Src,
      accessor<DestT, DestDims, DestMode, DestTgt, IsDestPlaceholder> Dest,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.require(Src);
          CGH.require(Dest);
          CGH.copy(Src, Dest);
        },
        CodeLoc);
  }
 
  template <typename T, int Dims, access_mode Mode, target Tgt,
            access::placeholder IsPlaceholder>
  event update_host(
      accessor<T, Dims, Mode, Tgt, IsPlaceholder> Acc,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.require(Acc);
          CGH.update_host(Acc);
        },
        CodeLoc);
  }
 
  template <typename T, int Dims, access_mode Mode, target Tgt,
            access::placeholder IsPlaceholder>
  event fill(
      accessor<T, Dims, Mode, Tgt, IsPlaceholder> Dest, const T &Src,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.require(Dest);
          CGH.fill<T>(Dest, Src);
        },
        CodeLoc);
  }
 
  bool ext_codeplay_supports_fusion() const;
 
// Clean KERNELFUNC macros.
#undef _KERNELFUNCPARAM
 
  event ext_oneapi_graph(
      ext::oneapi::experimental::command_graph<
          ext::oneapi::experimental::graph_state::executable>
          Graph,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit([&](handler &CGH) { CGH.ext_oneapi_graph(Graph); }, CodeLoc);
  }
 
  event ext_oneapi_graph(
      ext::oneapi::experimental::command_graph<
          ext::oneapi::experimental::graph_state::executable>
          Graph,
      event DepEvent,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.ext_oneapi_graph(Graph);
        },
        CodeLoc);
  }
 
  event ext_oneapi_graph(
      ext::oneapi::experimental::command_graph<
          ext::oneapi::experimental::graph_state::executable>
          Graph,
      const std::vector<event> &DepEvents,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.ext_oneapi_graph(Graph);
        },
        CodeLoc);
  }
 
  bool is_in_order() const;
 
  backend get_backend() const noexcept;
 
  bool ext_oneapi_empty() const;
 
  pi_native_handle getNative(int32_t &NativeHandleDesc) const;
 
  event ext_oneapi_get_last_event() const;
 
  void ext_oneapi_set_external_event(const event &external_event);
 
private:
  std::shared_ptr<detail::queue_impl> impl;
  queue(std::shared_ptr<detail::queue_impl> impl) : impl(impl) {}
 
  template <class Obj>
  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
  template <class T>
  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
 
  template <backend BackendName, class SyclObjectT>
  friend auto get_native(const SyclObjectT &Obj)
      -> backend_return_t<BackendName, SyclObjectT>;
 
#if __SYCL_USE_FALLBACK_ASSERT
  friend event detail::submitAssertCapture(queue &, event &, queue *,
                                           const detail::code_location &);
#endif
 
  event submit_impl(std::function<void(handler &)> CGH,
                    const detail::code_location &CodeLoc);
  event submit_impl(std::function<void(handler &)> CGH, queue secondQueue,
                    const detail::code_location &CodeLoc);
 
  event discard_or_return(const event &Event);
 
  // Function to postprocess submitted command
  // Arguments:
  // bool IsKernel - true if the submitted command was kernel, false otherwise
  // bool KernelUsesAssert - true if submitted kernel uses assert, only
  //                         meaningful when IsKernel is true
  // event &Event - event after which post processing should be executed
  using SubmitPostProcessF = std::function<void(bool, bool, event &)>;
 
  event submit_impl_and_postprocess(std::function<void(handler &)> CGH,
                                    const detail::code_location &CodeLoc,
                                    const SubmitPostProcessF &PostProcess);
  event submit_impl_and_postprocess(std::function<void(handler &)> CGH,
                                    queue secondQueue,
                                    const detail::code_location &CodeLoc,
                                    const SubmitPostProcessF &PostProcess);
 
  template <typename KernelName, int Dims, typename PropertiesT,
            typename... RestT>
  std::enable_if_t<
      detail::AreAllButLastReductions<RestT...>::value &&
          ext::oneapi::experimental::is_property_list<PropertiesT>::value,
      event>
  parallel_for_impl(range<Dims> Range, PropertiesT Properties,
                    RestT &&...Rest) {
    using KI = sycl::detail::KernelInfo<KernelName>;
    constexpr detail::code_location CodeLoc(
        KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
        KI::getColumnNumber());
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
        },
        CodeLoc);
  }
 
  template <typename KernelName, int Dims, typename... RestT>
  std::enable_if_t<detail::AreAllButLastReductions<RestT...>::value, event>
  parallel_for_impl(range<Dims> Range, RestT &&...Rest) {
    return parallel_for_impl<KernelName>(
        Range, ext::oneapi::experimental::empty_properties_t{}, Rest...);
  }
 
  template <typename KernelName, int Dims, typename PropertiesT,
            typename... RestT>
  std::enable_if_t<
      ext::oneapi::experimental::is_property_list<PropertiesT>::value, event>
  parallel_for_impl(range<Dims> Range, event DepEvent, PropertiesT Properties,
                    RestT &&...Rest) {
    using KI = sycl::detail::KernelInfo<KernelName>;
    constexpr detail::code_location CodeLoc(
        KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
        KI::getColumnNumber());
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
        },
        CodeLoc);
  }
 
  template <typename KernelName, int Dims, typename... RestT>
  event parallel_for_impl(range<Dims> Range, event DepEvent, RestT &&...Rest) {
    return parallel_for_impl<KernelName>(
        Range, DepEvent, ext::oneapi::experimental::empty_properties_t{},
        Rest...);
  }
 
  template <typename KernelName, int Dims, typename PropertiesT,
            typename... RestT>
  std::enable_if_t<
      ext::oneapi::experimental::is_property_list<PropertiesT>::value, event>
  parallel_for_impl(range<Dims> Range, const std::vector<event> &DepEvents,
                    PropertiesT Properties, RestT &&...Rest) {
    using KI = sycl::detail::KernelInfo<KernelName>;
    constexpr detail::code_location CodeLoc(
        KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
        KI::getColumnNumber());
    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
        },
        CodeLoc);
  }
 
  template <typename KernelName, int Dims, typename... RestT>
  event parallel_for_impl(range<Dims> Range,
                          const std::vector<event> &DepEvents,
                          RestT &&...Rest) {
    return parallel_for_impl<KernelName>(
        Range, DepEvents, ext::oneapi::experimental::empty_properties_t{},
        Rest...);
  }
 
  event memcpyToDeviceGlobal(void *DeviceGlobalPtr, const void *Src,
                             bool IsDeviceImageScope, size_t NumBytes,
                             size_t Offset,
                             const std::vector<event> &DepEvents);
  event memcpyFromDeviceGlobal(void *Dest, const void *DeviceGlobalPtr,
                               bool IsDeviceImageScope, size_t NumBytes,
                               size_t Offset,
                               const std::vector<event> &DepEvents);
};
 
} // namespace _V1
} // namespace sycl
 
namespace std {
template <> struct __SYCL_EXPORT hash<sycl::queue> {
  size_t operator()(const sycl::queue &Q) const;
};
} // namespace std
 
#if __SYCL_USE_FALLBACK_ASSERT
// Explicitly request format macros
#ifndef __STDC_FORMAT_MACROS
#define __STDC_FORMAT_MACROS 1
#endif
#include <cinttypes>
 
namespace sycl {
inline namespace _V1 {
 
namespace detail {
#define __SYCL_ASSERT_START 1
 
namespace __sycl_service_kernel__ {
class AssertInfoCopier;
} // namespace __sycl_service_kernel__
 
event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue,
                          const detail::code_location &CodeLoc) {
  buffer<detail::AssertHappened, 1> Buffer{1};
 
  event CopierEv, CheckerEv, PostCheckerEv;
  auto CopierCGF = [&](handler &CGH) {
    CGH.depends_on(Event);
 
    auto Acc = Buffer.get_access<access::mode::write>(CGH);
 
    CGH.single_task<__sycl_service_kernel__::AssertInfoCopier>([Acc] {
#if defined(__SYCL_DEVICE_ONLY__) && !defined(__NVPTX__)
      __devicelib_assert_read(&Acc[0]);
#else
      (void)Acc;
#endif // defined(__SYCL_DEVICE_ONLY__) && !defined(__NVPTX__)
    });
  };
  auto CheckerCGF = [&CopierEv, &Buffer](handler &CGH) {
    CGH.depends_on(CopierEv);
    using mode = access::mode;
    using target = access::target;
 
    auto Acc = Buffer.get_access<mode::read, target::host_buffer>(CGH);
 
    CGH.host_task([=] {
      const detail::AssertHappened *AH = &Acc[0];
 
      // Don't use assert here as msvc will insert reference to __imp__wassert
      // which won't be properly resolved in separate compile use-case
#ifndef NDEBUG
      if (AH->Flag == __SYCL_ASSERT_START)
        throw sycl::exception(
            make_error_code(errc::invalid),
            "Internal Error. Invalid value in assert description.");
#endif
 
      if (AH->Flag) {
        const char *Expr = AH->Expr[0] ? AH->Expr : "<unknown expr>";
        const char *File = AH->File[0] ? AH->File : "<unknown file>";
        const char *Func = AH->Func[0] ? AH->Func : "<unknown func>";
 
        fprintf(stderr,
                "%s:%d: %s: global id: [%" PRIu64 ",%" PRIu64 ",%" PRIu64
                "], local id: [%" PRIu64 ",%" PRIu64 ",%" PRIu64 "] "
                "Assertion `%s` failed.\n",
                File, AH->Line, Func, AH->GID0, AH->GID1, AH->GID2, AH->LID0,
                AH->LID1, AH->LID2, Expr);
        fflush(stderr);
        abort(); // no need to release memory as it's abort anyway
      }
    });
  };
 
  if (SecondaryQueue) {
    CopierEv = Self.submit_impl(CopierCGF, *SecondaryQueue, CodeLoc);
    CheckerEv = Self.submit_impl(CheckerCGF, *SecondaryQueue, CodeLoc);
  } else {
    CopierEv = Self.submit_impl(CopierCGF, CodeLoc);
    CheckerEv = Self.submit_impl(CheckerCGF, CodeLoc);
  }
 
  return CheckerEv;
}
#undef __SYCL_ASSERT_START
} // namespace detail
 
} // namespace _V1
} // namespace sycl
#endif // __SYCL_USE_FALLBACK_ASSERT