Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[ABI-Break][SYCL] Switch to new nd_item impl by default #13197

Merged
merged 1 commit into from
Apr 1, 2024
Merged

[ABI-Break][SYCL] Switch to new nd_item impl by default #13197

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
211 changes: 0 additions & 211 deletions sycl/include/sycl/nd_item.hpp
View file
Open in desktop
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
Open in desktop
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
Open in desktop

This file was deleted.

Loading