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.

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

Add test plan for oneapi_non_uniform_groups extension #866

Merged
Changes from 4 commits
Commits
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
78 changes: 54 additions & 24 deletions test_plans/non_uniform_groups.asciidoc
Original file line number Diff line number Diff line change
Expand Up @@ -92,22 +92,22 @@ equal to `N1` if the predicate was `true` or equal to `N2` if the predicate was
==== get_group_linear_id

Check that `get_group_linear_id()` return type is `linear_id_type` and return
value is equal to `ballot.get_group_id()` converted to `linear_id_type`.
value is equal to `get_group_id()` converted to `linear_id_type`.

==== get_local_linear_id

Check that `get_local_linear_id()` return type is `linear_id_type` and the
return value is equal to `ballot.get_local_id()` converted to `linear_id_type`.
return value is equal to `get_local_id()` converted to `linear_id_type`.

==== get_group_linear_range

Check that `get_group_linear_range()` return type is `linear_id_type` and return
value is equal to `ballot.get_group_range()` converted to `linear_id_type`.
value is equal to `get_group_range()` converted to `linear_id_type`.

==== get_local_linear_range

Check that `get_local_linear_range()` return type is `linear_id_type` and return
value is equal to `ballot.get_local_range()` converted to `linear_id_type`.
value is equal to `get_local_range()` converted to `linear_id_type`.

==== leader

Expand All @@ -119,8 +119,8 @@ Check that `leader()` return type is `bool` and return value is equal to
Let `N` be some power-of-two value greater than 1 that is expected to be a
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Are you excluding the case of 1. That seems interesting to test corner cases, otherwise.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

For the class API tests we only test for one group configuration. The intention here was to pick one that wasn't 1 to avoid some trivial cases and to test a more normal group size. We could potentially run the tests for multiple configurations, if that would be more interesting? We do it for the group algorithms.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Oh I thought you were running with several N and in that case testing also the boundaries would have made sense.

divisor of the sub-group size of most devices. The `get_fixed_size_group` is
called with the `sub_group` of the invocation and `N` as the partition size.
Let `M` be the size of the sub-group the given `fixed_size_group` was created
from.
Let `M` be the result of `get_local_range()` on the sub-group the given
`fixed_size_group` was created from.

==== Group traits

Expand Down Expand Up @@ -166,22 +166,23 @@ equal to `N`.
==== get_group_linear_id

Check that `get_group_linear_id()` return type is `linear_id_type` and return
value is equal to `ballot.get_group_id()` converted to `linear_id_type`.
value is equal to `get_group_id()` converted to `linear_id_type`.

==== get_local_linear_id

Check that `get_local_linear_id()` return type is `linear_id_type` and the
return value is equal to `ballot.get_local_id()` converted to `linear_id_type`.
return value is equal to `get_local_id()` converted to
`linear_id_type`.

==== get_group_linear_range

Check that `get_group_linear_range()` return type is `linear_id_type` and return
value is equal to `ballot.get_group_range()` converted to `linear_id_type`.
value is equal to `get_group_range()` converted to `linear_id_type`.

==== get_local_linear_range

Check that `get_local_linear_range()` return type is `linear_id_type` and return
value is equal to `ballot.get_local_range()` converted to `linear_id_type`.
value is equal to `get_local_range()` converted to `linear_id_type`.

==== leader

Expand All @@ -190,9 +191,11 @@ Check that `leader()` return type is `bool` and return value is equal to

=== The `tangle_group` class API

The `get_tangle_group` is called with the `sub_group` of the invocation. This
will only be called by the first `N` items of the sub-group, where `N` is
strictly less than the size of the sub-group.
The `get_tangle_group` is called with the `sub_group` of the invocation.
Let `M` be the result of `get_local_range()` on this `sub_group` and let `N` be
some value strictly less than `M`. `get_tangle_group` is called it two split
steffenlarsen marked this conversation as resolved.
Show resolved Hide resolved
control-flows in an if-else-statement, the if-branch with the first `N` items of
the sub-group and the else branch with the rest.

==== Group traits

Expand Down Expand Up @@ -238,22 +241,22 @@ equal to `N`.
==== get_group_linear_id

Check that `get_group_linear_id()` return type is `linear_id_type` and return
value is equal to `ballot.get_group_id()` converted to `linear_id_type`.
value is equal to `get_group_id()` converted to `linear_id_type`.

==== get_local_linear_id

Check that `get_local_linear_id()` return type is `linear_id_type` and the
return value is equal to `ballot.get_local_id()` converted to `linear_id_type`.
return value is equal to `get_local_id()` converted to `linear_id_type`.

==== get_group_linear_range

Check that `get_group_linear_range()` return type is `linear_id_type` and return
value is equal to `ballot.get_group_range()` converted to `linear_id_type`.
value is equal to `get_group_range()` converted to `linear_id_type`.

==== get_local_linear_range

Check that `get_local_linear_range()` return type is `linear_id_type` and return
value is equal to `ballot.get_local_range()` converted to `linear_id_type`.
value is equal to `get_local_range()` converted to `linear_id_type`.

==== leader

Expand All @@ -263,7 +266,7 @@ Check that `leader()` return type is `bool` and return value is equal to
=== The `opportunistic_group` class API

The `get_opportunistic_group` is called by all work items.
steffenlarsen marked this conversation as resolved.
Show resolved Hide resolved
Let `M` be the size of the sub-group of the invocation.
Let `M` be the result of `get_local_range()` on the sub-group of the invocation.

==== Group traits

Expand Down Expand Up @@ -309,22 +312,22 @@ less than or equal to `M`.
==== get_group_linear_id

Check that `get_group_linear_id()` return type is `linear_id_type` and return
value is equal to `ballot.get_group_id()` converted to `linear_id_type`.
value is equal to `get_group_id()` converted to `linear_id_type`.

==== get_local_linear_id

Check that `get_local_linear_id()` return type is `linear_id_type` and the
return value is equal to `ballot.get_local_id()` converted to `linear_id_type`.
return value is equal to `get_local_id()` converted to `linear_id_type`.

==== get_group_linear_range

Check that `get_group_linear_range()` return type is `linear_id_type` and return
value is equal to `ballot.get_group_range()` converted to `linear_id_type`.
value is equal to `get_group_range()` converted to `linear_id_type`.

==== get_local_linear_range

Check that `get_local_linear_range()` return type is `linear_id_type` and return
value is equal to `ballot.get_local_range()` converted to `linear_id_type`.
value is equal to `get_local_range()` converted to `linear_id_type`.

==== leader

Expand Down Expand Up @@ -366,5 +369,32 @@ The group algorithms

for `ballot_group`, `fixed_size_group`, `tangle_group`
and `opportunistic_group` are tested similar to how they are currently tested
with `group` and `sub_group` in the core CTS. The groups are constructed in the
same way as for the API testing described above.
with `group` and `sub_group` in the core CTS.

The groups are constructed as follows:

* `get_ballot_group` is called with a predicate that is `true` for the first `N`
work-items in the sub-group.
* `get_ballot_group` is called with a predicate that is `true` for work-items
with odd `sg.get_local_linear_id()` values, where `sg` is the sub-group.
* `get_ballot_group` is called with a predicate that is `true` for all
work-items in the sub-group.
* `get_ballot_group` is called with a predicate that is `false` for all
work-items in the sub-group.
* `get_fixed_size_group` is called with a partition-size of 1.
* `get_fixed_size_group` is called with a partition-size of 2, if 2 is greater
than or equal to the smallest supported sub-group size on the device.
* `get_fixed_size_group` is called with a partition-size of 4, if 4 is greater
than or equal to the smallest supported sub-group size on the device.
* `get_fixed_size_group` is called with a partition-size of 8, if 8 is greater
than or equal to the smallest supported sub-group size on the device.
* `get_tangle_group` is called in a branched control-flow with the first `N`
work-items in the sub-group.
* `get_tangle_group` is called in a branched control-flow with work-items with
odd `sg.get_local_linear_id()` values, where `sg` is the sub-group.
* `get_tangle_group` is called by all items in the sub-group.
* `get_opportunistic_group` is called in a branched control-flow with the first
`N` work-items in the sub-group.
* `get_opportunistic_group` is called in a branched control-flow with work-items
with odd `sg.get_local_linear_id()` values, where `sg` is the sub-group.
* `get_opportunistic_group` is called by all items in the sub-group.
Loading