32 #include <type_traits>
35 inline namespace _V1 {
40 namespace ext::oneapi::experimental {
48 template <
int Dimensions = 1>
class nd_item {
53 #ifdef __SYCL_DEVICE_ONLY__
54 return __spirv::initGlobalInvocationId<Dimensions, id<Dimensions>>();
72 LinId = Index[0] - Offset[0];
74 LinId = (Index[0] - Offset[0]) * Extent[1] + Index[1] - Offset[1];
76 LinId = (Index[0] - Offset[0]) * Extent[1] * Extent[2] +
77 (Index[1] - Offset[1]) * Extent[2] + Index[2] - Offset[2];
84 #ifdef __SYCL_DEVICE_ONLY__
85 return __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>();
104 LinId = Index[0] * Extent[1] + Index[1];
107 Index[0] * Extent[1] * Extent[2] + Index[1] * Extent[2] + Index[2];
135 LinId = Index[0] * Extent[1] + Index[1];
138 Index[0] * Extent[1] * Extent[2] + Index[1] * Extent[2] + Index[2];
145 #ifdef __SYCL_DEVICE_ONLY__
146 return __spirv::initNumWorkgroups<Dimensions, range<Dimensions>>();
159 #ifdef __SYCL_DEVICE_ONLY__
160 return __spirv::initGlobalSize<Dimensions, range<Dimensions>>();
173 #ifdef __SYCL_DEVICE_ONLY__
174 return __spirv::initWorkgroupSize<Dimensions, range<Dimensions>>();
188 #ifdef __SYCL_DEVICE_ONLY__
189 return __spirv::initGlobalOffset<Dimensions, id<Dimensions>>();
209 template <access::mode accessMode = access::mode::read_write>
212 typename
std::enable_if_t<accessMode == access::
mode::read ||
213 accessMode == access::
mode::write ||
216 accessSpace = access::
fence_space::global_and_local)
const {
234 template <
typename dataT>
236 std::enable_if_t<!detail::is_bool<dataT>::value,
240 size_t srcStride)
const {
246 numElements, srcStride, 0);
255 template <
typename dataT>
257 std::enable_if_t<!detail::is_bool<dataT>::value,
268 numElements, destStride, 0);
278 template <
typename DestDataT,
typename SrcDataT>
279 std::enable_if_t<!detail::is_bool<DestDataT>::value &&
280 std::is_same_v<std::remove_const_t<SrcDataT>, DestDataT>,
284 size_t srcStride)
const {
290 numElements, srcStride, 0);
300 template <
typename DestDataT,
typename SrcDataT>
301 std::enable_if_t<!detail::is_bool<DestDataT>::value &&
302 std::is_same_v<std::remove_const_t<SrcDataT>, DestDataT>,
306 size_t destStride)
const {
312 numElements, destStride, 0);
321 template <
typename T, access::address_space DestS, access::address_space SrcS>
324 detail::is_scalar_bool<T>::value,
332 size_t Stride)
const {
333 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
334 "Async copy to/from bool memory is not supported.");
336 reinterpret_cast<uint8_t *
>(Dest.get()));
338 reinterpret_cast<uint8_t *
>(Src.get()));
347 template <
typename T, access::address_space DestS, access::address_space SrcS>
350 detail::is_vector_bool<T>::value,
358 size_t Stride)
const {
359 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
360 "Async copy to/from bool memory is not supported.");
362 auto DestP = address_space_cast<DestS, access::decorated::legacy>(
363 reinterpret_cast<VecT *
>(Dest.get()));
364 auto SrcP = address_space_cast<SrcS, access::decorated::legacy>(
365 reinterpret_cast<VecT *
>(Src.get()));
376 std::enable_if_t<detail::is_scalar_bool<DestT>::value &&
377 std::is_same_v<std::remove_const_t<SrcT>, DestT>,
381 size_t NumElements,
size_t Stride)
const {
382 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
383 "Async copy to/from bool memory is not supported.");
385 std::conditional_t<std::is_const_v<SrcT>,
const uint8_t, uint8_t>;
389 Dest.get_decorated()));
393 Src.get_decorated()));
404 std::enable_if_t<detail::is_vector_bool<DestT>::value &&
405 std::is_same_v<std::remove_const_t<SrcT>, DestT>,
409 size_t NumElements,
size_t Stride)
const {
410 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
411 "Async copy to/from bool memory is not supported.");
414 std::conditional_t<std::is_const_v<SrcT>, std::add_const_t<VecT>, VecT>;
418 Dest.get_decorated()));
422 Src.get_decorated()));
431 template <
typename dataT>
435 size_t numElements)
const {
444 template <
typename dataT>
448 size_t numElements)
const {
458 template <
typename DestDataT,
typename SrcDataT>
459 typename std::enable_if_t<
460 std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>,
device_event>
463 size_t numElements)
const {
473 template <
typename DestDataT,
typename SrcDataT>
474 typename std::enable_if_t<
475 std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>,
device_event>
478 size_t numElements)
const {
482 template <
typename... eventTN>
void wait_for(eventTN... events)
const {
510 template <
typename T,
typename... Ts>
517 #ifdef __SYCL_DEVICE_ONLY__
518 return __spirv::initWorkgroupId<Dimensions, id<Dimensions>>();
static group< Dims > createGroup(const range< Dims > &Global, const range< Dims > &Local, const range< Dims > &Group, const id< Dims > &Index)
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.
Identifies an instance of the function object executing at each point in a range.
Identifies an instance of the function object executing at each point in an nd_range.
id< Dimensions > get_group_id() const
std::enable_if_t<!detail::is_bool< DestDataT >::value &&std::is_same_v< std::remove_const_t< SrcDataT >, DestDataT >, device_event > async_work_group_copy(decorated_local_ptr< DestDataT > dest, decorated_global_ptr< SrcDataT > src, size_t numElements, size_t srcStride) const
Asynchronously copies a number of elements specified by numElements from the source pointed by src to...
size_t get_local_linear_id() const
size_t get_local_range(int Dimension) const
void waitForHelper(T E, Ts... Es) const
std::enable_if_t< detail::is_scalar_bool< DestT >::value &&std::is_same_v< std::remove_const_t< SrcT >, DestT >, device_event > async_work_group_copy(multi_ptr< DestT, DestS, access::decorated::yes > Dest, multi_ptr< SrcT, SrcS, access::decorated::yes > Src, size_t NumElements, size_t Stride) const
Specialization for scalar bool type.
size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const
nd_item(const nd_item &rhs)=default
size_t get_global_range(int Dimension) const
sycl::ext::oneapi::experimental::root_group< Dimensions > ext_oneapi_get_root_group() const
id< Dimensions > get_local_id() const
nd_item(const item< Dimensions, true > &, const item< Dimensions, false > &, const group< Dimensions > &)
range< Dimensions > get_local_range() const
id< Dimensions > get_global_id() const
sub_group get_sub_group() const
void wait_for(eventTN... events) const
nd_item & operator=(const nd_item &rhs)=default
std::enable_if_t< detail::is_vector_bool< DestT >::value &&std::is_same_v< std::remove_const_t< SrcT >, DestT >, device_event > async_work_group_copy(multi_ptr< DestT, DestS, access::decorated::yes > Dest, multi_ptr< SrcT, SrcS, access::decorated::yes > Src, size_t NumElements, size_t Stride) const
Specialization for vector bool type.
std::enable_if_t<!detail::is_bool< dataT >::value, device_event > async_work_group_copy(local_ptr< dataT > dest, global_ptr< dataT > src, size_t numElements, size_t srcStride) const
Asynchronously copies a number of elements specified by numElements from the source pointed by src to...
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
Asynchronously copies a number of elements specified by numElements from the source pointed by src to...
std::enable_if_t<!detail::is_bool< DestDataT >::value &&std::is_same_v< std::remove_const_t< SrcDataT >, DestDataT >, device_event > async_work_group_copy(decorated_global_ptr< DestDataT > dest, decorated_local_ptr< SrcDataT > src, size_t numElements, size_t destStride) const
Asynchronously copies a number of elements specified by numElements from the source pointed by src to...
nd_range< Dimensions > get_nd_range() const
size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const
void waitForHelper() const
bool operator==(const nd_item &) const
size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const
bool operator!=(const nd_item &rhs) const
nd_item(nd_item &&rhs)=default
void waitForHelper(device_event Event) const
size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const
static constexpr int dimensions
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...
id< Dimensions > get_offset() const
group< Dimensions > get_group() const
nd_item & operator=(nd_item &&rhs)=default
size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const
range< Dimensions > get_global_range() const
range< Dimensions > get_group_range() const
size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const
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
Asynchronously copies a number of elements specified by numElements from the source pointed by src to...
void barrier(access::fence_space accessSpace=access::fence_space::global_and_local) const
Defines the iteration domain of both the work-groups and the overall dispatch.
#define __SYCL_ASSUME_INT(x)
#define __SYCL_ALWAYS_INLINE
typename change_base_type< T, B >::type change_base_type_t
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
decltype(convertToOpenCLType(std::declval< T >())) ConvertToOpenCLType_t
class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
signed char __SYCL2020_DEPRECATED
constexpr mode_tag_t< access_mode::read_write > read_write
std::conditional_t< is_decorated, decorated_type *, std::add_pointer_t< value_type > > pointer
__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
__SYCL_CONVERGENT__ __DPCPP_SYCL_EXTERNAL void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept
__SYCL_CONVERGENT__ __DPCPP_SYCL_EXTERNAL void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept
__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