Skip to content

Commit 7379a88

Browse files
AndreiZibrovAndreiZibrov
committed
[SYCL] Fixed-size tangle group renamed to "tangle" + formatting
1 parent a19b66f commit 7379a88

File tree

13 files changed

+44
-44
lines changed

13 files changed

+44
-44
lines changed

‎llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -73,7 +73,7 @@ def AspectExt_intel_esimd : Aspect<"ext_intel_esimd">;
7373
def AspectExt_oneapi_ballot_group : Aspect<"ext_oneapi_ballot_group">;
7474
def AspectExt_oneapi_chunk : Aspect<"ext_oneapi_chunk">;
7575
def AspectExt_oneapi_opportunistic_group : Aspect<"ext_oneapi_opportunistic_group">;
76-
def AspectExt_oneapi_tangle_group : Aspect<"ext_oneapi_tangle_group">;
76+
def AspectExt_oneapi_tangle : Aspect<"ext_oneapi_tangle">;
7777
def AspectExt_intel_matrix : Aspect<"ext_intel_matrix">;
7878
def AspectExt_oneapi_is_composite : Aspect<"ext_oneapi_is_composite">;
7979
def AspectExt_oneapi_is_component : Aspect<"ext_oneapi_is_component">;
@@ -145,7 +145,7 @@ def : TargetInfo<"__TestAspectList",
145145
AspectExt_oneapi_bindless_sampled_image_fetch_3d,
146146
AspectExt_intel_esimd,
147147
AspectExt_oneapi_ballot_group, AspectExt_oneapi_chunk, AspectExt_oneapi_opportunistic_group,
148-
AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component,
148+
AspectExt_oneapi_tangle, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component,
149149
AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph, AspectExt_oneapi_private_alloca,
150150
AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group,
151151
AspectExt_intel_fpga_task_sequence,
@@ -164,7 +164,7 @@ defvar IntelCpuAspects = [
164164
AspectExt_oneapi_srgb, AspectExt_oneapi_native_assert,
165165
AspectExt_intel_legacy_image, AspectExt_oneapi_ballot_group,
166166
AspectExt_oneapi_chunk, AspectExt_oneapi_opportunistic_group,
167-
AspectExt_oneapi_tangle_group, AspectExt_oneapi_private_alloca
167+
AspectExt_oneapi_tangle, AspectExt_oneapi_private_alloca
168168
] # AllUSMAspects;
169169

170170
def : TargetInfo<"spir64", [], [], "", "", 1>;

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_tangle.asciidoc‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -318,7 +318,7 @@ branches to safely communicate between all work-items executing the same
318318
control flow.
319319

320320
NOTE: This differs from the `fragment` returned by `get_opportunistic_group()`
321-
because a `tangle_group` requires the implementation to track group membership.
321+
because a `tangle` requires the implementation to track group membership.
322322
Which group type to use will depend on a combination of
323323
implementation/backend/device and programmer preference.
324324

‎sycl/include/sycl/detail/spirv.hpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -28,7 +28,7 @@ namespace experimental {
2828
template <typename ParentGroup> class ballot_group;
2929
template <size_t ChunkSize, typename ParentGroup> class chunk;
3030
template <int Dimensions> class root_group;
31-
template <typename ParentGroup> class tangle_group;
31+
template <typename ParentGroup> class tangle;
3232
class opportunistic_group;
3333
} // namespace experimental
3434
} // namespace oneapi
@@ -62,7 +62,7 @@ struct is_tangle_or_opportunistic_group : std::false_type {};
6262

6363
template <typename ParentGroup>
6464
struct is_tangle_or_opportunistic_group<
65-
sycl::ext::oneapi::experimental::tangle_group<ParentGroup>>
65+
sycl::ext::oneapi::experimental::tangle<ParentGroup>>
6666
: std::true_type {};
6767

6868
template <>
@@ -112,7 +112,7 @@ struct group_scope<
112112
};
113113

114114
template <typename ParentGroup>
115-
struct group_scope<sycl::ext::oneapi::experimental::tangle_group<ParentGroup>> {
115+
struct group_scope<sycl::ext::oneapi::experimental::tangle<ParentGroup>> {
116116
static constexpr __spv::Scope::Flag value = group_scope<ParentGroup>::value;
117117
};
118118

@@ -184,7 +184,7 @@ bool GroupAll(ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
184184
static_cast<uint32_t>(pred), ChunkSize);
185185
}
186186
template <typename ParentGroup>
187-
bool GroupAll(ext::oneapi::experimental::tangle_group<ParentGroup>, bool pred) {
187+
bool GroupAll(ext::oneapi::experimental::tangle<ParentGroup>, bool pred) {
188188
return __spirv_GroupNonUniformAll(group_scope<ParentGroup>::value, pred);
189189
}
190190

@@ -219,7 +219,7 @@ bool GroupAny(ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
219219
static_cast<uint32_t>(pred), ChunkSize);
220220
}
221221
template <typename ParentGroup>
222-
bool GroupAny(ext::oneapi::experimental::tangle_group<ParentGroup>, bool pred) {
222+
bool GroupAny(ext::oneapi::experimental::tangle<ParentGroup>, bool pred) {
223223
return __spirv_GroupNonUniformAny(group_scope<ParentGroup>::value, pred);
224224
}
225225
bool GroupAny(const ext::oneapi::experimental::opportunistic_group &,
@@ -347,7 +347,7 @@ GroupBroadcast(ext::oneapi::experimental::chunk<ChunkSize, ParentGroup> g, T x,
347347
}
348348
template <typename ParentGroup, typename T, typename IdT>
349349
EnableIfNativeBroadcast<T, IdT>
350-
GroupBroadcast(ext::oneapi::experimental::tangle_group<ParentGroup> g, T x,
350+
GroupBroadcast(ext::oneapi::experimental::tangle<ParentGroup> g, T x,
351351
IdT local_id) {
352352
// Remap local_id to its original numbering in ParentGroup.
353353
auto LocalId = detail::IdToMaskPosition(g, local_id);

‎sycl/include/sycl/ext/oneapi/experimental/non_uniform_groups.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -74,7 +74,7 @@ namespace ext::oneapi::experimental {
7474
// Forward declarations of non-uniform group types for algorithm definitions
7575
template <typename ParentGroup> class ballot_group;
7676
template <size_t ChunkSize, typename ParentGroup> class chunk;
77-
template <typename ParentGroup> class tangle_group;
77+
template <typename ParentGroup> class tangle;
7878
class opportunistic_group;
7979

8080
} // namespace ext::oneapi::experimental

‎sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp‎ renamed to ‎sycl/include/sycl/ext/oneapi/experimental/tangle.hpp‎

Lines changed: 16 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//==------ tangle_group.hpp --- SYCL extension for non-uniform groups ------==//
1+
//==------ tangle.hpp --- SYCL extension for non-uniform groups ------==//
22
//
33
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
44
// See https://llvm.org/LICENSE.txt for license information.
@@ -25,17 +25,17 @@ namespace sycl {
2525
inline namespace _V1 {
2626
namespace ext::oneapi::experimental {
2727

28-
template <typename ParentGroup> class tangle_group;
28+
template <typename ParentGroup> class tangle;
2929

3030
template <typename Group>
3131
#ifdef __SYCL_DEVICE_ONLY__
32-
[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_tangle_group)]]
32+
[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_tangle)]]
3333
#endif
3434
inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
3535
std::is_same_v<Group, sycl::sub_group>,
36-
tangle_group<Group>> get_tangle_group(Group group);
36+
tangle<Group>> get_tangle(Group group);
3737

38-
template <typename ParentGroup> class tangle_group {
38+
template <typename ParentGroup> class tangle {
3939
public:
4040
using id_type = id<1>;
4141
using range_type = range<1>;
@@ -128,19 +128,19 @@ template <typename ParentGroup> class tangle_group {
128128
protected:
129129
sub_group_mask Mask;
130130

131-
tangle_group(sub_group_mask m) : Mask(m) {}
131+
tangle(sub_group_mask m) : Mask(m) {}
132132

133-
friend tangle_group<ParentGroup> get_tangle_group<ParentGroup>(ParentGroup);
133+
friend tangle<ParentGroup> get_tangle<ParentGroup>(ParentGroup);
134134

135-
friend sub_group_mask sycl::detail::GetMask<tangle_group<ParentGroup>>(
136-
tangle_group<ParentGroup> Group);
135+
friend sub_group_mask sycl::detail::GetMask<tangle<ParentGroup>>(
136+
tangle<ParentGroup> Group);
137137
};
138138

139139
template <typename Group>
140140
inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
141141
std::is_same_v<Group, sycl::sub_group>,
142-
tangle_group<Group>>
143-
get_tangle_group(Group group) {
142+
tangle<Group>>
143+
get_tangle(Group group) {
144144
(void)group;
145145
#ifdef __SYCL_DEVICE_ONLY__
146146
#if defined(__SPIR__) || defined(__SPIRV__)
@@ -149,12 +149,12 @@ get_tangle_group(Group group) {
149149
// We store the mask here because it is required to calculate IDs, not
150150
// because it is required to construct the group.
151151
sub_group_mask mask = sycl::ext::oneapi::group_ballot(group, true);
152-
return tangle_group<sycl::sub_group>(mask);
152+
return tangle<sycl::sub_group>(mask);
153153
#elif defined(__NVPTX__)
154154
// TODO: Construct from compiler-generated mask. Return an invalid group in
155-
// in the meantime. CUDA devices will report false for the tangle_group
155+
// in the meantime. CUDA devices will report false for the tangle
156156
// support aspect so kernels launch should ensure this is never run.
157-
return tangle_group<sycl::sub_group>(0);
157+
return tangle<sycl::sub_group>(0);
158158
#endif
159159
#else
160160
throw exception(make_error_code(errc::runtime),
@@ -164,12 +164,12 @@ get_tangle_group(Group group) {
164164
} // namespace this_kernel
165165

166166
template <typename ParentGroup>
167-
struct is_user_constructed_group<tangle_group<ParentGroup>> : std::true_type {};
167+
struct is_user_constructed_group<tangle<ParentGroup>> : std::true_type {};
168168

169169
} // namespace ext::oneapi::experimental
170170

171171
template <typename ParentGroup>
172-
struct is_group<ext::oneapi::experimental::tangle_group<ParentGroup>>
172+
struct is_group<ext::oneapi::experimental::tangle<ParentGroup>>
173173
: std::true_type {};
174174

175175
} // namespace _V1

‎sycl/include/sycl/info/aspects.def‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -48,7 +48,7 @@ __SYCL_ASPECT(ext_intel_esimd, 53)
4848
__SYCL_ASPECT(ext_oneapi_ballot_group, 54)
4949
__SYCL_ASPECT(ext_oneapi_chunk, 55)
5050
__SYCL_ASPECT(ext_oneapi_opportunistic_group, 56)
51-
__SYCL_ASPECT(ext_oneapi_tangle_group, 57)
51+
__SYCL_ASPECT(ext_oneapi_tangle, 57)
5252
__SYCL_ASPECT(ext_intel_matrix, 58)
5353
__SYCL_ASPECT(ext_oneapi_is_composite, 59)
5454
__SYCL_ASPECT(ext_oneapi_is_component, 60)

‎sycl/include/sycl/sycl.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -103,7 +103,7 @@
103103
#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
104104
#include <sycl/ext/oneapi/experimental/reduction_properties.hpp>
105105
#include <sycl/ext/oneapi/experimental/root_group.hpp>
106-
#include <sycl/ext/oneapi/experimental/tangle_group.hpp>
106+
#include <sycl/ext/oneapi/experimental/tangle.hpp>
107107
#include <sycl/ext/oneapi/experimental/work_group_memory.hpp>
108108
#include <sycl/ext/oneapi/filter_selector.hpp>
109109
#include <sycl/ext/oneapi/free_function_queries.hpp>

‎sycl/source/detail/device_impl.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -664,8 +664,8 @@ bool device_impl::has(aspect Aspect) const {
664664
(this->getBackend() == backend::opencl) ||
665665
(this->getBackend() == backend::ext_oneapi_cuda);
666666
}
667-
case aspect::ext_oneapi_tangle_group: {
668-
// TODO: tangle_group is not currently supported for CUDA devices. Add when
667+
case aspect::ext_oneapi_tangle: {
668+
// TODO: tangle is not currently supported for CUDA devices. Add when
669669
// implemented.
670670
return (this->getBackend() == backend::ext_oneapi_level_zero) ||
671671
(this->getBackend() == backend::opencl);

‎sycl/test-e2e/NonUniformGroups/is_user_constructed.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@
44
#include <sycl/ext/oneapi/experimental/ballot_group.hpp>
55
#include <sycl/ext/oneapi/experimental/chunk.hpp>
66
#include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>
7-
#include <sycl/ext/oneapi/experimental/tangle_group.hpp>
7+
#include <sycl/ext/oneapi/experimental/tangle.hpp>
88
namespace syclex = sycl::ext::oneapi::experimental;
99

1010
static_assert(
@@ -14,5 +14,5 @@ static_assert(
1414
static_assert(
1515
syclex::is_user_constructed_group_v<syclex::chunk<2, sycl::sub_group>>);
1616
static_assert(
17-
syclex::is_user_constructed_group_v<syclex::tangle_group<sycl::sub_group>>);
17+
syclex::is_user_constructed_group_v<syclex::tangle<sycl::sub_group>>);
1818
static_assert(syclex::is_user_constructed_group_v<syclex::opportunistic_group>);

‎sycl/test-e2e/NonUniformGroups/tangle_group.cpp‎ renamed to ‎sycl/test-e2e/NonUniformGroups/tangle.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
// UNSUPPORTED: cuda || hip
99

1010
#include <sycl/detail/core.hpp>
11-
#include <sycl/ext/oneapi/experimental/tangle_group.hpp>
11+
#include <sycl/ext/oneapi/experimental/tangle.hpp>
1212
#include <vector>
1313
namespace syclex = sycl::ext::oneapi::experimental;
1414

@@ -45,7 +45,7 @@ int main() {
4545
// Branches deliberately duplicated to test impact of optimizations.
4646
// This only reliably works with optimizations disabled right now.
4747
if (item.get_global_id() % 2 == 0) {
48-
auto TangleGroup = syclex::get_tangle_group(SG);
48+
auto TangleGroup = syclex::get_tangle(SG);
4949

5050
bool Match = true;
5151
Match &= (TangleGroup.get_group_id() == 0);
@@ -56,7 +56,7 @@ int main() {
5656
MatchAcc[WI] = Match;
5757
LeaderAcc[WI] = TangleGroup.leader();
5858
} else {
59-
auto TangleGroup = syclex::get_tangle_group(SG);
59+
auto TangleGroup = syclex::get_tangle(SG);
6060

6161
bool Match = true;
6262
Match &= (TangleGroup.get_group_id() == 0);

0 commit comments

Comments
 (0)