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

[subgroups][non_uniform_broadcast] Fix broadcasting index generation #1680

Merged
merged 6 commits into from
Mar 12, 2024
Merged

[subgroups][non_uniform_broadcast] Fix broadcasting index generation #1680

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
9bbab53
[subgroups][non_uniform_broadcast] Fix broadcasting index generation
Nuullll Mar 23, 2023
082f5dd
Merge branch 'master' into fix-sub_group_non_uniform_broadcast
Nuullll Sep 27, 2023
2dcf4af
Revert "[subgroups][non_uniform_broadcast] Fix broadcasting index gen…
Nuullll Sep 27, 2023
cf67385
[subgroups][non_uniform_broadcast] Fix broadcasting index generation
Nuullll Sep 27, 2023
050239c
Merge branch 'main' into fix-sub_group_non_uniform_broadcast
Nuullll Feb 28, 2024
a4ff687
Apply suggestion
Nuullll Feb 28, 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
32 changes: 12 additions & 20 deletions test_conformance/subgroups/subgroup_common_templates.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@
// subgroup takes only one value from only one chosen (the smallest subgroup ID)
// work_item
// sub_group_non_uniform_broadcast - same as type 0 but
// only 4 work_items from subgroup enter the code (are active)
// only half of work_items from subgroup enter the code (are active)
template <typename Ty, SubgroupsBroadcastOp operation> struct BC
{
static void log_test(const WorkGroupParams &test_params,
Expand Down Expand Up @@ -78,24 +78,16 @@ template <typename Ty, SubgroupsBroadcastOp operation> struct BC
int bcast_elseif = 0;
int bcast_index = (int)(genrand_int32(gMTdata) & 0x7fffffff)
% (d > n ? n : d);
int num_of_active_items = n >> 1;
Nuullll marked this conversation as resolved.
Show resolved Hide resolved
Nuullll marked this conversation as resolved.
Show resolved Hide resolved
// l - calculate subgroup local id from which value will be
// broadcasted (one the same value for whole subgroup)
if (operation != SubgroupsBroadcastOp::broadcast)
{
// reduce brodcasting index in case of non_uniform and
// last workgroup last subgroup
if (last_subgroup_size && j == nj - 1
&& last_subgroup_size < NR_OF_ACTIVE_WORK_ITEMS)
{
bcast_if = bcast_index % last_subgroup_size;
bcast_elseif = bcast_if;
}
else
{
bcast_if = bcast_index % NR_OF_ACTIVE_WORK_ITEMS;
bcast_elseif = NR_OF_ACTIVE_WORK_ITEMS
+ bcast_index % (n - NR_OF_ACTIVE_WORK_ITEMS);
}
if (num_of_active_items != 0)
bcast_if = bcast_index % num_of_active_items;
if (num_of_active_items != n)
bcast_elseif = num_of_active_items
+ bcast_index % (n - num_of_active_items);
}

for (i = 0; i < n; ++i)
Expand All @@ -107,7 +99,7 @@ template <typename Ty, SubgroupsBroadcastOp operation> struct BC
}
else
{
if (i < NR_OF_ACTIVE_WORK_ITEMS)
if (i < num_of_active_items)
{
// index of the third
// element int the vector.
Expand Down Expand Up @@ -182,15 +174,15 @@ template <typename Ty, SubgroupsBroadcastOp operation> struct BC
}

// Check result
int num_of_active_items = n >> 1;
if (operation == SubgroupsBroadcastOp::broadcast_first)
{
int lowest_active_id = -1;
for (i = 0; i < n; ++i)
{

lowest_active_id = i < NR_OF_ACTIVE_WORK_ITEMS
? 0
: NR_OF_ACTIVE_WORK_ITEMS;
lowest_active_id =
i < num_of_active_items ? 0 : num_of_active_items;
// findout if broadcasted
// value is the same
tr = mx[ii + lowest_active_id];
Expand Down Expand Up @@ -221,7 +213,7 @@ template <typename Ty, SubgroupsBroadcastOp operation> struct BC
}
else
{
if (i < NR_OF_ACTIVE_WORK_ITEMS)
if (i < num_of_active_items)
{ // take index of array where info
// which work_item will be
// broadcast its value is stored
Expand Down
4 changes: 0 additions & 4 deletions test_conformance/subgroups/subhelpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,8 +28,6 @@
#include <regex>
#include <map>

#define NR_OF_ACTIVE_WORK_ITEMS 4

extern MTdata gMTdata;
typedef std::bitset<128> bs128;
extern cl_half_rounding_mode g_rounding_mode;
Expand Down Expand Up @@ -1474,8 +1472,6 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test

Fns::log_test(test_params, "");

kernel_sstr << "#define NR_OF_ACTIVE_WORK_ITEMS ";
kernel_sstr << NR_OF_ACTIVE_WORK_ITEMS << "\n";
// Make sure a test of type Ty is supported by the device
if (!TypeManager<Ty>::type_supported(device))
{
Expand Down
4 changes: 2 additions & 2 deletions test_conformance/subgroups/test_subgroup_ballot.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -767,7 +767,7 @@ __kernel void test_sub_group_non_uniform_broadcast(const __global Type *in, __gl
int gid = get_global_id(0);
XY(xy,gid);
Type x = in[gid];
if (xy[gid].x < NR_OF_ACTIVE_WORK_ITEMS) {
if (xy[gid].x < (get_sub_group_size() >> 1)) {
out[gid] = sub_group_non_uniform_broadcast(x, xy[gid].z);
} else {
out[gid] = sub_group_non_uniform_broadcast(x, xy[gid].w);
Expand All @@ -779,7 +779,7 @@ __kernel void test_sub_group_broadcast_first(const __global Type *in, __global i
int gid = get_global_id(0);
XY(xy,gid);
Type x = in[gid];
if (xy[gid].x < NR_OF_ACTIVE_WORK_ITEMS) {
if (xy[gid].x < (get_sub_group_size() >> 1)) {
out[gid] = sub_group_broadcast_first(x);;
} else {
out[gid] = sub_group_broadcast_first(x);;
Expand Down
Loading