[SYCL] Align Non-Uniform Groups with PR#14604:

[AndreiZibrov]

This PR implements the non-uniform groups API redesign proposed in PR 14604

1. Removed redundant _group suffix from class names
2. Used more descriptive function names that clearly indicate the partitioning behaviour
3. United similar functionality to reduce cognitive overhead

• fixed_size_group → chunk
• tangle_group → tangle
• ballot_group → fragment
• opportunistic_group → merged into fragment
• get_ballot_group() → binary_partition()
• get_fixed_size_group() → chunked_partition()
• get_tangle_group() → entangle()

• is_user_constructed_updated and HasExtensionWordBoundary tests

[AndreiZibrov]

each introduced entity has own aspect combination in opencl, so checking needed

UPD: left for separate PR

[Pennycook]

I'm surprised this logic still works now that we can create a fragment using either binary_partition or get_opportunistic_group.

Does this work because we still know that our implementation has at most two groups? Do you agree that if we added some way to create more than two fragments from a function, we'd need to rewrite this as a loop?

[AlexeySachkov]

Does this work because we still know that our implementation has at most two groups? Do you agree that if we added some way to create more than two fragments from a function, we'd need to rewrite this as a loop?

fragment can be further subdivided by binary_partition, but each fragment instance will only "know" about its immediate parent.

BTW, why do we need the else branch here in the first place?

[Pennycook]

fragment can be further subdivided by binary_partition, but each fragment instance will only "know" about its immediate parent.

Right, so at each call to binary_partition we're getting back a fragment with at most 2 groups.

I'm just thinking aloud. If we added something like a labeled_partition that returned more than 2 conceptual fragments, I don't think the current code would work. We might want to add a TODO or a note to that effect, is what I'm thinking. A ballot_group could only represent the result of a binary split, which is why it was different before.

BTW, why do we need the else branch here in the first place?

This was the only way I could find to trick Clang into generating code with the correct semantics. I think if we want to replace this with something that's simpler and easier to understand, then we need a new SPIR-V extension.

Here's a very simple example of the sort of code that we have to worry about:

int x = ...;
auto fragment = syclx::binary_partition(sg, sg.get_local_id() % 2);
auto sum = sycl::reduce_over_group(fragment, x, sycl::plus<>());

Each work-item gets its own fragment object, and the information stored inside that object describes which other work-items are part of the same conceptual fragment. But we haven't actually changed anything about the execution model: the hardware will continue to execute all the work-items in the sub-group "together" (typically in lock-step).

If we implement reduce_over_group with no branch, our only option today would be to write something like this:

float reduce_over_group(fragment, float x, BinaryOperation BOp) {
  return __spirv_OpGroupNonUniformFAdd(x, BOp);
}

...which wouldn't work. All the work-items in the same sub-group will execute the same "dynamic instance" of the GroupNonUniformFAdd instruction, and so we'd end up reducing across the whole sub-group, instead of the two individual partitions.

The only way I could find to reduce over separate partitions using existing SPIR-V instructions like GroupNonUniformFAdd was to convert the partitions into control flow. Each partition encounters a unique dynamic instance of the SPIR-V instruction, and the reduction works as expected.

If we had SPIR-V instructions that expected a mask, we could shift that complexity to the device compiler, and the headers would just have to say:

float reduce_over_group(fragment f, float x, BinaryOperation BOp) {
  return __spirv_OpGroupMaskedFAdd(x, BOp, f.Mask);
}

[AndreiZibrov]

added TODOs, please confirm it looks as assumed

[AlexeySachkov]

Thanks, @Pennycook, I understand why do we need the condition, but I'm still confused about the else branch.

int x = ...;
auto fragment = syclx::binary_partition(sg, sg.get_local_id() % 2);
auto val = sycl::all_of(fragment, x > 42);

What we have right now is:

if (g.get_group_id() == 1) {                                               
  return __spirv_GroupNonUniformAll(group_scope<ParentGroup>::value, pred);
} else {                                                                   
  return __spirv_GroupNonUniformAll(group_scope<ParentGroup>::value, pred);
}                                                                          

Shouldn't it be something like?

if (current work-item is active within the current fragment) {                                               
  return __spirv_GroupNonUniformAll(group_scope<ParentGroup>::value, pred);
} 

I.e. we should check mask embedded into the fragment and work-items which aren't a part of it just do nothing.

[Pennycook]

It's really subtle, because SPMD is weird. It took me a long time to wrap my head around this (or, at least, to convince myself that I understood what was happening -- I might still be wrong 😆).

Each fragment object here represents two different groups with different masks. But if we wrote it as:

if (current work-item is active within the current fragment) {
  return __spirv_GroupNonUniformAll(group_scope<ParentGroup>::value, pred);
}

...then there's only one branch, based on one condition. It might help to think of it like this:

bool condition = /* current work-item is active within the current fragment */;
if (condition) {
  return __spirv_GroupNonUniformAll(group_scope<ParentGroup>::value, pred);
}

The compiler isn't smart enough to see that different work-items within the sub-group are computing condition by comparing against different mask values. All it sees is a branch on condition, and condition always evaluates to true (because each work-item will always be in the mask that represents which part of the sub-group it is in).

So, a branch like this effectively doesn't do anything, and we end up reducing over the whole sub-group.

Does that make sense?

[AlexeySachkov]

Does this work because we still know that our implementation has at most two groups? Do you agree that if we added some way to create more than two fragments from a function, we'd need to rewrite this as a loop?

fragment can be further subdivided by binary_partition, but each fragment instance will only "know" about its immediate parent.

BTW, why do we need the else branch here in the first place?

[AlexeySachkov]

We use NonUniformBitwiseAnd and it seems like we need cl_khr_subgroup_non_uniform_arithmetic for it

[AndreiZibrov]

ok, so leaving for next PR?

[AndreiZibrov]

as requested

[AndreiZibrov]

as requested

[AndreiZibrov]

as requested

[AndreiZibrov]

as requested

[steffenlarsen]

Nit;

Suggested change

	} // namespace sycl
	} // namespace sycl

[AndreiZibrov]

I'm a right that clang-format shouldn't pass this?

[steffenlarsen]

We don't have any customizations to our clang-format to not pass these cases, but we normally require it anyway. I wouldn't be against a customization to our clang-format to enforce it though.

[steffenlarsen]

#14604 relaxed this to allow ParentGroup to also allow chunk and fragment. Will this be done in a follow-up?

[AndreiZibrov]

tbd in next commit

[steffenlarsen]

This should also allow chunk as a ParentGroup.

[AndreiZibrov]

tbd in next commit

[steffenlarsen]

I don't see why the fragment class is available here and inside the definition of chunk? Could this definition be moved into the body of fragment above?

Also, seems like we need to include fragment.hpp instead of doing a forward declaration. I suspect this works because we coincidentally end up including it at top-level or because instantiations are evaluated after the classes are made available.

[AndreiZibrov]

tbd in next feature commit. currently dealing with CI compilation erros

[AndreiZibrov]

It looks like there is something mismatching with namespaces building via includes between Windows part(passed Build + LIT) and Linux (Build + Lit):
https://github.com/intel/llvm/actions/runs/16032949268/job/45237784966?pr=19238#:~:text=sycl/ext/oneapi/experimental/cuda/non_uniform_algorithms.hpp%3A103%3A46%3A%20error%3A%20use,%23%20%7C%20%20%20103%20%7C%20inline%20__SYCL_ALWAYS_INLINE%20std%3A%3Aenable_if_t%3Csyclex%3A%3Ais_chunk_v%3CGroup%3E%2C%20T%3E it says that the ext::oneapi has no experimental in it. I've tried relative/system path type for includes and different combinations in latest commits... so leaving for later

[steffenlarsen]

Suggested change

	using syclex = ::sycl::ext::oneapi::experimental;
	namespace syclex = sycl::ext::oneapi::experimental;

I suspect that will fix the build failures.

[steffenlarsen]

On a side note, I am not too keen on defining these sorts of aliases in public headers. Granted, it is in the detail namespace, I fear it will cause us problems down the line.

Also, it's a change that isn't really related to the other changes, so it takes away from the main goal of this patch.

[AndreiZibrov]

only linux CI(only) jobs failed are saying that 'experimental' is not part of 'oneapi'
I was just trying to change the path of aimed namespace just at least to reach it on CI

I know it's not related - I'd replace it by working solution. Currently still I haven't done/reproduced such error locally

[steffenlarsen]

only linux CI(only) jobs failed are saying that 'experimental' is not part of 'oneapi'

More specifically, it says "error: no type named 'experimental' in namespace 'sycl::ext::oneapi'" which is true, there's no experimental type in that namespace. 😉

Even if some compilers allow this, we need to use the variants that all our supported compilers allow. I also don't see this kind of using being allowed by the using-directives, but I may also just be reading the reference wrong.