Skip to content

Commit

Permalink
[ABI-Break][SYCL] Switch to new nd_item impl by default (#13197)
Browse files Browse the repository at this point in the history
  • Loading branch information
@aelovikov-intel
aelovikov-intel authored Apr 1, 2024
1 parent 220a130 commit 2da97a5
Show file tree
Hide file tree
Showing 3 changed files with 7 additions and 358 deletions.
211 changes: 0 additions & 211 deletions sycl/include/sycl/nd_item.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -41,7 +41,6 @@ namespace ext::oneapi::experimental {
template <int Dimensions> class root_group;
}

#if __INTEL_PREVIEW_BREAKING_CHANGES
/// Identifies an instance of the function object executing at each point in an
/// nd_range.
///
Expand Down Expand Up @@ -536,216 +535,6 @@ template <int Dimensions = 1> class nd_item {
#endif
}
};
#else
/// Identifies an instance of the function object executing at each point in an
/// nd_range.
///
/// \ingroup sycl_api
template <int Dimensions = 1> class nd_item {
public:
static constexpr int dimensions = Dimensions;

nd_item() = delete;

id<Dimensions> get_global_id() const { return globalItem.get_id(); }

size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const {
size_t Id = globalItem.get_id(Dimension);
__SYCL_ASSUME_INT(Id);
return Id;
}

size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const {
size_t Id = globalItem.get_linear_id();
__SYCL_ASSUME_INT(Id);
return Id;
}

id<Dimensions> get_local_id() const { return localItem.get_id(); }

size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const {
size_t Id = localItem.get_id(Dimension);
__SYCL_ASSUME_INT(Id);
return Id;
}

size_t get_local_linear_id() const {
size_t Id = localItem.get_linear_id();
__SYCL_ASSUME_INT(Id);
return Id;
}

group<Dimensions> get_group() const { return Group; }

sub_group get_sub_group() const { return sub_group(); }

size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const {
size_t Id = Group[Dimension];
__SYCL_ASSUME_INT(Id);
return Id;
}

size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const {
size_t Id = Group.get_linear_id();
__SYCL_ASSUME_INT(Id);
return Id;
}

range<Dimensions> get_group_range() const { return Group.get_group_range(); }

size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const {
size_t Range = Group.get_group_range(Dimension);
__SYCL_ASSUME_INT(Range);
return Range;
}

range<Dimensions> get_global_range() const { return globalItem.get_range(); }

size_t get_global_range(int Dimension) const {
return globalItem.get_range(Dimension);
}

range<Dimensions> get_local_range() const { return localItem.get_range(); }

size_t get_local_range(int Dimension) const {
return localItem.get_range(Dimension);
}

__SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
id<Dimensions> get_offset() const { return globalItem.get_offset(); }

nd_range<Dimensions> get_nd_range() const {
return nd_range<Dimensions>(get_global_range(), get_local_range(),
get_offset());
}

void barrier(access::fence_space accessSpace =
access::fence_space::global_and_local) const {
uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace);
__spirv_ControlBarrier(__spv::Scope::Workgroup, __spv::Scope::Workgroup,
flags);
}

/// Executes a work-group mem-fence with memory ordering on the local address
/// space, global address space or both based on the value of \p accessSpace.
template <access::mode accessMode = access::mode::read_write>
__SYCL2020_DEPRECATED("use sycl::atomic_fence() free function instead")
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 {
(void)accessSpace;
Group.mem_fence();
}

template <typename dataT>
__SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
device_event
async_work_group_copy(local_ptr<dataT> dest, global_ptr<dataT> src,
size_t numElements) const {
return Group.async_work_group_copy(dest, src, numElements);
}

template <typename dataT>
__SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
device_event
async_work_group_copy(global_ptr<dataT> dest, local_ptr<dataT> src,
size_t numElements) const {
return Group.async_work_group_copy(dest, src, numElements);
}

template <typename dataT>
__SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
device_event
async_work_group_copy(local_ptr<dataT> dest, global_ptr<dataT> src,
size_t numElements, size_t srcStride) const {

return Group.async_work_group_copy(dest, src, numElements, srcStride);
}

template <typename dataT>
__SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
device_event
async_work_group_copy(global_ptr<dataT> dest, local_ptr<dataT> src,
size_t numElements, size_t destStride) const {
return Group.async_work_group_copy(dest, src, numElements, destStride);
}

template <typename DestDataT, typename SrcDataT>
typename 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 {
return Group.async_work_group_copy(dest, src, numElements);
}

template <typename DestDataT, typename SrcDataT>
typename 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 {
return Group.async_work_group_copy(dest, src, numElements);
}

template <typename DestDataT, typename SrcDataT>
typename 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,
size_t srcStride) const {

return Group.async_work_group_copy(dest, src, numElements, srcStride);
}

template <typename DestDataT, typename SrcDataT>
typename 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,
size_t destStride) const {
return Group.async_work_group_copy(dest, src, numElements, destStride);
}

template <typename... eventTN> void wait_for(eventTN... events) const {
Group.wait_for(events...);
}

sycl::ext::oneapi::experimental::root_group<Dimensions>
ext_oneapi_get_root_group() const {
return sycl::ext::oneapi::experimental::root_group<Dimensions>{*this};
}

nd_item(const nd_item &rhs) = default;

nd_item(nd_item &&rhs) = default;

nd_item &operator=(const nd_item &rhs) = default;

nd_item &operator=(nd_item &&rhs) = default;

bool operator==(const nd_item &rhs) const {
return (rhs.localItem == this->localItem) &&
(rhs.globalItem == this->globalItem) && (rhs.Group == this->Group);
}

bool operator!=(const nd_item &rhs) const { return !((*this) == rhs); }

protected:
friend class detail::Builder;
nd_item(const item<Dimensions, true> &GL, const item<Dimensions, false> &L,
const group<Dimensions> &GR)
: globalItem(GL), localItem(L), Group(GR) {}

private:
item<Dimensions, true> globalItem;
item<Dimensions, false> localItem;
group<Dimensions> Group;
};
#endif

template <int Dims>
__SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_nd_item() instead")
Expand Down
47 changes: 7 additions & 40 deletions sycl/test/abi/layout_array.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,13 @@ SYCL_EXTERNAL void range(sycl::range<2>) {}

//----------------------------

SYCL_EXTERNAL void nd_item(sycl::nd_item<2>) {}
// CHECK: 0 | class sycl::nd_item<> (empty)
// CHECK-NEXT: | [sizeof=1, dsize=0, align=1,
// CHECK-NEXT: | nvsize=0, nvalign=1]

//----------------------------

SYCL_EXTERNAL void item(sycl::item<2>) {}

// CHECK: 0 | class sycl::item<2>
Expand All @@ -47,46 +54,6 @@ SYCL_EXTERNAL void item(sycl::item<2>) {}

//----------------------------

SYCL_EXTERNAL void nd_item(sycl::nd_item<2>) {}

// CHECK: 0 | class sycl::nd_item<2>
// CHECK-NEXT: 0 | class sycl::item<2> globalItem
// CHECK-NEXT: 0 | struct sycl::detail::ItemBase<2, true> MImpl
// CHECK-NEXT: 0 | class sycl::range<2> MExtent
// CHECK-NEXT: 0 | class sycl::detail::array<2> (base)
// CHECK-NEXT: 0 | size_t[2] common_array
// CHECK-NEXT: 16 | class sycl::id<2> MIndex
// CHECK-NEXT: 16 | class sycl::detail::array<2> (base)
// CHECK-NEXT: 16 | size_t[2] common_array
// CHECK-NEXT: 32 | class sycl::id<2> MOffset
// CHECK-NEXT: 32 | class sycl::detail::array<2> (base)
// CHECK-NEXT: 32 | size_t[2] common_array
// CHECK-NEXT: 48 | class sycl::item<2, false> localItem
// CHECK-NEXT: 48 | struct sycl::detail::ItemBase<2, false> MImpl
// CHECK-NEXT: 48 | class sycl::range<2> MExtent
// CHECK-NEXT: 48 | class sycl::detail::array<2> (base)
// CHECK-NEXT: 48 | size_t[2] common_array
// CHECK-NEXT: 64 | class sycl::id<2> MIndex
// CHECK-NEXT: 64 | class sycl::detail::array<2> (base)
// CHECK-NEXT: 64 | size_t[2] common_array
// CHECK-NEXT: 80 | class sycl::group<2> Group
// CHECK-NEXT: 80 | class sycl::range<2> globalRange
// CHECK-NEXT: 80 | class sycl::detail::array<2> (base)
// CHECK-NEXT: 80 | size_t[2] common_array
// CHECK-NEXT: 96 | class sycl::range<2> localRange
// CHECK-NEXT: 96 | class sycl::detail::array<2> (base)
// CHECK-NEXT: 96 | size_t[2] common_array
// CHECK-NEXT: 112 | class sycl::range<2> groupRange
// CHECK-NEXT: 112 | class sycl::detail::array<2> (base)
// CHECK-NEXT: 112 | size_t[2] common_array
// CHECK-NEXT: 128 | class sycl::id<2> index
// CHECK-NEXT: 128 | class sycl::detail::array<2> (base)
// CHECK-NEXT: 128 | size_t[2] common_array
// CHECK-NEXT: | [sizeof=144, dsize=144, align=8,
// CHECK-NEXT: | nvsize=144, nvalign=8]

//----------------------------

SYCL_EXTERNAL void nd_range(sycl::nd_range<2>) {}
// CHECK: 0 | class sycl::nd_range<2>
// CHECK-NEXT: 0 | class sycl::range<2> globalSize
Expand Down
107 changes: 0 additions & 107 deletions sycl/test/basic_tests/nd_item.cpp
View file

This file was deleted.

0 comments on commit 2da97a5

Please sign in to comment.