From 3cac56c33ff1a2da2f86a0b8f33db2b3313a34a6 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 6 Feb 2024 06:29:49 -0800 Subject: [PATCH 1/7] Add test plan for oneapi_non_uniform_groups extension This commit adds a test plan for the [oneapi_non_uniform_groups](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc) oneAPI extension. Signed-off-by: Larsen, Steffen --- test_plans/non_uniform_groups.asciidoc | 370 +++++++++++++++++++++++++ 1 file changed, 370 insertions(+) create mode 100644 test_plans/non_uniform_groups.asciidoc diff --git a/test_plans/non_uniform_groups.asciidoc b/test_plans/non_uniform_groups.asciidoc new file mode 100644 index 000000000..c24101cba --- /dev/null +++ b/test_plans/non_uniform_groups.asciidoc @@ -0,0 +1,370 @@ +:sectnums: +:xrefstyle: short + += Test plan for sycl_ext_oneapi_non_uniform_groups + +This is a test plan for the APIs described in +https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_groups.asciidoc[sycl_ext_oneapi_non_uniform_groups]. + + +== Testing scope + +=== Device coverage + +All of the tests described below are performed only on the default device that +is selected on the CTS command line. + +=== Feature test macro + +All of the tests should use `#ifdef SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS` so they +can be skipped if feature is not supported. + +== Tests + +=== The `is_fixed_topology_group` trait for existing types + +Check the following: + +* `is_fixed_topology_group>::value` is `true`. +* `is_fixed_topology_group_v>` is `true`. +* `is_fixed_topology_group>::value` is `true`. +* `is_fixed_topology_group_v>` is `true`. +* `is_fixed_topology_group>::value` is `true`. +* `is_fixed_topology_group_v>` is `true`. +* `is_fixed_topology_group::value` is `true`. +* `is_fixed_topology_group_v` is `true`. + +If `SYCL_EXT_ONEAPI_ROOT_GROUP` is defined, check the following: + +* `is_fixed_topology_group::value` is `true`. +* `is_fixed_topology_group_v` is `true`. + +=== The `ballot_group` class API + +The `get_ballot_group` is called with the `sub_group` of the invocation and a +predicate splitting the work-items of the sub-group into uneven groups. Let `N1` +be the size of the group created with `true` predicate and let `N2` be the size +of the group created with `false` predicate. + +==== Group traits + +Check the following: + +* `is_group>::value` is `true`. +* `is_group_v>` is `true`. +* `is_user_constructed_group>::value` is `true`. +* `is_user_constructed_group_v>` is `true`. +* `is_fixed_topology_group>::value` is `false`. +* `is_fixed_topology_group_v>` is `false`. + +==== Members + +Check the following: + +* `id_type` is same as `id<1>`. +* `range_type` is same as `range<1>`. +* `linear_id_type` is same as `uint32_t`. +* `dimensions` is 1. +* `fence_scope` is equal to `sub_group::fence_scope`. + +==== get_group_id + +Check that `get_group_id()` return type is `id_type` and return value is +`1` if the predicate was `true` or `0` if the predicate was `false`. + +==== get_local_id + +Check that `get_local_id()` return type is `id_type` and return value is less +than `N1` if the predicate was `true` or less than `N2` if the predicate was +`false`. + +==== get_group_range + +Check that `get_group_range()` return type is `range_type` and return value is +equal to `2`. + +==== get_local_range + +Check that `get_local_range()` return type is `range_type` and return value is +equal to `N1` if the predicate was `true` or equal to `N2` if the predicate was +`false`. + +==== 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`. + +==== 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`. + +==== 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`. + +==== 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`. + +==== leader + +Check that `leader()` return type is `bool` and return value is equal to +`get_local_id() == 0`. + +=== The `fixed_size_group` class API + +Let `N` be some power-of-two value greater than 1 that is expected to be a +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. + +==== Group traits + +Check the following: + +* `is_group>::value` is `true`. +* `is_group_v>` is `true`. +* `is_user_constructed_group>::value` is `true`. +* `is_user_constructed_group>` is `true`. +* `is_fixed_topology_group_v>::value` is `false`. +* `is_fixed_topology_group_v>` is `false`. + +==== Members + +Check the following: + +* `id_type` is same as `id<1>`. +* `range_type` is same as `range<1>`. +* `linear_id_type` is same as `uint32_t`. +* `dimensions` is 1. +* `fence_scope` is equal to `sub_group::fence_scope`. + +==== get_group_id + +Check that `get_group_id()` return type is `id<1>` and return value is less than +`M/N`. + +==== get_local_id + +Check that `get_local_id()` return type is `id<1>` and return value is less than +`N`. + +==== get_group_range + +Check that `get_group_range()` return type is `range_type` and return value is +equal to `M/N`. + +==== get_local_range + +Check that `get_local_range()` return type is `range_type` and return value is +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`. + +==== 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`. + +==== 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`. + +==== 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`. + +==== leader + +Check that `leader()` return type is `bool` and return value is equal to +`get_local_id() == 0`. + +=== 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. + +==== Group traits + +Check the following: + +* `is_group>::value` is `true`. +* `is_group_v>` is `true`. +* `is_user_constructed_group>::value` is `true`. +* `is_user_constructed_group>` is `true`. +* `is_fixed_topology_group_v>::value` is `false`. +* `is_fixed_topology_group_v>` is `false`. + +==== Members + +Check the following: + +* `id_type` is same as `id<1>`. +* `range_type` is same as `range<1>`. +* `linear_id_type` is same as `uint32_t`. +* `dimensions` is 1. +* `fence_scope` is equal to `sub_group::fence_scope`. + +==== get_group_id + +Check that `get_group_id()` return type is `id_type` and return value is equal +to `0`. + +==== get_local_id + +Check that `get_local_id()` return type is `id_type` and return value is less +than `N`. + +==== get_group_range + +Check that `get_group_range()` return type is `range_type` and return value is +equal to `1`. + +==== get_local_range + +Check that `get_local_range()` return type is `range_type` and return value is +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`. + +==== 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`. + +==== 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`. + +==== 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`. + +==== leader + +Check that `leader()` return type is `bool` and return value is equal to +`get_local_id() == 0`. + +=== The `opportunistic_group` class API + +The `get_opportunistic_group` is called by all work items. +Let `M` be the size of the sub-group of the invocation. + +==== Group traits + +Check the following: + +* `is_group::value` is `true`. +* `is_group_v` is `true`. +* `is_user_constructed_group::value` is `true`. +* `is_user_constructed_group` is `true`. +* `is_fixed_topology_group_v::value` is `false`. +* `is_fixed_topology_group_v` is `false`. + +==== Members + +Check the following: + +* `id_type` is same as `id<1>`. +* `range_type` is same as `range<1>`. +* `linear_id_type` is same as `uint32_t`. +* `dimensions` is 1. +* `fence_scope` is equal to `sub_group::fence_scope`. + +==== get_group_id + +Check that `get_group_id()` return type is `id_type` and return value is equal +to `0`. + +==== get_local_id + +Check that `get_local_id()` return type is `id_type` and return value is less +than `get_local_range().size()`. + +==== get_group_range + +Check that `get_group_range()` return type is `range_type` and return value is +equal to `1`. + +==== get_local_range + +Check that `get_local_range()` return type is `range_type` and return value is +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`. + +==== 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`. + +==== 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`. + +==== 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`. + +==== leader + +Check that `leader()` return type is `bool` and return value is equal to +`get_local_id() == 0`. + +=== Group functions + +The group functions + +* `group_barrier` +* `group_broadcast` + +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. + +=== Group algorithms + +The group algorithms + +* `joint_any_of` +* `any_of_group` +* `joint_all_of` +* `all_of_group` +* `joint_none_of` +* `none_of_group` +* `shift_group_left` +* `shift_group_right` +* `permute_group_by_xor` +* `select_from_group` +* `joint_reduce` +* `reduce_over_group` +* `joint_exclusive_scan` +* `exclusive_scan_over_group` +* `joint_inclusive_scan` +* `inclusive_scan_over_group` + +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. From a29b71d752789829431299d25d06f1059255c79e Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 8 Feb 2024 00:13:36 -0800 Subject: [PATCH 2/7] Fix link Signed-off-by: Larsen, Steffen --- test_plans/non_uniform_groups.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test_plans/non_uniform_groups.asciidoc b/test_plans/non_uniform_groups.asciidoc index c24101cba..1565ab98e 100644 --- a/test_plans/non_uniform_groups.asciidoc +++ b/test_plans/non_uniform_groups.asciidoc @@ -4,7 +4,7 @@ = Test plan for sycl_ext_oneapi_non_uniform_groups This is a test plan for the APIs described in -https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_groups.asciidoc[sycl_ext_oneapi_non_uniform_groups]. +https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc[sycl_ext_oneapi_non_uniform_groups]. == Testing scope From 48afb567c5c71db9774bb98ab30886833ebabc96 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 20 Feb 2024 01:42:51 -0800 Subject: [PATCH 3/7] Remove ballot references Signed-off-by: Larsen, Steffen --- test_plans/non_uniform_groups.asciidoc | 33 +++++++++++++------------- 1 file changed, 17 insertions(+), 16 deletions(-) diff --git a/test_plans/non_uniform_groups.asciidoc b/test_plans/non_uniform_groups.asciidoc index 1565ab98e..8dec70351 100644 --- a/test_plans/non_uniform_groups.asciidoc +++ b/test_plans/non_uniform_groups.asciidoc @@ -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 @@ -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 @@ -238,22 +239,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 @@ -309,22 +310,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 From ac9a6f48776ddd2f7dbc2172b5266cdbae775374 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 20 Feb 2024 01:48:23 -0800 Subject: [PATCH 4/7] Specify 'size of sub-group' Signed-off-by: Larsen, Steffen --- test_plans/non_uniform_groups.asciidoc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/test_plans/non_uniform_groups.asciidoc b/test_plans/non_uniform_groups.asciidoc index 8dec70351..3251b4cb3 100644 --- a/test_plans/non_uniform_groups.asciidoc +++ b/test_plans/non_uniform_groups.asciidoc @@ -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 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 @@ -193,7 +193,7 @@ Check that `leader()` return type is `bool` and return value is equal to 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. +strictly less than the result of `get_local_range()` on the sub-group. ==== Group traits @@ -264,7 +264,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. -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 From b514678e20387bd0e72ea3652e601f8fd89fbacc Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 21 Feb 2024 05:21:31 -0800 Subject: [PATCH 5/7] Add specific cases for the test-plan Signed-off-by: Larsen, Steffen --- test_plans/non_uniform_groups.asciidoc | 31 ++++++++++++++++++++++++-- 1 file changed, 29 insertions(+), 2 deletions(-) diff --git a/test_plans/non_uniform_groups.asciidoc b/test_plans/non_uniform_groups.asciidoc index 3251b4cb3..8edc24f9b 100644 --- a/test_plans/non_uniform_groups.asciidoc +++ b/test_plans/non_uniform_groups.asciidoc @@ -367,5 +367,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. From 2fb29efe2a6f71756f5d8d37a7e9096a8c155072 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 21 Feb 2024 10:24:24 -0800 Subject: [PATCH 6/7] Specify API checks for tangle_group runs in if-else Signed-off-by: Larsen, Steffen --- test_plans/non_uniform_groups.asciidoc | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/test_plans/non_uniform_groups.asciidoc b/test_plans/non_uniform_groups.asciidoc index 8edc24f9b..1d649a5d3 100644 --- a/test_plans/non_uniform_groups.asciidoc +++ b/test_plans/non_uniform_groups.asciidoc @@ -191,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 result of `get_local_range()` on 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 +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 From 284624d58ce1f58ff42162bc8d0e5c7a577c58a5 Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Thu, 22 Feb 2024 09:18:28 +0100 Subject: [PATCH 7/7] Update test_plans/non_uniform_groups.asciidoc Co-authored-by: John Pennycook --- test_plans/non_uniform_groups.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test_plans/non_uniform_groups.asciidoc b/test_plans/non_uniform_groups.asciidoc index 1d649a5d3..c714b0dc8 100644 --- a/test_plans/non_uniform_groups.asciidoc +++ b/test_plans/non_uniform_groups.asciidoc @@ -193,7 +193,7 @@ Check that `leader()` return type is `bool` and return value is equal to 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 +some value strictly less than `M`. `get_tangle_group` is called in two split 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.