Skip to content

Latest commit

 

History

History
278 lines (185 loc) · 9.02 KB

group_functions.asciidoc

File metadata and controls

278 lines (185 loc) · 9.02 KB

Test plan for group functions and algorithms

This is a test plan for the APIs described in SYCL 2020 sections 4.17.3 "Group functions" and 4.17.4 "Group algorithms library".

Implementation time is estimated to be 3 person-days.

1. Testing scope

1.1. Backend coverage

All the tests described below are performed for any SYCL implementation.

1.2. Device coverage

All of the tests described below are performed only on the default device that is selected on the CTS command line.

2. Tests

Every function should be checked for accordance to spec:

  • Function existence

  • Availability for some specified types

  • Type of result

  • Result correctness

2.1. Types for functions

Fundamental type list for tests:

  • char

  • int

  • float

  • long long

For FULL conformance add:

  • signed char

  • unsigned char

  • short

  • unsigned short

  • unsigned int

  • long

  • unsigned long

  • unsigned long long

  • bool

In addition, if the device has aspect::fp64:

  • double

In addition, if the device has aspect::fp16:

  • half

If function accepts another types, the following list of types is added:

  • vec<int, 2>

  • vec<unsigned long long, 16>

  • marray<float, 5>

  • marray<short, 7>

  • struct { int a; bool b; }

Pointers are taken from the same list of types.

bool Predicate(T) functions are functions checking the ranges of T itself or the first element of vec, marray, struct type T. They are choosen in such a way that functions using them should return both true and false in different tests.

binary_op used in scans and reductions are addition and maximum of T itself or the first element of vec, marray, struct type T.

2.2. Data for the functions

Data for the functions are sequences of 1, 2, 3, etc. corresponding to increasing linear_id.

2.3. Execution

Functions will be executed by a group using sycl::parallel_for with nd_range<Dims> corresponding to one group of maximal group size available on device with Dims = 1,2,3 (by querying info::device::max_work_item_sizes<Dims> and info::device::max_work_group_size). Sub-group variants will execute functions on all the subgroups or, if this is not possible, on the first subgroup of this large group only.

2.4. Traits

2.4.1. is_group

For both group and sub_group template parameter T check that is_group<T> is std::true_type. For some other template parameter T check that is_group<T> is std::false_type.

2.4.2. is_group_v

For both group and sub_group template parameter T check that is_group_v<T> is true. For some other template parameter T check that is_group_v<T> is false.

2.5. Functions

2.5.1. group_broadcast

For each of the following use-cases:

  • using T group_broadcast(Group g, T x);

  • using T group_broadcast(Group g, T x, Group::linear_id_type local_linear_id);

  • using T group_broadcast(Group g, T x, Group::id_type local_id);

for some scalar, vector, marray, and user defined trivially copyable types T and both group and sub_group check returned type and value.

2.5.2. group_barrier

For each memory_scope wider than Group::fence_scope in the list

  • sycl::memory_scope::sub_group

  • sycl::memory_scope::work_group

  • sycl::memory_scope::device

  • sycl::memory_scope::system

If supported by the hardware (from query of info::device::atomic_fence_scope_capabilities) check that invocation of void group_barrier(Group g, memory_scope fence_scope) for both group and sub_group behaves as expected: when each of workitems writes 1 into zero-initialized local memory array (for sub_group and work_group memory scope) or zero-initialized global memory array (for device and system memory scope) after barrier all array values red by workitems in reverse order are equal to 1.

2.5.3. joint_any_of

With several different pointer types Ptr and bool Predicate(*Ptr) check returned type and value of bool joint_any_of(Group g, Ptr first, Ptr last, Predicate pred); for both group and sub_group.

2.5.4. joint_all_of

With several different pointer types Ptr and bool Predicate(*Ptr) check returned type and value of bool joint_all_of(Group g, Ptr first, Ptr last, Predicate pred); for both group and sub_group.

2.5.5. joint_none_of

With several different pointer types Ptr and bool Predicate(*Ptr) check returned type and value of bool joint_none_of(Group g, Ptr first, Ptr last, Predicate pred); for both group and sub_group.

2.5.6. any_of_group

For each of the following use-cases:

  • using bool any_of_group(Group g, T x, Predicate pred); with several different types T and bool Predicate(T)

  • using bool any_of_group(Group g, bool pred);

check returned type and value for both group and sub_group.

2.5.7. all_of_group

For each of the following use-cases:

  • using bool all_of_group(Group g, T x, Predicate pred); with several different types T and bool Predicate(T)

  • using bool all_of_group(Group g, bool pred);

check returned type and value for both group and sub_group.

2.5.8. none_of_group

For each of the following use-cases:

  • using bool none_of_group(Group g, T x, Predicate pred); with several different types T and bool Predicate(T)

  • using bool none_of_group(Group g, bool pred);

check returned type and value for both group and sub_group.

2.5.9. shift_group_left

For some scalar, vector, marray, and user defined trivially copyable types T and both group and sub_group with and without delta check returned type and value of T shift_group_left(Group g, T x, Group::linear_id_type delta).

2.5.10. shift_group_right

For some scalar, vector, marray, and user defined trivially copyable types T and both group and sub_group with and without delta check returned type and value of T shift_group_right(Group g, T x, Group::linear_id_type delta).

2.5.11. permute_group_by_xor

For some scalar, vector, marray, and user defined trivially copyable types T and both group and sub_group check returned type and value of T permute_group_by_xor(Group g, T x, Group::linear_id_type mask); with several masks.

2.5.12. select_from_group

For some scalar, vector, marray, and user defined trivially copyable types T and both group and sub_group check returned type and value of T select_from_group(Group g, T x, Group::id_type remote_local_id);.

2.5.13. joint_reduce

For some pointers to a fundamental type Ptr (using V = std::iterator_traits<Ptr>::value_type) for each of the following use-cases:

  • using V joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op); with some V binary_op(V, V)

  • using T joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op); with some T binary_op(T, V)

check returned type and value for both group and sub_group.

2.5.14. reduce_over_group

For some fundamental types T and V for each of the following use-cases:

  • using T reduce_over_group(Group g, T x, BinaryOperation binary_op); with some T binary_op(T, T)

  • using T reduce_over_group(Group g, V x, T init, BinaryOperation binary_op); with some T binary_op(T, V)

check returned type and value for both group and sub_group.

2.5.15. joint_exclusive_scan

For some pointers to a fundamental type InPtr and OutPtr (using I = std::iterator_traits<InPtr>::value_type and O = std::iterator_traits<OutPtr>::value_type) for each of the following use-cases:

  • using OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op); with some O binary_op(I, I)

  • using OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, BinaryOperation binary_op); with some fundamental type T and T binary_op(T, I)

check returned type and value for both group and sub_group.

2.5.16. joint_inclusive_scan

For some pointers to a fundamental type InPtr and OutPtr (using I = std::iterator_traits<InPtr>::value_type and O = std::iterator_traits<OutPtr>::value_type) for each of the following use-cases:

  • using OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op); with some O binary_op(I, I)

  • using OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op, T init); with some fundamental type T and T binary_op(T, I)

check returned type and value for both group and sub_group.

2.5.17. exclusive_scan_over_group

For some fundamental types T and V for each of the following use-cases:

  • using T exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op); with some T binary_op(T, T)

  • using T exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op); with some T binary_op(T, V)

check returned type and value for both group and sub_group.

2.5.18. inclusive_scan_over_group

For some fundamental types T and V for each of the following use-cases:

  • using T inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op); with some T binary_op(T, T)

  • using T inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init); with some T binary_op(T, V)

check returned type and value for both group and sub_group.