DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
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>
15 #include <CL/sycl/group.hpp>
16 #include <CL/sycl/id.hpp>
17 #include <CL/sycl/item.hpp>
18 #include <CL/sycl/nd_range.hpp>
19 #include <CL/sycl/range.hpp>
20 #include <CL/sycl/sub_group.hpp>
21 
22 #include <cstddef>
23 #include <stdexcept>
24 #include <type_traits>
25 
27 namespace sycl {
28 namespace detail {
29 class Builder;
30 }
31 
36 template <int dimensions = 1> class nd_item {
37 public:
38  nd_item() = delete;
39 
40  id<dimensions> get_global_id() const { return globalItem.get_id(); }
41 
42  size_t __SYCL_ALWAYS_INLINE get_global_id(int dimension) const {
43  size_t Id = globalItem.get_id(dimension);
45  return Id;
46  }
47 
49  size_t Id = globalItem.get_linear_id();
51  return Id;
52  }
53 
54  id<dimensions> get_local_id() const { return localItem.get_id(); }
55 
56  size_t __SYCL_ALWAYS_INLINE get_local_id(int dimension) const {
57  size_t Id = localItem.get_id(dimension);
59  return Id;
60  }
61 
62  size_t get_local_linear_id() const {
63  size_t Id = localItem.get_linear_id();
65  return Id;
66  }
67 
68  group<dimensions> get_group() const { return Group; }
69 
70  sub_group get_sub_group() const { return sub_group(); }
71 
72  size_t __SYCL_ALWAYS_INLINE get_group(int dimension) const {
73  size_t Size = Group[dimension];
74  __SYCL_ASSUME_INT(Size);
75  return Size;
76  }
77 
79  size_t Id = Group.get_linear_id();
81  return Id;
82  }
83 
84  range<dimensions> get_group_range() const { return Group.get_group_range(); }
85 
86  size_t __SYCL_ALWAYS_INLINE get_group_range(int dimension) const {
87  size_t Range = Group.get_group_range(dimension);
88  __SYCL_ASSUME_INT(Range);
89  return Range;
90  }
91 
92  range<dimensions> get_global_range() const { return globalItem.get_range(); }
93 
94  size_t get_global_range(int dimension) const {
95  return globalItem.get_range(dimension);
96  }
97 
98  range<dimensions> get_local_range() const { return localItem.get_range(); }
99 
100  size_t get_local_range(int dimension) const {
101  return localItem.get_range(dimension);
102  }
103 
104  __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
105  id<dimensions> get_offset() const { return globalItem.get_offset(); }
106 
108  return nd_range<dimensions>(get_global_range(), get_local_range(),
109  get_offset());
110  }
111 
112  void barrier(access::fence_space accessSpace =
113  access::fence_space::global_and_local) const {
114  uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace);
116  flags);
117  }
118 
121  template <access::mode accessMode = access::mode::read_write>
122  __SYCL2020_DEPRECATED("use sycl::atomic_fence() free function instead")
123  void mem_fence(
124  typename detail::enable_if_t<accessMode == access::mode::read ||
125  accessMode == access::mode::write ||
126  accessMode == access::mode::read_write,
127  access::fence_space>
128  accessSpace = access::fence_space::global_and_local) const {
129  (void)accessSpace;
130  Group.mem_fence();
131  }
132 
133  template <typename dataT>
135  global_ptr<dataT> src,
136  size_t numElements) const {
137  return Group.async_work_group_copy(dest, src, numElements);
138  }
139 
140  template <typename dataT>
142  local_ptr<dataT> src,
143  size_t numElements) const {
144  return Group.async_work_group_copy(dest, src, numElements);
145  }
146 
147  template <typename dataT>
149  global_ptr<dataT> src, size_t numElements,
150  size_t srcStride) const {
151 
152  return Group.async_work_group_copy(dest, src, numElements, srcStride);
153  }
154 
155  template <typename dataT>
157  local_ptr<dataT> src, size_t numElements,
158  size_t destStride) const {
159  return Group.async_work_group_copy(dest, src, numElements, destStride);
160  }
161 
162  template <typename... eventTN> void wait_for(eventTN... events) const {
163  Group.wait_for(events...);
164  }
165 
166  nd_item(const nd_item &rhs) = default;
167 
168  nd_item(nd_item &&rhs) = default;
169 
170  nd_item &operator=(const nd_item &rhs) = default;
171 
172  nd_item &operator=(nd_item &&rhs) = default;
173 
174  bool operator==(const nd_item &rhs) const {
175  return (rhs.localItem == this->localItem) &&
176  (rhs.globalItem == this->globalItem) && (rhs.Group == this->Group);
177  }
178 
179  bool operator!=(const nd_item &rhs) const { return !((*this) == rhs); }
180 
181 protected:
182  friend class detail::Builder;
184  const group<dimensions> &GR)
185  : globalItem(GL), localItem(L), Group(GR) {}
186 
187 private:
188  item<dimensions, true> globalItem;
189  item<dimensions, false> localItem;
190  group<dimensions> Group;
191 };
192 
193 namespace detail {
194 template <int Dims> nd_item<Dims> store_nd_item(const nd_item<Dims> *nd_i) {
195  return get_or_store(nd_i);
196 }
197 } // namespace detail
198 
199 template <int Dims>
200 __SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_nd_item() instead")
202 #ifdef __SYCL_DEVICE_ONLY__
203  return detail::Builder::getElement(detail::declptr<nd_item<Dims>>());
204 #else
205  return detail::store_nd_item<Dims>(nullptr);
206 #endif
207 }
208 
209 namespace ext {
210 namespace oneapi {
211 namespace experimental {
212 template <int Dims> nd_item<Dims> this_nd_item() {
213 #ifdef __SYCL_DEVICE_ONLY__
214  return sycl::detail::Builder::getElement(detail::declptr<nd_item<Dims>>());
215 #else
216  return sycl::detail::store_nd_item<Dims>(nullptr);
217 #endif
218 }
219 } // namespace experimental
220 } // namespace oneapi
221 } // namespace ext
222 } // namespace sycl
223 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::nd_item::get_nd_range
nd_range< dimensions > get_nd_range() const
Definition: nd_item.hpp:107
cl::sycl::nd_range
Defines the iteration domain of both the work-groups and the overall dispatch.
Definition: nd_range.hpp:23
spirv_ops.hpp
__spirv_ControlBarrier
__SYCL_CONVERGENT__ SYCL_EXTERNAL void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:26
cl::sycl::detail::Builder
Definition: helpers.hpp:68
__spv::Scope::Workgroup
@ Workgroup
Definition: spirv_types.hpp:30
cl::sycl::nd_item::get_sub_group
sub_group get_sub_group() const
Definition: nd_item.hpp:70
cl::sycl::nd_item::async_work_group_copy
device_event async_work_group_copy(global_ptr< dataT > dest, local_ptr< dataT > src, size_t numElements) const
Definition: nd_item.hpp:141
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:56
cl::sycl::item
Identifies an instance of the function object executing at each point in a range.
Definition: helpers.hpp:28
sub_group.hpp
item.hpp
cl::sycl::nd_item::async_work_group_copy
device_event async_work_group_copy(local_ptr< dataT > dest, global_ptr< dataT > src, size_t numElements) const
Definition: nd_item.hpp:134
cl::sycl::id
A unique identifier of an item in an index space.
Definition: array.hpp:17
__SYCL_DEPRECATED
#define __SYCL_DEPRECATED(message)
Definition: defines_elementary.hpp:47
cl::sycl::nd_item::get_local_linear_id
size_t get_local_linear_id() const
Definition: nd_item.hpp:62
cl::sycl::group
Encapsulates all functionality required to represent a particular work-group within a parallel execut...
Definition: helpers.hpp:29
cl::sycl::nd_item::get_group_linear_id
size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const
Definition: nd_item.hpp:78
cl::sycl::detail::declptr
T * declptr()
Definition: helpers.hpp:56
cl::sycl::nd_item::get_global_range
size_t get_global_range(int dimension) const
Definition: nd_item.hpp:94
helpers.hpp
cl::sycl::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:32
access.hpp
cl::sycl::nd_item::get_local_range
range< dimensions > get_local_range() const
Definition: nd_item.hpp:98
cl::sycl::nd_item::async_work_group_copy
device_event async_work_group_copy(global_ptr< dataT > dest, local_ptr< dataT > src, size_t numElements, size_t destStride) const
Definition: nd_item.hpp:156
group.hpp
cl::sycl::nd_item::barrier
void barrier(access::fence_space accessSpace=access::fence_space::global_and_local) const
Definition: nd_item.hpp:112
cl::sycl::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:110
id.hpp
cl::sycl::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: buffer.hpp:24
cl::sycl::nd_item::get_group_range
range< dimensions > get_group_range() const
Definition: nd_item.hpp:84
cl::sycl::access::fence_space
fence_space
Definition: access.hpp:37
nd_range.hpp
cl::sycl::detail::getSPIRVMemorySemanticsMask
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
Definition: helpers.hpp:200
cl::sycl::nd_item::wait_for
void wait_for(eventTN... events) const
Definition: nd_item.hpp:162
cl::sycl::ext::oneapi::experimental::this_nd_item
nd_item< Dims > this_nd_item()
Definition: nd_item.hpp:212
__SYCL_ALWAYS_INLINE
#define __SYCL_ALWAYS_INLINE
Definition: defines_elementary.hpp:29
cl::sycl::detail::store_nd_item
nd_item< Dims > store_nd_item(const nd_item< Dims > *nd_i)
Definition: nd_item.hpp:194
cl::sycl::nd_item::operator==
bool operator==(const nd_item &rhs) const
Definition: nd_item.hpp:174
cl::sycl::nd_item::get_global_range
range< dimensions > get_global_range() const
Definition: nd_item.hpp:92
range.hpp
defines.hpp
cl::sycl::nd_item::get_local_id
id< dimensions > get_local_id() const
Definition: nd_item.hpp:54
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::nd_item::operator!=
bool operator!=(const nd_item &rhs) const
Definition: nd_item.hpp:179
cl::sycl::nd_item::get_group
group< dimensions > get_group() const
Definition: nd_item.hpp:68
__SYCL_ASSUME_INT
#define __SYCL_ASSUME_INT(x)
Definition: defines.hpp:18
cl::sycl::nd_item::get_global_id
size_t __SYCL_ALWAYS_INLINE get_global_id(int dimension) const
Definition: nd_item.hpp:42
cl::sycl::nd_item::get_global_linear_id
size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const
Definition: nd_item.hpp:48
cl::sycl::ext::oneapi::sub_group
Definition: sub_group.hpp:108
cl::sycl::nd_item::get_global_id
id< dimensions > get_global_id() const
Definition: nd_item.hpp:40
cl::sycl::device_event
Encapsulates a single SYCL device event which is available only within SYCL kernel functions and can ...
Definition: device_event.hpp:22
cl::sycl::detail::get_or_store
T get_or_store(const T *obj)
Definition: helpers.hpp:60
cl::sycl::nd_item::get_local_range
size_t get_local_range(int dimension) const
Definition: nd_item.hpp:100
cl::sycl::nd_item
Identifies an instance of the function object executing at each point in an nd_range.
Definition: helpers.hpp:32
cl::sycl::nd_item::get_local_id
size_t __SYCL_ALWAYS_INLINE get_local_id(int dimension) const
Definition: nd_item.hpp:56
cl::sycl::nd_item::get_group_range
size_t __SYCL_ALWAYS_INLINE get_group_range(int dimension) const
Definition: nd_item.hpp:86
cl::sycl::nd_item::nd_item
nd_item(const item< dimensions, true > &GL, const item< dimensions, false > &L, const group< dimensions > &GR)
Definition: nd_item.hpp:183
cl::sycl::nd_item::async_work_group_copy
device_event async_work_group_copy(local_ptr< dataT > dest, global_ptr< dataT > src, size_t numElements, size_t srcStride) const
Definition: nd_item.hpp:148
cl::sycl::access::mode
mode
Definition: access.hpp:28
cl::sycl::nd_item::get_group
size_t __SYCL_ALWAYS_INLINE get_group(int dimension) const
Definition: nd_item.hpp:72
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12