From fbb2ca746a581e02848da27a978d4f85acad564d Mon Sep 17 00:00:00 2001 From: Jie Zhang Date: Mon, 9 Oct 2023 13:49:52 +0800 Subject: [PATCH 01/12] Split group_reduce and group_scan cases by types --- tests/group_functions/CMakeLists.txt | 34 ++++++ ...t_reduce.cpp => group_joint_reduce.cpp.in} | 33 +++--- .../group_joint_reduce_fp16.cpp | 84 --------------- .../group_joint_reduce_fp64.cpp | 86 --------------- .../group_joint_reduce_fp64_fp16.cpp | 53 --------- ...joint_scan.cpp => group_joint_scan.cpp.in} | 20 ++-- .../group_functions/group_joint_scan_fp16.cpp | 91 ---------------- .../group_functions/group_joint_scan_fp64.cpp | 91 ---------------- .../group_joint_scan_fp64_fp16.cpp | 102 ------------------ tests/group_functions/group_reduce.h | 56 ++++++++++ ...oup.cpp => group_reduce_over_group.cpp.in} | 31 +++--- .../group_reduce_over_group_fp16.cpp | 62 ----------- .../group_reduce_over_group_fp64.cpp | 62 ----------- .../group_reduce_over_group_fp64_fp16.cpp | 53 --------- tests/group_functions/group_scan.h | 84 +++++++++++++++ ...group.cpp => group_scan_over_group.cpp.in} | 14 ++- .../group_scan_over_group_fp16.cpp | 58 ---------- .../group_scan_over_group_fp64.cpp | 58 ---------- .../group_scan_over_group_fp64_fp16.cpp | 58 ---------- 19 files changed, 229 insertions(+), 901 deletions(-) rename tests/group_functions/{group_joint_reduce.cpp => group_joint_reduce.cpp.in} (57%) delete mode 100644 tests/group_functions/group_joint_reduce_fp16.cpp delete mode 100644 tests/group_functions/group_joint_reduce_fp64.cpp delete mode 100644 tests/group_functions/group_joint_reduce_fp64_fp16.cpp rename tests/group_functions/{group_joint_scan.cpp => group_joint_scan.cpp.in} (80%) delete mode 100644 tests/group_functions/group_joint_scan_fp16.cpp delete mode 100644 tests/group_functions/group_joint_scan_fp64.cpp delete mode 100644 tests/group_functions/group_joint_scan_fp64_fp16.cpp rename tests/group_functions/{group_reduce_over_group.cpp => group_reduce_over_group.cpp.in} (62%) delete mode 100644 tests/group_functions/group_reduce_over_group_fp16.cpp delete mode 100644 tests/group_functions/group_reduce_over_group_fp64.cpp delete mode 100644 tests/group_functions/group_reduce_over_group_fp64_fp16.cpp rename tests/group_functions/{group_scan_over_group.cpp => group_scan_over_group.cpp.in} (76%) delete mode 100644 tests/group_functions/group_scan_over_group_fp16.cpp delete mode 100644 tests/group_functions/group_scan_over_group_fp64.cpp delete mode 100644 tests/group_functions/group_scan_over_group_fp64_fp16.cpp diff --git a/tests/group_functions/CMakeLists.txt b/tests/group_functions/CMakeLists.txt index 82f462065..2f30733f9 100644 --- a/tests/group_functions/CMakeLists.txt +++ b/tests/group_functions/CMakeLists.txt @@ -1,3 +1,37 @@ +function(configure_test_case) + cmake_parse_arguments(CTS + "" "TYPE;IN_FILENAME;OUT_FILENAME;TEST_LIST" "" ${ARGN}) + set(CTS_TYPE_NAME ${CTS_TYPE}) + configure_file(${CTS_IN_FILENAME} ${CTS_OUT_FILENAME}) + list(APPEND ${CTS_TEST_LIST} "${CMAKE_CURRENT_BINARY_DIR}/${CTS_OUT_FILENAME}") + set(${CTS_TEST_LIST} ${${CTS_TEST_LIST}} PARENT_SCOPE) +endfunction() + +list(APPEND TEMPLATE_LIST + "group_joint_scan" + "group_scan_over_group" + "group_joint_reduce" + "group_reduce_over_group" +) +set(TYPE_LIST "") +get_std_type(TYPE_LIST) + file(GLOB test_cases_list *.cpp) +foreach(TEMP IN LISTS TEMPLATE_LIST) + foreach(TY IN LISTS TYPE_LIST) + if(TY STREQUAL "bool") + continue() + endif() + set(OUT_FILE "${TEMP}_${TY}.cpp") + STRING(REGEX REPLACE ":" "_" OUT_FILE ${OUT_FILE}) + STRING(REGEX REPLACE " " "_" OUT_FILE ${OUT_FILE}) + configure_test_case( + TYPE "${TY}" + IN_FILENAME "${TEMP}.cpp.in" + OUT_FILENAME ${OUT_FILE} + TEST_LIST test_cases_list) + endforeach() +endforeach() + add_cts_test(${test_cases_list}) diff --git a/tests/group_functions/group_joint_reduce.cpp b/tests/group_functions/group_joint_reduce.cpp.in similarity index 57% rename from tests/group_functions/group_joint_reduce.cpp rename to tests/group_functions/group_joint_reduce.cpp.in index 397decf13..fde841e69 100644 --- a/tests/group_functions/group_joint_reduce.cpp +++ b/tests/group_functions/group_joint_reduce.cpp.in @@ -20,32 +20,31 @@ #include "group_reduce.h" +// clang-format off +#cmakedefine CTS_TYPE @CTS_TYPE@ +#cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") +// clang-format on using ReduceTypes = Types; -// 2-dim Cartesian product of type lists -using prod2 = product::type; - // hipSYCL has no implementation over sub-groups -TEMPLATE_LIST_TEST_CASE("Group and sub-group joint reduce functions", - "[group_func][type_list][dim]", ReduceTypes) { +TEST_CASE(CTS_TYPE_NAME + " group and sub-group joint reduce functions", + "[group_func][type_list][dim]") { auto queue = once_per_unit::get_queue(); - // Get binary operators from TestType - const auto Operators = get_op_types(); - const auto Type = unnamed_type_pack(); - for_all_combinations(Dims, Type, Operators, queue); + const auto Operators = get_op_types(); + const auto RetType = unnamed_type_pack(); + for_all_combinations(Dims, RetType, Operators, + queue); } // hipSYCL has problems with 16-bit types for Ptr -TEMPLATE_LIST_TEST_CASE("Group and sub-group joint reduce functions with init", - "[group_func][type_list][dim]", prod2) { +TEMPLATE_LIST_TEST_CASE( + CTS_TYPE_NAME + " group and sub-group joint reduce functions with init", + "[group_func][type_list][dim]", ReduceTypes) { auto queue = once_per_unit::get_queue(); - using T = std::tuple_element_t<0, TestType>; - using U = std::tuple_element_t<1, TestType>; - // Get binary operators from T - const auto Operators = get_op_types(); - const auto RetType = unnamed_type_pack(); - const auto ReducedType = unnamed_type_pack(); + const auto Operators = get_op_types(); + const auto RetType = unnamed_type_pack(); + const auto ReducedType = unnamed_type_pack(); // check all work group dimensions for_all_combinations( Dims, RetType, ReducedType, Operators, queue); diff --git a/tests/group_functions/group_joint_reduce_fp16.cpp b/tests/group_functions/group_joint_reduce_fp16.cpp deleted file mode 100644 index 2f6725823..000000000 --- a/tests/group_functions/group_joint_reduce_fp16.cpp +++ /dev/null @@ -1,84 +0,0 @@ -/******************************************************************************* -// -// 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 "group_reduce.h" - -using ReduceTypes = Types; - -using HalfExtendedTypes = concatenation::type; -// 2-dim Cartesian product of type lists -using prod2 = product::type; - -// hipSYCL has no implementation over sub-groups -TEST_CASE("Group and sub-group joint reduce functions", - "[group_func][fp16][dim]") { - auto queue = once_per_unit::get_queue(); - // FIXME: hipSYCL omission -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - WARN( - "hipSYCL has no implementation of " - "std::iterator_traits::value_type joint_reduce(sub_group g, " - "Ptr first, Ptr last, BinaryOperation binary_op) over sub-groups. " - "Skipping the test case."); -#endif - - if (queue.get_device().has(sycl::aspect::fp16)) { - // Get binary operators from TestType - const auto Operators = get_op_types(); - const auto Type = unnamed_type_pack(); - for_all_combinations(Dims, Type, Operators, - queue); - } else { - WARN("Device does not support half precision floating point operations."); - } -} - -TEMPLATE_LIST_TEST_CASE("Group and sub-group joint reduce functions with init", - "[group_func][type_list][fp16][dim]", prod2) { - auto queue = once_per_unit::get_queue(); - using T = std::tuple_element_t<0, TestType>; - using U = std::tuple_element_t<1, TestType>; - - // check types to only print warning once - if constexpr (std::is_same_v && std::is_same_v) { - // FIXME: hipSYCL omission -#if SYCL_CTS_COMPILING_WITH_HIPSYCL - WARN( - "hipSYCL has no implementation of T joint_reduce(sub_group g, Ptr " - "first, Ptr last, T init, " - "BinaryOperation binary_op) over sub-groups. Skipping the test case."); -#endif - } - - if (queue.get_device().has(sycl::aspect::fp16)) { - if constexpr (std::is_same_v || - std::is_same_v) { - // Get binary operators from T - const auto Operators = get_op_types(); - const auto RetType = unnamed_type_pack(); - const auto ReducedType = unnamed_type_pack(); - // check all work group dimensions - for_all_combinations( - Dims, RetType, ReducedType, Operators, queue); - } - } else { - WARN("Device does not support half precision floating point operations."); - } -} diff --git a/tests/group_functions/group_joint_reduce_fp64.cpp b/tests/group_functions/group_joint_reduce_fp64.cpp deleted file mode 100644 index 2eecebe29..000000000 --- a/tests/group_functions/group_joint_reduce_fp64.cpp +++ /dev/null @@ -1,86 +0,0 @@ -/******************************************************************************* -// -// 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 "group_reduce.h" - -using ReduceTypes = Types; - -using DoubleExtendedTypes = concatenation::type; -// 2-dim Cartesian product of type lists -using prod2 = - product::type; - -// hipSYCL has no implementation over sub-groups -TEST_CASE("Group and sub-group joint reduce functions", - "[group_func][fp64][dim]") { - auto queue = once_per_unit::get_queue(); - // FIXME: hipSYCL omission -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - WARN( - "hipSYCL has no implementation of " - "std::iterator_traits::value_type joint_reduce(sub_group g, " - "Ptr first, Ptr last, BinaryOperation binary_op) over sub-groups. " - "Skipping the test case."); -#endif - - if (queue.get_device().has(sycl::aspect::fp64)) { - // Get binary operators from TestType - const auto Operators = get_op_types(); - const auto Type = unnamed_type_pack(); - for_all_combinations(Dims, Type, Operators, - queue); - } else { - WARN("Device does not support double precision floating point operations."); - } -} - -TEMPLATE_LIST_TEST_CASE("Group and sub-group joint reduce functions with init", - "[group_func][type_list][fp64][dim]", prod2) { - auto queue = once_per_unit::get_queue(); - using T = std::tuple_element_t<0, TestType>; - using U = std::tuple_element_t<1, TestType>; - - // check types to only print warning once - if constexpr (std::is_same_v && std::is_same_v) { - // FIXME: hipSYCL omission -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - WARN( - "hipSYCL has no implementation of T joint_reduce(sub_group g, Ptr " - "first, Ptr last, T init, " - "BinaryOperation binary_op) over sub-groups. Skipping the test case."); -#endif - } - - if (queue.get_device().has(sycl::aspect::fp64)) { - if constexpr (std::is_same_v || std::is_same_v) { - // Get binary operators from T - const auto Operators = get_op_types(); - const auto RetType = unnamed_type_pack(); - const auto ReducedType = unnamed_type_pack(); - // check all work group dimensions - for_all_combinations( - Dims, RetType, ReducedType, Operators, queue); - } - } else { - WARN( - "Device does not support double precision floating point " - "operations."); - } -} diff --git a/tests/group_functions/group_joint_reduce_fp64_fp16.cpp b/tests/group_functions/group_joint_reduce_fp64_fp16.cpp deleted file mode 100644 index e96a64e87..000000000 --- a/tests/group_functions/group_joint_reduce_fp64_fp16.cpp +++ /dev/null @@ -1,53 +0,0 @@ -/******************************************************************************* -// -// 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. -// -*******************************************************************************/ - -// need also double tests enabled -#ifdef SYCL_CTS_ENABLE_DOUBLE_TESTS - -#include "group_reduce.h" - -using ReduceTypes = unnamed_type_pack; - -// 2-dim Cartesian product of type lists -using prod2 = product::type; - -TEMPLATE_LIST_TEST_CASE("Group and sub-group joint reduce functions with init", - "[group_func][fp16][fp64][dim]", prod2) { - auto queue = once_per_unit::get_queue(); - using T = std::tuple_element_t<0, TestType>; - using U = std::tuple_element_t<1, TestType>; - - if (queue.get_device().has(sycl::aspect::fp16) && - queue.get_device().has(sycl::aspect::fp64)) { - // Get binary operators from T - const auto Operators = get_op_types(); - const auto RetType = unnamed_type_pack(); - const auto ReducedType = unnamed_type_pack(); - // check all work group dimensions - for_all_combinations( - Dims, RetType, ReducedType, Operators, queue); - } else { - WARN( - "Device does not support half and double precision floating point " - "operations simultaneously."); - } -} - -#endif diff --git a/tests/group_functions/group_joint_scan.cpp b/tests/group_functions/group_joint_scan.cpp.in similarity index 80% rename from tests/group_functions/group_joint_scan.cpp rename to tests/group_functions/group_joint_scan.cpp.in index 5d1740f1b..47c92f3ca 100644 --- a/tests/group_functions/group_joint_scan.cpp +++ b/tests/group_functions/group_joint_scan.cpp.in @@ -22,13 +22,19 @@ #if !SYCL_CTS_COMPILING_WITH_HIPSYCL #include "group_scan.h" +// clang-format off +#cmakedefine CTS_TYPE @CTS_TYPE@ +#cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") +// clang-format on +using TestType = unnamed_type_pack; using ScanTypes = Types; static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); #endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL // FIXME: known_identity is not impemented yet for hipSYCL. DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group joint scan functions", "[group_func][type_list][dim]")({ +(CTS_TYPE_NAME + " group and sub-group joint scan functions", + "[group_func][type_list][dim]")({ #if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) WARN( "hipSYCL cannot handle cases of different types for InPtr and OutPtr. " @@ -42,17 +48,17 @@ DISABLED_FOR_TEST_CASE(hipSYCL) auto queue = once_per_unit::get_queue(); // FIXME: hipSYCL cannot handle cases of different types #if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - for_all_combinations(Dims, ScanTypes{}, + for_all_combinations(Dims, TestType{}, queue); #else - for_all_combinations(Dims, ScanTypes{}, ScanTypes{}, + for_all_combinations(Dims, TestType{}, ScanTypes{}, queue); #endif }); // FIXME: known_identity is not impemented yet for hipSYCL. DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group joint scan functions with init", +(CTS_TYPE_NAME + " group and sub-group joint scan functions with init", "[group_func][type_list][dim]")({ #if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) WARN( @@ -67,10 +73,10 @@ DISABLED_FOR_TEST_CASE(hipSYCL) auto queue = once_per_unit::get_queue(); // FIXME: hipSYCL cannot handle cases of different types #if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - for_all_combinations( - Dims, ScanTypes{}, queue); + for_all_combinations(Dims, TestType{}, + queue); #else for_all_combinations( - Dims, ScanTypes{}, ScanTypes{}, ScanTypes{}, queue); + Dims, TestType{}, ScanTypes{}, ScanTypes{}, queue); #endif }); diff --git a/tests/group_functions/group_joint_scan_fp16.cpp b/tests/group_functions/group_joint_scan_fp16.cpp deleted file mode 100644 index 81bfef660..000000000 --- a/tests/group_functions/group_joint_scan_fp16.cpp +++ /dev/null @@ -1,91 +0,0 @@ -/******************************************************************************* -// -// 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/disabled_for_test_case.h" -#include "type_coverage.h" -#if !SYCL_CTS_COMPILING_WITH_HIPSYCL -#include "group_scan.h" - -using ScanTypes = Types; - -using HalfType = unnamed_type_pack; -using HalfExtendedTypes = concatenation::type; - -static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); -#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL -// FIXME: known_identity is not impemented yet for hipSYCL. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group joint scan functions", - "[group_func][type_list][fp16][dim]")({ -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - WARN( - "hipSYCL cannot handle cases of different types for InPtr and OutPtr. " - "Skipping such test cases."); - WARN( - "hipSYCL joint_exclusive_scan and joint_inclusive_scan cannot process " - "over several sub-groups simultaneously. Using one sub-group only."); - WARN("hipSYCL does not support sycl::known_identity_v yet."); -#endif - - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp16)) { - // FIXME: hipSYCL cannot handle cases of different types -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - for_all_combinations(Dims, HalfType{}, - queue); -#else - for_all_combinations_with( - Dims, HalfExtendedTypes{}, HalfExtendedTypes{}, queue); -#endif - } else { - WARN("Device does not support half precision floating point operations."); - } -}); - -// FIXME: known_identity is not impemented yet for hipSYCL. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group joint scan functions with init", - "[group_func][type_list][fp16][dim]")({ -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - WARN( - "hipSYCL cannot handle cases of different types for T, *InPtr and " - "*OutPtr. Skipping such test cases."); - WARN( - "hipSYCL joint_exclusive_scan and joint_inclusive_scan with init values " - "cannot process over several sub-groups simultaneously. Using one " - "sub-group only."); -#endif - - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp16)) { - // FIXME: hipSYCL cannot handle cases of different types -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - for_all_combinations( - Dims, HalfType{}, queue); -#else - for_all_combinations_with( - Dims, HalfExtendedTypes{}, HalfExtendedTypes{}, HalfExtendedTypes{}, - queue); -#endif - } else { - WARN("Device does not support half precision floating point operations."); - } -}); diff --git a/tests/group_functions/group_joint_scan_fp64.cpp b/tests/group_functions/group_joint_scan_fp64.cpp deleted file mode 100644 index 71c79fdd8..000000000 --- a/tests/group_functions/group_joint_scan_fp64.cpp +++ /dev/null @@ -1,91 +0,0 @@ -/******************************************************************************* -// -// 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/disabled_for_test_case.h" -#include "type_coverage.h" -#if !SYCL_CTS_COMPILING_WITH_HIPSYCL -#include "group_scan.h" - -using ScanTypes = Types; - -using DoubleType = unnamed_type_pack; -using DoubleExtendedTypes = concatenation::type; - -static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); -#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL -// FIXME: known_identity is not impemented yet for hipSYCL. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group joint scan functions", - "[group_func][type_list][fp64][dim]")({ -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - WARN( - "hipSYCL cannot handle cases of different types for InPtr and OutPtr. " - "Skipping such test cases."); - WARN( - "hipSYCL joint_exclusive_scan and joint_inclusive_scan cannot process " - "over several sub-groups simultaneously. Using one sub-group only."); - WARN("hipSYCL does not support sycl::known_identity_v yet."); -#endif - - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp64)) { - // FIXME: hipSYCL cannot handle cases of different types -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - for_all_combinations(Dims, DoubleType{}, - queue); -#else - for_all_combinations_with( - Dims, DoubleExtendedTypes{}, DoubleExtendedTypes{}, queue); -#endif - } else { - WARN("Device does not support double precision floating point operations."); - } -}); - -// FIXME: known_identity is not impemented yet for hipSYCL. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group joint scan functions with init", - "[group_func][type_list][fp64][dim]")({ -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - WARN( - "hipSYCL cannot handle cases of different types for T, *InPtr and " - "*OutPtr. Skipping such test cases."); - WARN( - "hipSYCL joint_exclusive_scan and joint_inclusive_scan with init values " - "cannot process over several sub-groups simultaneously. Using one " - "sub-group only."); -#endif - - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp64)) { - // FIXME: hipSYCL cannot handle cases of different types -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - for_all_combinations( - Dims, DoubleType{}, queue); -#else - for_all_combinations_with( - Dims, DoubleExtendedTypes{}, DoubleExtendedTypes{}, - DoubleExtendedTypes{}, queue); -#endif - } else { - WARN("Device does not support double precision floating point operations."); - } -}); diff --git a/tests/group_functions/group_joint_scan_fp64_fp16.cpp b/tests/group_functions/group_joint_scan_fp64_fp16.cpp deleted file mode 100644 index 4996ef5a8..000000000 --- a/tests/group_functions/group_joint_scan_fp64_fp16.cpp +++ /dev/null @@ -1,102 +0,0 @@ -/******************************************************************************* -// -// 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. -// -*******************************************************************************/ - -// need also double tests enabled -#ifdef SYCL_CTS_ENABLE_DOUBLE_TESTS - -#include "../common/common.h" -#include "../common/disabled_for_test_case.h" -#include "type_coverage.h" -#if !SYCL_CTS_COMPILING_WITH_HIPSYCL -#include "group_scan.h" - -using ScanTypes = Types; - -using DoubleHalfTypes = unnamed_type_pack; -using DoubleHalfExtendedTypes = concatenation::type; - -static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); -#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL -// FIXME: known_identity is not impemented yet for hipSYCL. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group joint scan functions", "[group_func][fp16][fp64][dim]")({ -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - WARN( - "hipSYCL cannot handle cases of different types for InPtr and OutPtr. " - "Skipping the test."); - WARN( - "hipSYCL joint_exclusive_scan and joint_inclusive_scan cannot process " - "over several sub-groups simultaneously. Using one sub-group only."); - WARN("hipSYCL does not support sycl::known_identity_v yet."); -#endif - - // FIXME: hipSYCL cannot handle cases of different types -#if SYCL_CTS_COMPILING_WITH_HIPSYCL - return; -#else - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp16) && - queue.get_device().has(sycl::aspect::fp64)) { - // here DoubleHalfTypes can be used as only half+double combinations are - // untested - for_all_combinations_with_two( - Dims, DoubleHalfTypes{}, DoubleHalfTypes{}, queue); - } else { - WARN( - "Device does not support half and double precision floating point " - "operations simultaneously."); - } -#endif -}); - -// FIXME: known_identity is not impemented yet for hipSYCL. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group joint scan functions with init", - "[group_func][type_list][fp16][fp64][dim]")({ -#if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) - WARN( - "hipSYCL cannot handle cases of different types for T, *InPtr and " - "*OutPtr. Skipping the test."); - WARN( - "hipSYCL joint_exclusive_scan and joint_inclusive_scan with init values " - "cannot process over several sub-groups simultaneously. Using one " - "sub-group only."); -#endif - - // FIXME: hipSYCL cannot handle cases of different types -#if SYCL_CTS_COMPILING_WITH_HIPSYCL - return; -#else - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp16) && - queue.get_device().has(sycl::aspect::fp64)) { - for_all_combinations_with_two(Dims, DoubleHalfExtendedTypes{}, - DoubleHalfExtendedTypes{}, - DoubleHalfExtendedTypes{}, queue); - } else { - WARN( - "Device does not support half and double precision floating point " - "operations simultaneously."); - } -#endif -}); - -#endif diff --git a/tests/group_functions/group_reduce.h b/tests/group_functions/group_reduce.h index 5d150e76c..da09bd0dc 100644 --- a/tests/group_functions/group_reduce.h +++ b/tests/group_functions/group_reduce.h @@ -182,6 +182,20 @@ class invoke_joint_reduce_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + if constexpr (type_traits::group_algorithms::is_legal_operator_v< T, OperatorT>) { joint_reduce_group(queue, op_name); @@ -292,6 +306,20 @@ class invoke_init_joint_reduce_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + if constexpr (type_traits::group_algorithms::is_legal_operator_v< RetT, OperatorT>) { init_joint_reduce_group(queue, op_name); @@ -408,6 +436,20 @@ class invoke_reduce_over_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + if constexpr (type_traits::group_algorithms::is_legal_operator_v< T, OperatorT>) { reduce_over_group(queue, op_name); @@ -529,6 +571,20 @@ class invoke_init_reduce_over_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + if constexpr (type_traits::group_algorithms::is_legal_operator_v< RetT, OperatorT>) { init_reduce_over_group(queue, op_name); diff --git a/tests/group_functions/group_reduce_over_group.cpp b/tests/group_functions/group_reduce_over_group.cpp.in similarity index 62% rename from tests/group_functions/group_reduce_over_group.cpp rename to tests/group_functions/group_reduce_over_group.cpp.in index 6fc0620ef..14a29e6e1 100644 --- a/tests/group_functions/group_reduce_over_group.cpp +++ b/tests/group_functions/group_reduce_over_group.cpp.in @@ -20,30 +20,31 @@ #include "group_reduce.h" +// clang-format off +#cmakedefine CTS_TYPE @CTS_TYPE@ +#cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") +// clang-format on using ReduceTypes = Types; -// 2-dim Cartesian product of type lists -using prod2 = product::type; - -TEMPLATE_LIST_TEST_CASE("Group and sub-group reduce functions", - "[group_func][type_list][dim]", ReduceTypes) { +TEST_CASE(CTS_TYPE_NAME + " group and sub-group reduce functions", + "[group_func][type_list][dim]") { auto queue = once_per_unit::get_queue(); // Get binary operators from TestType - const auto Operators = get_op_types(); - const auto Type = unnamed_type_pack(); - for_all_combinations(Dims, Type, Operators, queue); + const auto Operators = get_op_types(); + const auto RetType = unnamed_type_pack(); + for_all_combinations(Dims, RetType, Operators, + queue); } -TEMPLATE_LIST_TEST_CASE("Group and sub-group reduce functions with init", - "[group_func][type_list][dim]", prod2) { +TEMPLATE_LIST_TEST_CASE(CTS_TYPE_NAME + + " group and sub-group reduce functions with init", + "[group_func][type_list][dim]", ReduceTypes) { auto queue = once_per_unit::get_queue(); - using T = std::tuple_element_t<0, TestType>; - using U = std::tuple_element_t<1, TestType>; // Get binary operators from T - const auto Operators = get_op_types(); - const auto RetType = unnamed_type_pack(); - const auto ReducedType = unnamed_type_pack(); + const auto Operators = get_op_types(); + const auto RetType = unnamed_type_pack(); + const auto ReducedType = unnamed_type_pack(); // check all work group dimensions for_all_combinations( Dims, RetType, ReducedType, Operators, queue); diff --git a/tests/group_functions/group_reduce_over_group_fp16.cpp b/tests/group_functions/group_reduce_over_group_fp16.cpp deleted file mode 100644 index 480773539..000000000 --- a/tests/group_functions/group_reduce_over_group_fp16.cpp +++ /dev/null @@ -1,62 +0,0 @@ -/******************************************************************************* -// -// 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 "group_reduce.h" - -using ReduceTypes = Types; - -using HalfExtendedTypes = concatenation::type; -// 2-dim Cartesian product of type lists -using prod2 = product::type; - -TEST_CASE("Group and sub-group reduce functions", "[group_func][fp16][dim]") { - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp16)) { - // Get binary operators from TestType - const auto Operators = get_op_types(); - const auto Type = unnamed_type_pack(); - for_all_combinations(Dims, Type, Operators, - queue); - } else { - WARN("Device does not support half precision floating point operations."); - } -} - -TEMPLATE_LIST_TEST_CASE("Group and sub-group reduce functions with init", - "[group_func][type_list][fp16][dim]", prod2) { - auto queue = once_per_unit::get_queue(); - using T = std::tuple_element_t<0, TestType>; - using U = std::tuple_element_t<1, TestType>; - - if (queue.get_device().has(sycl::aspect::fp16)) { - if constexpr (std::is_same_v || - std::is_same_v) { - // Get binary operators from T - const auto Operators = get_op_types(); - const auto RetType = unnamed_type_pack(); - const auto ReducedType = unnamed_type_pack(); - // check all work group dimensions - for_all_combinations( - Dims, RetType, ReducedType, Operators, queue); - } - } else { - WARN("Device does not support half precision floating point operations."); - } -} diff --git a/tests/group_functions/group_reduce_over_group_fp64.cpp b/tests/group_functions/group_reduce_over_group_fp64.cpp deleted file mode 100644 index fa8da3d36..000000000 --- a/tests/group_functions/group_reduce_over_group_fp64.cpp +++ /dev/null @@ -1,62 +0,0 @@ -/******************************************************************************* -// -// 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 "group_reduce.h" - -using ReduceTypes = Types; - -using DoubleExtendedTypes = concatenation::type; -// 2-dim Cartesian product of type lists -using prod2 = - product::type; - -TEST_CASE("Group and sub-group reduce functions", "[group_func][fp64][dim]") { - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp64)) { - // Get binary operators from TestType - const auto Operators = get_op_types(); - const auto Type = unnamed_type_pack(); - for_all_combinations(Dims, Type, Operators, - queue); - } else { - WARN("Device does not support double precision floating point operations."); - } -} - -TEMPLATE_LIST_TEST_CASE("Group and sub-group reduce functions with init", - "[group_func][type_list][fp64][dim]", prod2) { - auto queue = once_per_unit::get_queue(); - using T = std::tuple_element_t<0, TestType>; - using U = std::tuple_element_t<1, TestType>; - - if (queue.get_device().has(sycl::aspect::fp64)) { - if constexpr (std::is_same_v || std::is_same_v) { - // Get binary operators from T - const auto Operators = get_op_types(); - const auto RetType = unnamed_type_pack(); - const auto ReducedType = unnamed_type_pack(); - // check all work group dimensions - for_all_combinations( - Dims, RetType, ReducedType, Operators, queue); - } - } else { - WARN("Device does not support double precision floating point operations."); - } -} diff --git a/tests/group_functions/group_reduce_over_group_fp64_fp16.cpp b/tests/group_functions/group_reduce_over_group_fp64_fp16.cpp deleted file mode 100644 index c8077c63d..000000000 --- a/tests/group_functions/group_reduce_over_group_fp64_fp16.cpp +++ /dev/null @@ -1,53 +0,0 @@ -/******************************************************************************* -// -// 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. -// -*******************************************************************************/ - -// need also double tests enabled -#ifdef SYCL_CTS_ENABLE_DOUBLE_TESTS - -#include "group_reduce.h" - -using ReduceTypes = unnamed_type_pack; - -// 2-dim Cartesian product of type lists -using prod2 = product::type; - -TEMPLATE_LIST_TEST_CASE("Group and sub-group reduce functions with init", - "[group_func][fp16][fp64][dim]", prod2) { - auto queue = once_per_unit::get_queue(); - using T = std::tuple_element_t<0, TestType>; - using U = std::tuple_element_t<1, TestType>; - - if (queue.get_device().has(sycl::aspect::fp16) && - queue.get_device().has(sycl::aspect::fp64)) { - // Get binary operators from T - const auto Operators = get_op_types(); - const auto RetType = unnamed_type_pack(); - const auto ReducedType = unnamed_type_pack(); - // check all work group dimensions - for_all_combinations( - Dims, RetType, ReducedType, Operators, queue); - } else { - WARN( - "Device does not support half and double precision floating point " - "operations simultaneously."); - } -} - -#endif diff --git a/tests/group_functions/group_scan.h b/tests/group_functions/group_scan.h index 889098dbd..4b0421dcc 100644 --- a/tests/group_functions/group_scan.h +++ b/tests/group_functions/group_scan.h @@ -219,6 +219,20 @@ template class invoke_joint_scan_group { public: void operator()(sycl::queue& queue) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + const auto operators = get_op_types(); for_all_combinations(operators, queue); } @@ -229,6 +243,20 @@ template class invoke_joint_scan_group_same_type { public: void operator()(sycl::queue& queue) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + const auto operators = get_op_types(); for_all_combinations(operators, queue); } @@ -275,6 +303,20 @@ template class invoke_init_joint_scan_group { public: void operator()(sycl::queue& queue) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + const auto operators = get_op_types(); for_all_combinations(operators, queue); @@ -286,6 +328,20 @@ template class invoke_init_joint_scan_group_same_type { public: void operator()(sycl::queue& queue) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + const auto operators = get_op_types(); for_all_combinations(operators, queue); @@ -477,6 +533,20 @@ template class invoke_scan_over_group { public: void operator()(sycl::queue& queue) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + const auto operators = get_op_types(); for_all_combinations(operators, queue); } @@ -516,6 +586,20 @@ template class invoke_init_scan_over_group { public: void operator()(sycl::queue& queue) { + if (!queue.get_device().has(sycl::aspect::fp16) && + std::is_same_v, sycl::half>) { + WARN( + "Device does not support half precision floating point " + "operations."); + return; + } else if (!queue.get_device().has(sycl::aspect::fp64) && + std::is_same_v, double>) { + WARN( + "Device does not support double precision floating point " + "operations."); + return; + } + const auto operators = get_op_types(); for_all_combinations(operators, queue); diff --git a/tests/group_functions/group_scan_over_group.cpp b/tests/group_functions/group_scan_over_group.cpp.in similarity index 76% rename from tests/group_functions/group_scan_over_group.cpp rename to tests/group_functions/group_scan_over_group.cpp.in index dd9e587e5..671cf8273 100644 --- a/tests/group_functions/group_scan_over_group.cpp +++ b/tests/group_functions/group_scan_over_group.cpp.in @@ -22,6 +22,11 @@ #if !SYCL_CTS_COMPILING_WITH_HIPSYCL #include "group_scan.h" +// clang-format off +#cmakedefine CTS_TYPE @CTS_TYPE@ +#cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") +// clang-format on +using TestType = unnamed_type_pack; using ScanTypes = Types; static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); @@ -29,17 +34,18 @@ static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); // FIXME: known_identity is not impemented yet for hipSYCL. DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group scan functions", "[group_func][type_list][dim]")({ +(CTS_TYPE_NAME + " group and sub-group scan functions", + "[group_func][type_list][dim]")({ auto queue = once_per_unit::get_queue(); - for_all_combinations(Dims, ScanTypes{}, queue); + for_all_combinations(Dims, TestType{}, queue); }); // FIXME: hipSYCL has wrong arguments order for inclusive_scan_over_group: init // and op are interchanged. known_identity is not implemented yet. DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group scan functions with init", +(CTS_TYPE_NAME + " group and sub-group scan functions with init", "[group_func][type_list][dim]")({ auto queue = once_per_unit::get_queue(); - for_all_combinations(Dims, ScanTypes{}, + for_all_combinations(Dims, TestType{}, ScanTypes{}, queue); }); diff --git a/tests/group_functions/group_scan_over_group_fp16.cpp b/tests/group_functions/group_scan_over_group_fp16.cpp deleted file mode 100644 index 1fdaba919..000000000 --- a/tests/group_functions/group_scan_over_group_fp16.cpp +++ /dev/null @@ -1,58 +0,0 @@ -/******************************************************************************* -// -// 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/disabled_for_test_case.h" -#include "type_coverage.h" -#if !SYCL_CTS_COMPILING_WITH_HIPSYCL -#include "group_scan.h" - -using ScanTypes = Types; - -using HalfType = unnamed_type_pack; -using HalfExtendedTypes = concatenation::type; - -static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); -#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL - -// FIXME: known_identity is not impemented yet for hipSYCL. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group scan functions", "[group_func][fp16][dim]")({ - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp16)) { - for_all_combinations(Dims, HalfType{}, queue); - } else { - WARN("Device does not support half precision floating point operations."); - } -}); - -// FIXME: hipSYCL has wrong arguments order for inclusive_scan_over_group: init -// and op are interchanged. known_identity is not impemented yet. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group scan functions with init", - "[group_func][type_list][fp16][dim]")({ - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp16)) { - for_all_combinations_with( - Dims, HalfExtendedTypes{}, HalfExtendedTypes{}, queue); - } else { - WARN("Device does not support half precision floating point operations."); - } -}); diff --git a/tests/group_functions/group_scan_over_group_fp64.cpp b/tests/group_functions/group_scan_over_group_fp64.cpp deleted file mode 100644 index 11b26be57..000000000 --- a/tests/group_functions/group_scan_over_group_fp64.cpp +++ /dev/null @@ -1,58 +0,0 @@ -/******************************************************************************* -// -// 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/disabled_for_test_case.h" -#include "type_coverage.h" -#if !SYCL_CTS_COMPILING_WITH_HIPSYCL -#include "group_scan.h" - -using ScanTypes = Types; - -using DoubleType = unnamed_type_pack; -using DoubleExtendedTypes = concatenation::type; - -static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); -#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL - -// FIXME: known_identity is not impemented yet for hipSYCL. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group scan functions", "[group_func][fp64][dim]")({ - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp64)) { - for_all_combinations(Dims, DoubleType{}, queue); - } else { - WARN("Device does not support double precision floating point operations."); - } -}); - -// FIXME: hipSYCL has wrong arguments order for inclusive_scan_over_group: init -// and op are interchanged. known_identity is not impemented yet. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group scan functions with init", - "[group_func][type_list][fp64][dim]")({ - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp64)) { - for_all_combinations_with( - Dims, DoubleExtendedTypes{}, DoubleExtendedTypes{}, queue); - } else { - WARN("Device does not support double precision floating point operations."); - } -}); diff --git a/tests/group_functions/group_scan_over_group_fp64_fp16.cpp b/tests/group_functions/group_scan_over_group_fp64_fp16.cpp deleted file mode 100644 index 846f81310..000000000 --- a/tests/group_functions/group_scan_over_group_fp64_fp16.cpp +++ /dev/null @@ -1,58 +0,0 @@ -/******************************************************************************* -// -// 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. -// -*******************************************************************************/ - -// need also double tests enabled -#ifdef SYCL_CTS_ENABLE_DOUBLE_TESTS - -#include "../common/common.h" -#include "../common/disabled_for_test_case.h" -#include "type_coverage.h" -#if !SYCL_CTS_COMPILING_WITH_HIPSYCL -#include "group_scan.h" - -using ScanTypes = Types; - -using DoubleHalfTypes = unnamed_type_pack; -using DoubleHalfExtendedTypes = concatenation::type; - -static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); -#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL - -// FIXME: hipSYCL has wrong arguments order for inclusive_scan_over_group: init -// and op are interchanged. known_identity is not impemented yet. -DISABLED_FOR_TEST_CASE(hipSYCL) -("Group and sub-group scan functions with init", - "[group_func][fp16][fp64][dim]")({ - auto queue = once_per_unit::get_queue(); - if (queue.get_device().has(sycl::aspect::fp16) && - queue.get_device().has(sycl::aspect::fp64)) { - // here DoubleHalfTypes can be used as only half+double combinations are - // untested - for_all_combinations_with_two(Dims, DoubleHalfTypes{}, - DoubleHalfTypes{}, queue); - } else { - WARN( - "Device does not support half and double precision floating point " - "operations simultaneously."); - } -}); - -#endif From 7eef73d4103bc1afcbc9e101667a31552029c693 Mon Sep 17 00:00:00 2001 From: JieZ Zhang Date: Wed, 11 Oct 2023 13:14:01 +0800 Subject: [PATCH 02/12] Remove fp16 and fp64 from 'get_std_type' if cmake options are set --- tests/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index eeab437af..1729e994f 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -41,7 +41,7 @@ function(get_std_type OUT_LIST) "unsigned long long" ) endif() - + half_double_filter(STD_TYPE_LIST) set(${OUT_LIST} ${${OUT_LIST}} ${STD_TYPE_LIST} PARENT_SCOPE) endfunction() From 633e0bb264921fd19355514a0ddb091ed21e8105 Mon Sep 17 00:00:00 2001 From: JieZ Zhang Date: Mon, 16 Oct 2023 16:42:52 +0800 Subject: [PATCH 03/12] Make CTS_TYPE visible to hipSYCL --- tests/group_functions/group_joint_scan.cpp.in | 7 +++++-- tests/group_functions/group_scan_over_group.cpp.in | 6 ++++-- 2 files changed, 9 insertions(+), 4 deletions(-) diff --git a/tests/group_functions/group_joint_scan.cpp.in b/tests/group_functions/group_joint_scan.cpp.in index 47c92f3ca..8e7039e9e 100644 --- a/tests/group_functions/group_joint_scan.cpp.in +++ b/tests/group_functions/group_joint_scan.cpp.in @@ -19,18 +19,21 @@ *******************************************************************************/ #include "../common/disabled_for_test_case.h" -#if !SYCL_CTS_COMPILING_WITH_HIPSYCL -#include "group_scan.h" // clang-format off #cmakedefine CTS_TYPE @CTS_TYPE@ #cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") // clang-format on + +#if !SYCL_CTS_COMPILING_WITH_HIPSYCL +#include "group_scan.h" + using TestType = unnamed_type_pack; using ScanTypes = Types; static const auto Dims = integer_pack<1, 2, 3>::generate_unnamed(); #endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL + // FIXME: known_identity is not impemented yet for hipSYCL. DISABLED_FOR_TEST_CASE(hipSYCL) (CTS_TYPE_NAME + " group and sub-group joint scan functions", diff --git a/tests/group_functions/group_scan_over_group.cpp.in b/tests/group_functions/group_scan_over_group.cpp.in index 671cf8273..f16f1fe60 100644 --- a/tests/group_functions/group_scan_over_group.cpp.in +++ b/tests/group_functions/group_scan_over_group.cpp.in @@ -19,13 +19,15 @@ *******************************************************************************/ #include "../common/disabled_for_test_case.h" -#if !SYCL_CTS_COMPILING_WITH_HIPSYCL -#include "group_scan.h" // clang-format off #cmakedefine CTS_TYPE @CTS_TYPE@ #cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") // clang-format on + +#if !SYCL_CTS_COMPILING_WITH_HIPSYCL +#include "group_scan.h" + using TestType = unnamed_type_pack; using ScanTypes = Types; From da51497a44049bcc97f3c788b7c38d14848944b6 Mon Sep 17 00:00:00 2001 From: Jie Zhang Date: Wed, 18 Oct 2023 14:38:58 +0800 Subject: [PATCH 04/12] Skip unsupported tests with SKIP --- tests/group_functions/group_reduce.h | 24 +++++++------------ tests/group_functions/group_scan.h | 36 ++++++++++------------------ 2 files changed, 20 insertions(+), 40 deletions(-) diff --git a/tests/group_functions/group_reduce.h b/tests/group_functions/group_reduce.h index da09bd0dc..e9c2253d3 100644 --- a/tests/group_functions/group_reduce.h +++ b/tests/group_functions/group_reduce.h @@ -184,16 +184,14 @@ class invoke_joint_reduce_group { void operator()(sycl::queue& queue, const std::string& op_name) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } if constexpr (type_traits::group_algorithms::is_legal_operator_v< @@ -308,16 +306,14 @@ class invoke_init_joint_reduce_group { void operator()(sycl::queue& queue, const std::string& op_name) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } if constexpr (type_traits::group_algorithms::is_legal_operator_v< @@ -438,16 +434,14 @@ class invoke_reduce_over_group { void operator()(sycl::queue& queue, const std::string& op_name) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } if constexpr (type_traits::group_algorithms::is_legal_operator_v< @@ -573,16 +567,14 @@ class invoke_init_reduce_over_group { void operator()(sycl::queue& queue, const std::string& op_name) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } if constexpr (type_traits::group_algorithms::is_legal_operator_v< diff --git a/tests/group_functions/group_scan.h b/tests/group_functions/group_scan.h index 4b0421dcc..80449b2b4 100644 --- a/tests/group_functions/group_scan.h +++ b/tests/group_functions/group_scan.h @@ -221,16 +221,14 @@ class invoke_joint_scan_group { void operator()(sycl::queue& queue) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } const auto operators = get_op_types(); @@ -245,16 +243,14 @@ class invoke_joint_scan_group_same_type { void operator()(sycl::queue& queue) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } const auto operators = get_op_types(); @@ -305,16 +301,14 @@ class invoke_init_joint_scan_group { void operator()(sycl::queue& queue) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } const auto operators = get_op_types(); @@ -330,16 +324,14 @@ class invoke_init_joint_scan_group_same_type { void operator()(sycl::queue& queue) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } const auto operators = get_op_types(); @@ -535,16 +527,14 @@ class invoke_scan_over_group { void operator()(sycl::queue& queue) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } const auto operators = get_op_types(); @@ -588,16 +578,14 @@ class invoke_init_scan_over_group { void operator()(sycl::queue& queue) { if (!queue.get_device().has(sycl::aspect::fp16) && std::is_same_v, sycl::half>) { - WARN( + SKIP( "Device does not support half precision floating point " "operations."); - return; } else if (!queue.get_device().has(sycl::aspect::fp64) && std::is_same_v, double>) { - WARN( + SKIP( "Device does not support double precision floating point " "operations."); - return; } const auto operators = get_op_types(); From dd7e58ca60743788d586c182fc5e48780f9ae845 Mon Sep 17 00:00:00 2001 From: Jie Zhang Date: Wed, 18 Oct 2023 16:38:23 +0800 Subject: [PATCH 05/12] Move aspect query into a seperate line --- tests/group_functions/group_reduce.h | 80 +++++++++--------- tests/group_functions/group_scan.h | 120 +++++++++++++-------------- 2 files changed, 100 insertions(+), 100 deletions(-) diff --git a/tests/group_functions/group_reduce.h b/tests/group_functions/group_reduce.h index e9c2253d3..9b71303d1 100644 --- a/tests/group_functions/group_reduce.h +++ b/tests/group_functions/group_reduce.h @@ -182,16 +182,16 @@ class invoke_joint_reduce_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } if constexpr (type_traits::group_algorithms::is_legal_operator_v< @@ -304,16 +304,16 @@ class invoke_init_joint_reduce_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } if constexpr (type_traits::group_algorithms::is_legal_operator_v< @@ -432,16 +432,16 @@ class invoke_reduce_over_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } if constexpr (type_traits::group_algorithms::is_legal_operator_v< @@ -565,16 +565,16 @@ class invoke_init_reduce_over_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } if constexpr (type_traits::group_algorithms::is_legal_operator_v< diff --git a/tests/group_functions/group_scan.h b/tests/group_functions/group_scan.h index 80449b2b4..698e70ea1 100644 --- a/tests/group_functions/group_scan.h +++ b/tests/group_functions/group_scan.h @@ -219,16 +219,16 @@ template class invoke_joint_scan_group { public: void operator()(sycl::queue& queue) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } const auto operators = get_op_types(); @@ -241,16 +241,16 @@ template class invoke_joint_scan_group_same_type { public: void operator()(sycl::queue& queue) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } const auto operators = get_op_types(); @@ -299,16 +299,16 @@ template class invoke_init_joint_scan_group { public: void operator()(sycl::queue& queue) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } const auto operators = get_op_types(); @@ -322,16 +322,16 @@ template class invoke_init_joint_scan_group_same_type { public: void operator()(sycl::queue& queue) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } const auto operators = get_op_types(); @@ -525,16 +525,16 @@ template class invoke_scan_over_group { public: void operator()(sycl::queue& queue) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } const auto operators = get_op_types(); @@ -576,16 +576,16 @@ template class invoke_init_scan_over_group { public: void operator()(sycl::queue& queue) { - if (!queue.get_device().has(sycl::aspect::fp16) && - std::is_same_v, sycl::half>) { - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (!queue.get_device().has(sycl::aspect::fp64) && - std::is_same_v, double>) { - SKIP( - "Device does not support double precision floating point " - "operations."); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); } const auto operators = get_op_types(); From 0453e6e98f996313bee7ef9607113dca494867f3 Mon Sep 17 00:00:00 2001 From: Jie Zhang Date: Wed, 8 Nov 2023 16:11:35 +0800 Subject: [PATCH 06/12] Unify indentation for group_function/CMakeLists --- tests/group_functions/CMakeLists.txt | 44 ++++++++++++++-------------- 1 file changed, 22 insertions(+), 22 deletions(-) diff --git a/tests/group_functions/CMakeLists.txt b/tests/group_functions/CMakeLists.txt index 2f30733f9..d473a8859 100644 --- a/tests/group_functions/CMakeLists.txt +++ b/tests/group_functions/CMakeLists.txt @@ -1,17 +1,17 @@ -function(configure_test_case) - cmake_parse_arguments(CTS - "" "TYPE;IN_FILENAME;OUT_FILENAME;TEST_LIST" "" ${ARGN}) - set(CTS_TYPE_NAME ${CTS_TYPE}) - configure_file(${CTS_IN_FILENAME} ${CTS_OUT_FILENAME}) - list(APPEND ${CTS_TEST_LIST} "${CMAKE_CURRENT_BINARY_DIR}/${CTS_OUT_FILENAME}") - set(${CTS_TEST_LIST} ${${CTS_TEST_LIST}} PARENT_SCOPE) +function(configure_test_case) + cmake_parse_arguments(CTS + "" "TYPE;IN_FILENAME;OUT_FILENAME;TEST_LIST" "" ${ARGN}) + set(CTS_TYPE_NAME ${CTS_TYPE}) + configure_file(${CTS_IN_FILENAME} ${CTS_OUT_FILENAME}) + list(APPEND ${CTS_TEST_LIST} "${CMAKE_CURRENT_BINARY_DIR}/${CTS_OUT_FILENAME}") + set(${CTS_TEST_LIST} ${${CTS_TEST_LIST}} PARENT_SCOPE) endfunction() list(APPEND TEMPLATE_LIST - "group_joint_scan" - "group_scan_over_group" - "group_joint_reduce" - "group_reduce_over_group" + "group_joint_scan" + "group_scan_over_group" + "group_joint_reduce" + "group_reduce_over_group" ) set(TYPE_LIST "") get_std_type(TYPE_LIST) @@ -20,17 +20,17 @@ file(GLOB test_cases_list *.cpp) foreach(TEMP IN LISTS TEMPLATE_LIST) foreach(TY IN LISTS TYPE_LIST) - if(TY STREQUAL "bool") - continue() - endif() - set(OUT_FILE "${TEMP}_${TY}.cpp") - STRING(REGEX REPLACE ":" "_" OUT_FILE ${OUT_FILE}) - STRING(REGEX REPLACE " " "_" OUT_FILE ${OUT_FILE}) - configure_test_case( - TYPE "${TY}" - IN_FILENAME "${TEMP}.cpp.in" - OUT_FILENAME ${OUT_FILE} - TEST_LIST test_cases_list) + if(TY STREQUAL "bool") + continue() + endif() + set(OUT_FILE "${TEMP}_${TY}.cpp") + STRING(REGEX REPLACE ":" "_" OUT_FILE ${OUT_FILE}) + STRING(REGEX REPLACE " " "_" OUT_FILE ${OUT_FILE}) + configure_test_case( + TYPE "${TY}" + IN_FILENAME "${TEMP}.cpp.in" + OUT_FILENAME ${OUT_FILE} + EST_LIST test_cases_list) endforeach() endforeach() From eddb5d5e5ef3955a5a763b489cb929f535cfa6d0 Mon Sep 17 00:00:00 2001 From: Jie Zhang Date: Fri, 10 Nov 2023 17:35:09 +0800 Subject: [PATCH 07/12] Remove trailing whitespace --- tests/group_functions/CMakeLists.txt | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/group_functions/CMakeLists.txt b/tests/group_functions/CMakeLists.txt index d473a8859..d1275bd05 100644 --- a/tests/group_functions/CMakeLists.txt +++ b/tests/group_functions/CMakeLists.txt @@ -1,4 +1,4 @@ -function(configure_test_case) +function(configure_test_case) cmake_parse_arguments(CTS "" "TYPE;IN_FILENAME;OUT_FILENAME;TEST_LIST" "" ${ARGN}) set(CTS_TYPE_NAME ${CTS_TYPE}) @@ -7,8 +7,8 @@ function(configure_test_case) set(${CTS_TEST_LIST} ${${CTS_TEST_LIST}} PARENT_SCOPE) endfunction() -list(APPEND TEMPLATE_LIST - "group_joint_scan" +list(APPEND TEMPLATE_LIST + "group_joint_scan" "group_scan_over_group" "group_joint_reduce" "group_reduce_over_group" From 83ce19d07a98f328e4109748ad73a9c04d3aac18 Mon Sep 17 00:00:00 2001 From: jiezzhang Date: Thu, 16 Nov 2023 17:01:02 +0800 Subject: [PATCH 08/12] Fix typo Co-authored-by: Marcos Maronas --- tests/group_functions/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/group_functions/CMakeLists.txt b/tests/group_functions/CMakeLists.txt index d1275bd05..66a3fa653 100644 --- a/tests/group_functions/CMakeLists.txt +++ b/tests/group_functions/CMakeLists.txt @@ -30,7 +30,7 @@ foreach(TEMP IN LISTS TEMPLATE_LIST) TYPE "${TY}" IN_FILENAME "${TEMP}.cpp.in" OUT_FILENAME ${OUT_FILE} - EST_LIST test_cases_list) + TEST_LIST test_cases_list) endforeach() endforeach() From 2cb959b43c185e91ee8f4d2eab0856733fd077fd Mon Sep 17 00:00:00 2001 From: Jie Zhang Date: Thu, 16 Nov 2023 19:57:15 +0800 Subject: [PATCH 09/12] Remove fp16 and fp64 for group_functions only This reverts commit 7eef73d4103bc1afcbc9e101667a31552029c693. --- .vscode/settings.json | 5 +++++ tests/CMakeLists.txt | 2 +- tests/group_functions/CMakeLists.txt | 1 + 3 files changed, 7 insertions(+), 1 deletion(-) create mode 100644 .vscode/settings.json diff --git a/.vscode/settings.json b/.vscode/settings.json new file mode 100644 index 000000000..4fe957006 --- /dev/null +++ b/.vscode/settings.json @@ -0,0 +1,5 @@ +{ + "files.associations": { + "algorithm": "cpp" + } +} \ No newline at end of file diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 1729e994f..eeab437af 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -41,7 +41,7 @@ function(get_std_type OUT_LIST) "unsigned long long" ) endif() - half_double_filter(STD_TYPE_LIST) + set(${OUT_LIST} ${${OUT_LIST}} ${STD_TYPE_LIST} PARENT_SCOPE) endfunction() diff --git a/tests/group_functions/CMakeLists.txt b/tests/group_functions/CMakeLists.txt index 66a3fa653..b6f7fe381 100644 --- a/tests/group_functions/CMakeLists.txt +++ b/tests/group_functions/CMakeLists.txt @@ -15,6 +15,7 @@ list(APPEND TEMPLATE_LIST ) set(TYPE_LIST "") get_std_type(TYPE_LIST) +half_double_filter(TYPE_LIST) file(GLOB test_cases_list *.cpp) From 05aa248a45e5de72aa96d24aff6a5c63aaef0ce4 Mon Sep 17 00:00:00 2001 From: Jie Zhang Date: Thu, 16 Nov 2023 20:49:05 +0800 Subject: [PATCH 10/12] Remove misconfiged files --- .vscode/settings.json | 5 ----- 1 file changed, 5 deletions(-) delete mode 100644 .vscode/settings.json diff --git a/.vscode/settings.json b/.vscode/settings.json deleted file mode 100644 index 4fe957006..000000000 --- a/.vscode/settings.json +++ /dev/null @@ -1,5 +0,0 @@ -{ - "files.associations": { - "algorithm": "cpp" - } -} \ No newline at end of file From 3cc9e85c3b4b467fde1ec4b82d31e03c12728e32 Mon Sep 17 00:00:00 2001 From: Jie Zhang Date: Fri, 17 Nov 2023 16:19:07 +0800 Subject: [PATCH 11/12] Move FP16 and FP64 aspect check into templates --- .../group_functions/group_joint_reduce.cpp.in | 26 +++++++ tests/group_functions/group_joint_scan.cpp.in | 24 +++++++ tests/group_functions/group_reduce.h | 48 ------------- .../group_reduce_over_group.cpp.in | 26 +++++++ tests/group_functions/group_scan.h | 72 ------------------- .../group_scan_over_group.cpp.in | 22 ++++++ 6 files changed, 98 insertions(+), 120 deletions(-) diff --git a/tests/group_functions/group_joint_reduce.cpp.in b/tests/group_functions/group_joint_reduce.cpp.in index fde841e69..ff8c652a2 100644 --- a/tests/group_functions/group_joint_reduce.cpp.in +++ b/tests/group_functions/group_joint_reduce.cpp.in @@ -32,6 +32,19 @@ TEST_CASE(CTS_TYPE_NAME + " group and sub-group joint reduce functions", auto queue = once_per_unit::get_queue(); const auto Operators = get_op_types(); const auto RetType = unnamed_type_pack(); + + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); + } + for_all_combinations(Dims, RetType, Operators, queue); } @@ -45,6 +58,19 @@ TEMPLATE_LIST_TEST_CASE( const auto Operators = get_op_types(); const auto RetType = unnamed_type_pack(); const auto ReducedType = unnamed_type_pack(); + + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); + } + // check all work group dimensions for_all_combinations( Dims, RetType, ReducedType, Operators, queue); diff --git a/tests/group_functions/group_joint_scan.cpp.in b/tests/group_functions/group_joint_scan.cpp.in index 8e7039e9e..7fbee0ae6 100644 --- a/tests/group_functions/group_joint_scan.cpp.in +++ b/tests/group_functions/group_joint_scan.cpp.in @@ -49,6 +49,18 @@ DISABLED_FOR_TEST_CASE(hipSYCL) #endif auto queue = once_per_unit::get_queue(); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); + } + // FIXME: hipSYCL cannot handle cases of different types #if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) for_all_combinations(Dims, TestType{}, @@ -74,6 +86,18 @@ DISABLED_FOR_TEST_CASE(hipSYCL) #endif auto queue = once_per_unit::get_queue(); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); + } + // FIXME: hipSYCL cannot handle cases of different types #if defined(SYCL_CTS_COMPILING_WITH_HIPSYCL) for_all_combinations(Dims, TestType{}, diff --git a/tests/group_functions/group_reduce.h b/tests/group_functions/group_reduce.h index 9b71303d1..5d150e76c 100644 --- a/tests/group_functions/group_reduce.h +++ b/tests/group_functions/group_reduce.h @@ -182,18 +182,6 @@ class invoke_joint_reduce_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - if constexpr (type_traits::group_algorithms::is_legal_operator_v< T, OperatorT>) { joint_reduce_group(queue, op_name); @@ -304,18 +292,6 @@ class invoke_init_joint_reduce_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - if constexpr (type_traits::group_algorithms::is_legal_operator_v< RetT, OperatorT>) { init_joint_reduce_group(queue, op_name); @@ -432,18 +408,6 @@ class invoke_reduce_over_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - if constexpr (type_traits::group_algorithms::is_legal_operator_v< T, OperatorT>) { reduce_over_group(queue, op_name); @@ -565,18 +529,6 @@ class invoke_init_reduce_over_group { public: void operator()(sycl::queue& queue, const std::string& op_name) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - if constexpr (type_traits::group_algorithms::is_legal_operator_v< RetT, OperatorT>) { init_reduce_over_group(queue, op_name); diff --git a/tests/group_functions/group_reduce_over_group.cpp.in b/tests/group_functions/group_reduce_over_group.cpp.in index 14a29e6e1..226b1655b 100644 --- a/tests/group_functions/group_reduce_over_group.cpp.in +++ b/tests/group_functions/group_reduce_over_group.cpp.in @@ -32,6 +32,19 @@ TEST_CASE(CTS_TYPE_NAME + " group and sub-group reduce functions", // Get binary operators from TestType const auto Operators = get_op_types(); const auto RetType = unnamed_type_pack(); + + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); + } + for_all_combinations(Dims, RetType, Operators, queue); } @@ -45,6 +58,19 @@ TEMPLATE_LIST_TEST_CASE(CTS_TYPE_NAME + const auto Operators = get_op_types(); const auto RetType = unnamed_type_pack(); const auto ReducedType = unnamed_type_pack(); + + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); + } + // check all work group dimensions for_all_combinations( Dims, RetType, ReducedType, Operators, queue); diff --git a/tests/group_functions/group_scan.h b/tests/group_functions/group_scan.h index 698e70ea1..889098dbd 100644 --- a/tests/group_functions/group_scan.h +++ b/tests/group_functions/group_scan.h @@ -219,18 +219,6 @@ template class invoke_joint_scan_group { public: void operator()(sycl::queue& queue) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - const auto operators = get_op_types(); for_all_combinations(operators, queue); } @@ -241,18 +229,6 @@ template class invoke_joint_scan_group_same_type { public: void operator()(sycl::queue& queue) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - const auto operators = get_op_types(); for_all_combinations(operators, queue); } @@ -299,18 +275,6 @@ template class invoke_init_joint_scan_group { public: void operator()(sycl::queue& queue) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - const auto operators = get_op_types(); for_all_combinations(operators, queue); @@ -322,18 +286,6 @@ template class invoke_init_joint_scan_group_same_type { public: void operator()(sycl::queue& queue) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - const auto operators = get_op_types(); for_all_combinations(operators, queue); @@ -525,18 +477,6 @@ template class invoke_scan_over_group { public: void operator()(sycl::queue& queue) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - const auto operators = get_op_types(); for_all_combinations(operators, queue); } @@ -576,18 +516,6 @@ template class invoke_init_scan_over_group { public: void operator()(sycl::queue& queue) { - if constexpr (std::is_same_v, sycl::half>) { - if (!queue.get_device().has(sycl::aspect::fp16)) - SKIP( - "Device does not support half precision floating point " - "operations."); - } else if (std::is_same_v, double>) { - if (!queue.get_device().has(sycl::aspect::fp64)) - SKIP( - "Device does not support double precision floating point " - "operations."); - } - const auto operators = get_op_types(); for_all_combinations(operators, queue); diff --git a/tests/group_functions/group_scan_over_group.cpp.in b/tests/group_functions/group_scan_over_group.cpp.in index f16f1fe60..39308d9e5 100644 --- a/tests/group_functions/group_scan_over_group.cpp.in +++ b/tests/group_functions/group_scan_over_group.cpp.in @@ -39,6 +39,17 @@ DISABLED_FOR_TEST_CASE(hipSYCL) (CTS_TYPE_NAME + " group and sub-group scan functions", "[group_func][type_list][dim]")({ auto queue = once_per_unit::get_queue(); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); + } for_all_combinations(Dims, TestType{}, queue); }); @@ -48,6 +59,17 @@ DISABLED_FOR_TEST_CASE(hipSYCL) (CTS_TYPE_NAME + " group and sub-group scan functions with init", "[group_func][type_list][dim]")({ auto queue = once_per_unit::get_queue(); + if constexpr (std::is_same_v, sycl::half>) { + if (!queue.get_device().has(sycl::aspect::fp16)) + SKIP( + "Device does not support half precision floating point " + "operations."); + } else if (std::is_same_v, double>) { + if (!queue.get_device().has(sycl::aspect::fp64)) + SKIP( + "Device does not support double precision floating point " + "operations."); + } for_all_combinations(Dims, TestType{}, ScanTypes{}, queue); }); From 8e91479f35d8ee1043f33d4692a43d3f4c385b66 Mon Sep 17 00:00:00 2001 From: Jie Zhang Date: Mon, 20 Nov 2023 16:46:24 +0800 Subject: [PATCH 12/12] Remove redundant comments in group_joint_reduce --- tests/group_functions/group_joint_reduce.cpp.in | 2 -- 1 file changed, 2 deletions(-) diff --git a/tests/group_functions/group_joint_reduce.cpp.in b/tests/group_functions/group_joint_reduce.cpp.in index ff8c652a2..c0e214c5f 100644 --- a/tests/group_functions/group_joint_reduce.cpp.in +++ b/tests/group_functions/group_joint_reduce.cpp.in @@ -26,7 +26,6 @@ // clang-format on using ReduceTypes = Types; -// hipSYCL has no implementation over sub-groups TEST_CASE(CTS_TYPE_NAME + " group and sub-group joint reduce functions", "[group_func][type_list][dim]") { auto queue = once_per_unit::get_queue(); @@ -49,7 +48,6 @@ TEST_CASE(CTS_TYPE_NAME + " group and sub-group joint reduce functions", queue); } -// hipSYCL has problems with 16-bit types for Ptr TEMPLATE_LIST_TEST_CASE( CTS_TYPE_NAME + " group and sub-group joint reduce functions with init", "[group_func][type_list][dim]", ReduceTypes) {