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>
12 #include <sycl/access/access.hpp>
13 #include <sycl/detail/defines.hpp>
14 #include <sycl/detail/helpers.hpp>
15 #include <sycl/group.hpp>
16 #include <sycl/id.hpp>
17 #include <sycl/item.hpp>
18 #include <sycl/nd_range.hpp>
19 #include <sycl/range.hpp>
20 #include <sycl/sub_group.hpp>
21 
22 #include <cstddef>
23 #include <stdexcept>
24 #include <type_traits>
25 
26 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 template <int Dims>
194 __SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_nd_item() instead")
196 #ifdef __SYCL_DEVICE_ONLY__
197  return detail::Builder::getElement(detail::declptr<nd_item<Dims>>());
198 #else
199  throw sycl::exception(
200  sycl::make_error_code(sycl::errc::feature_not_supported),
201  "Free function calls are not supported on host device");
202 #endif
203 }
204 
205 namespace ext {
206 namespace oneapi {
207 namespace experimental {
208 template <int Dims> nd_item<Dims> this_nd_item() {
209 #ifdef __SYCL_DEVICE_ONLY__
210  return sycl::detail::Builder::getElement(
212 #else
213  throw sycl::exception(
214  sycl::make_error_code(sycl::errc::feature_not_supported),
215  "Free function calls are not supported on host device");
216 #endif
217 }
218 } // namespace experimental
219 } // namespace oneapi
220 } // namespace ext
221 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
222 } // namespace sycl
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:32
Identifies an instance of the function object executing at each point in a range.
Definition: item.hpp:39
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: multi_ptr.hpp:78
Identifies an instance of the function object executing at each point in an nd_range.
Definition: nd_item.hpp:36
size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const
Definition: nd_item.hpp:48
nd_item(const item< dimensions, true > &GL, const item< dimensions, false > &L, const group< dimensions > &GR)
Definition: nd_item.hpp:183
size_t get_global_range(int dimension) const
Definition: nd_item.hpp:94
nd_item & operator=(nd_item &&rhs)=default
id< dimensions > get_local_id() const
Definition: nd_item.hpp:54
nd_item & operator=(const nd_item &rhs)=default
size_t __SYCL_ALWAYS_INLINE get_group(int dimension) const
Definition: nd_item.hpp:72
size_t __SYCL_ALWAYS_INLINE get_group_range(int dimension) const
Definition: nd_item.hpp:86
sub_group get_sub_group() const
Definition: nd_item.hpp:70
id< dimensions > get_global_id() const
Definition: nd_item.hpp:40
device_event async_work_group_copy(local_ptr< dataT > dest, global_ptr< dataT > src, size_t numElements) const
Definition: nd_item.hpp:134
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
range< dimensions > get_group_range() const
Definition: nd_item.hpp:84
size_t __SYCL_ALWAYS_INLINE get_global_id(int dimension) const
Definition: nd_item.hpp:42
range< dimensions > get_global_range() const
Definition: nd_item.hpp:92
void barrier(access::fence_space accessSpace=access::fence_space::global_and_local) const
Definition: nd_item.hpp:112
size_t get_local_linear_id() const
Definition: nd_item.hpp:62
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
nd_range< dimensions > get_nd_range() const
Definition: nd_item.hpp:107
size_t get_local_range(int dimension) const
Definition: nd_item.hpp:100
range< dimensions > get_local_range() const
Definition: nd_item.hpp:98
size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const
Definition: nd_item.hpp:78
size_t __SYCL_ALWAYS_INLINE get_local_id(int dimension) const
Definition: nd_item.hpp:56
void wait_for(eventTN... events) const
Definition: nd_item.hpp:162
bool operator!=(const nd_item &rhs) const
Definition: nd_item.hpp:179
device_event async_work_group_copy(global_ptr< dataT > dest, local_ptr< dataT > src, size_t numElements) const
Definition: nd_item.hpp:141
nd_item(const nd_item &rhs)=default
nd_item(nd_item &&rhs)=default
bool operator==(const nd_item &rhs) const
Definition: nd_item.hpp:174
group< dimensions > get_group() const
Definition: nd_item.hpp:68
Defines the iteration domain of both the work-groups and the overall dispatch.
Definition: nd_range.hpp:23
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: range.hpp:24
#define __SYCL_ASSUME_INT(x)
Definition: defines.hpp:18
#define __SYCL_INLINE_VER_NAMESPACE(X)
#define __SYCL_DEPRECATED(message)
#define __SYCL2020_DEPRECATED(message)
#define __SYCL_ALWAYS_INLINE
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
Definition: helpers.hpp:199
typename std::enable_if< B, T >::type enable_if_t
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:91
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
__SYCL_CONVERGENT__ SYCL_EXTERNAL void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:26