diff --git a/tests/common/invoke.h b/tests/common/invoke.h index a537820bc..454950866 100644 --- a/tests/common/invoke.h +++ b/tests/common/invoke.h @@ -18,7 +18,7 @@ namespace { * @brief Functor to invoke kernels with nd_item in use * @tparam dim Dimension to use * @tparam kernelT Type to use as the kernel name - */ + */ template struct invoke_nd_item { static constexpr int dimensions = dim; @@ -33,18 +33,15 @@ struct invoke_nd_item { * @param kernelBody Kernel body to call */ template - void operator()(sycl::handler &cgh, - sycl::range numWorkItems, - sycl::range workGroupSize, - kernelBodyT kernelBody) { - - cgh.parallel_for( - sycl::nd_range(numWorkItems, workGroupSize), - [=](sycl::nd_item ndItem) { - const size_t index = ndItem.get_global_linear_id(); - - kernelBody(ndItem, index); - }); + void operator()(sycl::handler& cgh, sycl::range numWorkItems, + sycl::range workGroupSize, kernelBodyT kernelBody) { + cgh.parallel_for(sycl::nd_range(numWorkItems, workGroupSize), + [=](sycl::nd_item ndItem) { + const size_t index = + ndItem.get_global_linear_id(); + + kernelBody(ndItem, index); + }); } }; @@ -52,7 +49,7 @@ struct invoke_nd_item { * @brief Functor to invoke kernels with group in use * @tparam dim Dimension to use * @tparam kernelT Type to use as the kernel name - */ + */ template struct invoke_group { static constexpr int dimensions = dim; @@ -67,31 +64,49 @@ struct invoke_group { * @param kernelBody Kernel body to call */ template - void operator()(sycl::handler &cgh, - sycl::range numWorkItems, - sycl::range workGroupSize, - kernelBodyT kernelBody) { + void operator()(sycl::handler& cgh, sycl::range numWorkItems, + sycl::range workGroupSize, kernelBodyT kernelBody) { sycl::range numWorkGroups = numWorkItems / workGroupSize; cgh.parallel_for_work_group( - numWorkGroups, workGroupSize, - [=](sycl::group group) { - const size_t index = group.get_linear_id(); + numWorkGroups, workGroupSize, [=](sycl::group group) { + const size_t index = group.get_linear_id(); - kernelBody(group, index); - }); + kernelBody(group, index); + }); + } +}; + +template +struct invoke_sub_group { + static constexpr int dimensions = dim; + using parameterT = sycl::sub_group; + + template + void operator()(sycl::handler& cgh, sycl::range numWorkItems, + sycl::range workGroupSize, kernelBodyT kernelBody) { + sycl::range numWorkGroups = numWorkItems / workGroupSize; + + cgh.parallel_for( + sycl::nd_range(numWorkItems, workGroupSize), + [=](sycl::nd_item<3> item) { + const size_t index = item.get_global_linear_id(); + sycl::sub_group sub_group = item.get_sub_group(); + + kernelBody(sub_group, index); + }); } }; /** - * @brief Generate and store the given number of nd_item/group/h_item instances + * @brief Generate and store the given number of nd_item/group/h_item/sub_group + * instances * @retval Array of instances * @tparam numItems Number of instances to store * @tparam kernelInvokeT Invocation functor to use */ template -std::array store_instances() -{ +std::array store_instances() { constexpr auto numDims = kernelInvokeT::dimensions; using item_t = typename kernelInvokeT::parameterT; using item_array_t = std::array; @@ -104,25 +119,21 @@ std::array store_instances() sycl_cts::util::get_cts_object::range::get(numItems, 1, 1); { - sycl::buffer itemBuf(items.data(), - sycl::range<1>(items.size())); + sycl::buffer itemBuf(items.data(), sycl::range<1>(items.size())); auto queue = sycl_cts::util::get_cts_object::queue(); queue.submit([&](sycl::handler& cgh) { - auto itemAcc = - itemBuf.template get_access(cgh); + auto itemAcc = itemBuf.template get_access(cgh); kernelInvokeT{}( cgh, itemRange, oneElemRange, - [=](item_t& item, const size_t index) { - itemAcc[index] = item; - }); + [=](item_t& item, const size_t index) { itemAcc[index] = item; }); }); queue.wait_and_throw(); } return items; } -} // namespace +} // namespace -#endif // __SYCLCTS_TESTS_COMMON_INVOKE_H +#endif // __SYCLCTS_TESTS_COMMON_INVOKE_H diff --git a/tests/common/semantics_by_value.h b/tests/common/semantics_by_value.h index 8d542a23c..ab84f1dd5 100644 --- a/tests/common/semantics_by_value.h +++ b/tests/common/semantics_by_value.h @@ -33,7 +33,7 @@ namespace common_by_value_semantics { * @brief Provides a safe index for checking an operation */ enum class current_check : size_t { - reflexivity_equal_self, + reflexivity_equal_self = 0, reflexivity_not_equal_self, equal_copy, equal_copy_symmetry, diff --git a/tests/sub_group/sub_group_semantics.cpp b/tests/sub_group/sub_group_semantics.cpp new file mode 100644 index 000000000..cb3080536 --- /dev/null +++ b/tests/sub_group/sub_group_semantics.cpp @@ -0,0 +1,146 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2023 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" +#include "../common/invoke.h" +#include "../common/once_per_unit.h" +#include "../common/semantics_by_value.h" + +namespace sub_group_semantics { + +enum class op_codes : size_t { + ctor_copy = 0, + ctor_move, + assign_copy, + assign_move, + code_count +}; + +constexpr size_t error_count = to_integral(op_codes::code_count); + +static constexpr size_t sizes[] = {16, 32, 64}; + +static const std::array error_strings{ + "sub_group with sub_group was not constructed correctly", + "sub_group with sub_group was not move constructed correctly", + "sub_group with sub_group was not copy assigned correctly", + "sub_group with sub_group was not move assigned correctly", +}; + +template +void set_success_operation(ResultArray& result, bool success) { + int index = to_integral(Code); + result[index] = success; +} + +std::string get_error_string(int code) { return error_strings[code]; } + +struct sub_group_semantics_kernel; +struct sub_group_equality_kernel; +struct setup_kernel; + +bool check_equality_by_id(const sycl::sub_group& actual, + sycl::id<1>* expected_ids) { + return actual.get_group_id() == expected_ids[0] && + actual.get_local_id() == expected_ids[1]; +} + +template +void check_by_value_semantics(sycl::sub_group& sub_group, ResultArray& result) { + sycl::id<1> expected_ids[] = {sub_group.get_group_id(), + sub_group.get_local_id()}; + // Check copy constructor + sycl::sub_group copied(sub_group); + set_success_operation( + result, check_equality_by_id(copied, expected_ids)); + + // Check copy assignment + sycl::sub_group copy_assigned(sub_group); + copy_assigned = sub_group; + set_success_operation( + result, check_equality_by_id(copy_assigned, expected_ids)); + + // Check move constructor; invalidates sub_group + sycl::sub_group moved(std::move(sub_group)); + set_success_operation( + result, check_equality_by_id(moved, expected_ids)); + + // Check move assignment + sycl::sub_group move_assigned(copy_assigned); + move_assigned = std::move(copy_assigned); + set_success_operation( + result, check_equality_by_id(move_assigned, expected_ids)); +} + +TEST_CASE("sub_group by-value semantics", "[sub_group]") { + bool result[error_count]; + std::fill(result, result + error_count, false); + { + sycl::buffer res_buf(result, sycl::range(error_count)); + + sycl::queue queue = once_per_unit::get_queue(); + const sycl::range<3> r{1, 1, 1}; + sycl::nd_range<3> nd_range(r, r); + queue + .submit([&](sycl::handler& cgh) { + auto res_acc = res_buf.get_access(cgh); + cgh.parallel_for( + nd_range, [=](sycl::nd_item<3> nd_item) { + sycl::sub_group sub_group = nd_item.get_sub_group(); + check_by_value_semantics(sub_group, res_acc); + }); + }) + .wait_and_throw(); + } + for (int i = 0; i < error_count; ++i) { + INFO(get_error_string(i)); + CHECK(result[i]); + } +} + +// FIXME: re-enable when == and != operators is implemented +#if !SYCL_CTS_COMPILING_WITH_HIPSYCL && !SYCL_CTS_COMPILING_WITH_COMPUTECPP +TEST_CASE("Check sycl::sub_group equality", "[sub_group]") { + size_t code_count = + to_integral(common_by_value_semantics::current_check::size); + bool result[code_count]; + std::fill(result, result + code_count, false); + auto items = store_instances<2, invoke_sub_group<3, setup_kernel>>(); + { + sycl::buffer res_buf(result, sycl::range(code_count)); + auto queue = once_per_unit::get_queue(); + queue + .submit([&](sycl::handler& cgh) { + auto res_acc = res_buf.get_access(cgh); + cgh.single_task([=] { + common_by_value_semantics::check_equality(items[0], items[1], + res_acc); + }); + }) + .wait_and_throw(); + } + for (int i = 0; i < code_count; ++i) { + INFO(common_by_value_semantics::get_error_string(i)); + CHECK(result[i]); + } +} +#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL && + // !SYCL_CTS_COMPILING_WITH_COMPUTECPP +} // namespace sub_group_semantics