17 #ifdef __SYCL_DEVICE_ONLY__
24 #include <type_traits>
29 inline namespace _V1 {
32 template <
int Dims,
bool WithOffset>
class item;
33 template <
int Dims>
class group;
34 template <
int Dims>
class range;
35 template <
int Dims>
class id;
36 template <
int Dims>
class nd_item;
37 template <
int Dims>
class h_item;
38 template <
typename Type, std::
size_t NumElements>
class marray;
47 __SYCL_EXPORT std::vector<sycl::detail::pi::PiEvent>
49 std::shared_ptr<sycl::detail::context_impl> Context);
51 __SYCL_EXPORT
void waitEvents(std::vector<sycl::event> DepEvents);
56 template <
typename T> T *
declptr() {
return static_cast<T *
>(
nullptr); }
61 static thread_local
auto stored = *
obj;
83 return group<Dims>(Global, Local, Global / Local, Index);
86 template <
class ResType,
typename BitsType>
88 return ResType(Bits, BitsNum);
91 template <
int Dims,
bool WithOffset>
92 static std::enable_if_t<WithOffset, item<Dims, WithOffset>>
98 template <
int Dims,
bool WithOffset>
99 static std::enable_if_t<!WithOffset, item<Dims, WithOffset>>
124 template <
int Dims,
bool WithOffset>
127 Item.MImpl.MIndex = NextIndex;
130 #ifdef __SYCL_DEVICE_ONLY__
133 using is_valid_dimensions = std::integral_constant<bool, (N > 0) && (N < 4)>;
136 static_assert(is_valid_dimensions<Dims>::value,
"invalid dimensions");
137 return __spirv::initGlobalInvocationId<Dims, id<Dims>>();
141 static_assert(is_valid_dimensions<Dims>::value,
"invalid dimensions");
142 range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
143 range<Dims> LocalSize{__spirv::initWorkgroupSize<Dims, range<Dims>>()};
144 range<Dims> GroupRange{__spirv::initNumWorkgroups<Dims, range<Dims>>()};
145 id<Dims> GroupId{__spirv::initWorkgroupId<Dims, id<Dims>>()};
146 return createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
149 template <
int Dims,
bool WithOffset>
150 static std::enable_if_t<WithOffset, const item<Dims, WithOffset>> getItem() {
151 static_assert(is_valid_dimensions<Dims>::value,
"invalid dimensions");
152 id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
153 range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
154 id<Dims> GlobalOffset{__spirv::initGlobalOffset<Dims, id<Dims>>()};
155 return createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
158 template <
int Dims,
bool WithOffset>
159 static std::enable_if_t<!WithOffset, const item<Dims, WithOffset>> getItem() {
160 static_assert(is_valid_dimensions<Dims>::value,
"invalid dimensions");
161 id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
162 range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
163 return createItem<Dims, false>(GlobalSize, GlobalId);
166 template <
int Dims>
static const nd_item<Dims> getElement(nd_item<Dims> *) {
167 static_assert(is_valid_dimensions<Dims>::value,
"invalid dimensions");
168 range<Dims> GlobalSize{__spirv::initGlobalSize<Dims, range<Dims>>()};
169 range<Dims> LocalSize{__spirv::initWorkgroupSize<Dims, range<Dims>>()};
170 range<Dims> GroupRange{__spirv::initNumWorkgroups<Dims, range<Dims>>()};
171 id<Dims> GroupId{__spirv::initWorkgroupId<Dims, id<Dims>>()};
172 id<Dims> GlobalId{__spirv::initGlobalInvocationId<Dims, id<Dims>>()};
173 id<Dims> LocalId{__spirv::initLocalInvocationId<Dims, id<Dims>>()};
174 id<Dims> GlobalOffset{__spirv::initGlobalOffset<Dims, id<Dims>>()};
176 createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
177 item<Dims, true> GlobalItem =
178 createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
179 item<Dims, false> LocalItem = createItem<Dims, false>(LocalSize, LocalId);
180 return createNDItem<Dims>(GlobalItem, LocalItem, Group);
183 template <
int Dims,
bool WithOffset>
184 static auto getElement(item<Dims, WithOffset> *)
185 -> decltype(getItem<Dims, WithOffset>()) {
186 return getItem<Dims, WithOffset>();
190 static auto getNDItem() -> decltype(getElement(
declptr<nd_item<Dims>>())) {
191 return getElement(
declptr<nd_item<Dims>>());
202 inline constexpr uint32_t
232 ?
static_cast<uint32_t
>(
236 ?
static_cast<uint32_t
>(
241 static_cast<uint32_t
>(
248 template <
size_t... Inds,
class F>
249 void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
250 (f(std::integral_constant<size_t, Inds>{}), ...);
253 template <
size_t count,
class F>
void loop(F &&f) {
254 loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
static std::enable_if_t< WithOffset, item< Dims, WithOffset > > createItem(const range< Dims > &Extent, const id< Dims > &Index, const id< Dims > &Offset)
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const range< Dims > &Group, const id< Dims > &Index)
static ResType createSubGroupMask(BitsType Bits, size_t BitsNum)
static h_item< Dims > createHItem(const item< Dims, false > &Global, const item< Dims, false > &Local, const range< Dims > &Flex)
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const id< Dims > &Index)
static void updateItemIndex(sycl::item< Dims, WithOffset > &Item, const id< Dims > &NextIndex)
static std::enable_if_t<!WithOffset, item< Dims, WithOffset > > createItem(const range< Dims > &Extent, const id< Dims > &Index)
static h_item< Dims > createHItem(const item< Dims, false > &Global, const item< Dims, false > &Local)
static nd_item< Dims > createNDItem(const item< Dims, true > &Global, const item< Dims, false > &Local, const group< Dims > &Group)
Identifies an instance of a group::parallel_for_work_item function object executing at each point in ...
A unique identifier of an item in an index space.
Identifies an instance of the function object executing at each point in a range.
Provides a cross-platform math array class template that works on SYCL devices as well as in host C++...
Identifies an instance of the function object executing at each point in an nd_range.
Defines the iteration domain of either a single work-group in a parallel dispatch,...
std::vector< sycl::detail::pi::PiEvent > getOrWaitEvents(std::vector< sycl::event > DepEvents, std::shared_ptr< sycl::detail::context_impl > Context)
T get_or_store(const T *obj)
void markBufferAsInternal(const std::shared_ptr< buffer_impl > &BufImpl)
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
void loop_impl(std::integer_sequence< size_t, Inds... >, F &&f)
void waitEvents(std::vector< sycl::event > DepEvents)
C++ wrapper of extern "C" PI interfaces.