Adds test for the non_uniform_groups oneAPI extension #875
Conversation
This commit adds tests for the [sycl_ext_oneapi_non_uniform_groups](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc) extension in accordance with the test plan currently in review here: KhronosGroup#866
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
| size_t split = sub_group_size / 3; | ||
| bool is_left = sg.get_local_linear_id() < split; | ||
|
|
||
| auto ballot = oneapi_ext::get_ballot_group(sg, is_left); |
There was a problem hiding this comment.
Should this be checked explicitly?
Constraints: Available only if `sycl::is_group_v<std::decay_t<Group>> &&
std::is_same_v<Group, sycl::sub_group>` is true.
There was a problem hiding this comment.
I don't believe the CTS normally does negative testing for availability like this. The test will fail if it isn't correctly available.
|
|
||
| template <size_t PartitionSize> | ||
| struct NonUniformGroupHelper< | ||
| oneapi_ext::fixed_size_group<PartitionSize, sycl::sub_group>> { |
There was a problem hiding this comment.
Should this be checked explicitly?
`PartitionSize` must be positive and a power of 2.
There was a problem hiding this comment.
The specification doesn't specify what happens if it isn't, so we cannot check this.
| size_t sub_group_size = sg.get_local_range().size(); | ||
|
|
||
| auto fixed_size = | ||
| oneapi_ext::get_fixed_size_group<partition_size>(sg); |
There was a problem hiding this comment.
Should there be a negative test?
Preconditions: `PartitionSize` must be less than or equal to the result of `group.get_max_local_range()`. `group.get_local_linear_range()` must be evenly divisible by `PartitionSize`.
There was a problem hiding this comment.
Same as in #875 (comment).
| size_t sub_group_size = it.get_sub_group().get_local_range().size(); | ||
|
|
||
| auto opportunistic = | ||
| oneapi_ext::this_kernel::get_opportunistic_group(); |
There was a problem hiding this comment.
Should this be checked?
Each call to `get_opportunistic_group()` returns a different group.
There was a problem hiding this comment.
That is a good question. @Pennycook - What is the intention here? Do the groups have common by-value semantics and if so, does this make the guarantee that get_opportunistic_group() != get_opportunistic_group(), which in turn means the groups will never have the same members values?
There was a problem hiding this comment.
That is a very good question.
The intention of the wording was really to say that the members may be different, and that in cases where the members are different, the group objects may be different. Without thinking of any specific implementation: if the group members are different, the representation (e.g., as a mask) will probably be different; even if the group members are the same, I can imagine situations where the different group objects would still have dedicated scratchpad memory or barriers.
I think the simplest thing to do for now is to guarantee that equality returns false. It's probably more useful to have a comparison that is always useless than to try and define the cases in which equality should/could return true. I'd prefer not to provide a comparison operator at all, but I've heard suggestions recently that this would break a lot of things.
There was a problem hiding this comment.
I think the simplest thing to do for now is to guarantee that equality returns false. It's probably more useful to have a comparison that is always useless than to try and define the cases in which equality should/could return true. I'd prefer not to provide a comparison operator at all, but I've heard suggestions recently that this would break a lot of things.
Are you saying that opportunistic_group needs to have a comparison operator? If you are saying that things would break if there is no comparison operator, it seems surprising that those things wouldn't also break if the comparison operator always returned false.
Are you saying that the comparison operator would return false even when comparing an object with a copy? For example:
opportunistic_group og1 = /*...*/;
opportunistic_group og2 = og1;
bool same = (og1 == og2); // returns false?
There was a problem hiding this comment.
Are you saying that
opportunistic_groupneeds to have a comparison operator? If you are saying that things would break if there is no comparison operator, it seems surprising that those things wouldn't also break if the comparison operator always returned false.
I don't have complete information here. This came up in a discussion we were having about some of the other types (e.g., nd_item), and my takeaway was that that if we were to remove certain operators then the classes wouldn't play nicely with certain type traits. I might have gotten the wrong end of the stick.
Are you saying that the comparison operator would return false even when comparing an object with a copy? For example:
opportunistic_group og1 = /*...*/; opportunistic_group og2 = og1; bool same = (og1 == og2); // returns false?
Yes, that's what I was suggesting.
I'm also open to defining a real comparison operator for groups, but I don't know what the criteria should be. I can see an argument that it should be based on which work-items are in the group (which would allow comparisons across groups of different types) but can also think of cases where that could lead to surprising results (e.g., just because two groups have the same work-items in them doesn't necessarily mean they can be used interchangably).
tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp
Outdated
Show resolved
Hide resolved
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in
Outdated
Show resolved
Hide resolved
tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp
Outdated
Show resolved
Hide resolved
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
|
Interesting discussion about defining group equality and usability. Related to also whether we can use them in another work-group or kernel by storing them into an intermediate buffer. |
This commit adds tests for the sycl_ext_oneapi_non_uniform_groups extension in accordance with the test plan currently in review here: #866