From 1ab967f093696964386f38bfa0f5c5c771e656d6 Mon Sep 17 00:00:00 2001 From: "Ipmagamabetov, AmirX" Date: Wed, 28 Jun 2023 14:38:08 +0300 Subject: [PATCH 1/4] sub_group by-value semantics --- tests/common/invoke.h | 25 ++++- tests/common/semantics_by_value.h | 2 +- tests/sub_group/sub_group_semantics.cpp | 143 ++++++++++++++++++++++++ 3 files changed, 168 insertions(+), 2 deletions(-) create mode 100644 tests/sub_group/sub_group_semantics.cpp diff --git a/tests/common/invoke.h b/tests/common/invoke.h index a537820bc..de1ac68dd 100644 --- a/tests/common/invoke.h +++ b/tests/common/invoke.h @@ -83,8 +83,31 @@ struct invoke_group { } }; +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 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..c23b2e389 --- /dev/null +++ b/tests/sub_group/sub_group_semantics.cpp @@ -0,0 +1,143 @@ +/******************************************************************************* +// +// 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]); + } +} + +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]); + } +} + +} // namespace sub_group_semantics From ad634b65aa150b5f4f79c1ebe4ed87f0721cf618 Mon Sep 17 00:00:00 2001 From: "Ipmagamabetov, AmirX" Date: Wed, 28 Jun 2023 17:10:27 +0300 Subject: [PATCH 2/4] apply new clang format --- tests/common/invoke.h | 70 ++++++++++++++++++------------------------- 1 file changed, 29 insertions(+), 41 deletions(-) diff --git a/tests/common/invoke.h b/tests/common/invoke.h index de1ac68dd..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,19 +64,16 @@ 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); + }); } }; @@ -88,11 +82,9 @@ 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) { + template + void operator()(sycl::handler& cgh, sycl::range numWorkItems, + sycl::range workGroupSize, kernelBodyT kernelBody) { sycl::range numWorkGroups = numWorkItems / workGroupSize; cgh.parallel_for( @@ -102,19 +94,19 @@ template 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/sub_group 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; @@ -127,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 From 71b8306a5534a5d8b28dbfc0f20fac43c252408b Mon Sep 17 00:00:00 2001 From: "Ipmagamabetov, AmirX" Date: Wed, 28 Jun 2023 17:21:56 +0300 Subject: [PATCH 3/4] disable for computecpp and hipsycl --- tests/sub_group/sub_group_semantics.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/tests/sub_group/sub_group_semantics.cpp b/tests/sub_group/sub_group_semantics.cpp index c23b2e389..cb3080536 100644 --- a/tests/sub_group/sub_group_semantics.cpp +++ b/tests/sub_group/sub_group_semantics.cpp @@ -115,6 +115,8 @@ TEST_CASE("sub_group by-value semantics", "[sub_group]") { } } +// 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); @@ -139,5 +141,6 @@ TEST_CASE("Check sycl::sub_group equality", "[sub_group]") { CHECK(result[i]); } } - +#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL && + // !SYCL_CTS_COMPILING_WITH_COMPUTECPP } // namespace sub_group_semantics From de5e9343978f7e906241755311badaccd4d290f6 Mon Sep 17 00:00:00 2001 From: Philip Salzmann Date: Mon, 17 Jul 2023 17:19:37 +0200 Subject: [PATCH 4/4] CI: Set ccache directory explicitly ccache uses various rules to determine the cache directory, which may result in different CI containers having different paths. Specify as absolute path instead of using "~", as that seems to be interpreted differently by cache action and ccache. --- .github/workflows/cts_ci.yml | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/.github/workflows/cts_ci.yml b/.github/workflows/cts_ci.yml index a2359d0b8..71accf0b6 100644 --- a/.github/workflows/cts_ci.yml +++ b/.github/workflows/cts_ci.yml @@ -143,24 +143,28 @@ jobs: - name: Set up ccache uses: actions/cache@v3 with: - path: ~/.ccache + path: ${{ env.container-workspace }}/.ccache key: ${{ matrix.sycl-impl }}-ccache-${{ github.sha }} restore-keys: | ${{ matrix.sycl-impl }}-ccache- # Use ccache's "depend mode" to work around DPC++ issue (see https://github.com/intel/llvm/issues/5260) # This requires compilation with -MD, which is enabled because we use the Ninja generator # Using this mode should not have any practical disadvantages + - name: Set ccache environment variables + run: | + echo "CCACHE_DEPEND=1" >> "$GITHUB_ENV" + echo "CCACHE_DIR=${{ env.container-workspace }}/.ccache" >> "$GITHUB_ENV" - name: Build 'oclmath' working-directory: ${{ env.container-workspace }}/build - run: CCACHE_DEPEND=1 cmake --build . --target oclmath + run: cmake --build . --target oclmath - name: Build 'util' working-directory: ${{ env.container-workspace }}/build - run: CCACHE_DEPEND=1 cmake --build . --target util + run: cmake --build . --target util - name: Build all supported test categories working-directory: ${{ env.container-workspace }}/build run: | TS_BEFORE=$(date +%s) - CCACHE_DEPEND=1 cmake --build . --target test_all --parallel ${{ env.parallel-build-jobs }} + cmake --build . --target test_all --parallel ${{ env.parallel-build-jobs }} TS_AFTER=$(date +%s) ELAPSED=$(($TS_AFTER - $TS_BEFORE)) sort --numeric-sort --reverse --output=build_times.log build_times.log