From 2e3cc3daa203aab30d5213928949b46f81a583fb Mon Sep 17 00:00:00 2001 From: "Ipmagamabetov, AmirX" Date: Fri, 9 Jun 2023 12:56:27 +0300 Subject: [PATCH 1/4] kernel attributes tests --- tests/kernel/kernel_attributes.h | 53 ++++ tests/kernel/kernel_attributes_vec_hint.cpp | 210 ++++++++++++++ tests/kernel/kernel_attributes_wg_hint.cpp | 303 ++++++++++++++++++++ 3 files changed, 566 insertions(+) create mode 100644 tests/kernel/kernel_attributes.h create mode 100644 tests/kernel/kernel_attributes_vec_hint.cpp create mode 100644 tests/kernel/kernel_attributes_wg_hint.cpp diff --git a/tests/kernel/kernel_attributes.h b/tests/kernel/kernel_attributes.h new file mode 100644 index 000000000..5bf912bd7 --- /dev/null +++ b/tests/kernel/kernel_attributes.h @@ -0,0 +1,53 @@ +/******************************************************************************* +// +// 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" + +namespace kernel_attributes { + +template +struct kernel_functor_st; +template +struct kernel_functor; +template +struct kernel_functor_wg; + +template +struct kernel_separate_lambda_st; +template +struct kernel_separate_lambda; +template +struct kernel_separate_lambda_wg; + +template +struct kernel_lambda_st; +template +struct kernel_lambda; +template +struct kernel_lambda_wg; + +template +inline constexpr int expected_val() { + return 40 + dims; +} + +const sycl::range<1> range(1); + +} // namespace kernel_attributes diff --git a/tests/kernel/kernel_attributes_vec_hint.cpp b/tests/kernel/kernel_attributes_vec_hint.cpp new file mode 100644 index 000000000..6bc08ba2f --- /dev/null +++ b/tests/kernel/kernel_attributes_vec_hint.cpp @@ -0,0 +1,210 @@ +/******************************************************************************* +// +// 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 "kernel_attributes.h" +#include "../common/disabled_for_test_case.h" + +using namespace kernel_attributes; + +// FIXME: enable when [[sycl::vec_type_hint()]] is implemented +#ifndef SYCL_CTS_COMPILING_WITH_DPCPP + +#define RUN_TEST(K_NAME1, K_NAME2, K_NAME3, VEC_T, FUNC1, FUNC2, FUNC3) \ + \ + { \ + auto queue = sycl_cts::util::get_cts_object::queue(); \ + VEC_T vec_st; \ + VEC_T vec; \ + VEC_T vec_wg; \ + { \ + sycl::buffer buf_st(&vec_st, range); \ + sycl::buffer buf(&vec, range); \ + sycl::buffer buf_wg(&vec_wg, range); \ + \ + queue.submit([&](sycl::handler& cgh) { \ + auto acc = buf_st.template get_access(cgh); \ + cgh.single_task(FUNC1(acc)); \ + }); \ + queue.submit([&](sycl::handler& cgh) { \ + auto acc = buf.template get_access(cgh); \ + cgh.parallel_for(sycl::nd_range{range, range}, FUNC2(acc)); \ + }); \ + queue.submit([&](sycl::handler& cgh) { \ + auto acc = buf_wg.template get_access(cgh); \ + cgh.parallel_for_work_group(range, range, FUNC3(acc)); \ + }); \ + queue.wait_and_throw(); \ + } \ + verify(vec_st); \ + verify(vec); \ + verify(vec_wg); \ + } + +template +class functor { + public: + using vector_t = vec_t; + using acc_t = + sycl::accessor; + + functor(acc_t _acc) : acc(_acc) {} + [[sycl::vec_type_hint(vec_t)]] void operator()() const { + acc[sycl::id<1>()] = expected_val<1>(); + } + + [[sycl::vec_type_hint(vec_t)]] void operator()( + sycl::nd_item<1> nd_item) const { + acc[nd_item.get_local_id()] = expected_val<1>(); + } + + [[sycl::vec_type_hint(vec_t)]] void operator()(sycl::group<1> group) const { + acc[group.get_group_id()] = expected_val<1>(); + } + + private: + acc_t acc; +}; + +template +void verify(vec_t& vec) { + const auto size = vec_t::size(); + bool res = true; + for (int i = 0; i < size; i++) { + res &= vec[i] == expected_val<1>(); + } + INFO( + "Check that kernel is executed without any " + "exception and have expected result using sycl::vec<" + << typeid(typename vec_t::element_type).name() << ", " << size << ">"); + CHECK(res); +} + +template +void check_functor() { + using k_name1 = kernel_functor_st; + using k_name2 = kernel_functor; + using k_name3 = kernel_functor_wg; + + RUN_TEST(k_name1, k_name2, k_name3, typename functor::vector_t, functor, + functor, functor); +} + +template +const auto get_lambda_st(acc_t& acc) { + return [=]() [[sycl::vec_type_hint(vec_t)]] { + acc[sycl::id<1>()] = expected_val<1>(); + }; +} + +template +const auto get_lambda(acc_t& acc) { + return [=](auto nd_item) [[sycl::vec_type_hint(vec_t)]] { + acc[nd_item.get_local_id()] = expected_val<1>(); + }; +} + +template +const auto get_lambda_wg(acc_t& acc) { + return [=](auto group) [[sycl::vec_type_hint(vec_t)]] { + acc[group.get_group_id()] = expected_val<1>(); + }; +} + +template +void check_separate_lambda() { + using vec_t = sycl::vec; + using k_name1 = kernel_separate_lambda_st; + using k_name2 = kernel_separate_lambda<0, N, T>; + using k_name3 = kernel_separate_lambda_wg<0, N, T>; + + RUN_TEST(k_name1, k_name2, k_name3, vec_t, get_lambda_st, + get_lambda, get_lambda_wg); +} + +template +void check_lambda() { + using vec_t = sycl::vec; + + auto queue = sycl_cts::util::get_cts_object::queue(); + vec_t vec_st; + vec_t vec; + vec_t vec_wg; + { + sycl::buffer buf_st(&vec_st, range); + sycl::buffer buf(&vec, range); + sycl::buffer buf_wg(&vec_wg, range); + + queue.submit([&](sycl::handler& cgh) { + auto acc = buf_st.template get_access(cgh); + cgh.single_task>( + [=] { acc[sycl::id<1>()] = expected_val<1>(); }); + }); + queue.submit([&](sycl::handler& cgh) { + auto acc = buf.template get_access(cgh); + cgh.parallel_for>( + sycl::nd_range{range, range}, [=](auto nd_item) { + acc[nd_item.get_local_id()] = expected_val<1>(); + }); + }); + queue.submit([&](sycl::handler& cgh) { + auto acc = buf_wg.template get_access(cgh); + cgh.parallel_for_work_group>( + range, range, + [=](auto group) { acc[group.get_group_id()] = expected_val<1>(); }); + }); + queue.wait_and_throw(); + } + verify(vec_st); + verify(vec); + verify(vec_wg); +} + +template +void run_tests_for_size() { + using vec_t = sycl::vec; + + check_functor>(); + check_separate_lambda(); + check_lambda(); +} + +template +void run_tests_for_type() { + run_tests_for_size(); + run_tests_for_size(); + run_tests_for_size(); + run_tests_for_size(); + run_tests_for_size(); + run_tests_for_size(); +} + +DISABLED_FOR_TEST_CASE(hipSYCL, ComputeCpp, DPCPP) +("Behavior of kernel attribute vec_type_hint", "[kernel]") ({ +#if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS + run_tests_for_type(); + run_tests_for_type(); + +#else + SKIP("Tests for deprecated features are disabled."); +#endif // SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS +}) + +#endif // SYCL_CTS_COMPILING_WITH_DPCPP diff --git a/tests/kernel/kernel_attributes_wg_hint.cpp b/tests/kernel/kernel_attributes_wg_hint.cpp new file mode 100644 index 000000000..0c596cabc --- /dev/null +++ b/tests/kernel/kernel_attributes_wg_hint.cpp @@ -0,0 +1,303 @@ +/******************************************************************************* +// +// 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 "kernel_attributes.h" + +using namespace kernel_attributes; + +static constexpr int size = 4; +static auto dev = sycl_cts::util::get_cts_object::device(); +static auto max_wg_size = pow(size * 2, 3); +static bool max_wg_size_info = + dev.get_info() < max_wg_size + ? false + : true; + +template +class functor { + public: + static constexpr int dims = Dimensions; + static constexpr int size = Size; + + using acc_t = + sycl::accessor; + + functor(acc_t _acc) : acc(_acc) {} + + [[sycl::work_group_size_hint(size)]] void operator()( + sycl::nd_item<1> nd_item) const { + acc[nd_item.get_local_id()] = expected_val<1>(); + } + + [[sycl::work_group_size_hint(size, size)]] void operator()( + sycl::nd_item<2> nd_item) const { + acc[nd_item.get_local_id()] = expected_val<2>(); + } + + [[sycl::work_group_size_hint(size, size, size)]] void operator()( + sycl::nd_item<3> nd_item) const { + acc[nd_item.get_local_id()] = expected_val<3>(); + } + + [[sycl::work_group_size_hint(size)]] void operator()( + sycl::group<1> group) const { + acc[group.get_group_id()] = expected_val<1>(); + } + + [[sycl::work_group_size_hint(size, size)]] void operator()( + sycl::group<2> group) const { + acc[group.get_group_id()] = expected_val<2>(); + } + + [[sycl::work_group_size_hint(size, size, size)]] void operator()( + sycl::group<3> group) const { + acc[group.get_group_id()] = expected_val<3>(); + } + + private: + acc_t acc; +}; + +template +const auto get_lambda(acc_t& acc) { + if constexpr (dims == 1) { + return [=](auto nd_item) [[sycl::work_group_size_hint(size)]] { + acc[nd_item.get_local_id()] = expected_val(); + }; + } else if constexpr (dims == 2) { + return [=](auto nd_item) [[sycl::work_group_size_hint(size, size)]] { + acc[nd_item.get_local_id()] = expected_val(); + }; + } else if constexpr (dims == 3) { + return [=](auto nd_item) [[sycl::work_group_size_hint(size, size, size)]] { + acc[nd_item.get_local_id()] = expected_val(); + }; + } else { + return [=](auto nd_item) {}; + } +} + +template +const auto get_lambda_wg(acc_t& acc) { + if constexpr (dims == 1) { + return [=](auto group) [[sycl::work_group_size_hint(size)]] { + acc[group.get_group_id()] = expected_val(); + }; + } else if constexpr (dims == 2) { + return [=](auto group) [[sycl::work_group_size_hint(size, size)]] { + acc[group.get_group_id()] = expected_val(); + }; + } else if constexpr (dims == 3) { + return [=](auto group) [[sycl::work_group_size_hint(size, size, size)]] { + acc[group.get_group_id()] = expected_val(); + }; + } else { + return [=](auto group) {}; + } +} + +template +void verify(arr_t arr, const char* msg) { + INFO( + "Check that kernel is executed without any " + "exception and has expected result using " + << msg); + CHECK(std::all_of(arr.cbegin(), arr.cend(), + [](const int i) { return i == expected_val(); })); +} + +template +void check_functor_and_sep_lambda() { + auto queue = sycl_cts::util::get_cts_object::queue(); + + constexpr int dims = functor::dims; + constexpr int size = functor::size; + constexpr int buffer_size = (dims == 3) ? size * size * size + : dims == 2 ? size * size + : size; + + std::array data_functor; + std::array data_sep_lambda; + + std::array data_functor_wg; + std::array data_sep_lambda_wg; + + const auto range = + sycl_cts::util::get_cts_object::range::get(size, size, size); + const auto range_wg = + sycl_cts::util::get_cts_object::range::get(1, 1, 1); + { + sycl::buffer buf1(data_functor.data(), range); + sycl::buffer buf2(data_sep_lambda.data(), range); + + sycl::buffer buf1_wg(data_functor_wg.data(), range); + sycl::buffer buf2_wg(data_sep_lambda_wg.data(), range); + + // functor + queue.submit([&](sycl::handler& cgh) { + auto acc1 = sycl::accessor(buf1, cgh, sycl::write_only); + cgh.parallel_for>(sycl::nd_range{range, range}, + functor{acc1}); + }); + queue.submit([&](sycl::handler& cgh) { + auto acc1_wg = sycl::accessor(buf1_wg, cgh, sycl::write_only); + cgh.parallel_for_work_group>(range, range_wg, + functor{acc1_wg}); + }); + + // separate lambda + queue.submit([&](sycl::handler& cgh) { + auto acc2 = sycl::accessor(buf2, cgh, sycl::write_only); + cgh.parallel_for>( + sycl::nd_range{range, range}, get_lambda(acc2)); + }); + queue.submit([&](sycl::handler& cgh) { + auto acc2_wg = sycl::accessor(buf2_wg, cgh, sycl::write_only); + cgh.parallel_for_work_group>( + range, range_wg, get_lambda_wg(acc2_wg)); + }); + queue.wait_and_throw(); + } + + verify(data_functor, "functor with nd_item"); + verify(data_functor_wg, "functor with group"); + verify(data_sep_lambda, "separate lambda with nd_item"); + verify(data_sep_lambda_wg, "separate lambda with group"); +} + +template +void run_tests_for_lambda() { + auto queue = sycl_cts::util::get_cts_object::queue(); + + constexpr int buffer_size_1d = size; + constexpr int buffer_size_2d = size * size; + constexpr int buffer_size_3d = size * size * size; + + std::array data_1d; + std::array data_2d; + std::array data_3d; + + std::array data_wg_1d; + std::array data_wg_2d; + std::array data_wg_3d; + + sycl::range<1> range_1d(size); + sycl::range<2> range_2d(size, size); + sycl::range<3> range_3d(size, size, size); + { + sycl::buffer buf_1d(data_1d.data(), range_1d); + sycl::buffer buf_2d(data_2d.data(), range_2d); + sycl::buffer buf_3d(data_3d.data(), range_3d); + + sycl::buffer buf_wg_1d(data_wg_1d.data(), range_1d); + sycl::buffer buf_wg_2d(data_wg_2d.data(), range_2d); + sycl::buffer buf_wg_3d(data_wg_3d.data(), range_3d); + + // lambda submission call + queue.submit([&](sycl::handler& cgh) { + auto acc_1d = sycl::accessor(buf_1d, cgh, sycl::write_only); + cgh.parallel_for>( + sycl::nd_range<1>{range_1d, range_1d}, + [=](auto nd_item) [[sycl::work_group_size_hint(size)]] { + acc_1d[nd_item.get_local_id()] = expected_val<1>(); + }); + }); + queue.submit([&](sycl::handler& cgh) { + auto acc_2d = sycl::accessor(buf_2d, cgh, sycl::write_only); + cgh.parallel_for>( + sycl::nd_range<2>{range_2d, range_2d}, + [=](auto nd_item) [[sycl::work_group_size_hint(size, size)]] { + acc_2d[nd_item.get_local_id()] = expected_val<2>(); + }); + }); + queue.submit([&](sycl::handler& cgh) { + auto acc_wg_1d = sycl::accessor(buf_wg_1d, cgh, sycl::write_only); + cgh.parallel_for_work_group>( + range_1d, sycl::range<1>{1}, + [=](auto group) [[sycl::work_group_size_hint(size)]] { + acc_wg_1d[group.get_group_id()] = expected_val<1>(); + }); + }); + queue.submit([&](sycl::handler& cgh) { + auto acc_wg_2d = sycl::accessor(buf_wg_2d, cgh, sycl::write_only); + cgh.parallel_for_work_group>( + range_2d, sycl::range<2>{1, 1}, + [=](auto group) [[sycl::work_group_size_hint(size, size)]] { + acc_wg_2d[group.get_group_id()] = expected_val<2>(); + }); + }); + + if (max_wg_size_info) { + queue.submit([&](sycl::handler& cgh) { + auto acc_3d = sycl::accessor(buf_3d, cgh, sycl::write_only); + cgh.parallel_for>( + sycl::nd_range<3>{range_3d, range_3d}, + [=](auto nd_item) [[sycl::work_group_size_hint(size, size, size)]] { + acc_3d[nd_item.get_local_id()] = expected_val<3>(); + }); + }); + queue.submit([&](sycl::handler& cgh) { + auto acc_wg_3d = sycl::accessor(buf_wg_3d, cgh, sycl::write_only); + cgh.parallel_for_work_group>( + range_3d, sycl::range<3>{1, 1, 1}, + [=](auto group) [[sycl::work_group_size_hint(size, size, size)]] { + acc_wg_3d[group.get_group_id()] = expected_val<3>(); + }); + }); + } else { + WARN("Device does not support work group size " << max_wg_size); + } + queue.wait_and_throw(); + } + + verify<1>(data_1d, "lambda with nd_item 1 dim"); + verify<2>(data_2d, "lambda with nd_item 2 dims"); + + verify<1>(data_wg_1d, "lambda with group 1 dim"); + verify<2>(data_wg_2d, "lambda with group 2 dims"); + + if (max_wg_size_info) { + verify<3>(data_wg_3d, "lambda with group 3 dims"); + verify<3>(data_3d, "lambda with nd_item 3 dims"); + } +} + +template +void run_tests_for_dim() { + check_functor_and_sep_lambda>(); + check_functor_and_sep_lambda>(); + + if (max_wg_size_info) { + check_functor_and_sep_lambda>(); + } else { + WARN("Device does not support work group size " << max_wg_size); + } +} + +TEST_CASE("Behavior of kernel attribute work_group_size_hint", "[kernel]") { + run_tests_for_dim<1>(); + run_tests_for_dim<2>(); + run_tests_for_dim<3>(); + + run_tests_for_lambda(); + run_tests_for_lambda(); + run_tests_for_lambda(); +} From 876ca5d9a55fce255fb8a8753314a41dac788c7b Mon Sep 17 00:00:00 2001 From: "Ipmagamabetov, AmirX" Date: Fri, 9 Jun 2023 12:57:51 +0300 Subject: [PATCH 2/4] clang format --- tests/kernel/kernel_attributes_vec_hint.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/kernel/kernel_attributes_vec_hint.cpp b/tests/kernel/kernel_attributes_vec_hint.cpp index 6bc08ba2f..83ec71ae7 100644 --- a/tests/kernel/kernel_attributes_vec_hint.cpp +++ b/tests/kernel/kernel_attributes_vec_hint.cpp @@ -19,8 +19,8 @@ *******************************************************************************/ #include "../common/common.h" -#include "kernel_attributes.h" #include "../common/disabled_for_test_case.h" +#include "kernel_attributes.h" using namespace kernel_attributes; @@ -197,7 +197,7 @@ void run_tests_for_type() { } DISABLED_FOR_TEST_CASE(hipSYCL, ComputeCpp, DPCPP) -("Behavior of kernel attribute vec_type_hint", "[kernel]") ({ +("Behavior of kernel attribute vec_type_hint", "[kernel]")({ #if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS run_tests_for_type(); run_tests_for_type(); From b661d942568603908091d0effbb905ced1c08849 Mon Sep 17 00:00:00 2001 From: "Ipmagamabetov, AmirX" Date: Tue, 13 Jun 2023 11:48:34 +0300 Subject: [PATCH 3/4] Link to issue --- tests/kernel/kernel_attributes_vec_hint.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tests/kernel/kernel_attributes_vec_hint.cpp b/tests/kernel/kernel_attributes_vec_hint.cpp index 83ec71ae7..06d5eaba7 100644 --- a/tests/kernel/kernel_attributes_vec_hint.cpp +++ b/tests/kernel/kernel_attributes_vec_hint.cpp @@ -25,6 +25,7 @@ using namespace kernel_attributes; // FIXME: enable when [[sycl::vec_type_hint()]] is implemented +// https://github.com/intel/llvm/issues/9836 #ifndef SYCL_CTS_COMPILING_WITH_DPCPP #define RUN_TEST(K_NAME1, K_NAME2, K_NAME3, VEC_T, FUNC1, FUNC2, FUNC3) \ @@ -196,6 +197,7 @@ void run_tests_for_type() { run_tests_for_size(); } +// Enable when https://github.com/intel/llvm/issues/9836 is fixed DISABLED_FOR_TEST_CASE(hipSYCL, ComputeCpp, DPCPP) ("Behavior of kernel attribute vec_type_hint", "[kernel]")({ #if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS From 90f5f068d59c50d959214e12d6f8ccc47339bd7a Mon Sep 17 00:00:00 2001 From: Amir <122785142+AmirIpma@users.noreply.github.com> Date: Thu, 15 Jun 2023 18:44:43 +0300 Subject: [PATCH 4/4] apply comment MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Peter Žužek --- tests/kernel/kernel_attributes_vec_hint.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/kernel/kernel_attributes_vec_hint.cpp b/tests/kernel/kernel_attributes_vec_hint.cpp index 06d5eaba7..bad989527 100644 --- a/tests/kernel/kernel_attributes_vec_hint.cpp +++ b/tests/kernel/kernel_attributes_vec_hint.cpp @@ -29,7 +29,6 @@ using namespace kernel_attributes; #ifndef SYCL_CTS_COMPILING_WITH_DPCPP #define RUN_TEST(K_NAME1, K_NAME2, K_NAME3, VEC_T, FUNC1, FUNC2, FUNC3) \ - \ { \ auto queue = sycl_cts::util::get_cts_object::queue(); \ VEC_T vec_st; \