DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
item.hpp
Go to the documentation of this file.
1 //==------------ item.hpp --- SYCL iteration 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 
15 #include <CL/sycl/id.hpp>
16 #include <CL/sycl/range.hpp>
17 
18 #include <cstddef>
19 
21 namespace sycl {
22 namespace detail {
23 class Builder;
24 template <typename TransformedArgType, int Dims, typename KernelType>
25 class RoundedRangeKernel;
26 template <typename TransformedArgType, int Dims, typename KernelType>
27 class RoundedRangeKernelWithKH;
28 }
29 template <int dimensions> class id;
30 template <int dimensions> class range;
31 
36 template <int dimensions = 1, bool with_offset = true> class item {
37 #ifndef __SYCL_DISABLE_ITEM_TO_INT_CONV__
38  /* Helper class for conversion operator. Void type is not suitable. User
39  * cannot even try to get address of the operator __private_class(). User
40  * may try to get an address of operator void() and will get the
41  * compile-time error */
42  class __private_class;
43 
44  template <bool B, typename T>
45  using EnableIfT = detail::conditional_t<B, T, __private_class>;
46 #endif // __SYCL_DISABLE_ITEM_TO_INT_CONV__
47 public:
48  item() = delete;
49 
50  id<dimensions> get_id() const { return MImpl.MIndex; }
51 
52  size_t __SYCL_ALWAYS_INLINE get_id(int dimension) const {
53  size_t Id = MImpl.MIndex[dimension];
55  return Id;
56  }
57 
58  size_t __SYCL_ALWAYS_INLINE operator[](int dimension) const {
59  size_t Id = MImpl.MIndex[dimension];
61  return Id;
62  }
63 
64  range<dimensions> get_range() const { return MImpl.MExtent; }
65 
66  size_t __SYCL_ALWAYS_INLINE get_range(int dimension) const {
67  size_t Id = MImpl.MExtent[dimension];
69  return Id;
70  }
71 #ifndef __SYCL_DISABLE_ITEM_TO_INT_CONV__
72  operator EnableIfT<dimensions == 1, std::size_t>() const { return get_id(0); }
73 #endif // __SYCL_DISABLE_ITEM_TO_INT_CONV__
74  template <bool has_offset = with_offset>
75  __SYCL2020_DEPRECATED("offsets are deprecated in SYCL2020")
76  detail::enable_if_t<has_offset, id<dimensions>> get_offset() const {
77  return MImpl.MOffset;
78  }
79 
80  template <bool has_offset = with_offset>
81  __SYCL2020_DEPRECATED("offsets are deprecated in SYCL2020")
82  detail::enable_if_t<has_offset, size_t> __SYCL_ALWAYS_INLINE
83  get_offset(int dimension) const {
84  size_t Id = MImpl.MOffset[dimension];
86  return Id;
87  }
88 
89  template <bool has_offset = with_offset>
91  return detail::Builder::createItem<dimensions, true>(
92  MImpl.MExtent, MImpl.MIndex, /*Offset*/ {});
93  }
94 
96  size_t Id = MImpl.get_linear_id();
98  return Id;
99  }
100 
101  item(const item &rhs) = default;
102 
103  item(item<dimensions, with_offset> &&rhs) = default;
104 
105  item &operator=(const item &rhs) = default;
106 
107  item &operator=(item &&rhs) = default;
108 
109  bool operator==(const item &rhs) const { return rhs.MImpl == MImpl; }
110 
111  bool operator!=(const item &rhs) const { return rhs.MImpl != MImpl; }
112 
113 protected:
114  template <bool has_offset = with_offset>
115  item(detail::enable_if_t<has_offset, const range<dimensions>> &extent,
116  const id<dimensions> &index, const id<dimensions> &offset)
117  : MImpl{extent, index, offset} {}
118 
119  template <bool has_offset = with_offset>
120  item(detail::enable_if_t<!has_offset, const range<dimensions>> &extent,
121  const id<dimensions> &index)
122  : MImpl{extent, index} {}
123 
124  friend class detail::Builder;
125 
126 private:
127  // Friend to get access to private method set_allowed_range().
128  template <typename, int, typename> friend class detail::RoundedRangeKernel;
129  template <typename, int, typename>
131  void set_allowed_range(const range<dimensions> rnwi) { MImpl.MExtent = rnwi; }
132 
134 };
135 
136 namespace detail {
137 template <int Dims> item<Dims> store_item(const item<Dims> *i) {
138  return get_or_store(i);
139 }
140 } // namespace detail
141 
142 template <int Dims>
143 __SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_item() instead")
144 item<Dims> this_item() {
145 #ifdef __SYCL_DEVICE_ONLY__
146  return detail::Builder::getElement(detail::declptr<item<Dims>>());
147 #else
148  return detail::store_item<Dims>(nullptr);
149 #endif
150 }
151 
152 namespace ext {
153 namespace oneapi {
154 namespace experimental {
155 template <int Dims> item<Dims> this_item() {
156 #ifdef __SYCL_DEVICE_ONLY__
157  return sycl::detail::Builder::getElement(detail::declptr<item<Dims>>());
158 #else
159  return sycl::detail::store_item<Dims>(nullptr);
160 #endif
161 }
162 } // namespace experimental
163 } // namespace oneapi
164 } // namespace ext
165 } // namespace sycl
166 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::item::operator==
bool operator==(const item &rhs) const
Definition: item.hpp:109
cl::sycl::detail::Builder
Definition: helpers.hpp:68
cl::sycl::item::get_range
size_t __SYCL_ALWAYS_INLINE get_range(int dimension) const
Definition: item.hpp:66
cl::sycl::item::item
item(detail::enable_if_t<!has_offset, const range< dimensions >> &extent, const id< dimensions > &index)
Definition: item.hpp:120
type_traits.hpp
__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
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::detail::declptr
T * declptr()
Definition: helpers.hpp:56
cl::sycl::item::get_id
size_t __SYCL_ALWAYS_INLINE get_id(int dimension) const
Definition: item.hpp:52
helpers.hpp
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::item::operator[]
size_t __SYCL_ALWAYS_INLINE operator[](int dimension) const
Definition: item.hpp:58
__SYCL_ALWAYS_INLINE
#define __SYCL_ALWAYS_INLINE
Definition: defines_elementary.hpp:29
cl::sycl::detail::RoundedRangeKernel
Definition: handler.hpp:203
range.hpp
cl::sycl::detail::store_item
item< Dims > store_item(const item< Dims > *i)
Definition: item.hpp:137
defines.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
__SYCL_ASSUME_INT
#define __SYCL_ASSUME_INT(x)
Definition: defines.hpp:18
cl::sycl::detail::RoundedRangeKernelWithKH
Definition: handler.hpp:221
cl::sycl::item::get_id
id< dimensions > get_id() const
Definition: item.hpp:50
cl::sycl::item::get_linear_id
size_t __SYCL_ALWAYS_INLINE get_linear_id() const
Definition: item.hpp:95
cl::sycl::detail::ItemBase< dimensions, with_offset >
cl::sycl::detail::get_or_store
T get_or_store(const T *obj)
Definition: helpers.hpp:60
item_base.hpp
cl::sycl::item::operator!=
bool operator!=(const item &rhs) const
Definition: item.hpp:111
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::item::get_range
range< dimensions > get_range() const
Definition: item.hpp:64
cl::sycl::item::item
item(detail::enable_if_t< has_offset, const range< dimensions >> &extent, const id< dimensions > &index, const id< dimensions > &offset)
Definition: item.hpp:115
cl::sycl::ext::oneapi::experimental::this_item
item< Dims > this_item()
Definition: item.hpp:155
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12