diff --git a/.github/workflows/cts_ci.yml b/.github/workflows/cts_ci.yml index 9b9fa5850..f0fd497ab 100644 --- a/.github/workflows/cts_ci.yml +++ b/.github/workflows/cts_ci.yml @@ -56,7 +56,7 @@ jobs: matrix: include: - sycl-impl: dpcpp - version: 1dbee22f9c8a3a825deb871bab76937e04fa26fc + version: b209b321b5a8540263af9ba317c89a1882f06120 - sycl-impl: hipsycl version: 3d8b1cd steps: @@ -114,7 +114,7 @@ jobs: matrix: include: - sycl-impl: dpcpp - version: 1dbee22f9c8a3a825deb871bab76937e04fa26fc + version: b209b321b5a8540263af9ba317c89a1882f06120 - sycl-impl: hipsycl version: 3d8b1cd env: diff --git a/CMakeLists.txt b/CMakeLists.txt index 1365767bc..fe4d33ab8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -83,6 +83,10 @@ add_cts_option(SYCL_CTS_ENABLE_EXT_ONEAPI_AUTO_LOCAL_RANGE_TESTS "Enable extension oneAPI auto_local_range tests" OFF FORCE_ON ${SYCL_CTS_ENABLE_EXT_ONEAPI_TESTS}) +add_cts_option(SYCL_CTS_ENABLE_EXT_ONEAPI_KERNEL_COMPILER_SPIRV_TESTS + "Enable extension oneAPI kernel_compiler_spirv tests" OFF + FORCE_ON ${SYCL_CTS_ENABLE_EXT_ONEAPI_TESTS}) + add_cts_option(SYCL_CTS_ENABLE_EXT_ONEAPI_NON_UNIFORM_GROUPS_TESTS "Enable extension oneAPI non_uniform_groups tests" OFF FORCE_ON ${SYCL_CTS_ENABLE_EXT_ONEAPI_TESTS}) diff --git a/test_plans/non_uniform_groups.asciidoc b/test_plans/non_uniform_groups.asciidoc new file mode 100644 index 000000000..c714b0dc8 --- /dev/null +++ b/test_plans/non_uniform_groups.asciidoc @@ -0,0 +1,400 @@ +: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/experimental/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 `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 `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 `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 `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 result of `get_local_range()` on 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 `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 `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 `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 `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. +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 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. + +==== 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 `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 `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 `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 `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 result of `get_local_range()` on 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 `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 `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 `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 `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 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. diff --git a/test_plans/oneapi_kernel_compiler_spirv.asciidoc b/test_plans/oneapi_kernel_compiler_spirv.asciidoc new file mode 100644 index 000000000..d21bc3703 --- /dev/null +++ b/test_plans/oneapi_kernel_compiler_spirv.asciidoc @@ -0,0 +1,75 @@ +:sectnums: +:xrefstyle: short + += Test plan for sycl_ext_oneapi_kernel_compiler_spirv + +This is a test plan for the API described in +https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc[sycl_ext_oneapi_kernel_compiler_spirv]. + +== 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_KERNEL_COMPILER_SPIRV` so +they can be skipped if feature is not supported. + +== Tests + +All of the following tests run SPIR-V kernels loaded in binary form. + +=== SPIR-V Source Language Test + +Load a simple SPIR-V kernel and run it to ensure it has the expected behavior. +This kernel should take two parameters: an input pointer and an output pointer. +Each of these pointers should have type *OpTypePointer*, with *CrossWorkgroup* +storage class pointing to *OpTypeInt* with width 32. For each work item, the +kernel computes the following expression: `out[id] = (a * in[id]) + b`, where +`a` and `b` are integer constants. The host code will create this kernel, pass +its USM pointer arguments, run it, and assert that the output has the expected +values. This test ensures that `create_kernel_bundle_from_source` can be used +with `source_language::spirv` to obtain a kernel, set its parameters, and run +it. + +=== Kernel API Test + +This test checks that `ext_oneapi_has_kernel` and `ext_oneapi_get_kernel` have +the expected behavior. Assert that `ext_oneapi_has_kernel` returns true and +`ext_oneapi_get_kernel` returns a kernel when the name parameter matches a +SPIR-V entrypoint. Also, assert `ext_oneapi_has_kernel` returns false and +`ext_oneapi_get_kernel` throws an exception with `errc::invalid` if the name is +not valid. + +=== Parameter Tests + +This test checks that kernels can accept parameters for all of the SPIR-V types +required by the extension. The required types are the following: + +- *OpTypeInt*, width 8, 16, 32, and 64. +- *OpTypeFloat*, width 16, 32 and 64. + +For each type `T`, define a kernel with parameters `T`, *OpTypePointer* with +*Workgroup* storage class pointing to `T`, and *OpTypePointer* with +*CrossWorkgroup* storage class pointing to `T`. This kernel should compute an +expression using all three parameters and store the result. The host code can +then check this result to ensure that the parameter types are working. + +For the *OpTypeFloat* parameters with width 16 or 64, the host code should use +`sycl::aspect::fp16` and `sycl::aspect::fp64` to determine if the test kernels +can be built and run, or if they should be skipped. + +=== Struct Parameter Tests + +This test checks that kernels can accept *OpTypeStruct* parameters that match +the constraints specified by the extension. Define a kernel that accepts an +input *OpTypePointer* with *Function* storage class pointing to an +*OpTypeStruct*, and an output *OpTypePointer* with *CrossWorkgroup* storage +class pointing to the same *OpTypeStruct*. The struct should contain +*OpTypeInt*, *OpTypeFloat*, and inner *OpTypeStruct* members. The kernel +computes an expression using each member from the input and stores the result in +the corresponding member in the output. The host code then checks the output +struct to ensure the members have the expected values. diff --git a/test_plans/sycl_ext_oneapi_composite_device.asciidoc b/test_plans/sycl_ext_oneapi_composite_device.asciidoc new file mode 100644 index 000000000..12e046341 --- /dev/null +++ b/test_plans/sycl_ext_oneapi_composite_device.asciidoc @@ -0,0 +1,148 @@ +:sectnums: +:xrefstyle: short + += Test plan for sycl_ext_oneapi_composite_device + +This is a test plan for the API described in the +https://github.com/intel/llvm/blob/e94b24718e60a7fa03ca1abbde4f7e37bbd0557d/sycl/doc/extensions/proposed/sycl_ext_oneapi_composite_device.asciidoc[`sycl_ext_oneapi_composite_device`] +extension specification. + +== Testing scope + +=== Device coverage + +All of the tests described below are performed on the default device that is +selected on the CTS command line, but some of the test cases may query for +additional devices from the same platform the default device belongs to. + +=== Feature test macro + +All of the tests should use `#ifdef SYCL_EXT_ONEAPI_COMPOSITE_DEVICE` so they +can be skipped if the extension is not supported. + +== Tests + +Note that some of the sections below closely follow the extension specification +document and describe tests with narrower scope, which are aimed to check +specific functionality. There are sections below which describe bigger test +cases which are intended to check how different elements of the extension work +together. + +=== Enumerating composite devices + +==== Through the new `get_composite_devices()` API + +This test should check that freestanding function `get_composite_devices()`: + +* has the correct signature +* does not throw an exception +* returns the same result on subsequent calls + +==== Through the new `platform::ext_oneapi_get_composite_devices()` API + +This test should take a platform the selected device belongs to and check that +`platform::ext_oneapi_get_composite_devices()` method: + +* has the correct signature +* does not throw an exception +* returns the same result on subsequent calls + +==== Composite devices are not considered as root devices + +The test should ensure that no composite devices are returned through the +`platform::get_devices()` method for a platform the selected device belongs to. + +=== New device information descriptors + +Composite devices are not considered as root devices and therefore they won't +ever be selected by the CTS. To increase coverage, tests in this subcategory +should be additionally performed on a composite device the selected device +belongs to (if the selected device is a component device). + +==== `info::device::component_devices` + +The test should check that a query for the information descriptor has correct +return type. + +The test should also check the value returned for the descriptor: + +* if a tested device is not a composite device, an empty vector is expected +* otherwise, returned vector is expected to have at least two elements and those + elements should be additionally checked: +** each returned device should have the `ext_oneapi_is_component` aspect +** each returned device should return the tested device through + `info::device::composite_device` query + +==== `info::device::composite_device` + +The test should check that a query for the information descriptor has correct +return type. + +The test should also check the value returned for the descriptor: + +* if a tested device is not a component device, then a synchronous + `exception` with `errc::invalid` error code is expected +* otherwise, returned device should have the `ext_oneapi_is_composite` aspect + +=== New device aspects + +==== `ext_oneapi_is_composite` + +The test should ensure that all of devices returned by +`get_composite_devices()` and `platform::ext_oneapi_get_composite_devices()` +have `ext_oneapi_is_composite` aspect. + +==== `ext_oneapi_is_component` + +The test should ensure that none of devices returned by +`get_composite_devices()` and `platform::ext_oneapi_get_composite_devices()` +have `ext_oneapi_is_component` aspect. + +If the CTS selected device is a component device, then test should then attempt +to partition a component device using one of partition methods it supports and +check that none of sub-devices have the `ext_oneapi_is_component` aspect. + +=== Impact on "descendent device" + +The test should first obtain both a component and a composite device and it is +skipped if that is not possible using the selected device. + +Having those two, the test should create a `context` using the _composite_ +device and then attempt to create a `queue` by passing the created context in +there *and* the _component_ device. The test checks that `queue` was created +successfully, i.e. no exceptions were thrown. + +=== More complex test cases + +Tests in this category are intended to represent miniature versions of potential +real-world applications to check how the extension works with other core SYCL +features. + +==== Basic tests for the composite device + +The test obtains a composite device using the selected device and it is skipped +if unable to do so. + +The composite device is then used to submit a series of kernels that perform +some computations and memory operations using a `queue` object associated with +that device. + +==== "Interoperability" between composite and component devices + +The test should first obtain both a component and a composite device and it is +skipped if that is not possible using the selected device. + +The test then creates a shared context for both devices, allocates a shared +memory and creates separate command queues for each device. Test performs some +computations and memory operations on that shared data using both devices. + +==== Sharing memory to a descendent device + +The test should first obtain both a component and a composite device and it is +skipped if that is not possible using the selected device. + +The test allocates a memory in a context a component device belongs to. It then +create a queue for a composite device using _the same_ context and checks that +memory allocated for a composite device is available to a descendent component +device as well. + diff --git a/tests/common/common.h b/tests/common/common.h index 6b00a492f..ec72c36ca 100644 --- a/tests/common/common.h +++ b/tests/common/common.h @@ -38,6 +38,7 @@ #include "string_makers.h" #include "value_operations.h" +#include #include #include #include @@ -850,4 +851,22 @@ inline sycl::id<3> unlinearize(sycl::range<3> range, size_t id) { return {id0, id1, id2}; } +/** @brief Checks that two vectors of devices have the exact same devices, + * ignoring order and repeated devices. + * @param lhs std::vector with sycl::device + * @param rhs std::vector with sycl::device + */ +inline bool have_same_devices(std::vector lhs, + std::vector rhs) { + // TODO: If SYCL devices are given well-defined ordering, this can be + // implemented using std::set_difference. + auto create_check_func = [](const std::vector& devices) { + return [&devices](const sycl::device& dev) { + return std::find(devices.cbegin(), devices.cend(), dev) != devices.cend(); + }; + }; + return std::all_of(lhs.cbegin(), lhs.cend(), create_check_func(rhs)) && + std::all_of(rhs.cbegin(), rhs.cend(), create_check_func(lhs)); +} + #endif // __SYCLCTS_TESTS_COMMON_COMMON_H diff --git a/tests/extension/oneapi_kernel_compiler_spirv/CMakeLists.txt b/tests/extension/oneapi_kernel_compiler_spirv/CMakeLists.txt new file mode 100644 index 000000000..eec2e11dc --- /dev/null +++ b/tests/extension/oneapi_kernel_compiler_spirv/CMakeLists.txt @@ -0,0 +1,12 @@ +if(SYCL_CTS_ENABLE_EXT_ONEAPI_KERNEL_COMPILER_SPIRV_TESTS) + file(GLOB test_cases_list *.cpp) + + file(TO_CMAKE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/kernels.spv" KERNELS_PATH) + file(TO_CMAKE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/kernels_fp16.spv" KERNELS_FP16_PATH) + file(TO_CMAKE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/kernels_fp64.spv" KERNELS_FP64_PATH) + add_compile_definitions(KERNELS_PATH="${KERNELS_PATH}") + add_compile_definitions(KERNELS_FP16_PATH="${KERNELS_FP16_PATH}") + add_compile_definitions(KERNELS_FP64_PATH="${KERNELS_FP64_PATH}") + + add_cts_test(${test_cases_list}) +endif() diff --git a/tests/extension/oneapi_kernel_compiler_spirv/kernel_compiler_spirv.cpp b/tests/extension/oneapi_kernel_compiler_spirv/kernel_compiler_spirv.cpp new file mode 100644 index 000000000..645000ed2 --- /dev/null +++ b/tests/extension/oneapi_kernel_compiler_spirv/kernel_compiler_spirv.cpp @@ -0,0 +1,248 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +*******************************************************************************/ + +#include "../../common/common.h" + +namespace kernel_compiler_spirv::tests { + +#ifdef SYCL_EXT_ONEAPI_AUTO_LOCAL_RANGE + +sycl::kernel_bundle loadKernelsFromFile( + sycl::queue& q, const std::string& file_name) { + namespace syclex = sycl::ext::oneapi::experimental; + + // Read the SPIR-V module from disk. + std::ifstream spv_stream(file_name, std::ios::binary); + if (!spv_stream.is_open()) { + throw std::runtime_error("Failed to open '" + file_name + "'"); + } + spv_stream.seekg(0, std::ios::end); + size_t sz = spv_stream.tellg(); + spv_stream.seekg(0); + std::vector spv(sz); + spv_stream.read(reinterpret_cast(spv.data()), sz); + + // Create a kernel bundle from the binary SPIR-V. + sycl::kernel_bundle kb_src = + syclex::create_kernel_bundle_from_source( + q.get_context(), syclex::source_language::spirv, spv); + + // Build the SPIR-V module for our device. + sycl::kernel_bundle kb_exe = + syclex::build(kb_src); + return kb_exe; +} + +void testSimpleKernel(sycl::queue& q, const sycl::kernel& kernel, + int multiplier, int added) { + const auto num_args = kernel.get_info(); + REQUIRE(num_args == 2); + + constexpr int N = 4; + std::array input_array{0, 1, 2, 3}; + + sycl::buffer input_buffer{input_array.data(), sycl::range<1>(N)}; + sycl::buffer output_buffer{sycl::range<1>(N)}; + + q.submit([&](sycl::handler& cgh) { + cgh.set_args(sycl::accessor{input_buffer, cgh, sycl::read_only}, + sycl::accessor{output_buffer, cgh, sycl::write_only}); + cgh.parallel_for(sycl::range<1>{N}, kernel); + }); + + sycl::host_accessor out{output_buffer}; + for (int i = 0; i < N; i++) { + CHECK(out[i] == ((i * multiplier) + added)); + } +} + +/* +For each type T, the matching SPIR-V kernel takes four parameters: + + 1. [in] a: T + 2. [in] b: OpTypePointer(CrossWorkgroup) to T + 3. [in/out] tmp: OpTypePointer(Workgroup) to T + 4. [out] c: OpTypePointer(CrossWorkgroup) to T + +The kernel computes the following expressions: + + 1. *tmp = (a * a); + 2. *c = (*tmp) + ((*b) * (*b)); + +This test case sets the four parameters, runs the kernel, and asserts that +output c has the expected value. +*/ +template +void testParam(sycl::queue& q, const sycl::kernel& kernel) { + const auto num_args = kernel.get_info(); + REQUIRE(num_args == 4); + + // Kernel computes sum of squared inputs. + const T a = 2; + const T b = 5; + const T expected = (a * a) + (b * b); + + sycl::buffer a_buffer(&a, sycl::range<1>(1)); + + T* const b_ptr = sycl::malloc_shared(1, q); + b_ptr[0] = b; + + T output{}; + sycl::buffer output_buffer(&output, sycl::range<1>(1)); + + q.submit([&](sycl::handler& cgh) { + sycl::local_accessor local(1, cgh); + // Pass T for scalar parameter. + cgh.set_arg(0, a); + // Pass USM pointer for OpTypePointer(CrossWorkgroup) parameter. + cgh.set_arg(1, b_ptr); + // Pass sycl::accessor for OpTypePointer(CrossWorkgroup) parameter. + cgh.set_arg(2, sycl::accessor{output_buffer, cgh, sycl::write_only}); + // Pass sycl::local_accessor for OpTypePointer(Workgroup) parameter. + cgh.set_arg(3, local); + cgh.parallel_for(sycl::range<1>{1}, kernel); + }); + + sycl::host_accessor out{output_buffer}; + CHECK(out[0] == expected); + sycl::free(b_ptr, q); +} + +void testStruct(sycl::queue& q, const sycl::kernel& kernel) { + const auto num_args = kernel.get_info(); + REQUIRE(num_args == 2); + + // This definition must match the one used in the kernel. + struct S { + std::int32_t i; + cl_float f; + std::int32_t* p; + struct Inner { + std::int32_t i; + float f; + std::int32_t* p; + } inner; + }; + + // Any constants can be used to initialize this input. + std::int32_t* const in_p0 = sycl::malloc_shared(1, q); + std::int32_t* const in_p1 = sycl::malloc_shared(1, q); + *in_p0 = 3; + *in_p1 = 6; + S input{1, 2.0f, in_p0, S::Inner{4, 5.0f, in_p1}}; + + std::int32_t* const out_p0 = sycl::malloc_shared(1, q); + std::int32_t* const out_p1 = sycl::malloc_shared(1, q); + *out_p0 = 0; + *out_p1 = 0; + S* output = sycl::malloc_shared(1, q); + *output = S{0, 0, out_p0, S::Inner{0, 0, out_p1}}; + + q.submit([&](sycl::handler& cgh) { + cgh.set_args(input, output); + cgh.parallel_for(sycl::range<1>{1}, kernel); + }).wait(); + + // For each scalar struct member, output == (2 * input). For pointer members, + // *output == (2 * (*input)). + CHECK(output->i == input.i * 2); + CHECK(output->f == input.f * 2); + CHECK(*output->p == (*input.p) * 2); + CHECK(output->inner.i == input.inner.i * 2); + CHECK(output->inner.f == input.inner.f * 2); + CHECK(*output->inner.p == (*input.inner.p) * 2); + + sycl::free(output, q); + sycl::free(in_p0, q); + sycl::free(in_p1, q); + sycl::free(out_p0, q); + sycl::free(out_p1, q); +} + +void testKernelsFromSpvFile(std::string kernels_file, + std::string fp16_kernel_file, + std::string fp64_kernel_file) { + const auto getKernel = + [](sycl::kernel_bundle& bundle, + const std::string& name) { + return bundle.ext_oneapi_get_kernel(name); + }; + + sycl::queue q; + auto bundle = loadKernelsFromFile(q, kernels_file); + + // Test kernel retrieval functions. + { + CHECK(bundle.ext_oneapi_has_kernel("my_kernel")); + CHECK(bundle.ext_oneapi_has_kernel("OpTypeInt8")); + CHECK(bundle.ext_oneapi_has_kernel("OpTypeInt16")); + CHECK(bundle.ext_oneapi_has_kernel("OpTypeInt32")); + CHECK(bundle.ext_oneapi_has_kernel("OpTypeInt64")); + CHECK(bundle.ext_oneapi_has_kernel("OpTypeInt9") == false); + CHECK(bundle.ext_oneapi_has_kernel("") == false); + + CHECK_NOTHROW(bundle.ext_oneapi_get_kernel("my_kernel")); + CHECK_NOTHROW(bundle.ext_oneapi_get_kernel("OpTypeInt8")); + CHECK_NOTHROW(bundle.ext_oneapi_get_kernel("OpTypeInt16")); + CHECK_NOTHROW(bundle.ext_oneapi_get_kernel("OpTypeInt32")); + CHECK_NOTHROW(bundle.ext_oneapi_get_kernel("OpTypeInt64")); + CHECK_THROWS_AS(bundle.ext_oneapi_get_kernel("OpTypeInt9"), + sycl::exception); + CHECK_THROWS_AS(bundle.ext_oneapi_get_kernel(""), sycl::exception); + } + + // Test simple kernel. + testSimpleKernel(q, getKernel(bundle, "my_kernel"), 2, 100); + + // Test parameters. + testParam(q, getKernel(bundle, "OpTypeInt8")); + testParam(q, getKernel(bundle, "OpTypeInt16")); + testParam(q, getKernel(bundle, "OpTypeInt32")); + testParam(q, getKernel(bundle, "OpTypeInt64")); + testParam(q, getKernel(bundle, "OpTypeFloat32")); + + // Test OpTypeFloat16 parameters. + if (q.get_device().has(sycl::aspect::fp16)) { + auto fp16_bundle = loadKernelsFromFile(q, fp16_kernel_file); + testParam(q, getKernel(fp16_bundle, "OpTypeFloat16")); + } + + // Test OpTypeFloat64 parameters. + if (q.get_device().has(sycl::aspect::fp64)) { + auto fp64_bundle = loadKernelsFromFile(q, fp64_kernel_file); + testParam(q, getKernel(fp64_bundle, "OpTypeFloat64")); + } + + // Test OpTypeStruct parameters. + testStruct(q, getKernel(bundle, "OpTypeStruct")); +} + +#endif + +TEST_CASE("Test case for \"Kernel Compiler SPIR-V\" extension", + "[oneapi_kernel_compiler_spirv]") { +#ifndef SYCL_EXT_ONEAPI_KERNEL_COMPILER_SPIRV + SKIP("SYCL_EXT_ONEAPI_KERNEL_COMPILER_SPIRV is not defined"); +#else + testKernelsFromSpvFile(KERNELS_PATH, KERNELS_FP16_PATH, KERNELS_FP64_PATH); +#endif +} + +} // namespace kernel_compiler_spirv::tests diff --git a/tests/extension/oneapi_kernel_compiler_spirv/kernels.spv b/tests/extension/oneapi_kernel_compiler_spirv/kernels.spv new file mode 100644 index 000000000..7a4d2d1df Binary files /dev/null and b/tests/extension/oneapi_kernel_compiler_spirv/kernels.spv differ diff --git a/tests/extension/oneapi_kernel_compiler_spirv/kernels_fp16.spv b/tests/extension/oneapi_kernel_compiler_spirv/kernels_fp16.spv new file mode 100644 index 000000000..6a132a0d9 Binary files /dev/null and b/tests/extension/oneapi_kernel_compiler_spirv/kernels_fp16.spv differ diff --git a/tests/extension/oneapi_kernel_compiler_spirv/kernels_fp64.spv b/tests/extension/oneapi_kernel_compiler_spirv/kernels_fp64.spv new file mode 100644 index 000000000..37aa8c025 Binary files /dev/null and b/tests/extension/oneapi_kernel_compiler_spirv/kernels_fp64.spv differ diff --git a/tests/group_functions/group_broadcast.h b/tests/group_functions/group_broadcast.h index b8b120378..06207e980 100644 --- a/tests/group_functions/group_broadcast.h +++ b/tests/group_functions/group_broadcast.h @@ -131,87 +131,132 @@ void broadcast_sub_group(sycl::queue& queue) { sycl::range work_group_range = sycl_cts::util::work_group_range(queue); // array to return results - T res[test_matrix + 1] = {splat_init(0)}; + T origin_values[test_matrix] = {splat_init(0)}; + T broadcasted_values[test_matrix] = {splat_init(0)}; { - sycl::buffer res_sycl(res, sycl::range<1>(test_matrix + 1)); + sycl::buffer origin_values_buf(origin_values, + sycl::range<1>(test_matrix)); + sycl::buffer broadcasted_values_buf(broadcasted_values, + sycl::range<1>(test_matrix)); queue.submit([&](sycl::handler& cgh) { - auto res_acc = - res_sycl.template get_access(cgh); + auto origin_values_acc = + origin_values_buf.template get_access( + cgh); + auto broadcasted_values_acc = + broadcasted_values_buf + .template get_access(cgh); sycl::nd_range executionRange(work_group_range, work_group_range); + // Values computed in a kernel depend on global linear id. We need to make + // sure that there are no overflows + REQUIRE(executionRange.get_global_range().size() < + std::numeric_limits::max() / 100); cgh.parallel_for>(executionRange, [=](sycl::nd_item item) { sycl::sub_group sub_group = item.get_sub_group(); + // Each work-item computes a unique value + T value_to_broadcast(splat_init(static_cast( + item.get_global_linear_id() * 100 + sub_group.get_local_id()))); T local_var(splat_init(0)); + // To simplify the test, we are only checking the first sub-group if (sub_group.get_group_id()[0] == 0) { - // find local id of last group item - sycl::id<1> last_item = sub_group.get_local_range(); - --last_item[0]; - - // broadcast from the first workitem - local_var = splat_init(item.get_global_linear_id() + 1); + // Find local id of the leader, last and some other work-item in the + // sub-group. They will be used to check different combinations of + // broadcasting and receiving work-items + sycl::id<1> first_id = 0; + sycl::id<1> mid_id = sub_group.get_local_range() / 2; + sycl::id<1> last_id = sub_group.get_local_range() - 1; + + // Broadcast from the first work-item ASSERT_RETURN_TYPE( - T, sycl::group_broadcast(sub_group, local_var), + T, sycl::group_broadcast(sub_group, value_to_broadcast), "Return type of group_broadcast(sub_group g, T x) is wrong\n"); - local_var = sycl::group_broadcast(sub_group, local_var); - if (sub_group.get_local_linear_id() == - sub_group.get_local_linear_range() - 1) - res_acc[0] = local_var; - - // broadcast from the last workitem 1 - local_var = splat_init(item.get_global_linear_id() + 1); + if (sub_group.leader()) { + // Work-item which does the broadcast stores value to broadcast to + // use it later as a reference + origin_values_acc[0] = value_to_broadcast; + } + auto broadcasted_value = + sycl::group_broadcast(sub_group, value_to_broadcast); + // We read broadcasted value in another work-item + if (sub_group.get_local_id() == last_id) + broadcasted_values_acc[0] = broadcasted_value; + + // Broadcast from the last work-item, we specifically don't use + // sycl::id as the third argument to check overload with + // sub_group::linear_id_type as the third argument ASSERT_RETURN_TYPE( - T, sycl::group_broadcast(sub_group, local_var, last_item), + T, + sycl::group_broadcast(sub_group, value_to_broadcast, + sub_group.get_local_linear_range() - 1), "Return type of group_broadcast(sub_group g, T x, " "sub_group::linear_id_type local_linear_id) is wrong\n"); - local_var = sycl::group_broadcast( - sub_group, local_var, sub_group.get_local_linear_range() - 1); - if (sub_group.get_local_linear_id() == 0) res_acc[1] = local_var; + if (sub_group.get_local_id() == last_id) { + // Work-item which does the broadcast stores value to broadcast to + // use it later as a reference + origin_values_acc[1] = value_to_broadcast; + } - // broadcast from the last workitem 2 - local_var = splat_init(item.get_global_linear_id() + 1); + broadcasted_value = + sycl::group_broadcast(sub_group, value_to_broadcast, + sub_group.get_local_linear_range() - 1); + // We read broadcasted value in another work-item + if (sub_group.get_local_id() == mid_id) + broadcasted_values_acc[1] = broadcasted_value; + + // Broadcast from a mid work-item. This is similar to the test case + // above, but it checks overload which accepts sub_group::id_type as + // the last argument ASSERT_RETURN_TYPE( - T, sycl::group_broadcast(sub_group, local_var, last_item), + T, sycl::group_broadcast(sub_group, value_to_broadcast, mid_id), "Return type of group_broadcast(sub_group g, T x, " "sub_group::id_type local_id) is wrong\n"); - local_var = sycl::group_broadcast(sub_group, local_var, last_item); - if (sub_group.get_local_linear_id() == 0) res_acc[2] = local_var; + if (sub_group.get_local_id() == mid_id) { + // Work-item which does the broadcast stores value to broadcast to + // use it later as a reference + origin_values_acc[2] = value_to_broadcast; + } + broadcasted_value = + sycl::group_broadcast(sub_group, value_to_broadcast, mid_id); + // We read broadcasted value in another work-item + if (sub_group.get_local_id() == first_id) + broadcasted_values_acc[2] = broadcasted_value; - // select from the last workitem - local_var = splat_init(item.get_global_linear_id() + 1); + // Select from the first work-item ASSERT_RETURN_TYPE( - T, sycl::select_from_group(sub_group, local_var, last_item), + T, + sycl::select_from_group(sub_group, value_to_broadcast, first_id), "Return type of select_from_group(sub_group g, T x, " "sub_group::id_type local_id) is wrong\n"); - local_var = sycl::select_from_group(sub_group, local_var, last_item); - if (sub_group.get_local_linear_id() == 0) res_acc[3] = local_var; - - // Return the sub-group size when possible or just its parity - if (sub_group.get_local_linear_id() == 0) { - if constexpr (std::is_same_v) - res_acc[4] = sub_group.get_local_linear_range() % 2; - else - res_acc[4] = sub_group.get_local_linear_range(); + if (sub_group.get_local_id() == first_id) { + // Work-item which does the broadcast stores value to broadcast to + // use it later as a reference + origin_values_acc[3] = value_to_broadcast; } + broadcasted_value = + sycl::select_from_group(sub_group, value_to_broadcast, first_id); + // We read broadcasted value in another work-item + if (sub_group.get_local_id() == mid_id) + broadcasted_values_acc[3] = broadcasted_value; } }); }); } - T expected[test_matrix] = {splat_init(1), res[4], res[4], res[4]}; for (int i = 0; i < test_matrix; ++i) { std::string work_group = sycl_cts::util::work_group_print(work_group_range); CAPTURE(D, work_group); INFO("Return value of " << test_names[i] << " with T = " << type_name() << " is " - << (equal(res[i], expected[i]) ? "right" : "wrong")); - CHECK(equal(res[i], expected[i])); + << (equal(broadcasted_values[i], origin_values[i]) ? "right" + : "wrong")); + CHECK(equal(broadcasted_values[i], origin_values[i])); } } diff --git a/tests/kernel_bundle/sycl_build_verify_kernel_invoked_and_kernel_in_result_bundle.cpp b/tests/kernel_bundle/sycl_build_verify_kernel_invoked_and_kernel_in_result_bundle.cpp index e308525a3..2794b9e93 100644 --- a/tests/kernel_bundle/sycl_build_verify_kernel_invoked_and_kernel_in_result_bundle.cpp +++ b/tests/kernel_bundle/sycl_build_verify_kernel_invoked_and_kernel_in_result_bundle.cpp @@ -59,7 +59,7 @@ void verify_results( if (kernel_bundle.get_context() != ctx) { FAIL(log, "Kernel bundle's context does not equal to provided context"); } - if (kernel_bundle.get_devices() != dev_vector) { + if (!have_same_devices(kernel_bundle.get_devices(), dev_vector)) { FAIL(log, "Devices from kernel bundle not equal to provided devices"); } } @@ -118,7 +118,7 @@ struct verify_that_bundles_are_same { */ void run_verification(util::logger &log, sycl::queue &q) { auto ctx = q.get_context(); - std::vector dev_vector{ctx.get_devices()[0]}; + std::vector dev_vector{ctx.get_devices()}; const auto first_simple_kernel_id = sycl::get_kernel_id(); diff --git a/tests/kernel_bundle/sycl_join_kernel_bundle_with_empty_one.cpp b/tests/kernel_bundle/sycl_join_kernel_bundle_with_empty_one.cpp index e33633724..e0a744042 100644 --- a/tests/kernel_bundle/sycl_join_kernel_bundle_with_empty_one.cpp +++ b/tests/kernel_bundle/sycl_join_kernel_bundle_with_empty_one.cpp @@ -30,13 +30,12 @@ template void run_verification(util::logger &log) { auto queue = util::get_cts_object::queue(); const auto ctx = queue.get_context(); - const auto dev = queue.get_device(); auto kb = sycl::get_kernel_bundle(ctx); // Selector that always returns false. Used to get empty kernel_bundle auto false_selector = [](const sycl::device_image &) { return false; }; - auto empty_kb = sycl::get_kernel_bundle(ctx, {dev}, false_selector); + auto empty_kb = sycl::get_kernel_bundle(ctx, false_selector); // Check joined bundles in such order: (kernel_bundle, empty_kernel_bundle) { diff --git a/tests/kernel_bundle/sycl_link_verify_kernel_invoked_and_kernel_in_result_bundle.cpp b/tests/kernel_bundle/sycl_link_verify_kernel_invoked_and_kernel_in_result_bundle.cpp index 59c108810..0e27b27bb 100644 --- a/tests/kernel_bundle/sycl_link_verify_kernel_invoked_and_kernel_in_result_bundle.cpp +++ b/tests/kernel_bundle/sycl_link_verify_kernel_invoked_and_kernel_in_result_bundle.cpp @@ -71,7 +71,7 @@ void verify_results( if (kernel_bundle.get_context() != ctx) { FAIL(log, "Kernel bundle's context does not equal to provided context"); } - if (kernel_bundle.get_devices() != dev_vector) { + if (!have_same_devices(kernel_bundle.get_devices(), dev_vector)) { FAIL(log, "Devices from kernel bundle not equal to provided devices"); } } @@ -117,7 +117,7 @@ void run_verification(util::logger &log, sycl::queue &queue) { kb_with_first_simple_kernel_from_input, kb_with_second_simple_kernel}; std::vector dev_vector{ctx.get_devices()}; - std::vector current_dev_vector{util::get_cts_object::device()}; + std::vector current_dev_vector{queue.get_device()}; log.note("Verify link(vector>, vector) overload"); verify_results(log, current_dev_vector, ctx, diff --git a/tests/marray/marray_constructor.h b/tests/marray/marray_constructor.h index 894d1e460..dfd4157d8 100644 --- a/tests/marray/marray_constructor.h +++ b/tests/marray/marray_constructor.h @@ -145,7 +145,6 @@ class run_marray_constructor_test { }) .wait_and_throw(); } - run_checks(check_results); for (size_t i = 0; i < num_test_cases; ++i) { INFO(check_names[i]); CHECK(check_results[i]); diff --git a/tests/nd_item/nd_item_constructors.cpp b/tests/nd_item/nd_item_constructors.cpp index f3e5126dc..ba7af3917 100644 --- a/tests/nd_item/nd_item_constructors.cpp +++ b/tests/nd_item/nd_item_constructors.cpp @@ -203,50 +203,6 @@ void test_constructors(util::logger& log) { CHECK(success_deprecated[to_integral(current_check::move_constructor)]); CHECK(success_deprecated[to_integral(current_check::move_assignment)]); #endif - - // nd_item is not default constructible, store two objects into the array - static constexpr size_t numItems = 2; - using setup_kernel_t = nd_item_setup_kernel; - auto items = - store_instances>(); - { - const auto& item = items[0]; - const auto& itemReadOnly = item; - state_storage expected(itemReadOnly); - - { // Check copy constructor - sycl::nd_item copied(itemReadOnly); - CHECK(expected.check_equality(copied)); -#if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS - CHECK(expected.check_equality_deprecated(copied)); -#endif - } - { // Check copy assignment - auto copied = itemReadOnly; - CHECK(expected.check_equality(copied)); -#if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS - CHECK(expected.check_equality_deprecated(copied)); -#endif - } - { // Check move constructor; invalidates item - sycl::nd_item moved(item); - CHECK(expected.check_equality(moved)); -#if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS - CHECK(expected.check_equality_deprecated(moved)); -#endif - } - } - { - const auto& item = items[1]; - state_storage expected(item); - - // Check move assignment; invalidates item - auto moved = std::move(item); - CHECK(expected.check_equality(moved)); -#if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS - CHECK(expected.check_equality_deprecated(moved)); -#endif - } } class TEST_NAME : public util::test_base { diff --git a/tests/nd_item/nd_item_equality.cpp b/tests/nd_item/nd_item_equality.cpp index bfc4724b3..abae83eed 100644 --- a/tests/nd_item/nd_item_equality.cpp +++ b/tests/nd_item/nd_item_equality.cpp @@ -28,9 +28,6 @@ namespace TEST_NAMESPACE { using namespace sycl_cts; -template -struct nd_item_setup_kernel; - template struct nd_item_equality_kernel; @@ -46,23 +43,41 @@ class TEST_NAME : public util::test_base { template void test_equality(util::logger& log) { + using item_t = sycl::nd_item; + using kernel_t = nd_item_equality_kernel; + + // Store comparison results from kernel into a success array + std::array + success; + std::fill(std::begin(success), std::end(success), true); + { - using item_t = sycl::nd_item; + sycl::buffer successBuf(success.data(), + sycl::range<1>(success.size())); + + const auto oneElemRange = + util::get_cts_object::range::get(1, 1, 1); - // nd_item is not default constructible, store two objects into the array - static constexpr size_t numItems = 2; - using setup_kernel_t = nd_item_setup_kernel; - auto items = - store_instances>(); + auto queue = util::get_cts_object::queue(); + queue + .submit([&](sycl::handler& cgh) { + auto successAcc = + successBuf.get_access(cgh); - // Check nd_item equality operator on the device side - common_by_value_semantics::on_device_checker::template run< - nd_item_equality_kernel>( - log, items, "nd_item " + std::to_string(numDims) + " (device)"); + cgh.parallel_for( + sycl::nd_range(oneElemRange, oneElemRange), + [=](item_t item) { + common_by_value_semantics::check_equality(item, successAcc); + }); + }) + .wait_and_throw(); + } - // Check nd_item equality operator on the host side - common_by_value_semantics::check_on_host( - log, items[0], "nd_item " + std::to_string(numDims) + " (host)"); + for (int i = 0; i < success.size(); ++i) { + INFO(std::string(TOSTRING(TEST_NAME)) + " is " + + common_by_value_semantics::get_error_string(i)); + CHECK(success[i]); } } diff --git a/tests/spec_constants/spec_constants_defined_various_ways.h b/tests/spec_constants/spec_constants_defined_various_ways.h index 50a314b08..d4d14e45c 100644 --- a/tests/spec_constants/spec_constants_defined_various_ways.h +++ b/tests/spec_constants/spec_constants_defined_various_ways.h @@ -27,6 +27,7 @@ void perform_test(util::logger &log, const std::string &type_name, auto queue = util::get_cts_object::queue(); const sycl::context ctx = queue.get_context(); const sycl::device dev = queue.get_device(); + bool has_target_kernel = true; if constexpr (via_kb::value) { if (!dev.has(sycl::aspect::online_compiler)) { @@ -50,6 +51,7 @@ void perform_test(util::logger &log, const std::string &type_name, sycl::get_kernel_bundle(ctx, {dev}, {kernelId}); if (!k_bundle.has_kernel(kernelId)) { + has_target_kernel = false; log.note("kernel_bundle doesn't contain target kernel in case (" + case_hint + ") for " + type_name_string::get(type_name) + " (skipped)"); @@ -72,9 +74,12 @@ void perform_test(util::logger &log, const std::string &type_name, } }); } - if (!check_equal_values(ref, result)) - FAIL(log, - "case (" + case_hint + ") for " + type_name_string::get(type_name)); + if (has_target_kernel) { + // Check results only if target kernel is available + if (!check_equal_values(ref, result)) + FAIL(log, "case (" + case_hint + ") for " + + type_name_string::get(type_name)); + } } template diff --git a/tests/spec_constants/spec_constants_multiple.h b/tests/spec_constants/spec_constants_multiple.h index 1396f9567..68858d42d 100644 --- a/tests/spec_constants/spec_constants_multiple.h +++ b/tests/spec_constants/spec_constants_multiple.h @@ -47,6 +47,7 @@ class check_specialization_constants_multiple_for_type { auto queue = util::get_cts_object::queue(); const sycl::context ctx = queue.get_context(); const sycl::device dev = queue.get_device(); + bool has_target_kernel = true; if constexpr (via_kb::value) { if (!dev.has(sycl::aspect::online_compiler)) { @@ -77,6 +78,7 @@ class check_specialization_constants_multiple_for_type { sycl::get_kernel_bundle(ctx, {dev}, {kernelId}); if (!k_bundle.has_kernel(kernelId)) { + has_target_kernel = false; log.note( "kernel_bundle doesn't contain target kernel;" "multiple spec const for " + @@ -111,17 +113,20 @@ class check_specialization_constants_multiple_for_type { } }); } - if (!check_equal_values(ref1, result_vec[0].value) || - !check_equal_values(ref2, result_vec[1].value) || - !check_equal_values(ref3, result_vec[2].value) || - !check_equal_values(user_def_types::get_init_value(def_values[3]), - result_vec[3].value) || - !check_equal_values(user_def_types::get_init_value(def_values[4]), - result_vec[4].value) || - !check_equal_values(user_def_types::get_init_value(def_values[5]), - result_vec[5].value)) - FAIL(log, - "multiple spec const for " + type_name_string::get(type_name)); + if (has_target_kernel) { + // Check results only if target kernel is available + if (!check_equal_values(ref1, result_vec[0].value) || + !check_equal_values(ref2, result_vec[1].value) || + !check_equal_values(ref3, result_vec[2].value) || + !check_equal_values(user_def_types::get_init_value(def_values[3]), + result_vec[3].value) || + !check_equal_values(user_def_types::get_init_value(def_values[4]), + result_vec[4].value) || + !check_equal_values(user_def_types::get_init_value(def_values[5]), + result_vec[5].value)) + FAIL(log, + "multiple spec const for " + type_name_string::get(type_name)); + } } }; diff --git a/tests/spec_constants/spec_constants_same_name_inter_link.h b/tests/spec_constants/spec_constants_same_name_inter_link.h index 0f6b70376..e74bc9c3b 100644 --- a/tests/spec_constants/spec_constants_same_name_inter_link.h +++ b/tests/spec_constants/spec_constants_same_name_inter_link.h @@ -40,6 +40,7 @@ class check_specialization_constants_same_name_inter_link_for_type { auto queue = util::get_cts_object::queue(); const sycl::context ctx = queue.get_context(); const sycl::device dev = queue.get_device(); + bool has_target_kernel = true; if constexpr (via_kb::value) { if (!dev.has(sycl::aspect::online_compiler)) { @@ -67,6 +68,7 @@ class check_specialization_constants_same_name_inter_link_for_type { sycl::get_kernel_bundle(ctx, {dev}, {kernelId}); if (!k_bundle.has_kernel(kernelId)) { + has_target_kernel = false; log.note("kernel_bundle doesn't contain target kernel for " + type_name_string::get(type_name) + " (skipped)"); return; @@ -93,14 +95,17 @@ class check_specialization_constants_same_name_inter_link_for_type { } }); } - if (!check_equal_values(ref_def_value, def_value)) - FAIL(log, "Wrong linked spec const; (translation unit " + - std::to_string(TestConfig::tu) + ") for " + - type_name_string::get(type_name)); - if (!check_equal_values(ref, result)) - FAIL(log, "Wrong returned value; (translation unit " + - std::to_string(TestConfig::tu) + ") for " + - type_name_string::get(type_name)); + if (has_target_kernel) { + // Check results only if target kernel is available + if (!check_equal_values(ref_def_value, def_value)) + FAIL(log, "Wrong linked spec const; (translation unit " + + std::to_string(TestConfig::tu) + ") for " + + type_name_string::get(type_name)); + if (!check_equal_values(ref, result)) + FAIL(log, "Wrong returned value; (translation unit " + + std::to_string(TestConfig::tu) + ") for " + + type_name_string::get(type_name)); + } } } catch (...) { std::string message{"translation unit " + std::to_string(TestConfig::tu) + diff --git a/tests/spec_constants/spec_constants_same_name_stress.h b/tests/spec_constants/spec_constants_same_name_stress.h index fc03b0a2d..7cc020f56 100644 --- a/tests/spec_constants/spec_constants_same_name_stress.h +++ b/tests/spec_constants/spec_constants_same_name_stress.h @@ -47,6 +47,7 @@ class check_specialization_constants_same_name_stress_for_type { auto queue = util::get_cts_object::queue(); const sycl::context ctx = queue.get_context(); const sycl::device dev = queue.get_device(); + bool has_target_kernel = true; if constexpr (via_kb::value) { if (!dev.has(sycl::aspect::online_compiler)) { @@ -149,6 +150,7 @@ class check_specialization_constants_same_name_stress_for_type { auto k_bundle = sycl::get_kernel_bundle( ctx, {dev}, {kernelId}); if (!k_bundle.has_kernel(kernelId)) { + has_target_kernel = false; log.note("kernel_bundle doesn't contain target kernel for " + type_name_string::get(type_name) + " (skipped)"); return; @@ -160,14 +162,17 @@ class check_specialization_constants_same_name_stress_for_type { } }); } - for (int i = 0; i < size; ++i) { - if (!check_equal_values(def_values_arr[i], ref_def_values_arr[i])) { - FAIL(log, "Wrong default value for spec const defined in " + - get_hint(i) + " for type " + type_name); - } - if (!check_equal_values(result_arr[i], ref_arr[i])) { - FAIL(log, "Wrong result value for spec const defined in " + - get_hint(i) + "for type " + type_name); + if (has_target_kernel) { + // Check results only if target kernel is available + for (int i = 0; i < size; ++i) { + if (!check_equal_values(def_values_arr[i], ref_def_values_arr[i])) { + FAIL(log, "Wrong default value for spec const defined in " + + get_hint(i) + " for type " + type_name); + } + if (!check_equal_values(result_arr[i], ref_arr[i])) { + FAIL(log, "Wrong result value for spec const defined in " + + get_hint(i) + "for type " + type_name); + } } } } catch (...) { diff --git a/tests/std_classes/CMakeLists.txt b/tests/std_classes/CMakeLists.txt deleted file mode 100644 index 82f462065..000000000 --- a/tests/std_classes/CMakeLists.txt +++ /dev/null @@ -1,3 +0,0 @@ -file(GLOB test_cases_list *.cpp) - -add_cts_test(${test_cases_list}) diff --git a/tests/std_classes/std_classes.cpp b/tests/std_classes/std_classes.cpp deleted file mode 100644 index 35e171b5e..000000000 --- a/tests/std_classes/std_classes.cpp +++ /dev/null @@ -1,132 +0,0 @@ -/******************************************************************************* -// -// SYCL 2020 Conformance Test Suite -// -// Copyright (c) 2017-2022 Codeplay Software LTD. All Rights Reserved. -// Copyright (c) 2022 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -*******************************************************************************/ - -#include "../common/common.h" - -#define TEST_NAME std_classes - -namespace std_classes__ { -using namespace sycl_cts; - -/** check std::vector -*/ -template -using vectorClass = std::vector; - -/** check std::string -*/ -using stringClass = std::string; - -/** check std::function -*/ -template -using functionClass = std::function; - -/** check std::mutex -*/ -using mutexClass = std::mutex; - -/** check std::unique_ptr -*/ -template -using uniquePtrClass = std::unique_ptr; - -/** check std::shared_ptr -*/ -template -using sharedPtrClass = std::shared_ptr; - -/** check std::weak_ptr -*/ -template -using weakPtrClass = std::weak_ptr; - -/** check std::hash -*/ -template -using hashClass = std::hash; - -/** check std::exception_ptr -*/ -using exceptionPtrClass = std::exception_ptr; - -struct custom_deleter { - void operator()(int *p) const {}; -}; - -/** tests the availability of std classes -*/ -class TEST_NAME : public util::test_base { - public: - /** return information about this test - */ - void get_info(test_base::info &out) const override { - set_test_info(out, TOSTRING(TEST_NAME), TEST_FILE); - } - - /** execute this test - */ - void run(util::logger &log) override { - /* Try instantiating these classes */ - { - /** check std::vector - */ - std::vector vector; - - /** check std::string - */ - stringClass string; - - /** check std::function - */ - functionClass function; - - /** check std::mutex - */ - mutexClass mutex; - - /** check std::unique_ptr - */ - uniquePtrClass uniquePtr; - - /** check std::shared_ptr - */ - sharedPtrClass sharedPtr; - - /** check std::weak_ptr - */ - weakPtrClass weakPtr; - - /** check std::hash - */ - hashClass hash; - - /** check std::exception_ptr - */ - exceptionPtrClass exceptionPtr; - } - } -}; - -// register this test with the test_collection -util::test_proxy proxy; - -} /* namespace std_classes__ */ diff --git a/tests/vector_api/generate_vector_api.py b/tests/vector_api/generate_vector_api.py index 2920081c6..7f5802d37 100755 --- a/tests/vector_api/generate_vector_api.py +++ b/tests/vector_api/generate_vector_api.py @@ -34,6 +34,7 @@ vector_element_type_template = Template(""" CHECK(std::is_same_v::element_type, ${type}>); sycl::vec<${type}, ${size}> vec; + CHECK(std::is_same_v::value_type, ${type}>); CHECK(std::is_same_v())::element_type, ${type}>); """)