32 #include <type_traits>
35 inline namespace _V1 {
40 namespace ext::oneapi::experimental {
41 template <
int Dimensions>
class root_group;
48 template <
int Dimensions = 1>
class nd_item {
53 #ifdef __SYCL_DEVICE_ONLY__
54 return __spirv::initGlobalInvocationId<Dimensions, id<Dimensions>>();
57 "nd_item methods can't be invoked on the host");
74 LinId = Index[0] - Offset[0];
76 LinId = (Index[0] - Offset[0]) * Extent[1] + Index[1] - Offset[1];
78 LinId = (Index[0] - Offset[0]) * Extent[1] * Extent[2] +
79 (Index[1] - Offset[1]) * Extent[2] + Index[2] - Offset[2];
86 #ifdef __SYCL_DEVICE_ONLY__
87 return __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>();
90 "nd_item methods can't be invoked on the host");
108 LinId = Index[0] * Extent[1] + Index[1];
111 Index[0] * Extent[1] * Extent[2] + Index[1] * Extent[2] + Index[2];
139 LinId = Index[0] * Extent[1] + Index[1];
142 Index[0] * Extent[1] * Extent[2] + Index[1] * Extent[2] + Index[2];
149 #ifdef __SYCL_DEVICE_ONLY__
150 return __spirv::initNumWorkgroups<Dimensions, range<Dimensions>>();
153 "nd_item methods can't be invoked on the host");
165 #ifdef __SYCL_DEVICE_ONLY__
166 return __spirv::initGlobalSize<Dimensions, range<Dimensions>>();
169 "nd_item methods can't be invoked on the host");
181 #ifdef __SYCL_DEVICE_ONLY__
182 return __spirv::initWorkgroupSize<Dimensions, range<Dimensions>>();
185 "nd_item methods can't be invoked on the host");
198 #ifdef __SYCL_DEVICE_ONLY__
199 return __spirv::initGlobalOffset<Dimensions, id<Dimensions>>();
202 "nd_item methods can't be invoked on the host");
221 template <access::mode accessMode = access::mode::read_write>
224 typename
std::enable_if_t<accessMode == access::
mode::read ||
225 accessMode == access::
mode::write ||
228 accessSpace = access::
fence_space::global_and_local)
const {
246 template <
typename dataT>
248 std::enable_if_t<!detail::is_bool<dataT>::value,
252 size_t srcStride)
const {
258 numElements, srcStride, 0);
267 template <
typename dataT>
269 std::enable_if_t<!detail::is_bool<dataT>::value,
280 numElements, destStride, 0);
290 template <
typename DestDataT,
typename SrcDataT>
291 std::enable_if_t<!detail::is_bool<DestDataT>::value &&
292 std::is_same_v<std::remove_const_t<SrcDataT>, DestDataT>,
296 size_t srcStride)
const {
302 numElements, srcStride, 0);
312 template <
typename DestDataT,
typename SrcDataT>
313 std::enable_if_t<!detail::is_bool<DestDataT>::value &&
314 std::is_same_v<std::remove_const_t<SrcDataT>, DestDataT>,
318 size_t destStride)
const {
324 numElements, destStride, 0);
333 template <
typename T, access::address_space DestS, access::address_space SrcS>
336 detail::is_scalar_bool<T>::value,
344 size_t Stride)
const {
345 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
346 "Async copy to/from bool memory is not supported.");
348 reinterpret_cast<uint8_t *
>(Dest.get()));
350 reinterpret_cast<uint8_t *
>(Src.get()));
359 template <
typename T, access::address_space DestS, access::address_space SrcS>
362 detail::is_vector_bool<T>::value,
370 size_t Stride)
const {
371 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
372 "Async copy to/from bool memory is not supported.");
374 auto DestP = address_space_cast<DestS, access::decorated::legacy>(
375 reinterpret_cast<VecT *
>(Dest.get()));
376 auto SrcP = address_space_cast<SrcS, access::decorated::legacy>(
377 reinterpret_cast<VecT *
>(Src.get()));
388 std::enable_if_t<detail::is_scalar_bool<DestT>::value &&
389 std::is_same_v<std::remove_const_t<SrcT>, DestT>,
393 size_t NumElements,
size_t Stride)
const {
394 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
395 "Async copy to/from bool memory is not supported.");
397 std::conditional_t<std::is_const_v<SrcT>,
const uint8_t, uint8_t>;
401 Dest.get_decorated()));
405 Src.get_decorated()));
416 std::enable_if_t<detail::is_vector_bool<DestT>::value &&
417 std::is_same_v<std::remove_const_t<SrcT>, DestT>,
421 size_t NumElements,
size_t Stride)
const {
422 static_assert(
sizeof(
bool) ==
sizeof(uint8_t),
423 "Async copy to/from bool memory is not supported.");
426 std::conditional_t<std::is_const_v<SrcT>, std::add_const_t<VecT>, VecT>;
430 Dest.get_decorated()));
434 Src.get_decorated()));
443 template <
typename dataT>
447 size_t numElements)
const {
456 template <
typename dataT>
460 size_t numElements)
const {
470 template <
typename DestDataT,
typename SrcDataT>
471 typename std::enable_if_t<
472 std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>,
device_event>
475 size_t numElements)
const {
485 template <
typename DestDataT,
typename SrcDataT>
486 typename std::enable_if_t<
487 std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>,
device_event>
490 size_t numElements)
const {
494 template <
typename... eventTN>
void wait_for(eventTN... events)
const {
522 template <
typename T,
typename... Ts>
529 #ifdef __SYCL_DEVICE_ONLY__
530 return __spirv::initWorkgroupId<Dimensions, id<Dimensions>>();
533 "nd_item methods can't be invoked on the host");
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 __SYCL_TYPE(local_accessor) local_accessor 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
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
__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