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 <CL/sycl/backend_types.hpp>
#include <CL/sycl/detail/assert_happened.hpp>
#include <CL/sycl/detail/common.hpp>
#include <CL/sycl/detail/export.hpp>
#include <CL/sycl/detail/service_kernel_names.hpp>
#include <CL/sycl/device.hpp>
#include <CL/sycl/device_selector.hpp>
#include <CL/sycl/event.hpp>
#include <CL/sycl/exception_list.hpp>
#include <CL/sycl/handler.hpp>
#include <CL/sycl/info/info_desc.hpp>
#include <CL/sycl/property_list.hpp>
#include <CL/sycl/stl.hpp>
 
#include <inttypes.h>
#include <utility>
 
// having _TWO_ mid-param #ifdefs makes the functions very difficult to read.
// Here we simplify the &CodeLoc declaration to be _CODELOCPARAM(&CodeLoc) and
// _CODELOCARG(&CodeLoc) Similarly, the KernelFunc param is simplified to be
// _KERNELFUNCPARAM(KernelFunc) Once the queue kernel functions are defined,
// these macros are #undef immediately.
 
// replace _CODELOCPARAM(&CodeLoc) with nothing
// or :   , const detail::code_location &CodeLoc =
// detail::code_location::current()
// replace _CODELOCARG(&CodeLoc) with nothing
// or :  const detail::code_location &CodeLoc = {}
 
#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
#define _CODELOCONLYPARAM(a)                                                   \
  const detail::code_location a = detail::code_location::current()
#define _CODELOCPARAM(a)                                                       \
  , const detail::code_location a = detail::code_location::current()
 
#define _CODELOCARG(a)
#define _CODELOCFW(a) , a
#else
#define _CODELOCONLYPARAM(a)
#define _CODELOCPARAM(a)
 
#define _CODELOCARG(a) const detail::code_location a = {}
#define _CODELOCFW(a)
#endif
 
// 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
 
// Helper macro to identify if fallback assert is needed
// FIXME remove __NVPTX__ condition once devicelib supports CUDA
#if !defined(SYCL_DISABLE_FALLBACK_ASSERT) && !defined(__NVPTX__)
#define __SYCL_USE_FALLBACK_ASSERT 1
#else
#define __SYCL_USE_FALLBACK_ASSERT 0
#endif
 
__SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
 
// Forward declaration
class context;
class device;
class queue;
 
namespace detail {
class queue_impl;
#if __SYCL_USE_FALLBACK_ASSERT
static event submitAssertCapture(queue &, event &, queue *,
                                 const detail::code_location &);
#endif
} // namespace detail
 
class __SYCL_EXPORT queue {
public:
  explicit queue(const property_list &PropList = {})
      : queue(default_selector(), async_handler{}, PropList) {}
 
  queue(const async_handler &AsyncHandler, const property_list &PropList = {})
      : queue(default_selector(), AsyncHandler, PropList) {}
 
  queue(const device_selector &DeviceSelector,
        const property_list &PropList = {})
      : queue(DeviceSelector.select_device(), async_handler{}, PropList) {}
 
  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, async_handler{}, PropList) {}
 
  explicit queue(const device &SyclDevice, const async_handler &AsyncHandler,
                 const property_list &PropList = {});
 
  queue(const context &SyclContext, const device_selector &DeviceSelector,
        const property_list &PropList = {});
 
  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 = {});
 
  __SYCL2020_DEPRECATED("OpenCL interop APIs are deprecated")
  queue(cl_command_queue ClQueue, const context &SyclContext,
        const async_handler &AsyncHandler = {});
 
  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); }
 
  __SYCL2020_DEPRECATED("OpenCL interop APIs are deprecated")
  cl_command_queue get() const;
 
  context get_context() const;
 
  device get_device() const;
 
  bool is_host() const;
 
  template <info::queue param>
  typename info::param_traits<info::queue, param>::return_type get_info() const;
 
  // 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> event submit(T CGF _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
 
    event Event;
 
#if __SYCL_USE_FALLBACK_ASSERT
    if (!is_host()) {
      auto PostProcess = [this, &CodeLoc](bool IsKernel, bool KernelUsesAssert,
                                          event &E) {
        if (IsKernel && !device_has(aspect::ext_oneapi_native_assert) &&
            KernelUsesAssert) {
          // __devicelib_assert_fail isn't supported by Device-side Runtime
          // Linking against fallback impl of __devicelib_assert_fail is
          // performed by program manager class
          submitAssertCapture(*this, E, /* SecondaryQueue = */ nullptr,
                              CodeLoc);
        }
      };
 
      Event = submit_impl_and_postprocess(CGF, CodeLoc, PostProcess);
    } else
#endif // __SYCL_USE_FALLBACK_ASSERT
    {
      Event = submit_impl(CGF, CodeLoc);
    }
 
    return Event;
  }
 
  template <typename T>
  event submit(T CGF, queue &SecondaryQueue _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
 
    event Event;
 
#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) {
        // __devicelib_assert_fail isn't supported by Device-side Runtime
        // Linking against fallback impl of __devicelib_assert_fail is performed
        // by program manager class
        submitAssertCapture(*this, E, /* SecondaryQueue = */ nullptr, CodeLoc);
      }
    };
 
    Event =
        submit_impl_and_postprocess(CGF, SecondaryQueue, CodeLoc, PostProcess);
#else
    Event = submit_impl(CGF, SecondaryQueue, CodeLoc);
#endif // __SYCL_USE_FALLBACK_ASSERT
 
    return Event;
  }
 
  event ext_oneapi_submit_barrier(_CODELOCONLYPARAM(&CodeLoc)) {
    return submit(
        [=](handler &CGH) { CGH.ext_oneapi_barrier(); } _CODELOCFW(CodeLoc));
  }
 
  __SYCL2020_DEPRECATED("use 'ext_oneapi_submit_barrier' instead")
  event submit_barrier(_CODELOCONLYPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return ext_oneapi_submit_barrier(CodeLoc);
  }
 
  event ext_oneapi_submit_barrier(
      const std::vector<event> &WaitList _CODELOCPARAM(&CodeLoc)) {
    return submit([=](handler &CGH) {
      CGH.ext_oneapi_barrier(WaitList);
    } _CODELOCFW(CodeLoc));
  }
 
  __SYCL2020_DEPRECATED("use 'ext_oneapi_submit_barrier' instead")
  event
  submit_barrier(const std::vector<event> &WaitList _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return ext_oneapi_submit_barrier(WaitList, CodeLoc);
  }
 
  void wait(_CODELOCONLYPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
 
    wait_proxy(CodeLoc);
  }
 
  void wait_and_throw(_CODELOCONLYPARAM(&CodeLoc)) {
    _CODELOCARG(&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;
 
  template <typename PropertyT> PropertyT get_property() const;
 
  template <typename T> event fill(void *Ptr, const T &Pattern, size_t Count) {
    return submit([&](handler &CGH) { CGH.fill<T>(Ptr, Pattern, Count); });
  }
 
  template <typename T>
  event fill(void *Ptr, const T &Pattern, size_t Count, event DepEvent) {
    return submit([&](handler &CGH) {
      CGH.depends_on(DepEvent);
      CGH.fill<T>(Ptr, Pattern, Count);
    });
  }
 
  template <typename T>
  event fill(void *Ptr, const T &Pattern, size_t Count,
             const std::vector<event> &DepEvents) {
    return submit([&](handler &CGH) {
      CGH.depends_on(DepEvents);
      CGH.fill<T>(Ptr, Pattern, Count);
    });
  }
 
  event memset(void *Ptr, int Value, size_t Count);
 
  event memset(void *Ptr, int Value, size_t Count, event DepEvent);
 
  event memset(void *Ptr, int Value, size_t Count,
               const std::vector<event> &DepEvents);
 
  event memcpy(void *Dest, const void *Src, size_t Count);
 
  event memcpy(void *Dest, const void *Src, size_t Count, event DepEvent);
 
  event memcpy(void *Dest, const void *Src, size_t Count,
               const std::vector<event> &DepEvents);
 
  template <typename T> event copy(const T *Src, T *Dest, size_t Count) {
    return this->memcpy(Dest, Src, Count * sizeof(T));
  }
 
  template <typename T>
  event copy(const T *Src, T *Dest, size_t Count, event DepEvent) {
    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) {
    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);
 
  event mem_advise(const void *Ptr, size_t Length, int Advice);
 
  event mem_advise(const void *Ptr, size_t Length, int Advice, event DepEvent);
 
  event mem_advise(const void *Ptr, size_t Length, int Advice,
                   const std::vector<event> &DepEvents);
 
  event prefetch(const void *Ptr, size_t Count) {
    return submit([=](handler &CGH) { CGH.prefetch(Ptr, Count); });
  }
 
  event prefetch(const void *Ptr, size_t Count, event DepEvent) {
    return submit([=](handler &CGH) {
      CGH.depends_on(DepEvent);
      CGH.prefetch(Ptr, Count);
    });
  }
 
  event prefetch(const void *Ptr, size_t Count,
                 const std::vector<event> &DepEvents) {
    return submit([=](handler &CGH) {
      CGH.depends_on(DepEvents);
      CGH.prefetch(Ptr, Count);
    });
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event single_task(_KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
 
    return submit(
        [&](handler &CGH) {
          CGH.template single_task<KernelName, KernelType>(KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event single_task(event DepEvent,
                    _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.template single_task<KernelName, KernelType>(KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event single_task(const std::vector<event> &DepEvents,
                    _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.template single_task<KernelName, KernelType>(KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<1> NumWorkItems,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<2> NumWorkItems,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<3> NumWorkItems,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<1> NumWorkItems, event DepEvent,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return parallel_for_impl<KernelName>(NumWorkItems, DepEvent, KernelFunc,
                                         CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<2> NumWorkItems, event DepEvent,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return parallel_for_impl<KernelName>(NumWorkItems, DepEvent, KernelFunc,
                                         CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<3> NumWorkItems, event DepEvent,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return parallel_for_impl<KernelName>(NumWorkItems, DepEvent, KernelFunc,
                                         CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<1> NumWorkItems, const std::vector<event> &DepEvents,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return parallel_for_impl<KernelName>(NumWorkItems, DepEvents, KernelFunc,
                                         CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<2> NumWorkItems, const std::vector<event> &DepEvents,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return parallel_for_impl<KernelName>(NumWorkItems, DepEvents, KernelFunc,
                                         CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<3> NumWorkItems, const std::vector<event> &DepEvents,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return parallel_for_impl<KernelName>(NumWorkItems, DepEvents, KernelFunc,
                                         CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.template parallel_for<KernelName, KernelType>(
              NumWorkItems, WorkItemOffset, KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
                     event DepEvent,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.template parallel_for<KernelName, KernelType>(
              NumWorkItems, WorkItemOffset, KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
                     const std::vector<event> &DepEvents,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.template parallel_for<KernelName, KernelType>(
              NumWorkItems, WorkItemOffset, KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  event parallel_for(nd_range<Dims> ExecutionRange,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.template parallel_for<KernelName, KernelType>(ExecutionRange,
                                                            KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  event parallel_for(nd_range<Dims> ExecutionRange, event DepEvent,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.template parallel_for<KernelName, KernelType>(ExecutionRange,
                                                            KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  event parallel_for(nd_range<Dims> ExecutionRange,
                     const std::vector<event> &DepEvents,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.template parallel_for<KernelName, KernelType>(ExecutionRange,
                                                            KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims, typename Reduction>
  event parallel_for(nd_range<Dims> ExecutionRange, Reduction Redu,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
    _CODELOCARG(&CodeLoc);
    return submit(
        [&](handler &CGH) {
          CGH.template parallel_for<KernelName, KernelType, Dims, Reduction>(
              ExecutionRange, Redu, KernelFunc);
        },
        CodeLoc);
  }
 
// Clean up CODELOC and KERNELFUNC macros.
#undef _CODELOCPARAM
#undef _CODELOCONLYPARAM
#undef _CODELOCARG
#undef _CODELOCFW
#undef _KERNELFUNCPARAM
 
  bool is_in_order() const;
 
  backend get_backend() const noexcept;
 
  template <backend BackendName>
  __SYCL_DEPRECATED("Use SYCL 2020 sycl::get_native free function")
  auto get_native() const -> typename interop<BackendName, queue>::type {
    return reinterpret_cast<typename interop<BackendName, queue>::type>(
        getNative());
  }
 
private:
  pi_native_handle getNative() const;
 
  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);
 
#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);
 
  // 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 = detail::auto_name, typename KernelType,
            int Dims>
  event parallel_for_impl(
      range<Dims> NumWorkItems, KernelType KernelFunc,
      const detail::code_location &CodeLoc = detail::code_location::current()) {
    return submit(
        [&](handler &CGH) {
          CGH.template parallel_for<KernelName, KernelType>(NumWorkItems,
                                                            KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  event parallel_for_impl(range<Dims> NumWorkItems, event DepEvent,
                          KernelType KernelFunc,
                          const detail::code_location &CodeLoc) {
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvent);
          CGH.template parallel_for<KernelName, KernelType>(NumWorkItems,
                                                            KernelFunc);
        },
        CodeLoc);
  }
 
  template <typename KernelName = detail::auto_name, typename KernelType,
            int Dims>
  event parallel_for_impl(range<Dims> NumWorkItems,
                          const std::vector<event> &DepEvents,
                          KernelType KernelFunc,
                          const detail::code_location &CodeLoc) {
    return submit(
        [&](handler &CGH) {
          CGH.depends_on(DepEvents);
          CGH.template parallel_for<KernelName, KernelType>(NumWorkItems,
                                                            KernelFunc);
        },
        CodeLoc);
  }
 
  buffer<detail::AssertHappened, 1> &getAssertHappenedBuffer();
};
 
namespace detail {
#if __SYCL_USE_FALLBACK_ASSERT
#define __SYCL_ASSERT_START 1
 
event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue,
                          const detail::code_location &CodeLoc) {
  using AHBufT = buffer<detail::AssertHappened, 1>;
 
  AHBufT &Buffer = Self.getAssertHappenedBuffer();
 
  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] {
#ifdef __SYCL_DEVICE_ONLY__
      __devicelib_assert_read(&Acc[0]);
#else
      (void)Acc;
#endif // __SYCL_DEVICE_ONLY__
    });
  };
  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::runtime_error(
            "Internal Error. Invalid value in assert description.",
            PI_INVALID_VALUE);
#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);
        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
#endif // __SYCL_USE_FALLBACK_ASSERT
} // namespace detail
 
} // namespace sycl
} // __SYCL_INLINE_NAMESPACE(cl)
 
namespace std {
template <> struct hash<cl::sycl::queue> {
  size_t operator()(const cl::sycl::queue &Q) const {
    return std::hash<std::shared_ptr<cl::sycl::detail::queue_impl>>()(
        cl::sycl::detail::getSyclObjImpl(Q));
  }
};
} // namespace std
 
#undef __SYCL_USE_FALLBACK_ASSERT