Cold-Air-Networks/mirror-opencl-cts
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

subgroups: Fix maximum subgroup size for tail calculation (#2752)

Browse source 
Use get_max_sub_group_size() when calculating the size of the final
subgroup in a non-uniform workgroup.

get_sub_group_size() may already return the size of the current partial
subgroup. Using it as the modulus divisor can therefore calculate an
incorrect size, split the subgroup at the wrong lane, and pass
non-uniform indices to sub_group_non_uniform_broadcast.

Signed-off-by: Riccardo Zanetti <riccardo.zanetti@arm.com>

As approved on 28Jul26 call
This commit is contained in:
riccardo-zanetti-arm 2026-07-28 17:32:33 +01:00 committed by GitHub
parent 6874c1177c
commit 5013399741
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
1 changed files with 2 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

4
test_conformance/subgroups/test_subgroup_ballot.cpp
View file

@ -714,7 +714,7 @@ __kernel void test_sub_group_non_uniform_broadcast(const __global Type *in, __gl
uint sub_group_size = get_sub_group_size();
// If we are at the edge, calculate our own sub_group_size as it's implementation defined otherwise.
if (get_local_size(0) != get_enqueued_local_size(0) && get_sub_group_id() == get_num_sub_groups() - 1) {
uint new_sub_group_size = get_local_size(0) % sub_group_size;
uint new_sub_group_size = get_local_size(0) % get_max_sub_group_size();
if (new_sub_group_size != 0)
sub_group_size = new_sub_group_size;
}
@ -735,7 +735,7 @@ __kernel void test_sub_group_broadcast_first(const __global Type *in, __global i
uint sub_group_size = get_sub_group_size();
// If we are at the edge, calculate our own sub_group_size as it's implementation defined otherwise.
if (get_local_size(0) != get_enqueued_local_size(0) && get_sub_group_id() == get_num_sub_groups() - 1) {
uint new_sub_group_size = get_local_size(0) % sub_group_size;
uint new_sub_group_size = get_local_size(0) % get_max_sub_group_size();
if (new_sub_group_size != 0)
sub_group_size = new_sub_group_size;
}