sycl::sub_group by-value semantics tests

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 <int dim, typename kernelT>
 struct invoke_nd_item {
  static constexpr int dimensions = dim;
@@ -33,26 +33,23 @@ struct invoke_nd_item {
  * @param kernelBody Kernel body to call
  */
  template <typename kernelBodyT>
- void operator()(sycl::handler &cgh,
- sycl::range<dim> numWorkItems,
- sycl::range<dim> workGroupSize,
- kernelBodyT kernelBody) {
-
- cgh.parallel_for<kernelT>(
- sycl::nd_range<dim>(numWorkItems, workGroupSize),
- [=](sycl::nd_item<dim> ndItem) {
- const size_t index = ndItem.get_global_linear_id();
-
- kernelBody(ndItem, index);
- });
+ void operator()(sycl::handler& cgh, sycl::range<dim> numWorkItems,
+ sycl::range<dim> workGroupSize, kernelBodyT kernelBody) {
+ cgh.parallel_for<kernelT>(sycl::nd_range<dim>(numWorkItems, workGroupSize),
+ [=](sycl::nd_item<dim> ndItem) {
+ const size_t index =
+ ndItem.get_global_linear_id();
+
+ kernelBody(ndItem, index);
+ });
  }
 };
 
 /**
  * @brief Functor to invoke kernels with group in use
  * @tparam dim Dimension to use
  * @tparam kernelT Type to use as the kernel name
-  */
+ */
 template <int dim, typename kernelT>
 struct invoke_group {
  static constexpr int dimensions = dim;
@@ -67,31 +64,49 @@ struct invoke_group {
  * @param kernelBody Kernel body to call
  */
  template <typename kernelBodyT>
- void operator()(sycl::handler &cgh,
- sycl::range<dim> numWorkItems,
- sycl::range<dim> workGroupSize,
- kernelBodyT kernelBody) {
+ void operator()(sycl::handler& cgh, sycl::range<dim> numWorkItems,
+ sycl::range<dim> workGroupSize, kernelBodyT kernelBody) {
  sycl::range<dim> numWorkGroups = numWorkItems / workGroupSize;
 
  cgh.parallel_for_work_group<kernelT>(
- numWorkGroups, workGroupSize,
- [=](sycl::group<dim> group) {
- const size_t index = group.get_linear_id();
+ numWorkGroups, workGroupSize, [=](sycl::group<dim> group) {
+ const size_t index = group.get_linear_id();
 
- kernelBody(group, index);
- });
+ kernelBody(group, index);
+ });
+ }
+};
+
+template <int dim, typename kernelT>
+struct invoke_sub_group {
+ static constexpr int dimensions = dim;
+ using parameterT = sycl::sub_group;
+
+ template <typename kernelBodyT>
+ void operator()(sycl::handler& cgh, sycl::range<dim> numWorkItems,
+ sycl::range<dim> workGroupSize, kernelBodyT kernelBody) {
+ sycl::range<dim> numWorkGroups = numWorkItems / workGroupSize;
+
+ cgh.parallel_for<kernelT>(
+ sycl::nd_range<dim>(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 <size_t numItems, class kernelInvokeT>
-std::array<typename kernelInvokeT::parameterT, numItems> store_instances()
-{
+std::array<typename kernelInvokeT::parameterT, numItems> store_instances() {
  constexpr auto numDims = kernelInvokeT::dimensions;
  using item_t = typename kernelInvokeT::parameterT;
  using item_array_t = std::array<item_t, numItems>;
@@ -104,25 +119,21 @@ std::array<typename kernelInvokeT::parameterT, numItems> store_instances()
  sycl_cts::util::get_cts_object::range<numDims>::get(numItems, 1, 1);
 
  {
- sycl::buffer<item_t> itemBuf(items.data(),
- sycl::range<1>(items.size()));
+ sycl::buffer<item_t> 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<sycl::access_mode::write>(cgh);
+ auto itemAcc = itemBuf.template get_access<sycl::access_mode::write>(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

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,

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<std::string, error_count> 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 <op_codes Code, typename ResultArray>
+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 <typename ResultArray>
+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<op_codes::ctor_copy>(
+ 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<op_codes::assign_copy>(
+ 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<op_codes::ctor_move>(
+ 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<op_codes::assign_move>(
+ 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<bool, 1> 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<sycl::access_mode::read_write>(cgh);
+ cgh.parallel_for<sub_group_semantics_kernel>(
+ 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<bool, 1> 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<sub_group_equality_kernel>([=] {
+ 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