-
Notifications
You must be signed in to change notification settings - Fork 82
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Merge branch 'SYCL-2020' into steffen/test_conformance_target
- Loading branch information
Showing
1 changed file
with
45 additions
and
0 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,45 @@ | ||
:sectnums: | ||
:xrefstyle: short | ||
|
||
= Test plan for sycl_ext_oneapi_auto_local_range | ||
|
||
This is a test plan for the API described in | ||
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_auto_local_range.asciidoc[sycl_ext_oneapi_auto_local_range]. | ||
|
||
|
||
== 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_AUTO_LOCAL_RANGE` so they can be skipped | ||
if feature is not supported. | ||
|
||
== Tests | ||
|
||
* All following tests run with `Dimensions` = 1, 2, 3 | ||
|
||
=== auto_range function | ||
|
||
Check that `auto_range<Dimensions>()` return type is `range<Dimensions>` | ||
|
||
=== auto_range in parallel_for | ||
|
||
For following `parallel_for` functions: | ||
|
||
* `queue::parallel_for` | ||
* `handler::parallel_for` | ||
|
||
Check that a kernel launched using `auto_range<Dimensions>()` as the local range behaves as expected and can use group APIs: Create a local accumulator in each work item to sum values from an input buffer. Get the total input sum using `sycl::reduce_over_group` to accumulate the partial sums from the work items within the group. Check that this total has the expected value. Example kernel: | ||
|
||
``` | ||
int local_accumulator = 0; | ||
for (int i = g.get_local_id(); i < N; i += g.get_local_linear_range()) { | ||
local_accumulator += input[i]; | ||
} | ||
int total = sycl::reduce_over_group(g, local_accumulator, sycl::plus<>()); | ||
``` |