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

[SYCL] Fixed-size groups and partitions are renamed to "chunks" #16151

Draft
wants to merge 6 commits into
base: sycl
Choose a base branch
from
Draft

[SYCL] Fixed-size groups and partitions are renamed to "chunks" #16151

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
45bb32f
[SYCL] Fixed-size groups and partitions are renamed to "chunks"
AndreiZibrov Nov 21, 2024
4abd474
Merge branch 'intel:sycl' into sycl
AndreiZibrov Nov 22, 2024
fb1d33d
SYCL] Fixed-size groups and partitions are renamed to "chunks"
AndreiZibrov Nov 24, 2024
5e74e57
SYCL] Fixed-size groups and partitions are renamed to "chunks"
AndreiZibrov Nov 24, 2024
66bafee
SYCL] Fixed-size groups and partitions are renamed to "chunks" + miss…
AndreiZibrov Nov 24, 2024
359fc93
Merge branch 'intel:sycl' into sycl
AndreiZibrov Nov 29, 2024
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
8 changes: 4 additions & 4 deletions llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,7 +71,7 @@ def AspectExt_oneapi_bindless_images_sample_1d_usm : Aspect<"ext_oneapi_bindless
def AspectExt_oneapi_bindless_images_sample_2d_usm : Aspect<"ext_oneapi_bindless_images_sample_2d_usm">;
def AspectExt_intel_esimd : Aspect<"ext_intel_esimd">;
def AspectExt_oneapi_ballot_group : Aspect<"ext_oneapi_ballot_group">;
def AspectExt_oneapi_fixed_size_group : Aspect<"ext_oneapi_fixed_size_group">;
def AspectExt_oneapi_chunk : Aspect<"ext_oneapi_chunk">;
def AspectExt_oneapi_opportunistic_group : Aspect<"ext_oneapi_opportunistic_group">;
def AspectExt_oneapi_tangle_group : Aspect<"ext_oneapi_tangle_group">;
def AspectExt_intel_matrix : Aspect<"ext_intel_matrix">;
Expand Down Expand Up @@ -144,7 +144,7 @@ def : TargetInfo<"__TestAspectList",
AspectExt_oneapi_bindless_sampled_image_fetch_2d_usm, AspectExt_oneapi_bindless_sampled_image_fetch_2d,
AspectExt_oneapi_bindless_sampled_image_fetch_3d,
AspectExt_intel_esimd,
AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
AspectExt_oneapi_ballot_group, AspectExt_oneapi_chunk, AspectExt_oneapi_opportunistic_group,
AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component,
AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph, AspectExt_oneapi_private_alloca,
AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group,
Expand All @@ -163,7 +163,7 @@ defvar IntelCpuAspects = [
AspectCpu, AspectFp16, AspectFp64, AspectQueue_profiling, AspectAtomic64,
AspectExt_oneapi_srgb, AspectExt_oneapi_native_assert,
AspectExt_intel_legacy_image, AspectExt_oneapi_ballot_group,
AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
AspectExt_oneapi_chunk, AspectExt_oneapi_opportunistic_group,
AspectExt_oneapi_tangle_group, AspectExt_oneapi_private_alloca
] # AllUSMAspects;

Expand Down Expand Up @@ -231,7 +231,7 @@ class CudaTargetInfo<string targetName, list<Aspect> aspectList, int subGroupSiz
defvar CudaMinAspects = !listconcat(AllUSMAspects, [AspectGpu, AspectFp64, AspectOnline_compiler, AspectOnline_linker,
AspectQueue_profiling, AspectExt_intel_pci_address, AspectExt_intel_max_mem_bandwidth, AspectExt_intel_memory_bus_width,
AspectExt_intel_device_info_uuid, AspectExt_oneapi_native_assert, AspectExt_intel_free_memory, AspectExt_intel_device_id,
AspectExt_intel_memory_clock_rate, AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group,
AspectExt_intel_memory_clock_rate, AspectExt_oneapi_ballot_group, AspectExt_oneapi_chunk,
AspectExt_oneapi_opportunistic_group, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]);
// Bindless images aspects are partially supported on CUDA and disabled by default at the moment.
defvar CudaBindlessImagesAspects = [AspectExt_oneapi_bindless_images, AspectExt_oneapi_bindless_images_shared_usm,
Expand Down
66 changes: 31 additions & 35 deletions sycl/include/sycl/detail/spirv.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ namespace oneapi {
struct sub_group;
namespace experimental {
template <typename ParentGroup> class ballot_group;
template <size_t PartitionSize, typename ParentGroup> class fixed_size_group;
template <size_t ChunkSize, typename ParentGroup> class chunk;
template <int Dimensions> class root_group;
template <typename ParentGroup> class tangle_group;
class opportunistic_group;
Expand Down Expand Up @@ -76,11 +76,11 @@ struct is_ballot_group<
sycl::ext::oneapi::experimental::ballot_group<ParentGroup>>
: std::true_type {};

template <typename Group> struct is_fixed_size_group : std::false_type {};
template <typename Group> struct is_chunk : std::false_type {};

template <size_t PartitionSize, typename ParentGroup>
struct is_fixed_size_group<sycl::ext::oneapi::experimental::fixed_size_group<
PartitionSize, ParentGroup>> : std::true_type {};
template <size_t ChunkSize, typename ParentGroup>
struct is_chunk<sycl::ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>>
: std::true_type {};

template <typename Group> struct group_scope {};

Expand All @@ -105,9 +105,9 @@ struct group_scope<sycl::ext::oneapi::experimental::ballot_group<ParentGroup>> {
static constexpr __spv::Scope::Flag value = group_scope<ParentGroup>::value;
};

template <size_t PartitionSize, typename ParentGroup>
struct group_scope<sycl::ext::oneapi::experimental::fixed_size_group<
PartitionSize, ParentGroup>> {
template <size_t ChunkSize, typename ParentGroup>
struct group_scope<
sycl::ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>> {
static constexpr __spv::Scope::Flag value = group_scope<ParentGroup>::value;
};

Expand Down Expand Up @@ -174,15 +174,14 @@ bool GroupAll(ext::oneapi::experimental::ballot_group<ParentGroup> g,
return __spirv_GroupNonUniformAll(group_scope<ParentGroup>::value, pred);
}
}
template <size_t PartitionSize, typename ParentGroup>
bool GroupAll(
ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>,
bool pred) {
template <size_t ChunkSize, typename ParentGroup>
bool GroupAll(ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
bool pred) {
// GroupNonUniformAll doesn't support cluster size, so use a reduction
return __spirv_GroupNonUniformBitwiseAnd(
group_scope<ParentGroup>::value,
static_cast<uint32_t>(__spv::GroupOperation::ClusteredReduce),
static_cast<uint32_t>(pred), PartitionSize);
static_cast<uint32_t>(pred), ChunkSize);
}
template <typename ParentGroup>
bool GroupAll(ext::oneapi::experimental::tangle_group<ParentGroup>, bool pred) {
Expand Down Expand Up @@ -210,15 +209,14 @@ bool GroupAny(ext::oneapi::experimental::ballot_group<ParentGroup> g,
return __spirv_GroupNonUniformAny(group_scope<ParentGroup>::value, pred);
}
}
template <size_t PartitionSize, typename ParentGroup>
bool GroupAny(
ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>,
bool pred) {
template <size_t ChunkSize, typename ParentGroup>
bool GroupAny(ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
bool pred) {
// GroupNonUniformAny doesn't support cluster size, so use a reduction
return __spirv_GroupNonUniformBitwiseOr(
group_scope<ParentGroup>::value,
static_cast<uint32_t>(__spv::GroupOperation::ClusteredReduce),
static_cast<uint32_t>(pred), PartitionSize);
static_cast<uint32_t>(pred), ChunkSize);
}
template <typename ParentGroup>
bool GroupAny(ext::oneapi::experimental::tangle_group<ParentGroup>, bool pred) {
Expand Down Expand Up @@ -327,12 +325,12 @@ GroupBroadcast(sycl::ext::oneapi::experimental::ballot_group<ParentGroup> g,
WideOCLX, OCLId);
}
}
template <size_t PartitionSize, typename ParentGroup, typename T, typename IdT>
EnableIfNativeBroadcast<T, IdT> GroupBroadcast(
ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup> g,
T x, IdT local_id) {
template <size_t ChunkSize, typename ParentGroup, typename T, typename IdT>
EnableIfNativeBroadcast<T, IdT>
GroupBroadcast(ext::oneapi::experimental::chunk<ChunkSize, ParentGroup> g, T x,
IdT local_id) {
// Remap local_id to its original numbering in ParentGroup
auto LocalId = g.get_group_linear_id() * PartitionSize + local_id;
auto LocalId = g.get_group_linear_id() * ChunkSize + local_id;

// TODO: Refactor to avoid duplication after design settles.
auto GroupLocalId = static_cast<typename GroupId<ParentGroup>::type>(LocalId);
Expand All @@ -341,9 +339,9 @@ EnableIfNativeBroadcast<T, IdT> GroupBroadcast(
auto OCLId = detail::convertToOpenCLType(GroupLocalId);

// NonUniformBroadcast requires Id to be dynamically uniform, which does not
// hold here; each partition is broadcasting a separate index. We could
// hold here; each chunk is broadcasting a separate index. We could
// fallback to either NonUniformShuffle or a NonUniformBroadcast per
// partition, and it's unclear which will be faster in practice.
// chunk, and it's unclear which will be faster in practice.
return __spirv_GroupNonUniformShuffle(group_scope<ParentGroup>::value,
WideOCLX, OCLId);
}
Expand Down Expand Up @@ -888,7 +886,7 @@ inline uint32_t MapShuffleID(GroupT g, id<1> local_id) {
if constexpr (is_tangle_or_opportunistic_group<GroupT>::value ||
is_ballot_group<GroupT>::value)
return detail::IdToMaskPosition(g, local_id);
else if constexpr (is_fixed_size_group<GroupT>::value)
else if constexpr (is_chunk<GroupT>::value)
return g.get_group_linear_id() * g.get_local_range().size() + local_id;
else
return local_id.get(0);
Expand Down Expand Up @@ -983,7 +981,7 @@ EnableIfNativeShuffle<T> ShuffleXor(GroupT g, T x, id<1> mask) {
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
GroupT>) {
auto MemberMask = detail::ExtractMask(detail::GetMask(g))[0];
if constexpr (is_fixed_size_group_v<GroupT>) {
if constexpr (is_chunk_v<GroupT>) {
return cuda_shfl_sync_bfly_i32(MemberMask, x,
static_cast<uint32_t>(mask.get(0)), 0x1f);

Expand Down Expand Up @@ -1031,7 +1029,7 @@ EnableIfNativeShuffle<T> ShuffleDown(GroupT g, T x, uint32_t delta) {
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
GroupT>) {
auto MemberMask = detail::ExtractMask(detail::GetMask(g))[0];
if constexpr (is_fixed_size_group_v<GroupT>) {
if constexpr (is_chunk_v<GroupT>) {
return cuda_shfl_sync_down_i32(MemberMask, x, delta, 31);
} else {
unsigned localSetBit = g.get_local_id()[0] + 1;
Expand Down Expand Up @@ -1075,7 +1073,7 @@ EnableIfNativeShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
GroupT>) {
auto MemberMask = detail::ExtractMask(detail::GetMask(g))[0];
if constexpr (is_fixed_size_group_v<GroupT>) {
if constexpr (is_chunk_v<GroupT>) {
return cuda_shfl_sync_up_i32(MemberMask, x, delta, 0);
} else {
unsigned localSetBit = g.get_local_id()[0] + 1;
Expand Down Expand Up @@ -1298,12 +1296,10 @@ ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
} \
} \
\
template <__spv::GroupOperation Op, size_t PartitionSize, \
typename ParentGroup, typename T> \
template <__spv::GroupOperation Op, size_t ChunkSize, typename ParentGroup, \
typename T> \
inline T Group##Instruction( \
ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup> \
g, \
T x) { \
ext::oneapi::experimental::chunk<ChunkSize, ParentGroup> g, T x) { \
using ConvertedT = detail::ConvertToOpenCLType_t<T>; \
\
using OCLT = std::conditional_t< \
Expand All @@ -1321,7 +1317,7 @@ ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
constexpr auto OpInt = \
static_cast<unsigned int>(__spv::GroupOperation::ClusteredReduce); \
return __spirv_GroupNonUniform##Instruction(Scope, OpInt, Arg, \
PartitionSize); \
ChunkSize); \
} else { \
T tmp; \
for (size_t Cluster = 0; Cluster < g.get_group_linear_range(); \
Expand Down
7 changes: 3 additions & 4 deletions sycl/include/sycl/detail/type_traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,10 +20,9 @@
namespace sycl {
inline namespace _V1 {
namespace detail {
template <class T> struct is_fixed_size_group : std::false_type {};
template <class T> struct is_chunk : std::false_type {};

template <class T>
inline constexpr bool is_fixed_size_group_v = is_fixed_size_group<T>::value;
template <class T> inline constexpr bool is_chunk_v = is_chunk<T>::value;

template <typename VecT, typename OperationLeftT, typename OperationRightT,
template <typename> class OperationCurrentT, int... Indexes>
Expand Down Expand Up @@ -157,7 +156,7 @@ template <typename T, int N> struct get_elem_type_unqual<vec<T, N>> {
template <typename VecT, typename OperationLeftT, typename OperationRightT,
template <typename> class OperationCurrentT, int... Indexes>
struct get_elem_type_unqual<SwizzleOp<VecT, OperationLeftT, OperationRightT,
OperationCurrentT, Indexes...>> {
OperationCurrentT, Indexes...>> {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Strictly speaking, there should be no unrelated formatting changes, but in this particular case it is more of a nitpicking and can be ignored.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

the formatting task failed until I called clang-format -i per changed file

strictly:

cd ../llvm &&
git diff --cached --name-only --diff-filter=ACM | grep '\.[ch]pp$'|xargs -I{} clang-format -i```

using type = typename get_elem_type_unqual<std::remove_cv_t<VecT>>::type;
};

Expand Down
71 changes: 35 additions & 36 deletions .../oneapi/experimental/fixed_size_group.hpp → ...de/sycl/ext/oneapi/experimental/chunk.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
//==--- fixed_size_group.hpp --- SYCL extension for non-uniform groups -----==//
//==--- chunk.hpp --- SYCL extension for non-uniform groups -----==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
Expand All @@ -10,7 +10,7 @@

#include <sycl/aspects.hpp>
#include <sycl/detail/spirv.hpp>
#include <sycl/detail/type_traits.hpp> // for is_fixed_size_group, is_group
#include <sycl/detail/type_traits.hpp> // for is_chunk, is_group
#include <sycl/exception.hpp>
#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
#include <sycl/ext/oneapi/sub_group_mask.hpp> // for sub_group_mask
Expand All @@ -26,28 +26,31 @@ namespace sycl {
inline namespace _V1 {
namespace ext::oneapi::experimental {

template <size_t PartitionSize, typename ParentGroup> class fixed_size_group;
template <size_t ChunkSize, typename ParentGroup> class chunk;

template <size_t PartitionSize, typename Group>
template <size_t ChunkSize, typename Group>
#ifdef __SYCL_DEVICE_ONLY__
[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_fixed_size_group)]]
[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_chunk)]]
#endif
inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
std::is_same_v<Group, sycl::sub_group>,
fixed_size_group<PartitionSize, Group>>
get_fixed_size_group(Group group);
chunk<ChunkSize, Group>> chunked_partition(Group group);

template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
template <size_t ChunkSize, typename ParentGroup> class chunk {
public:
using id_type = id<1>;
using range_type = range<1>;
using linear_id_type = typename ParentGroup::linear_id_type;
static constexpr int dimensions = 1;
static constexpr sycl::memory_scope fence_scope = ParentGroup::fence_scope;
/* ToDo:

fragment implementation to be done:
operator fragment<ParentGroup>() const;
*/
id_type get_group_id() const {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv_SubgroupLocalInvocationId() / PartitionSize;
return __spirv_SubgroupLocalInvocationId() / ChunkSize;
#else
throw exception(make_error_code(errc::runtime),
"Non-uniform groups are not supported on host.");
Expand All @@ -56,7 +59,7 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {

id_type get_local_id() const {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv_SubgroupLocalInvocationId() % PartitionSize;
return __spirv_SubgroupLocalInvocationId() % ChunkSize;
#else
throw exception(make_error_code(errc::runtime),
"Non-uniform groups are not supported on host.");
Expand All @@ -65,7 +68,7 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {

range_type get_group_range() const {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv_SubgroupSize() / PartitionSize;
return __spirv_SubgroupSize() / ChunkSize;
#else
throw exception(make_error_code(errc::runtime),
"Non-uniform groups are not supported on host.");
Expand All @@ -74,7 +77,7 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {

range_type get_local_range() const {
#ifdef __SYCL_DEVICE_ONLY__
return PartitionSize;
return ChunkSize;
#else
throw exception(make_error_code(errc::runtime),
"Non-uniform groups are not supported on host.");
Expand Down Expand Up @@ -132,62 +135,58 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
#endif

#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
fixed_size_group(ext::oneapi::sub_group_mask mask) : Mask(mask) {}
chunk(ext::oneapi::sub_group_mask mask) : Mask(mask) {}
#else
fixed_size_group() {}
chunk() {}
#endif

friend fixed_size_group<PartitionSize, ParentGroup>
get_fixed_size_group<PartitionSize, ParentGroup>(ParentGroup g);
friend chunk<ChunkSize, ParentGroup>
chunked_partition<ChunkSize, ParentGroup>(ParentGroup g);

friend sub_group_mask
sycl::detail::GetMask<fixed_size_group<PartitionSize, ParentGroup>>(
fixed_size_group<PartitionSize, ParentGroup> Group);
friend sub_group_mask sycl::detail::GetMask<chunk<ChunkSize, ParentGroup>>(
chunk<ChunkSize, ParentGroup> Group);
};

template <size_t PartitionSize, typename Group>
template <size_t ChunkSize, typename Group>
inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
std::is_same_v<Group, sycl::sub_group>,
fixed_size_group<PartitionSize, Group>>
get_fixed_size_group(Group group) {
chunk<ChunkSize, Group>>
chunked_partition(Group group) {
(void)group;
#ifdef __SYCL_DEVICE_ONLY__
#if defined(__NVPTX__)
uint32_t loc_id = group.get_local_linear_id();
uint32_t loc_size = group.get_local_linear_range();
uint32_t bits = PartitionSize == 32
? 0xffffffff
: ((1 << PartitionSize) - 1)
<< ((loc_id / PartitionSize) * PartitionSize);
uint32_t bits = ChunkSize == 32 ? 0xffffffff
: ((1 << ChunkSize) - 1)
<< ((loc_id / ChunkSize) * ChunkSize);

return fixed_size_group<PartitionSize, sycl::sub_group>(
return chunk<ChunkSize, sycl::sub_group>(
sycl::detail::Builder::createSubGroupMask<ext::oneapi::sub_group_mask>(
bits, loc_size));
#else
return fixed_size_group<PartitionSize, sycl::sub_group>();
return chunk<ChunkSize, sycl::sub_group>();
#endif
#else
throw exception(make_error_code(errc::runtime),
"Non-uniform groups are not supported on host.");
#endif
}

template <size_t PartitionSize, typename ParentGroup>
struct is_user_constructed_group<fixed_size_group<PartitionSize, ParentGroup>>
template <size_t ChunkSize, typename ParentGroup>
struct is_user_constructed_group<chunk<ChunkSize, ParentGroup>>
: std::true_type {};

} // namespace ext::oneapi::experimental

namespace detail {
template <size_t PartitionSize, typename ParentGroup>
struct is_fixed_size_group<
ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>>
template <size_t ChunkSize, typename ParentGroup>
struct is_chunk<ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>>
: std::true_type {};
} // namespace detail

template <size_t PartitionSize, typename ParentGroup>
struct is_group<
ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>>
template <size_t ChunkSize, typename ParentGroup>
struct is_group<ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>>
: std::true_type {};

} // namespace _V1
Expand Down
Loading
Loading