DPC++ Runtime
Runtime libraries for oneAPI DPC++
nd_item.hpp
Go to the documentation of this file.
1 //==--------- nd_item.hpp --- SYCL iteration nd_item -----------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <CL/__spirv/spirv_ops.hpp> // for __spirv_ControlBarrier
12 #include <CL/__spirv/spirv_types.hpp> // for Scope
13 #include <CL/__spirv/spirv_vars.hpp> // for initLocalInvocationId
14 #include <sycl/access/access.hpp> // for mode, fence_space
15 #include <sycl/detail/defines.hpp> // for __SYCL_ASSUME_INT
16 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED, __SY...
17 #include <sycl/detail/generic_type_traits.hpp> // for ConvertToOpenCLType_t
18 #include <sycl/detail/helpers.hpp> // for getSPIRVMemorySemanticsMask
19 #include <sycl/detail/type_traits.hpp> // for is_bool, change_base_...
20 #include <sycl/device_event.hpp> // for device_event
21 #include <sycl/exception.hpp> // for make_error_code, errc, exce...
22 #include <sycl/group.hpp> // for group
23 #include <sycl/id.hpp> // for id
24 #include <sycl/item.hpp> // for item
25 #include <sycl/nd_range.hpp> // for nd_range
26 #include <sycl/pointers.hpp> // for decorated_global_ptr, decor...
27 #include <sycl/range.hpp> // for range
28 #include <sycl/sub_group.hpp> // for sub_group
29 
30 #include <cstddef> // for size_t
31 #include <stdint.h> // for uint32_t
32 #include <type_traits> // for enable_if_t, remove_const_t
33 
34 namespace sycl {
35 inline namespace _V1 {
36 namespace detail {
37 class Builder;
38 }
39 
40 namespace ext::oneapi::experimental {
41 template <int Dimensions> class root_group;
42 }
43 
44 #if __INTEL_PREVIEW_BREAKING_CHANGES
49 template <int Dimensions = 1> class nd_item {
50 public:
51  static constexpr int dimensions = Dimensions;
52 
53  id<Dimensions> get_global_id() const {
54 #ifdef __SYCL_DEVICE_ONLY__
55  return __spirv::initGlobalInvocationId<Dimensions, id<Dimensions>>();
56 #else
57  throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
58  "nd_item methods can't be invoked on the host");
59  return {};
60 #endif
61  }
62 
63  size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const {
64  size_t Id = get_global_id()[Dimension];
66  return Id;
67  }
68 
70  size_t LinId = 0;
71  id<Dimensions> Index = get_global_id();
72  range<Dimensions> Extent = get_global_range();
73  id<Dimensions> Offset = get_offset();
74  if (1 == Dimensions) {
75  LinId = Index[0] - Offset[0];
76  } else if (2 == Dimensions) {
77  LinId = (Index[0] - Offset[0]) * Extent[1] + Index[1] - Offset[1];
78  } else {
79  LinId = (Index[0] - Offset[0]) * Extent[1] * Extent[2] +
80  (Index[1] - Offset[1]) * Extent[2] + Index[2] - Offset[2];
81  }
82  __SYCL_ASSUME_INT(LinId);
83  return LinId;
84  }
85 
86  id<Dimensions> get_local_id() const {
87 #ifdef __SYCL_DEVICE_ONLY__
88  return __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>();
89 #else
90  throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
91  "nd_item methods can't be invoked on the host");
92  return {};
93 #endif
94  }
95 
96  size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const {
97  size_t Id = get_local_id()[Dimension];
99  return Id;
100  }
101 
102  size_t get_local_linear_id() const {
103  size_t LinId = 0;
104  id<Dimensions> Index = get_local_id();
105  range<Dimensions> Extent = get_local_range();
106  if (1 == Dimensions) {
107  LinId = Index[0];
108  } else if (2 == Dimensions) {
109  LinId = Index[0] * Extent[1] + Index[1];
110  } else {
111  LinId =
112  Index[0] * Extent[1] * Extent[2] + Index[1] * Extent[2] + Index[2];
113  }
114  __SYCL_ASSUME_INT(LinId);
115  return LinId;
116  }
117 
118  group<Dimensions> get_group() const {
119  // TODO: ideally Group object should be stateless and have a contructor with
120  // no arguments.
122  get_group_range(), get_group_id());
123  }
124 
125  sub_group get_sub_group() const { return sub_group(); }
126 
127  size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const {
128  size_t Id = get_group_id()[Dimension];
129  __SYCL_ASSUME_INT(Id);
130  return Id;
131  }
132 
134  size_t LinId = 0;
135  id<Dimensions> Index = get_group_id();
136  range<Dimensions> Extent = get_group_range();
137  if (1 == Dimensions) {
138  LinId = Index[0];
139  } else if (2 == Dimensions) {
140  LinId = Index[0] * Extent[1] + Index[1];
141  } else {
142  LinId =
143  Index[0] * Extent[1] * Extent[2] + Index[1] * Extent[2] + Index[2];
144  }
145  __SYCL_ASSUME_INT(LinId);
146  return LinId;
147  }
148 
149  range<Dimensions> get_group_range() const {
150 #ifdef __SYCL_DEVICE_ONLY__
151  return __spirv::initNumWorkgroups<Dimensions, range<Dimensions>>();
152 #else
153  throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
154  "nd_item methods can't be invoked on the host");
155  return {};
156 #endif
157  }
158 
159  size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const {
160  size_t Range = get_group_range()[Dimension];
161  __SYCL_ASSUME_INT(Range);
162  return Range;
163  }
164 
165  range<Dimensions> get_global_range() const {
166 #ifdef __SYCL_DEVICE_ONLY__
167  return __spirv::initGlobalSize<Dimensions, range<Dimensions>>();
168 #else
169  throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
170  "nd_item methods can't be invoked on the host");
171  return {};
172 #endif
173  }
174 
175  size_t get_global_range(int Dimension) const {
176  size_t Val = get_global_range()[Dimension];
177  __SYCL_ASSUME_INT(Val);
178  return Val;
179  }
180 
181  range<Dimensions> get_local_range() const {
182 #ifdef __SYCL_DEVICE_ONLY__
183  return __spirv::initWorkgroupSize<Dimensions, range<Dimensions>>();
184 #else
185  throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
186  "nd_item methods can't be invoked on the host");
187  return {};
188 #endif
189  }
190 
191  size_t get_local_range(int Dimension) const {
192  size_t Id = get_local_range()[Dimension];
193  __SYCL_ASSUME_INT(Id);
194  return Id;
195  }
196 
197  __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
198  id<Dimensions> get_offset() const {
199 #ifdef __SYCL_DEVICE_ONLY__
200  return __spirv::initGlobalOffset<Dimensions, id<Dimensions>>();
201 #else
202  throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
203  "nd_item methods can't be invoked on the host");
204  return {};
205 #endif
206  }
207 
208  nd_range<Dimensions> get_nd_range() const {
209  return nd_range<Dimensions>(get_global_range(), get_local_range(),
210  get_offset());
211  }
212 
213  void barrier(access::fence_space accessSpace =
215  uint32_t flags = _V1::detail::getSPIRVMemorySemanticsMask(accessSpace);
217  flags);
218  }
219 
222  template <access::mode accessMode = access::mode::read_write>
223  __SYCL2020_DEPRECATED("use sycl::atomic_fence() free function instead")
224  void mem_fence(
225  typename std::enable_if_t<accessMode == access::mode::read ||
226  accessMode == access::mode::write ||
227  accessMode == access::mode::read_write,
228  access::fence_space>
229  accessSpace = access::fence_space::global_and_local) const {
230  uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace);
231  // TODO: currently, there is no good way in SPIR-V to set the memory
232  // barrier only for load operations or only for store operations.
233  // The full read-and-write barrier is used and the template parameter
234  // 'accessMode' is ignored for now. Either SPIR-V or SYCL spec may be
235  // changed to address this discrepancy between SPIR-V and SYCL,
236  // or if we decide that 'accessMode' is the important feature then
237  // we can fix this later, for example, by using OpenCL 1.2 functions
238  // read_mem_fence() and write_mem_fence().
240  }
241 
247  template <typename dataT>
248  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
249  std::enable_if_t<!detail::is_bool<dataT>::value,
250  device_event> async_work_group_copy(local_ptr<dataT> dest,
251  global_ptr<dataT> src,
252  size_t numElements,
253  size_t srcStride) const {
254  using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
255  using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
256 
258  __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
259  numElements, srcStride, 0);
260  return device_event(E);
261  }
262 
268  template <typename dataT>
269  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
270  std::enable_if_t<!detail::is_bool<dataT>::value,
271  device_event> async_work_group_copy(global_ptr<dataT> dest,
272  local_ptr<dataT> src,
273  size_t numElements,
274  size_t destStride)
275  const {
276  using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
277  using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
278 
280  __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
281  numElements, destStride, 0);
282  return device_event(E);
283  }
284 
291  template <typename DestDataT, typename SrcDataT>
292  std::enable_if_t<!detail::is_bool<DestDataT>::value &&
293  std::is_same_v<std::remove_const_t<SrcDataT>, DestDataT>,
294  device_event>
295  async_work_group_copy(decorated_local_ptr<DestDataT> dest,
296  decorated_global_ptr<SrcDataT> src, size_t numElements,
297  size_t srcStride) const {
298  using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
299  using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
300 
302  __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
303  numElements, srcStride, 0);
304  return device_event(E);
305  }
306 
313  template <typename DestDataT, typename SrcDataT>
314  std::enable_if_t<!detail::is_bool<DestDataT>::value &&
315  std::is_same_v<std::remove_const_t<SrcDataT>, DestDataT>,
316  device_event>
317  async_work_group_copy(decorated_global_ptr<DestDataT> dest,
318  decorated_local_ptr<SrcDataT> src, size_t numElements,
319  size_t destStride) const {
320  using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
321  using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
322 
324  __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
325  numElements, destStride, 0);
326  return device_event(E);
327  }
328 
334  template <typename T, access::address_space DestS, access::address_space SrcS>
335  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
336  std::enable_if_t<
337  detail::is_scalar_bool<T>::value,
338  device_event> async_work_group_copy(multi_ptr<T, DestS,
339  access::decorated::legacy>
340  Dest,
341  multi_ptr<T, SrcS,
342  access::decorated::legacy>
343  Src,
344  size_t NumElements,
345  size_t Stride) const {
346  static_assert(sizeof(bool) == sizeof(uint8_t),
347  "Async copy to/from bool memory is not supported.");
348  auto DestP = multi_ptr<uint8_t, DestS, access::decorated::legacy>(
349  reinterpret_cast<uint8_t *>(Dest.get()));
350  auto SrcP = multi_ptr<uint8_t, SrcS, access::decorated::legacy>(
351  reinterpret_cast<uint8_t *>(Src.get()));
352  return async_work_group_copy(DestP, SrcP, NumElements, Stride);
353  }
354 
360  template <typename T, access::address_space DestS, access::address_space SrcS>
361  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
362  std::enable_if_t<
363  detail::is_vector_bool<T>::value,
364  device_event> async_work_group_copy(multi_ptr<T, DestS,
365  access::decorated::legacy>
366  Dest,
367  multi_ptr<T, SrcS,
368  access::decorated::legacy>
369  Src,
370  size_t NumElements,
371  size_t Stride) const {
372  static_assert(sizeof(bool) == sizeof(uint8_t),
373  "Async copy to/from bool memory is not supported.");
374  using VecT = detail::change_base_type_t<T, uint8_t>;
375  auto DestP = address_space_cast<DestS, access::decorated::legacy>(
376  reinterpret_cast<VecT *>(Dest.get()));
377  auto SrcP = address_space_cast<SrcS, access::decorated::legacy>(
378  reinterpret_cast<VecT *>(Src.get()));
379  return async_work_group_copy(DestP, SrcP, NumElements, Stride);
380  }
381 
387  template <typename DestT, access::address_space DestS, typename SrcT,
389  std::enable_if_t<detail::is_scalar_bool<DestT>::value &&
390  std::is_same_v<std::remove_const_t<SrcT>, DestT>,
391  device_event>
392  async_work_group_copy(multi_ptr<DestT, DestS, access::decorated::yes> Dest,
393  multi_ptr<SrcT, SrcS, access::decorated::yes> Src,
394  size_t NumElements, size_t Stride) const {
395  static_assert(sizeof(bool) == sizeof(uint8_t),
396  "Async copy to/from bool memory is not supported.");
397  using QualSrcT =
398  std::conditional_t<std::is_const_v<SrcT>, const uint8_t, uint8_t>;
399  auto DestP = multi_ptr<uint8_t, DestS, access::decorated::yes>(
400  detail::cast_AS<typename multi_ptr<uint8_t, DestS,
401  access::decorated::yes>::pointer>(
402  Dest.get_decorated()));
403  auto SrcP = multi_ptr<QualSrcT, SrcS, access::decorated::yes>(
404  detail::cast_AS<typename multi_ptr<QualSrcT, SrcS,
405  access::decorated::yes>::pointer>(
406  Src.get_decorated()));
407  return async_work_group_copy(DestP, SrcP, NumElements, Stride);
408  }
409 
415  template <typename DestT, access::address_space DestS, typename SrcT,
417  std::enable_if_t<detail::is_vector_bool<DestT>::value &&
418  std::is_same_v<std::remove_const_t<SrcT>, DestT>,
419  device_event>
420  async_work_group_copy(multi_ptr<DestT, DestS, access::decorated::yes> Dest,
421  multi_ptr<SrcT, SrcS, access::decorated::yes> Src,
422  size_t NumElements, size_t Stride) const {
423  static_assert(sizeof(bool) == sizeof(uint8_t),
424  "Async copy to/from bool memory is not supported.");
425  using VecT = detail::change_base_type_t<DestT, uint8_t>;
426  using QualSrcVecT =
427  std::conditional_t<std::is_const_v<SrcT>, std::add_const_t<VecT>, VecT>;
428  auto DestP = multi_ptr<VecT, DestS, access::decorated::yes>(
431  Dest.get_decorated()));
432  auto SrcP = multi_ptr<QualSrcVecT, SrcS, access::decorated::yes>(
433  detail::cast_AS<typename multi_ptr<QualSrcVecT, SrcS,
434  access::decorated::yes>::pointer>(
435  Src.get_decorated()));
436  return async_work_group_copy(DestP, SrcP, NumElements, Stride);
437  }
438 
444  template <typename dataT>
445  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
446  device_event
447  async_work_group_copy(local_ptr<dataT> dest, global_ptr<dataT> src,
448  size_t numElements) const {
449  return async_work_group_copy(dest, src, numElements, 1);
450  }
451 
457  template <typename dataT>
458  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
459  device_event
460  async_work_group_copy(global_ptr<dataT> dest, local_ptr<dataT> src,
461  size_t numElements) const {
462  return async_work_group_copy(dest, src, numElements, 1);
463  }
464 
471  template <typename DestDataT, typename SrcDataT>
472  typename std::enable_if_t<
473  std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
474  async_work_group_copy(decorated_local_ptr<DestDataT> dest,
475  decorated_global_ptr<SrcDataT> src,
476  size_t numElements) const {
477  return async_work_group_copy(dest, src, numElements, 1);
478  }
479 
486  template <typename DestDataT, typename SrcDataT>
487  typename std::enable_if_t<
488  std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
489  async_work_group_copy(decorated_global_ptr<DestDataT> dest,
490  decorated_local_ptr<SrcDataT> src,
491  size_t numElements) const {
492  return async_work_group_copy(dest, src, numElements, 1);
493  }
494 
495  template <typename... eventTN> void wait_for(eventTN... events) const {
496  waitForHelper(events...);
497  }
498 
500  ext_oneapi_get_root_group() const {
502  }
503 
504  nd_item(const nd_item &rhs) = default;
505  nd_item(nd_item &&rhs) = default;
506 
507  nd_item &operator=(const nd_item &rhs) = default;
508  nd_item &operator=(nd_item &&rhs) = default;
509 
510  bool operator==(const nd_item &) const { return true; }
511  bool operator!=(const nd_item &rhs) const { return !((*this) == rhs); }
512 
513 protected:
514  friend class detail::Builder;
515  nd_item() {}
516  nd_item(const item<Dimensions, true> &, const item<Dimensions, false> &,
517  const group<Dimensions> &) {}
518 
519  void waitForHelper() const {}
520 
521  void waitForHelper(device_event Event) const { Event.wait(); }
522 
523  template <typename T, typename... Ts>
524  void waitForHelper(T E, Ts... Es) const {
525  waitForHelper(E);
526  waitForHelper(Es...);
527  }
528 
529  id<Dimensions> get_group_id() const {
530 #ifdef __SYCL_DEVICE_ONLY__
531  return __spirv::initWorkgroupId<Dimensions, id<Dimensions>>();
532 #else
533  throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
534  "nd_item methods can't be invoked on the host");
535  return {};
536 #endif
537  }
538 };
539 #else
544 template <int Dimensions = 1> class nd_item {
545 public:
546  static constexpr int dimensions = Dimensions;
547 
548  nd_item() = delete;
549 
550  id<Dimensions> get_global_id() const { return globalItem.get_id(); }
551 
552  size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const {
553  size_t Id = globalItem.get_id(Dimension);
554  __SYCL_ASSUME_INT(Id);
555  return Id;
556  }
557 
559  size_t Id = globalItem.get_linear_id();
560  __SYCL_ASSUME_INT(Id);
561  return Id;
562  }
563 
564  id<Dimensions> get_local_id() const { return localItem.get_id(); }
565 
566  size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const {
567  size_t Id = localItem.get_id(Dimension);
568  __SYCL_ASSUME_INT(Id);
569  return Id;
570  }
571 
572  size_t get_local_linear_id() const {
573  size_t Id = localItem.get_linear_id();
574  __SYCL_ASSUME_INT(Id);
575  return Id;
576  }
577 
578  group<Dimensions> get_group() const { return Group; }
579 
580  sub_group get_sub_group() const { return sub_group(); }
581 
582  size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const {
583  size_t Id = Group[Dimension];
584  __SYCL_ASSUME_INT(Id);
585  return Id;
586  }
587 
589  size_t Id = Group.get_linear_id();
590  __SYCL_ASSUME_INT(Id);
591  return Id;
592  }
593 
594  range<Dimensions> get_group_range() const { return Group.get_group_range(); }
595 
596  size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const {
597  size_t Range = Group.get_group_range(Dimension);
598  __SYCL_ASSUME_INT(Range);
599  return Range;
600  }
601 
602  range<Dimensions> get_global_range() const { return globalItem.get_range(); }
603 
604  size_t get_global_range(int Dimension) const {
605  return globalItem.get_range(Dimension);
606  }
607 
608  range<Dimensions> get_local_range() const { return localItem.get_range(); }
609 
610  size_t get_local_range(int Dimension) const {
611  return localItem.get_range(Dimension);
612  }
613 
614  __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
615  id<Dimensions> get_offset() const { return globalItem.get_offset(); }
616 
619  get_offset());
620  }
621 
622  void barrier(access::fence_space accessSpace =
624  uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace);
626  flags);
627  }
628 
631  template <access::mode accessMode = access::mode::read_write>
632  __SYCL2020_DEPRECATED("use sycl::atomic_fence() free function instead")
633  void mem_fence(
634  typename std::enable_if_t<accessMode == access::mode::read ||
635  accessMode == access::mode::write ||
636  accessMode == access::mode::read_write,
637  access::fence_space>
638  accessSpace = access::fence_space::global_and_local) const {
639  (void)accessSpace;
640  Group.mem_fence();
641  }
642 
643  template <typename dataT>
644  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
646  async_work_group_copy(local_ptr<dataT> dest, global_ptr<dataT> src,
647  size_t numElements) const {
648  return Group.async_work_group_copy(dest, src, numElements);
649  }
650 
651  template <typename dataT>
652  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
654  async_work_group_copy(global_ptr<dataT> dest, local_ptr<dataT> src,
655  size_t numElements) const {
656  return Group.async_work_group_copy(dest, src, numElements);
657  }
658 
659  template <typename dataT>
660  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
662  async_work_group_copy(local_ptr<dataT> dest, global_ptr<dataT> src,
663  size_t numElements, size_t srcStride) const {
664 
665  return Group.async_work_group_copy(dest, src, numElements, srcStride);
666  }
667 
668  template <typename dataT>
669  __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
671  async_work_group_copy(global_ptr<dataT> dest, local_ptr<dataT> src,
672  size_t numElements, size_t destStride) const {
673  return Group.async_work_group_copy(dest, src, numElements, destStride);
674  }
675 
676  template <typename DestDataT, typename SrcDataT>
677  typename std::enable_if_t<
678  std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
681  size_t numElements) const {
682  return Group.async_work_group_copy(dest, src, numElements);
683  }
684 
685  template <typename DestDataT, typename SrcDataT>
686  typename std::enable_if_t<
687  std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
690  size_t numElements) const {
691  return Group.async_work_group_copy(dest, src, numElements);
692  }
693 
694  template <typename DestDataT, typename SrcDataT>
695  typename std::enable_if_t<
696  std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
698  decorated_global_ptr<SrcDataT> src, size_t numElements,
699  size_t srcStride) const {
700 
701  return Group.async_work_group_copy(dest, src, numElements, srcStride);
702  }
703 
704  template <typename DestDataT, typename SrcDataT>
705  typename std::enable_if_t<
706  std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
708  decorated_local_ptr<SrcDataT> src, size_t numElements,
709  size_t destStride) const {
710  return Group.async_work_group_copy(dest, src, numElements, destStride);
711  }
712 
713  template <typename... eventTN> void wait_for(eventTN... events) const {
714  Group.wait_for(events...);
715  }
716 
720  }
721 
722  nd_item(const nd_item &rhs) = default;
723 
724  nd_item(nd_item &&rhs) = default;
725 
726  nd_item &operator=(const nd_item &rhs) = default;
727 
728  nd_item &operator=(nd_item &&rhs) = default;
729 
730  bool operator==(const nd_item &rhs) const {
731  return (rhs.localItem == this->localItem) &&
732  (rhs.globalItem == this->globalItem) && (rhs.Group == this->Group);
733  }
734 
735  bool operator!=(const nd_item &rhs) const { return !((*this) == rhs); }
736 
737 protected:
738  friend class detail::Builder;
740  const group<Dimensions> &GR)
741  : globalItem(GL), localItem(L), Group(GR) {}
742 
743 private:
744  item<Dimensions, true> globalItem;
745  item<Dimensions, false> localItem;
746  group<Dimensions> Group;
747 };
748 #endif
749 
750 template <int Dims>
751 __SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_nd_item() instead")
753 #ifdef __SYCL_DEVICE_ONLY__
754  return detail::Builder::getElement(detail::declptr<nd_item<Dims>>());
755 #else
756  throw sycl::exception(
757  sycl::make_error_code(sycl::errc::feature_not_supported),
758  "Free function calls are not supported on host");
759 #endif
760 }
761 
762 namespace ext::oneapi::experimental {
763 template <int Dims> nd_item<Dims> this_nd_item() {
764 #ifdef __SYCL_DEVICE_ONLY__
765  return sycl::detail::Builder::getElement(
767 #else
768  throw sycl::exception(
769  sycl::make_error_code(sycl::errc::feature_not_supported),
770  "Free function calls are not supported on host");
771 #endif
772 }
773 } // namespace ext::oneapi::experimental
774 } // namespace _V1
775 } // namespace sycl
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const range< Dims > &Group, const id< Dims > &Index)
Definition: helpers.hpp:74
Encapsulates a single SYCL device event which is available only within SYCL kernel functions and can ...
A unique identifier of an item in an index space.
Definition: id.hpp:36
Identifies an instance of the function object executing at each point in a range.
Definition: item.hpp:37
id< Dimensions > get_id() const
Definition: item.hpp:55
range< Dimensions > get_range() const
Definition: item.hpp:69
size_t __SYCL_ALWAYS_INLINE get_linear_id() const
Definition: item.hpp:100
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: multi_ptr.hpp:83
std::conditional_t< is_decorated, decorated_type *, std::add_pointer_t< value_type > > pointer
Definition: multi_ptr.hpp:95
Identifies an instance of the function object executing at each point in an nd_range.
Definition: nd_item.hpp:544
size_t get_local_linear_id() const
Definition: nd_item.hpp:572
size_t get_local_range(int Dimension) const
Definition: nd_item.hpp:610
size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const
Definition: nd_item.hpp:588
nd_item(const nd_item &rhs)=default
size_t get_global_range(int Dimension) const
Definition: nd_item.hpp:604
sycl::ext::oneapi::experimental::root_group< Dimensions > ext_oneapi_get_root_group() const
Definition: nd_item.hpp:718
id< Dimensions > get_local_id() const
Definition: nd_item.hpp:564
range< Dimensions > get_local_range() const
Definition: nd_item.hpp:608
id< Dimensions > get_global_id() const
Definition: nd_item.hpp:550
sub_group get_sub_group() const
Definition: nd_item.hpp:580
void wait_for(eventTN... events) const
Definition: nd_item.hpp:713
nd_item(const item< Dimensions, true > &GL, const item< Dimensions, false > &L, const group< Dimensions > &GR)
Definition: nd_item.hpp:739
nd_item & operator=(const nd_item &rhs)=default
std::enable_if_t< std::is_same_v< DestDataT, std::remove_const_t< SrcDataT > >, device_event > async_work_group_copy(decorated_global_ptr< DestDataT > dest, decorated_local_ptr< SrcDataT > src, size_t numElements) const
Definition: nd_item.hpp:688
nd_range< Dimensions > get_nd_range() const
Definition: nd_item.hpp:617
size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const
Definition: nd_item.hpp:552
size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const
Definition: nd_item.hpp:558
bool operator!=(const nd_item &rhs) const
Definition: nd_item.hpp:735
nd_item(nd_item &&rhs)=default
bool operator==(const nd_item &rhs) const
Definition: nd_item.hpp:730
size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const
Definition: nd_item.hpp:596
static constexpr int dimensions
Definition: nd_item.hpp:546
void mem_fence(typename std::enable_if_t< accessMode==access::mode::read||accessMode==access::mode::write||accessMode==access::mode::read_write, access::fence_space > accessSpace=access::fence_space::global_and_local) const
Executes a work-group mem-fence with memory ordering on the local address space, global address space...
Definition: nd_item.hpp:633
id< Dimensions > get_offset() const
Definition: nd_item.hpp:615
group< Dimensions > get_group() const
Definition: nd_item.hpp:578
std::enable_if_t< std::is_same_v< DestDataT, std::remove_const_t< SrcDataT > >, device_event > async_work_group_copy(decorated_local_ptr< DestDataT > dest, decorated_global_ptr< SrcDataT > src, size_t numElements, size_t srcStride) const
Definition: nd_item.hpp:697
std::enable_if_t< std::is_same_v< DestDataT, std::remove_const_t< SrcDataT > >, device_event > async_work_group_copy(decorated_global_ptr< DestDataT > dest, decorated_local_ptr< SrcDataT > src, size_t numElements, size_t destStride) const
Definition: nd_item.hpp:707
nd_item & operator=(nd_item &&rhs)=default
size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const
Definition: nd_item.hpp:566
range< Dimensions > get_global_range() const
Definition: nd_item.hpp:602
range< Dimensions > get_group_range() const
Definition: nd_item.hpp:594
friend class detail::Builder
Definition: nd_item.hpp:738
device_event async_work_group_copy(local_ptr< dataT > dest, global_ptr< dataT > src, size_t numElements) const
Definition: nd_item.hpp:646
size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const
Definition: nd_item.hpp:582
std::enable_if_t< std::is_same_v< DestDataT, std::remove_const_t< SrcDataT > >, device_event > async_work_group_copy(decorated_local_ptr< DestDataT > dest, decorated_global_ptr< SrcDataT > src, size_t numElements) const
Definition: nd_item.hpp:679
void barrier(access::fence_space accessSpace=access::fence_space::global_and_local) const
Definition: nd_item.hpp:622
Defines the iteration domain of both the work-groups and the overall dispatch.
Definition: nd_range.hpp:22
#define __SYCL_ASSUME_INT(x)
Definition: defines.hpp:17
#define __SYCL_ALWAYS_INLINE
ToT cast_AS(FromT from)
Definition: access.hpp:330
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
Definition: helpers.hpp:198
decltype(convertToOpenCLType(std::declval< T >())) ConvertToOpenCLType_t
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
Definition: accessor.hpp:3233
signed char __SYCL2020_DEPRECATED
Definition: aliases.hpp:94
multi_ptr< ElementType, access::address_space::global_space, IsDecorated > global_ptr
Definition: pointers.hpp:30
constexpr mode_tag_t< access_mode::read_write > read_write
Definition: access.hpp:85
nd_item< Dims > this_nd_item()
Definition: nd_item.hpp:752
__SYCL_DEPRECATED("This is a deprecated argument type for SYCL nan built-in function.") std
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:94
multi_ptr< ElementType, access::address_space::local_space, IsDecorated > local_ptr
Definition: pointers.hpp:52
Definition: access.hpp:18
__SYCL_CONVERGENT__ __ocl_event_t __SYCL_OpGroupAsyncCopyGlobalToLocal(__spv::Scope::Flag, dataT *Dest, const dataT *Src, size_t NumElements, size_t Stride, __ocl_event_t) noexcept
Definition: spirv_ops.hpp:1265
__SYCL_CONVERGENT__ __DPCPP_SYCL_EXTERNAL void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:26
__SYCL_CONVERGENT__ __DPCPP_SYCL_EXTERNAL void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:36
__SYCL_CONVERGENT__ __ocl_event_t __SYCL_OpGroupAsyncCopyLocalToGlobal(__spv::Scope::Flag, dataT *Dest, const dataT *Src, size_t NumElements, size_t Stride, __ocl_event_t) noexcept
Definition: spirv_ops.hpp:1277
void * __ocl_event_t