30 #ifndef __STDC_FORMAT_MACROS
31 #define __STDC_FORMAT_MACROS 1
34 #include <type_traits>
43 #ifdef __SYCL_NONCONST_FUNCTOR__
44 #define _KERNELFUNCPARAM(a) KernelType a
46 #define _KERNELFUNCPARAM(a) const KernelType &a
51 #if defined(SYCL_FALLBACK_ASSERT)
52 #define __SYCL_USE_FALLBACK_ASSERT SYCL_FALLBACK_ASSERT
54 #define __SYCL_USE_FALLBACK_ASSERT 0
65 template <backend BackendName,
class SyclObjectT>
67 -> backend_return_t<BackendName, SyclObjectT>;
72 #if __SYCL_USE_FALLBACK_ASSERT
73 static event submitAssertCapture(queue &, event &, queue *,
74 const detail::code_location &);
104 :
queue(default_selector(), AsyncHandler, PropList) {}
112 template <
typename DeviceSelector,
114 detail::EnableIfSYCL2020DeviceSelectorInvocable<DeviceSelector>>
115 explicit queue(
const DeviceSelector &deviceSelector,
125 template <
typename DeviceSelector,
127 detail::EnableIfSYCL2020DeviceSelectorInvocable<DeviceSelector>>
128 explicit queue(
const DeviceSelector &deviceSelector,
139 template <
typename DeviceSelector,
141 detail::EnableIfSYCL2020DeviceSelectorInvocable<DeviceSelector>>
143 const DeviceSelector &deviceSelector,
155 template <
typename DeviceSelector,
157 detail::EnableIfSYCL2020DeviceSelectorInvocable<DeviceSelector>>
159 const DeviceSelector &deviceSelector,
163 AsyncHandler, propList) {}
171 "use SYCL 2020 device selectors instead.")
184 "use SYCL 2020 device selectors instead.")
186 const
async_handler &AsyncHandler, const property_list &PropList = {})
187 : queue(DeviceSelector.select_device(), AsyncHandler, PropList) {}
193 explicit queue(
const device &SyclDevice,
const property_list &PropList = {})
202 explicit queue(
const device &SyclDevice,
const async_handler &AsyncHandler,
203 const property_list &PropList = {});
212 "use SYCL 2020 device selectors instead.")
225 "use SYCL 2020 device selectors instead.")
245 queue(
const context &SyclContext,
const device &SyclDevice,
246 const async_handler &AsyncHandler,
const property_list &PropList = {});
256 #ifdef __SYCL_INTERNAL_API
257 queue(cl_command_queue ClQueue,
const context &SyclContext,
261 queue(
const queue &RHS) =
default;
263 queue(queue &&RHS) =
default;
265 queue &
operator=(
const queue &RHS) =
default;
275 #ifdef __SYCL_INTERNAL_API
276 cl_command_queue
get()
const;
283 device get_device()
const;
287 "is_host() is deprecated as the host device is no longer supported.")
288 bool is_host() const;
293 template <typename Param>
294 typename detail::is_queue_info_desc<Param>::return_type get_info() const;
311 #if __SYCL_USE_FALLBACK_ASSERT
312 auto PostProcess = [
this, &CodeLoc](
bool IsKernel,
bool KernelUsesAssert,
314 if (IsKernel && !
device_has(aspect::ext_oneapi_native_assert) &&
315 KernelUsesAssert && !
device_has(aspect::accelerator)) {
320 submitAssertCapture(*
this, E,
nullptr, CodeLoc);
324 auto Event = submit_impl_and_postprocess(CGF, CodeLoc, PostProcess);
325 return discard_or_return(Event);
327 auto Event = submit_impl(CGF, CodeLoc);
328 return discard_or_return(Event);
329 #endif // __SYCL_USE_FALLBACK_ASSERT
343 template <
typename T>
347 #if __SYCL_USE_FALLBACK_ASSERT
348 auto PostProcess = [
this, &SecondaryQueue, &CodeLoc](
349 bool IsKernel,
bool KernelUsesAssert,
event &E) {
350 if (IsKernel && !
device_has(aspect::ext_oneapi_native_assert) &&
351 KernelUsesAssert && !
device_has(aspect::accelerator)) {
358 submitAssertCapture(*
this, E, &SecondaryQueue, CodeLoc);
363 submit_impl_and_postprocess(CGF, SecondaryQueue, CodeLoc, PostProcess);
364 return discard_or_return(Event);
366 auto Event = submit_impl(CGF, SecondaryQueue, CodeLoc);
367 return discard_or_return(Event);
368 #endif // __SYCL_USE_FALLBACK_ASSERT
393 return ext_oneapi_submit_barrier(CodeLoc);
406 const std::vector<event> &WaitList
_CODELOCPARAM(&CodeLoc)) {
407 return submit([=](
handler &CGH) {
425 return ext_oneapi_submit_barrier(WaitList, CodeLoc);
450 wait_and_throw_proxy(CodeLoc);
465 void throw_asynchronous();
469 template <
typename PropertyT>
bool has_property() const noexcept;
474 template <typename PropertyT> PropertyT
get_property() const;
483 template <typename T>
event fill(
void *Ptr, const T &Pattern,
size_t Count) {
486 return submit([&](
handler &CGH) { CGH.
fill<T>(Ptr, Pattern, Count); },
498 template <
typename T>
499 event fill(
void *Ptr,
const T &Pattern,
size_t Count,
event DepEvent) {
500 return submit([&](
handler &CGH) {
502 CGH.
fill<T>(Ptr, Pattern, Count);
515 template <
typename T>
516 event fill(
void *Ptr,
const T &Pattern,
size_t Count,
517 const std::vector<event> &DepEvents) {
518 return submit([&](
handler &CGH) {
520 CGH.
fill<T>(Ptr, Pattern, Count);
533 event memset(
void *Ptr,
int Value,
size_t Count);
545 event memset(
void *Ptr,
int Value,
size_t Count,
event DepEvent);
558 event memset(
void *Ptr,
int Value,
size_t Count,
559 const std::vector<event> &DepEvents);
572 event memcpy(
void *Dest,
const void *Src,
size_t Count);
586 event memcpy(
void *Dest,
const void *Src,
size_t Count,
event DepEvent);
601 event memcpy(
void *Dest,
const void *Src,
size_t Count,
602 const std::vector<event> &DepEvents);
616 template <
typename T>
620 return this->
memcpy(Dest, Src, Count *
sizeof(T));
636 template <
typename T>
637 event copy(
const T *Src, T *Dest,
size_t Count,
641 return this->
memcpy(Dest, Src, Count *
sizeof(T), DepEvent);
657 template <
typename T>
658 event copy(
const T *Src, T *Dest,
size_t Count,
659 const std::vector<event> &DepEvents
_CODELOCPARAM(&CodeLoc)) {
662 return this->
memcpy(Dest, Src, Count *
sizeof(T), DepEvents);
682 event mem_advise(const
void *Ptr,
size_t Length,
int Advice);
692 event mem_advise(const
void *Ptr,
size_t Length,
int Advice,
event DepEvent);
703 event mem_advise(const
void *Ptr,
size_t Length,
int Advice,
704 const
std::vector<
event> &DepEvents);
713 event prefetch(const
void *Ptr,
size_t Count) {
716 return submit([=](
handler &CGH) { CGH.
prefetch(Ptr, Count); }, CodeLoc);
728 return submit([=](
handler &CGH) {
744 const std::vector<event> &DepEvents) {
745 return submit([=](
handler &CGH) {
769 template <
typename T =
unsigned char,
770 typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
772 size_t SrcPitch,
size_t Width,
774 return submit([=](
handler &CGH) {
798 template <
typename T =
unsigned char,
799 typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
801 size_t SrcPitch,
size_t Width,
size_t Height,
803 return submit([=](
handler &CGH) {
829 template <
typename T =
unsigned char,
830 typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
832 size_t SrcPitch,
size_t Width,
size_t Height,
833 const std::vector<event> &DepEvents
835 return submit([=](
handler &CGH) {
856 template <
typename T>
858 size_t DestPitch,
size_t Width,
860 return submit([=](
handler &CGH) {
881 template <
typename T>
883 size_t DestPitch,
size_t Width,
size_t Height,
885 return submit([=](
handler &CGH) {
908 template <
typename T>
910 size_t DestPitch,
size_t Width,
size_t Height,
911 const std::vector<event> &DepEvents
913 return submit([=](
handler &CGH) {
935 template <
typename T =
unsigned char,
936 typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
940 return submit([=](
handler &CGH) {
962 template <
typename T =
unsigned char,
963 typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
965 size_t Width,
size_t Height,
967 return submit([=](
handler &CGH) {
991 template <
typename T =
unsigned char,
992 typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
994 void *Dest,
size_t DestPitch,
int Value,
size_t Width,
size_t Height,
995 const std::vector<event> &DepEvents
_CODELOCPARAM(&CodeLoc)) {
996 return submit([=](
handler &CGH) {
1015 template <
typename T>
1018 return submit([=](
handler &CGH) {
1037 template <
typename T>
1039 size_t Width,
size_t Height,
1041 return submit([=](
handler &CGH) {
1062 template <
typename T>
1064 size_t Width,
size_t Height,
1065 const std::vector<event> &DepEvents
1067 return submit([=](
handler &CGH) {
1085 template <
typename T,
typename PropertyListT>
1087 const void *Src,
size_t NumBytes,
size_t Offset,
1088 const std::vector<event> &DepEvents) {
1089 if (
sizeof(T) < Offset + NumBytes)
1091 "Copy to device_global is out of bounds.");
1093 constexpr
bool IsDeviceImageScoped = PropertyListT::template
has_property<
1095 return memcpyToDeviceGlobal(&Dest, Src, IsDeviceImageScoped, NumBytes,
1111 template <
typename T,
typename PropertyListT>
1113 const void *Src,
size_t NumBytes,
size_t Offset,
1115 return this->
memcpy(Dest, Src, NumBytes, Offset,
1116 std::vector<event>{DepEvent});
1129 template <
typename T,
typename PropertyListT>
1131 const void *Src,
size_t NumBytes =
sizeof(T),
1132 size_t Offset = 0) {
1133 return this->
memcpy(Dest, Src, NumBytes, Offset, std::vector<event>{});
1148 template <
typename T,
typename PropertyListT>
1152 size_t NumBytes,
size_t Offset,
const std::vector<event> &DepEvents) {
1153 if (
sizeof(T) < Offset + NumBytes)
1155 "Copy from device_global is out of bounds.");
1157 constexpr
bool IsDeviceImageScoped = PropertyListT::template
has_property<
1159 return memcpyFromDeviceGlobal(Dest, &Src, IsDeviceImageScoped, NumBytes,
1175 template <
typename T,
typename PropertyListT>
1179 size_t NumBytes,
size_t Offset,
event DepEvent) {
1180 return this->
memcpy(Dest, Src, NumBytes, Offset,
1181 std::vector<event>{DepEvent});
1194 template <
typename T,
typename PropertyListT>
1198 size_t NumBytes =
sizeof(T),
size_t Offset = 0) {
1199 return this->
memcpy(Dest, Src, NumBytes, Offset, std::vector<event>{});
1215 template <
typename T,
typename PropertyListT>
1216 event copy(
const std::remove_all_extents_t<T> *Src,
1218 size_t Count,
size_t StartIndex,
1219 const std::vector<event> &DepEvents) {
1220 return this->
memcpy(Dest, Src, Count *
sizeof(std::remove_all_extents_t<T>),
1221 StartIndex *
sizeof(std::remove_all_extents_t<T>),
1238 template <
typename T,
typename PropertyListT>
1239 event copy(
const std::remove_all_extents_t<T> *Src,
1241 size_t Count,
size_t StartIndex,
event DepEvent) {
1242 return this->
memcpy(Dest, Src, Count *
sizeof(std::remove_all_extents_t<T>),
1243 StartIndex *
sizeof(std::remove_all_extents_t<T>),
1258 template <
typename T,
typename PropertyListT>
1259 event copy(
const std::remove_all_extents_t<T> *Src,
1261 size_t Count =
sizeof(T) /
sizeof(std::remove_all_extents_t<T>),
1262 size_t StartIndex = 0) {
1263 return this->
memcpy(Dest, Src, Count *
sizeof(std::remove_all_extents_t<T>),
1264 StartIndex *
sizeof(std::remove_all_extents_t<T>));
1280 template <
typename T,
typename PropertyListT>
1283 std::remove_all_extents_t<T> *Dest,
size_t Count,
size_t StartIndex,
1284 const std::vector<event> &DepEvents) {
1285 return this->
memcpy(Dest, Src, Count *
sizeof(std::remove_all_extents_t<T>),
1286 StartIndex *
sizeof(std::remove_all_extents_t<T>),
1303 template <
typename T,
typename PropertyListT>
1306 std::remove_all_extents_t<T> *Dest,
size_t Count,
size_t StartIndex,
1308 return this->
memcpy(Dest, Src, Count *
sizeof(std::remove_all_extents_t<T>),
1309 StartIndex *
sizeof(std::remove_all_extents_t<T>),
1324 template <
typename T,
typename PropertyListT>
1327 std::remove_all_extents_t<T> *Dest,
1328 size_t Count =
sizeof(T) /
sizeof(std::remove_all_extents_t<T>),
1329 size_t StartIndex = 0) {
1330 return this->
memcpy(Dest, Src, Count *
sizeof(std::remove_all_extents_t<T>),
1331 StartIndex *
sizeof(std::remove_all_extents_t<T>));
1340 typename PropertiesT>
1349 void(kernel_handler)>::value),
1350 "sycl::queue.single_task() requires a kernel instead of command group. "
1351 "Use queue.submit() instead");
1356 CGH.template single_task<KernelName, KernelType, PropertiesT>(
1366 template <
typename KernelName = detail::auto_name,
typename KernelType>
1368 return single_task<KernelName, KernelType>(
1380 typename PropertiesT>
1389 void(kernel_handler)>::value),
1390 "sycl::queue.single_task() requires a kernel instead of command group. "
1391 "Use queue.submit() instead");
1397 CGH.template single_task<KernelName, KernelType, PropertiesT>(
1408 template <
typename KernelName = detail::auto_name,
typename KernelType>
1411 return single_task<KernelName, KernelType>(
1424 typename PropertiesT>
1427 single_task(
const std::vector<event> &DepEvents, PropertiesT Properties,
1433 void(kernel_handler)>::value),
1434 "sycl::queue.single_task() requires a kernel instead of command group. "
1435 "Use queue.submit() instead");
1441 CGH.template single_task<KernelName, KernelType, PropertiesT>(
1453 template <
typename KernelName = detail::auto_name,
typename KernelType>
1456 return single_task<KernelName, KernelType>(
1469 return parallel_for_impl<KernelName>(Range, Rest...);
1480 return parallel_for_impl<KernelName>(Range, Rest...);
1491 return parallel_for_impl<KernelName>(Range, Rest...);
1503 return parallel_for_impl<KernelName>(Range, DepEvent, Rest...);
1515 return parallel_for_impl<KernelName>(Range, DepEvent, Rest...);
1527 return parallel_for_impl<KernelName>(Range, DepEvent, Rest...);
1541 return parallel_for_impl<KernelName>(Range, DepEvents, Rest...);
1555 return parallel_for_impl<KernelName>(Range, DepEvents, Rest...);
1569 return parallel_for_impl<KernelName>(Range, DepEvents, Rest...);
1585 const std::vector<event> &DepEvents,
1587 static_assert(1 <= Dim && Dim <= 3,
"Invalid number of dimensions");
1588 return parallel_for_impl<KernelName>(Range, WorkItemOffset, DepEvents,
1602 event parallel_for_impl(
range<Dims> Range,
id<Dims> WorkItemOffset,
1608 CGH.template parallel_for<KernelName>(Range, WorkItemOffset,
1625 event parallel_for_impl(
range<Dims> Range,
id<Dims> WorkItemOffset,
1632 CGH.template parallel_for<KernelName>(Range, WorkItemOffset,
1650 event parallel_for_impl(
range<Dims> Range,
id<Dims> WorkItemOffset,
1651 const
std::vector<
event> &DepEvents,
1658 CGH.template parallel_for<KernelName>(Range, WorkItemOffset,
1672 typename PropertiesT,
typename... RestT>
1678 using KI = sycl::detail::KernelInfo<KernelName>;
1680 KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
1681 KI::getColumnNumber());
1685 CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
1700 return parallel_for<KernelName>(
1715 using KI = sycl::detail::KernelInfo<KernelName>;
1717 KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
1718 KI::getColumnNumber());
1723 CGH.template parallel_for<KernelName>(Range, Rest...);
1740 using KI = sycl::detail::KernelInfo<KernelName>;
1742 KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
1743 KI::getColumnNumber());
1748 CGH.template parallel_for<KernelName>(Range, Rest...);
1763 return submit([&](
handler &CGH) {
1765 CGH.
copy(Src, Dest);
1775 template <
typename SrcT,
typename DestT,
int DestDims,
access_mode DestMode,
1777 event copy(std::shared_ptr<SrcT> Src,
1780 return submit([&](
handler &CGH) {
1782 CGH.
copy(Src, Dest);
1796 return submit([&](
handler &CGH) {
1798 CGH.
copy(Src, Dest);
1808 template <
typename SrcT,
typename DestT,
int DestDims,
access_mode DestMode,
1813 return submit([&](
handler &CGH) {
1815 CGH.
copy(Src, Dest);
1833 return submit([&](
handler &CGH) {
1836 CGH.
copy(Src, Dest);
1849 return submit([&](
handler &CGH) {
1865 return submit([&](
handler &CGH) {
1867 CGH.
fill<T>(Dest, Src);
1877 bool ext_codeplay_supports_fusion()
const;
1880 #undef _KERNELFUNCPARAM
1885 bool is_in_order()
const;
1890 backend get_backend() const noexcept;
1896 bool ext_oneapi_empty() const;
1901 std::shared_ptr<detail::queue_impl> impl;
1902 queue(
std::shared_ptr<detail::queue_impl> impl) : impl(impl) {}
1904 template <
class Obj>
1909 template <backend BackendName,
class SyclObjectT>
1910 friend auto get_native(
const SyclObjectT &Obj)
1911 -> backend_return_t<BackendName, SyclObjectT>;
1913 #if __SYCL_USE_FALLBACK_ASSERT
1914 friend event detail::submitAssertCapture(queue &, event &, queue *,
1915 const detail::code_location &);
1919 event submit_impl(std::function<
void(handler &)> CGH,
1920 const detail::code_location &CodeLoc);
1922 event submit_impl(std::function<
void(handler &)> CGH, queue secondQueue,
1923 const detail::code_location &CodeLoc);
1927 event discard_or_return(
const event &Event);
1935 using SubmitPostProcessF = std::function<void(
bool,
bool, event &)>;
1942 event submit_impl_and_postprocess(std::function<
void(handler &)> CGH,
1943 const detail::code_location &CodeLoc,
1944 const SubmitPostProcessF &PostProcess);
1951 event submit_impl_and_postprocess(std::function<
void(handler &)> CGH,
1953 const detail::code_location &CodeLoc,
1954 const SubmitPostProcessF &PostProcess);
1962 template <
typename KernelName,
int Dims,
typename PropertiesT,
1965 detail::AreAllButLastReductions<RestT...>::value &&
1966 ext::oneapi::experimental::is_property_list<PropertiesT>::value,
1968 parallel_for_impl(range<Dims> Range, PropertiesT Properties,
1970 using KI = sycl::detail::KernelInfo<KernelName>;
1971 constexpr detail::code_location CodeLoc(
1972 KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
1973 KI::getColumnNumber());
1974 detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1977 CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
1987 template <
typename KernelName,
int Dims,
typename... RestT>
1989 parallel_for_impl(range<Dims> Range, RestT &&...Rest) {
1990 return parallel_for_impl<KernelName>(
2002 template <
typename KernelName,
int Dims,
typename PropertiesT,
2005 ext::oneapi::experimental::is_property_list<PropertiesT>::value,
event>
2006 parallel_for_impl(range<Dims> Range, event DepEvent, PropertiesT Properties,
2008 using KI = sycl::detail::KernelInfo<KernelName>;
2009 constexpr detail::code_location CodeLoc(
2010 KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
2011 KI::getColumnNumber());
2012 detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
2015 CGH.depends_on(DepEvent);
2016 CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
2027 template <
typename KernelName,
int Dims,
typename... RestT>
2028 event parallel_for_impl(range<Dims> Range, event DepEvent, RestT &&...Rest) {
2029 return parallel_for_impl<KernelName>(
2042 template <
typename KernelName,
int Dims,
typename PropertiesT,
2045 ext::oneapi::experimental::is_property_list<PropertiesT>::value,
event>
2046 parallel_for_impl(range<Dims> Range,
const std::vector<event> &DepEvents,
2047 PropertiesT Properties, RestT &&...Rest) {
2048 using KI = sycl::detail::KernelInfo<KernelName>;
2049 constexpr detail::code_location CodeLoc(
2050 KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(),
2051 KI::getColumnNumber());
2052 detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
2055 CGH.depends_on(DepEvents);
2056 CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
2068 template <
typename KernelName,
int Dims,
typename... RestT>
2069 event parallel_for_impl(range<Dims> Range,
2070 const std::vector<event> &DepEvents,
2072 return parallel_for_impl<KernelName>(
2077 buffer<detail::AssertHappened, 1> &getAssertHappenedBuffer();
2079 event memcpyToDeviceGlobal(
void *DeviceGlobalPtr,
const void *Src,
2080 bool IsDeviceImageScope,
size_t NumBytes,
2082 const std::vector<event> &DepEvents);
2083 event memcpyFromDeviceGlobal(
void *Dest,
const void *DeviceGlobalPtr,
2084 bool IsDeviceImageScope,
size_t NumBytes,
2086 const std::vector<event> &DepEvents);
2090 #if __SYCL_USE_FALLBACK_ASSERT
2091 #define __SYCL_ASSERT_START 1
2103 event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue,
2104 const detail::code_location &CodeLoc) {
2105 using AHBufT = buffer<detail::AssertHappened, 1>;
2107 AHBufT &Buffer = Self.getAssertHappenedBuffer();
2109 event CopierEv, CheckerEv, PostCheckerEv;
2110 auto CopierCGF = [&](handler &CGH) {
2111 CGH.depends_on(Event);
2115 CGH.single_task<__sycl_service_kernel__::AssertInfoCopier>([Acc] {
2116 #if defined(__SYCL_DEVICE_ONLY__) && !defined(__NVPTX__)
2117 __devicelib_assert_read(&Acc[0]);
2120 #endif // defined(__SYCL_DEVICE_ONLY__) && !defined(__NVPTX__)
2123 auto CheckerCGF = [&CopierEv, &Buffer](handler &CGH) {
2124 CGH.depends_on(CopierEv);
2128 auto Acc = Buffer.get_access<mode::read, target::host_buffer>(CGH);
2131 const detail::AssertHappened *AH = &Acc[0];
2136 if (AH->Flag == __SYCL_ASSERT_START)
2137 throw sycl::runtime_error(
2138 "Internal Error. Invalid value in assert description.",
2139 PI_ERROR_INVALID_VALUE);
2143 const char *Expr = AH->Expr[0] ? AH->Expr :
"<unknown expr>";
2144 const char *File = AH->File[0] ? AH->File :
"<unknown file>";
2145 const char *Func = AH->Func[0] ? AH->Func :
"<unknown func>";
2148 "%s:%d: %s: global id: [%" PRIu64
",%" PRIu64
",%" PRIu64
2149 "], local id: [%" PRIu64
",%" PRIu64
",%" PRIu64
"] "
2150 "Assertion `%s` failed.\n",
2151 File, AH->Line, Func, AH->GID0, AH->GID1, AH->GID2, AH->LID0,
2152 AH->LID1, AH->LID2, Expr);
2159 if (SecondaryQueue) {
2160 CopierEv = Self.submit_impl(CopierCGF, *SecondaryQueue, CodeLoc);
2161 CheckerEv = Self.submit_impl(CheckerCGF, *SecondaryQueue, CodeLoc);
2163 CopierEv = Self.submit_impl(CopierCGF, CodeLoc);
2164 CheckerEv = Self.submit_impl(CheckerCGF, CodeLoc);
2169 #undef __SYCL_ASSERT_START
2170 #endif // __SYCL_USE_FALLBACK_ASSERT
2177 template <>
struct hash<
sycl::queue> {
2179 return std::hash<std::shared_ptr<sycl::detail::queue_impl>>()(
2185 #undef __SYCL_USE_FALLBACK_ASSERT