From 46917ee67e8f78465983b47e5616516a8c039139 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 5 Feb 2024 04:25:16 -0800 Subject: [PATCH 01/13] Adds test for the non_uniform_groups oneAPI extension This commit adds tests for the [sycl_ext_oneapi_non_uniform_groups](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc) extension in accordance with the test plan currently in review here: https://github.com/KhronosGroup/SYCL-CTS/pull/866 --- CMakeLists.txt | 4 + .../oneapi_non_uniform_groups/CMakeLists.txt | 40 ++ .../ballot_group_api.cpp | 175 ++++++ .../fixed_size_group_api.cpp | 183 ++++++ .../group_barrier.cpp | 37 ++ .../oneapi_non_uniform_groups/group_barrier.h | 215 +++++++ .../group_broadcast.cpp | 36 ++ .../group_broadcast.h | 203 ++++++ .../group_broadcast_fp16.cpp | 38 ++ .../group_broadcast_fp64.cpp | 36 ++ .../group_joint_reduce.cpp.in | 85 +++ .../group_joint_scan.cpp.in | 81 +++ .../oneapi_non_uniform_groups/group_of.cpp | 50 ++ .../oneapi_non_uniform_groups/group_of.h | 322 ++++++++++ .../group_permute.cpp | 36 ++ .../oneapi_non_uniform_groups/group_permute.h | 110 ++++ .../group_permute_fp16.cpp | 39 ++ .../group_permute_fp64.cpp | 38 ++ .../oneapi_non_uniform_groups/group_reduce.h | 560 +++++++++++++++++ .../group_reduce_over_group.cpp.in | 88 +++ .../oneapi_non_uniform_groups/group_scan.h | 594 ++++++++++++++++++ .../group_scan_over_group.cpp.in | 76 +++ .../oneapi_non_uniform_groups/group_shift.cpp | 36 ++ .../oneapi_non_uniform_groups/group_shift.h | 149 +++++ .../group_shift_fp16.cpp | 38 ++ .../group_shift_fp64.cpp | 38 ++ .../is_fixed_topology_group.cpp | 59 ++ .../non_uniform_group_common.h | 207 ++++++ .../opportunistic_group_api.cpp | 182 ++++++ .../tangle_group_api.cpp | 182 ++++++ 30 files changed, 3937 insertions(+) create mode 100644 tests/extension/oneapi_non_uniform_groups/CMakeLists.txt create mode 100644 tests/extension/oneapi_non_uniform_groups/ballot_group_api.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_barrier.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_barrier.h create mode 100644 tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_broadcast.h create mode 100644 tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in create mode 100644 tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in create mode 100644 tests/extension/oneapi_non_uniform_groups/group_of.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_of.h create mode 100644 tests/extension/oneapi_non_uniform_groups/group_permute.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_permute.h create mode 100644 tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_reduce.h create mode 100644 tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in create mode 100644 tests/extension/oneapi_non_uniform_groups/group_scan.h create mode 100644 tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in create mode 100644 tests/extension/oneapi_non_uniform_groups/group_shift.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_shift.h create mode 100644 tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/is_fixed_topology_group.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h create mode 100644 tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp create mode 100644 tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index c3161e9c1..1365767bc 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -83,6 +83,10 @@ add_cts_option(SYCL_CTS_ENABLE_EXT_ONEAPI_AUTO_LOCAL_RANGE_TESTS "Enable extension oneAPI auto_local_range tests" OFF FORCE_ON ${SYCL_CTS_ENABLE_EXT_ONEAPI_TESTS}) +add_cts_option(SYCL_CTS_ENABLE_EXT_ONEAPI_NON_UNIFORM_GROUPS_TESTS + "Enable extension oneAPI non_uniform_groups tests" OFF + FORCE_ON ${SYCL_CTS_ENABLE_EXT_ONEAPI_TESTS}) + # TODO: Deprecated - remove add_cts_option(SYCL_CTS_ENABLE_VERBOSE_LOG "Enable debug-level logs (deprecated)" OFF) diff --git a/tests/extension/oneapi_non_uniform_groups/CMakeLists.txt b/tests/extension/oneapi_non_uniform_groups/CMakeLists.txt new file mode 100644 index 000000000..79b819177 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/CMakeLists.txt @@ -0,0 +1,40 @@ +if(SYCL_CTS_ENABLE_EXT_ONEAPI_NON_UNIFORM_GROUPS_TESTS) + 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) + half_double_filter(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}) +endif() diff --git a/tests/extension/oneapi_non_uniform_groups/ballot_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/ballot_group_api.cpp new file mode 100644 index 000000000..a54fedb66 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/ballot_group_api.cpp @@ -0,0 +1,175 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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 non_uniform_groups::tests { + +struct checks { + enum { + get_group_id, + get_local_id, + get_group_range, + get_local_range, + get_group_linear_id, + get_local_linear_id, + get_group_linear_range, + get_local_linear_range, + leader, + COUNT, + }; +}; + +TEST_CASE("Test for ballot_group apis.", "[oneapi_non_uniform_groups]") { +#ifndef SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS + SKIP("SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS is not defined"); +#else + namespace oneapi_ext = sycl::ext::oneapi::experimental; + using ballot_group_t = oneapi_ext::ballot_group; + using CheckResults = bool[checks::COUNT]; + + constexpr size_t num_items = 64; + + sycl::buffer results_buffer{num_items}; + + auto q = sycl_cts::util::get_cts_object::queue(); + + if (!q.get_device().has(sycl::aspect::ext_oneapi_ballot_group)) { + SKIP("Device does not support ballot_group."); + } + + q.submit([&](sycl::handler& cgh) { + sycl::accessor acc{results_buffer, cgh, sycl::write_only}; + + sycl::nd_range<1> nd_range{sycl::range<1>{num_items}, + sycl::range<1>{num_items}}; + + cgh.parallel_for(nd_range, [=](sycl::nd_item<1> it) { + auto& results = acc[it.get_global_id()]; + + sycl::sub_group sg = it.get_sub_group(); + size_t sub_group_size = sg.get_local_range().size(); + size_t split = sub_group_size / 3; + bool is_left = sg.get_local_linear_id() < split; + + auto ballot = oneapi_ext::get_ballot_group(sg, is_left); + static_assert(std::is_same_v); + + // Since we make an uneven split, the group size will differ based on + // which side of the split this item is. + size_t expected_group_size = is_left ? split : sub_group_size - split; + + results[checks::get_group_id] = + ballot.get_group_id() == (is_left ? 1 : 0); + results[checks::get_local_id] = + ballot.get_local_id() < expected_group_size; + results[checks::get_group_range] = ballot.get_group_range().size() == 2; + results[checks::get_local_range] = + ballot.get_local_range().size() == expected_group_size; + results[checks::get_group_linear_id] = + ballot.get_group_linear_id() == ballot.get_group_id(); + results[checks::get_local_linear_id] = + ballot.get_local_linear_id() == ballot.get_local_id(); + results[checks::get_group_linear_range] = + ballot.get_group_linear_range() == ballot.get_group_range().size(); + results[checks::get_local_linear_range] = + ballot.get_local_linear_range() == ballot.get_local_range().size(); + results[checks::leader] = ballot.leader() == (ballot.get_local_id() == 0); + }); + }); + q.wait(); + + CheckResults results = {}; + sycl::accessor acc = results_buffer.get_host_access(); + for (size_t check = 0; check < checks::COUNT; check++) + results[check] = std::all_of(acc.cbegin(), acc.cend(), + [=](const auto& it) { return it[check]; }); + + // Group-category traits. + STATIC_CHECK(sycl::is_group::value); + STATIC_CHECK(sycl::is_group_v); + STATIC_CHECK(oneapi_ext::is_user_constructed_group::value); + STATIC_CHECK(oneapi_ext::is_user_constructed_group_v); + STATIC_CHECK(!oneapi_ext::is_fixed_topology_group::value); + STATIC_CHECK(!oneapi_ext::is_fixed_topology_group_v); + + // Aliases. + STATIC_CHECK(std::is_same_v>); + STATIC_CHECK(std::is_same_v>); + STATIC_CHECK(std::is_same_v); + + // Static constexpr members. + STATIC_CHECK(ballot_group_t::dimensions == 1); + STATIC_CHECK(ballot_group_t::fence_scope == sycl::sub_group::fence_scope); + + // get_group_id + CHECK(std::is_same_v().get_group_id()), + ballot_group_t::id_type>); + CHECK(results[checks::get_group_id]); + + // get_local_id + CHECK(std::is_same_v().get_local_id()), + ballot_group_t::id_type>); + CHECK(results[checks::get_local_id]); + + // get_group_range + CHECK( + std::is_same_v().get_group_range()), + ballot_group_t::range_type>); + CHECK(results[checks::get_group_range]); + + // get_local_range + CHECK( + std::is_same_v().get_local_range()), + ballot_group_t::range_type>); + CHECK(results[checks::get_local_range]); + + // get_group_linear_id + CHECK(std::is_same_v< + decltype(std::declval().get_group_linear_id()), + ballot_group_t::linear_id_type>); + CHECK(results[checks::get_group_linear_id]); + + // get_local_linear_id + CHECK(std::is_same_v< + decltype(std::declval().get_local_linear_id()), + ballot_group_t::linear_id_type>); + CHECK(results[checks::get_local_linear_id]); + + // get_group_linear_range + CHECK(std::is_same_v< + decltype(std::declval().get_group_linear_range()), + ballot_group_t::linear_id_type>); + CHECK(results[checks::get_group_linear_range]); + + // get_local_linear_range + CHECK(std::is_same_v< + decltype(std::declval().get_local_linear_range()), + ballot_group_t::linear_id_type>); + CHECK(results[checks::get_local_linear_range]); + + // leader + CHECK( + std::is_same_v().leader()), bool>); + CHECK(results[checks::leader]); +#endif +} + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp new file mode 100644 index 000000000..742cefc86 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp @@ -0,0 +1,183 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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 non_uniform_groups::tests { + +struct checks { + enum { + get_group_id, + get_local_id, + get_group_range, + get_local_range, + get_group_linear_id, + get_local_linear_id, + get_group_linear_range, + get_local_linear_range, + leader, + COUNT, + }; +}; + +constexpr size_t partition_size = 4; + +TEST_CASE("Test for fixed_size_group apis.", "[oneapi_non_uniform_groups]") { +#ifndef SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS + SKIP("SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS is not defined"); +#else + namespace oneapi_ext = sycl::ext::oneapi::experimental; + using fixed_size_group_t = + oneapi_ext::fixed_size_group; + using CheckResults = bool[checks::COUNT]; + + constexpr size_t num_items = 64; + + sycl::buffer results_buffer{num_items}; + + auto q = sycl_cts::util::get_cts_object::queue(); + + if (!q.get_device().has(sycl::aspect::ext_oneapi_fixed_size_group)) { + SKIP("Device does not support fixed_size_group."); + } + + q.submit([&](sycl::handler& cgh) { + sycl::accessor acc{results_buffer, cgh, sycl::write_only}; + + sycl::nd_range<1> nd_range{sycl::range<1>{num_items}, + sycl::range<1>{num_items}}; + + cgh.parallel_for( + nd_range, [=](sycl::nd_item<1> it) { + auto& results = acc[it.get_global_id()]; + + sycl::sub_group sg = it.get_sub_group(); + size_t sub_group_size = sg.get_local_range().size(); + + auto fixed_size = + oneapi_ext::get_fixed_size_group(sg); + static_assert( + std::is_same_v); + + results[checks::get_group_id] = + fixed_size.get_group_id() < sub_group_size / partition_size; + results[checks::get_local_id] = + fixed_size.get_local_id() < partition_size; + results[checks::get_group_range] = + fixed_size.get_group_range().size() == + sub_group_size / partition_size; + results[checks::get_local_range] = + fixed_size.get_local_range().size() == partition_size; + results[checks::get_group_linear_id] = + fixed_size.get_group_linear_id() == fixed_size.get_group_id(); + results[checks::get_local_linear_id] = + fixed_size.get_local_linear_id() == fixed_size.get_local_id(); + results[checks::get_group_linear_range] = + fixed_size.get_group_linear_range() == + fixed_size.get_group_range().size(); + results[checks::get_local_linear_range] = + fixed_size.get_local_linear_range() == + fixed_size.get_local_range().size(); + results[checks::leader] = + fixed_size.leader() == (fixed_size.get_local_id() == 0); + }); + }); + q.wait(); + + CheckResults results = {}; + sycl::accessor acc = results_buffer.get_host_access(); + for (size_t check = 0; check < checks::COUNT; check++) + results[check] = std::all_of(acc.cbegin(), acc.cend(), + [=](const auto& it) { return it[check]; }); + + // Group-category traits. + STATIC_CHECK(sycl::is_group::value); + STATIC_CHECK(sycl::is_group_v); + STATIC_CHECK( + oneapi_ext::is_user_constructed_group::value); + STATIC_CHECK(oneapi_ext::is_user_constructed_group_v); + STATIC_CHECK(!oneapi_ext::is_fixed_topology_group::value); + STATIC_CHECK(!oneapi_ext::is_fixed_topology_group_v); + + // Aliases. + STATIC_CHECK(std::is_same_v>); + STATIC_CHECK(std::is_same_v>); + STATIC_CHECK(std::is_same_v); + + // Static constexpr members. + STATIC_CHECK(fixed_size_group_t::dimensions == 1); + STATIC_CHECK(fixed_size_group_t::fence_scope == sycl::sub_group::fence_scope); + + // get_group_id + CHECK(std::is_same_v< + decltype(std::declval().get_group_id()), + fixed_size_group_t::id_type>); + CHECK(results[checks::get_group_id]); + + // get_local_id + CHECK(std::is_same_v< + decltype(std::declval().get_local_id()), + fixed_size_group_t::id_type>); + CHECK(results[checks::get_local_id]); + + // get_group_range + CHECK(std::is_same_v< + decltype(std::declval().get_group_range()), + fixed_size_group_t::range_type>); + CHECK(results[checks::get_group_range]); + + // get_local_range + CHECK(std::is_same_v< + decltype(std::declval().get_local_range()), + fixed_size_group_t::range_type>); + CHECK(results[checks::get_local_range]); + + // get_group_linear_id + CHECK(std::is_same_v< + decltype(std::declval().get_group_linear_id()), + fixed_size_group_t::linear_id_type>); + CHECK(results[checks::get_group_linear_id]); + + // get_local_linear_id + CHECK(std::is_same_v< + decltype(std::declval().get_local_linear_id()), + fixed_size_group_t::linear_id_type>); + CHECK(results[checks::get_local_linear_id]); + + // get_group_linear_range + CHECK(std::is_same_v< + decltype(std::declval().get_group_linear_range()), + fixed_size_group_t::linear_id_type>); + CHECK(results[checks::get_group_linear_range]); + + // get_local_linear_range + CHECK(std::is_same_v< + decltype(std::declval().get_local_linear_range()), + fixed_size_group_t::linear_id_type>); + CHECK(results[checks::get_local_linear_range]); + + // leader + CHECK(std::is_same_v().leader()), + bool>); + CHECK(results[checks::leader]); +#endif +} + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp b/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp new file mode 100644 index 000000000..6942a3327 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp @@ -0,0 +1,37 @@ +/******************************************************************************* +// +// 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 + +#include "group_barrier.h" + +template +class test_fence; + +TEST_CASE("Non-uniform-group barriers", + "[oneapi_non_uniform_groups][group_func]") { + auto queue = once_per_unit::get_queue(); + + non_uniform_group_barrier>(queue); + non_uniform_group_barrier>( + queue); + non_uniform_group_barrier>(queue); + non_uniform_group_barrier(queue); +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_barrier.h b/tests/extension/oneapi_non_uniform_groups/group_barrier.h new file mode 100644 index 000000000..652b9e0a4 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_barrier.h @@ -0,0 +1,215 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_functions/group_functions_common.h" +#include "non_uniform_group_common.h" + +template +class non_uniform_group_barrier_kernel; + +/** + * @brief Provides test for arbitraty non-uniform group barriers + * @tparam GroupT Type of the non-uniform group to test with + * @tparam T Type pointed by Ptr + */ +template +void non_uniform_group_barrier(sycl::queue& queue) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing group-of predicate function for " + group_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + std::vector supported_barriers = + queue.get_context() + .get_info(); + + using sms = std::tuple; + // indices of the tuple components + enum s { scope = 0, support = 1, test = 2 }; + + constexpr int non_uniform_group_barrier_variants = 5; + std::array + non_uniform_group_barriers{{{sycl::memory_scope::sub_group, true, true}, + {sycl::memory_scope::sub_group, true, true}, + {sycl::memory_scope::work_group, true, true}, + {sycl::memory_scope::device, true, true}, + {sycl::memory_scope::system, true, true}}}; + std::array + non_uniform_group_barriers_names{ + "default", "sycl::memory_scope::sub_group", + "sycl::memory_scope::work_group", "sycl::memory_scope::device", + "sycl::memory_scope::system"}; + for (auto& barrier : non_uniform_group_barriers) { + auto& sb = supported_barriers; + if (std::find(sb.begin(), sb.end(), std::get(barrier)) == + sb.end()) { + std::get(barrier) = false; + } + } + + using el_type = int32_t; + sycl::device device = queue.get_device(); + + // Check the maximum number elements of type "el_type" that can be + // placed in the device's global and local memory. Since the test + // tries to allocate local and global buffers with a size equal to + // the work group size, the latter must be limited by the allowed + // buffer size. + uint64_t global_mem_size_in_bytes = + device.get_info(); + uint64_t global_mem_size_in_elements = + global_mem_size_in_bytes / sizeof(el_type); + + uint64_t local_mem_size_in_bytes = + device.get_info(); + uint64_t local_mem_size_in_elements = + local_mem_size_in_bytes / sizeof(el_type); + + uint64_t work_items_limit = + std::min(global_mem_size_in_elements, local_mem_size_in_elements); + + sycl::range<1> work_group_range = + sycl_cts::util::work_group_range<1>(queue, work_items_limit); + size_t work_group_size = work_group_range.size(); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + std::vector v(work_group_size, 0); + sycl::buffer global_mem(v.data(), + sycl::range<1>(work_group_size)); + + sycl::buffer non_uniform_group_barriers_buf( + non_uniform_group_barriers.data(), sycl::range<1>(5)); + + queue.submit([&](sycl::handler& cgh) { + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + auto non_uniform_group_barriers_acc = + non_uniform_group_barriers_buf + .get_access(cgh); + + sycl::local_accessor local_acc( + sycl::range<1>(work_group_size), cgh); + sycl::accessor global_acc = + global_mem.get_access(cgh); + + cgh.parallel_for>( + executionRange, [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate(sub_group, + test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + + size_t llid = non_uniform_group.get_local_linear_id(); + size_t max_id = non_uniform_group.get_local_linear_range() - 1; + + ASSERT_RETURN_TYPE( + void, sycl::group_barrier(non_uniform_group), + "Return type of group_barrier(GroupT g) is wrong\n"); + ASSERT_RETURN_TYPE( + void, + sycl::group_barrier(non_uniform_group, + non_uniform_group.fence_scope), + "Return type of group_barrier(GroupT g, " + "memory_scope fence_scope) is wrong\n"); + + // test of default barrier + local_acc[llid] = llid; + sycl::group_barrier(non_uniform_group); + + if (local_acc[max_id - llid] != max_id - llid) + std::get(non_uniform_group_barriers_acc[0]) = false; + sycl::group_barrier(non_uniform_group); + + local_acc[llid] = 1; + sycl::group_barrier(non_uniform_group); + + if (local_acc[max_id - llid] != 1) + std::get(non_uniform_group_barriers_acc[0]) = false; + sycl::group_barrier(non_uniform_group); + + // tests for other barriers + for (int i = 1; i < non_uniform_group_barrier_variants; ++i) { + auto& barrier = non_uniform_group_barriers_acc[i]; + + if ((sub_group.get_group_linear_id() == 0) && + (non_uniform_group.get_group_linear_id() == + NonUniformGroupHelper< + GroupT>::preferred_single_worker_group_id(test_case)) && + std::get(barrier)) { + local_acc[llid] = llid; + global_acc[llid] = llid; + sycl::group_barrier(non_uniform_group); + + if (local_acc[max_id - llid] != max_id - llid || + global_acc[max_id - llid] != max_id - llid) + std::get(barrier) = false; + sycl::group_barrier(non_uniform_group); + + switch (std::get(barrier)) { + case sycl::memory_scope::sub_group: + case sycl::memory_scope::work_group: + local_acc[llid] = 1; + sycl::group_barrier(non_uniform_group, + std::get(barrier)); + + if (local_acc[max_id - llid] != 1) + std::get(barrier) = false; + sycl::group_barrier(non_uniform_group); + + [[fallthrough]]; + default: + global_acc[llid] = 1; + sycl::group_barrier(non_uniform_group, + std::get(barrier)); + + if (global_acc[max_id - llid] != 1) + std::get(barrier) = false; + sycl::group_barrier(non_uniform_group); + } + } + } + }); + }); + + for (int i = 0; i < non_uniform_group_barrier_variants; ++i) { + bool result = std::get(non_uniform_group_barriers[i]); + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group); + INFO("Result of group_barrier invocation for sub-group and " + << non_uniform_group_barriers_names[i] << " memory scope is " + << (result ? "right" : "wrong")); + CHECK(result); + } + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp new file mode 100644 index 000000000..630e01694 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp @@ -0,0 +1,36 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_broadcast.h" + +using BroadcastTypes = CustomTypes; + +TEMPLATE_LIST_TEST_CASE("Non-uniform group broadcast and select", + "[oneapi_non_uniform_groups][group_func][type_list]", + BroadcastTypes) { + auto queue = once_per_unit::get_queue(); + broadcast_non_uniform_group, + TestType>(queue); + broadcast_non_uniform_group, + TestType>(queue); + broadcast_non_uniform_group, + TestType>(queue); + broadcast_non_uniform_group(queue); +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast.h b/tests/extension/oneapi_non_uniform_groups/group_broadcast.h new file mode 100644 index 000000000..f3be91049 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast.h @@ -0,0 +1,203 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_functions/group_functions_common.h" +#include "non_uniform_group_common.h" + +template +class broadcast_non_uniform_group_kernel; + +/** + * @brief Provides test for arbitraty non-uniform group broadcast functions + * @tparam GroupT Type of the non-uniform group to test with + * @tparam T Type pointed by Ptr + */ +template +void broadcast_non_uniform_group(sycl::queue& queue) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing broadcast and select for " + group_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + // 4 functions + constexpr int test_matrix = 4; + const std::string test_names[test_matrix] = { + "T group_broadcast(GroupT g, T x)", + "T group_broadcast(GroupT g, T x, GroupT::linear_id_type " + "local_linear_id)", + "T group_broadcast(GroupT g, T x, GroupT::id_type local_id)", + "T select_from_group(GroupT g, T x, GroupT::id_type local_id)"}; + + sycl::range<1> work_group_range = sycl_cts::util::work_group_range<1>(queue); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + // array to return results + T origin_values[test_matrix] = {splat_init(0)}; + T broadcasted_values[test_matrix] = {splat_init(0)}; + { + sycl::buffer origin_values_buf(origin_values, + sycl::range<1>(test_matrix)); + sycl::buffer broadcasted_values_buf(broadcasted_values, + sycl::range<1>(test_matrix)); + + queue.submit([&](sycl::handler& cgh) { + auto origin_values_acc = + origin_values_buf + .template get_access(cgh); + auto broadcasted_values_acc = + broadcasted_values_buf + .template get_access(cgh); + + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + // Values computed in a kernel depend on global linear id. We need to + // make sure that there are no overflows + REQUIRE(executionRange.get_global_range().size() < + std::numeric_limits::max() / 100); + + cgh.parallel_for>( + executionRange, [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate(sub_group, + test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + + // Each work-item computes a unique value + T value_to_broadcast(splat_init( + static_cast(item.get_global_linear_id() * 100 + + non_uniform_group.get_local_id()))); + + T local_var(splat_init(0)); + + // To simplify the test, we are only checking the first group in + // the first sub-group. + if (item.get_sub_group().get_group_id()[0] == 0 && + non_uniform_group.get_group_id()[0] == 0) { + // Find local id of first, last and some third sub-group item in + // between. Will be used to check different combinations of + // broadcasting and receiving work-items + sycl::id<1> first_id = 0; + sycl::id<1> mid_id = non_uniform_group.get_local_range() / 2; + sycl::id<1> last_id = non_uniform_group.get_local_range(); + --last_id[0]; + + // Broadcast from the first work-item + ASSERT_RETURN_TYPE( + T, + sycl::group_broadcast(non_uniform_group, + value_to_broadcast), + "Return type of group_broadcast(GroupT g, T x) is wrong\n"); + + if (non_uniform_group.leader()) { + // Work-item which does the broadcast stores value to + // broadcast to use it later as a reference + origin_values_acc[0] = value_to_broadcast; + } + auto broadcasted_value = sycl::group_broadcast( + non_uniform_group, value_to_broadcast); + // We read broadcasted value in another work-item + if (non_uniform_group.get_local_id() == last_id) + broadcasted_values_acc[0] = broadcasted_value; + + // Broadcast from the last work-item + ASSERT_RETURN_TYPE( + T, + sycl::group_broadcast(non_uniform_group, value_to_broadcast, + last_id), + "Return type of group_broadcast(GroupT g, T x, " + "GroupT::linear_id_type local_linear_id) is wrong\n"); + + if (non_uniform_group.get_local_id() == last_id) { + // Work-item which does the broadcast stores value to + // broadcast to use it later as a reference + origin_values_acc[1] = value_to_broadcast; + } + + broadcasted_value = sycl::group_broadcast( + non_uniform_group, value_to_broadcast, + non_uniform_group.get_local_linear_range() - 1); + // We read broadcasted value in another work-item + if (non_uniform_group.get_local_id() == mid_id) + broadcasted_values_acc[1] = broadcasted_value; + + // Broadcast from a mid work-item + ASSERT_RETURN_TYPE( + T, + sycl::group_broadcast(non_uniform_group, value_to_broadcast, + mid_id), + "Return type of group_broadcast(GroupT g, T x, " + "GroupT::id_type local_id) is wrong\n"); + + if (non_uniform_group.get_local_id() == mid_id) { + // Work-item which does the broadcast stores value to + // broadcast to use it later as a reference + origin_values_acc[2] = value_to_broadcast; + } + broadcasted_value = sycl::group_broadcast( + non_uniform_group, value_to_broadcast, mid_id); + // We read broadcasted value in another work-item + if (non_uniform_group.get_local_id() == first_id) + broadcasted_values_acc[2] = broadcasted_value; + + // Select from the first work-item + ASSERT_RETURN_TYPE( + T, + sycl::select_from_group(non_uniform_group, + value_to_broadcast, first_id), + "Return type of select_from_group(GroupT g, T x, " + "GroupT::id_type local_id) is wrong\n"); + + if (non_uniform_group.get_local_id() == first_id) { + // Work-item which does the broadcast stores value to + // broadcast to use it later as a reference + origin_values_acc[3] = value_to_broadcast; + } + broadcasted_value = sycl::select_from_group( + non_uniform_group, value_to_broadcast, first_id); + // We read broadcasted value in another work-item + if (non_uniform_group.get_local_id() == mid_id) + broadcasted_values_acc[3] = broadcasted_value; + } + }); + }); + } + for (int i = 0; i < test_matrix; ++i) { + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group); + INFO("Return value of " + << test_names[i] << " with T = " << type_name() << " is " + << (equal(broadcasted_values[i], origin_values[i]) ? "right" + : "wrong")); + CHECK(equal(broadcasted_values[i], origin_values[i])); + } + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp new file mode 100644 index 000000000..c5b7103dd --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp @@ -0,0 +1,38 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_broadcast.h" + +TEST_CASE("Non-uniform group broadcast and select", + "[oneapi_non_uniform_groups][group_func][fp16][dim]") { + auto queue = once_per_unit::get_queue(); + if (queue.get_device().has(sycl::aspect::fp16)) { + broadcast_non_uniform_group, + sycl::half>(queue); + broadcast_non_uniform_group< + oneapi_ext::fixed_size_group<4, sycl::sub_group>, sycl::half>(queue); + broadcast_non_uniform_group, + sycl::half>(queue); + broadcast_non_uniform_group( + queue); + } else { + WARN("Device does not support half precision floating point operations."); + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp new file mode 100644 index 000000000..adc3a65ea --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp @@ -0,0 +1,36 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_broadcast.h" + +TEST_CASE("Non-uniform group broadcast and select", "[group_func][fp64][dim]") { + auto queue = once_per_unit::get_queue(); + if (queue.get_device().has(sycl::aspect::fp64)) { + broadcast_non_uniform_group, + double>(queue); + broadcast_non_uniform_group< + oneapi_ext::fixed_size_group<4, sycl::sub_group>, double>(queue); + broadcast_non_uniform_group, + double>(queue); + broadcast_non_uniform_group(queue); + } else { + WARN("Device does not support double precision floating point operations."); + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in new file mode 100644 index 000000000..22eac4e4b --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in @@ -0,0 +1,85 @@ +/******************************************************************************* +// +// 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" + +// clang-format off +#cmakedefine CTS_TYPE @CTS_TYPE@ +#cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") +// clang-format on +using ReduceTypes = Types; + +TEST_CASE(CTS_TYPE_NAME + " non-uniform group joint reduce functions", + "[oneapi_non_uniform_groups][group_func][type_list]") { + auto queue = once_per_unit::get_queue(); + const auto Operators = get_op_types(); + const auto RetType = unnamed_type_pack(); + const auto GroupTypes = unnamed_type_pack< + oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::tangle_group, + oneapi_ext::opportunistic_group>(); + + 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(GroupTypes, RetType, + Operators, queue); +} + +TEMPLATE_LIST_TEST_CASE( + CTS_TYPE_NAME + " non-uniform group joint reduce functions with init", + "[oneapi_non_uniform_groups][group_func][type_list]", ReduceTypes) { + auto queue = once_per_unit::get_queue(); + + const auto Operators = get_op_types(); + const auto RetType = unnamed_type_pack(); + const auto ReducedType = unnamed_type_pack(); + const auto GroupTypes = unnamed_type_pack< + oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::tangle_group, + oneapi_ext::opportunistic_group>(); + + 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( + GroupTypes, RetType, ReducedType, Operators, queue); +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in new file mode 100644 index 000000000..71af50f63 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in @@ -0,0 +1,81 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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. +// +*******************************************************************************/ + +// 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; +#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL + +TEST_CASE(CTS_TYPE_NAME + " non-uniform group joint scan functions", + "[oneapi_non_uniform_groups][group_func][type_list]"){ + auto queue = once_per_unit::get_queue(); + const auto GroupTypes = unnamed_type_pack< + oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::tangle_group, + oneapi_ext::opportunistic_group>(); + + 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(GroupTypes, TestType{}, + ScanTypes{}, queue); +}; + +TEST_CASE(CTS_TYPE_NAME + " non-uniform group joint scan functions with init", + "[oneapi_non_uniform_groups][group_func][type_list]"){ + auto queue = once_per_unit::get_queue(); + const auto GroupTypes = unnamed_type_pack< + oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::tangle_group, + oneapi_ext::opportunistic_group>(); + + 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( + GroupTypes, TestType{}, ScanTypes{}, ScanTypes{}, queue); +}; diff --git a/tests/extension/oneapi_non_uniform_groups/group_of.cpp b/tests/extension/oneapi_non_uniform_groups/group_of.cpp new file mode 100644 index 000000000..f0532e1e8 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_of.cpp @@ -0,0 +1,50 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_of.h" + +// use wide types to exclude truncation of init values +using WideTypes = std::tuple; + +TEMPLATE_LIST_TEST_CASE( + "Non-uniform group of bool functions with predicate functions", + "[oneapi_non_uniform_groups][group_func][type_list]", WideTypes) { + auto queue = once_per_unit::get_queue(); + predicate_function_of_non_uniform_group< + oneapi_ext::ballot_group, TestType>(queue); + predicate_function_of_non_uniform_group< + oneapi_ext::fixed_size_group<4, sycl::sub_group>, TestType>(queue); + predicate_function_of_non_uniform_group< + oneapi_ext::tangle_group, TestType>(queue); + predicate_function_of_non_uniform_group(queue); +} + +TEST_CASE("Non-uniform group of bool functions", + "[oneapi_non_uniform_groups][group_func]") { + auto queue = once_per_unit::get_queue(); + bool_function_of_non_uniform_group>( + queue); + bool_function_of_non_uniform_group< + oneapi_ext::fixed_size_group<4, sycl::sub_group>>(queue); + bool_function_of_non_uniform_group>( + queue); + bool_function_of_non_uniform_group(queue); +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_of.h b/tests/extension/oneapi_non_uniform_groups/group_of.h new file mode 100644 index 000000000..512fb56bd --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_of.h @@ -0,0 +1,322 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_functions/group_functions_common.h" +#include "non_uniform_group_common.h" + +template +class predicate_function_of_non_uniform_group_kernel; + +/** + * @brief Provides test for arbitraty non-uniform group bool of operations with + * predicate functions + * @tparam GroupT Type of the non-uniform group to test with + * @tparam T Type pointed by Ptr + */ +template +void predicate_function_of_non_uniform_group(sycl::queue& queue) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing group-of predicate function for " + group_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + // 3 functions * 4 predicates + constexpr int test_matrix = 3; + const std::string test_names[test_matrix] = { + "bool any_of_group(GroupT g, T x, Predicate pred)", + "bool all_of_group(GroupT g, T x, Predicate pred)", + "bool none_of_group(GroupT g, T x, Predicate pred)"}; + constexpr int test_cases = 4; + const std::string test_cases_names[test_cases] = {"none true", "one true", + "some true", "all true"}; + + sycl::range<1> work_group_range = sycl_cts::util::work_group_range<1>(queue); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + // array to return results: 4 predicates * 3 functions + constexpr int total_case_count = test_matrix * test_cases; + bool res[total_case_count]; + // Initially fill the results array with 'true'. Each non-uniform group test + // 'ands' with this to ensure every non-uniform group in the work-group + // returns the correct result. + std::fill(res, res + total_case_count, true); + { + sycl::buffer res_sycl(res, sycl::range<1>(total_case_count)); + + queue.submit([&](sycl::handler& cgh) { + auto res_acc = res_sycl.get_access(cgh); + + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + cgh.parallel_for>(executionRange, [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate(sub_group, + test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + + size_t size = non_uniform_group.get_local_linear_range(); + + // Use the non-uniform group local ID (plus 1) as a variable against + // which to test our predicates. Note that this has a well-defined set + // of values [1,2,...,N] where N is the non-uniform group size. Note + // that the non-uniform group could also just be of size 1. + T local_var(non_uniform_group.get_local_linear_id() + 1); + + // predicates + // The variable is never 1 for any member of the non-uniform group + auto none_true = [&](T i) { return i == 0; }; + // Exactly one member of the non-uniform group has value 1 (the first) + auto one_true = [&](T i) { return i == 1; }; + // Some (or all, for non-uniform groups of size 1) members of the + // non-uniform group have this value + auto some_true = [&](T i) { return i > size / 2; }; + // The variable is less than or equal to the non-uniform group size + // for all members of the non-uniform group. + auto all_true = [&](T i) { return i <= size; }; + + { + ASSERT_RETURN_TYPE( + bool, + sycl::any_of_group(non_uniform_group, local_var, none_true), + "Return type of any_of_group(GroupT g, bool pred) is wrong\n"); + res_acc[0] &= + !sycl::any_of_group(non_uniform_group, local_var, none_true); + res_acc[1] &= + sycl::any_of_group(non_uniform_group, local_var, one_true); + res_acc[2] &= + sycl::any_of_group(non_uniform_group, local_var, some_true); + res_acc[3] &= + sycl::any_of_group(non_uniform_group, local_var, all_true); + + ASSERT_RETURN_TYPE( + bool, + sycl::all_of_group(non_uniform_group, local_var, none_true), + "Return type of all_of_group(GroupT g, bool pred) is wrong\n"); + res_acc[4] &= + !sycl::all_of_group(non_uniform_group, local_var, none_true); + // Note that 'one_true' returns true for the first item. Thus in the + // case that the non-uniform group size is 1, check that all items + // match; otherwise check that not all items match. + res_acc[5] &= + sycl::all_of_group(non_uniform_group, local_var, one_true) ^ + (size != 1); + // Note that 'some_true' returns true for the first item if the + // non-uniform group size is 1. In that case, check that all items + // match; otherwise check that not all items match. + res_acc[6] &= + sycl::all_of_group(non_uniform_group, local_var, some_true) ^ + (size != 1); + res_acc[7] &= + sycl::all_of_group(non_uniform_group, local_var, all_true); + + ASSERT_RETURN_TYPE( + bool, + sycl::none_of_group(non_uniform_group, local_var, none_true), + "Return type of none_of_group(GroupT g, bool pred) is " + "wrong\n"); + res_acc[8] &= + sycl::none_of_group(non_uniform_group, local_var, none_true); + res_acc[9] &= + !sycl::none_of_group(non_uniform_group, local_var, one_true); + res_acc[10] &= + !sycl::none_of_group(non_uniform_group, local_var, some_true); + res_acc[11] &= + !sycl::none_of_group(non_uniform_group, local_var, all_true); + } + }); + }); + } + int index = 0; + for (int i = 0; i < test_matrix; ++i) + for (int j = 0; j < test_cases; ++j) { + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group); + INFO("Value of " << test_names[i] << " with " << test_cases_names[j] + << " predicate is " + << (res[index] ? "right" : "wrong")); + CHECK(res[index++]); + } + } +} + +template +class predicate_function_of_non_uniform_group_bool_kernel; + +/** + * @brief Provides test for group bool of operations + * @tparam GroupT Type of the non-uniform group to test with + */ +template +void bool_function_of_non_uniform_group(sycl::queue& queue) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing group-of bool function for " + group_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + // 3 functions * 4 predicates + constexpr int test_matrix = 3; + const std::string test_names[test_matrix] = { + "bool any_of_group(GroupT g, bool pred)", + "bool all_of_group(GroupT g, bool pred)", + "bool none_of_group(GroupT g, bool pred)"}; + constexpr int test_cases = 4; + const std::string test_cases_names[test_cases] = {"none true", "one true", + "some true", "all true"}; + + using T = size_t; + + sycl::range<1> work_group_range = sycl_cts::util::work_group_range<1>(queue); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + // array to return results: 4 predicates * 3 functions + constexpr int total_case_count = test_matrix * test_cases; + bool res[total_case_count]; + // Initially fill the results array with 'true'. Each non-uniform group test + // 'ands' with this to ensure every non-uniform group in the work-group + // returns the correct result. + std::fill(res, res + total_case_count, true); + { + sycl::buffer res_sycl(res, sycl::range<1>(total_case_count)); + + queue.submit([&](sycl::handler& cgh) { + auto res_acc = res_sycl.get_access(cgh); + + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + cgh.parallel_for>(executionRange, [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate(sub_group, + test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + + size_t size = non_uniform_group.get_local_linear_range(); + + // Use the non-uniform group local ID (plus 1) as a variable against + // which to test our predicates. Note that this has a well-defined set + // of values [1,2,...,N] where N is the non-uniform group size. Note + // that the non-uniform group could also just be of size 1. + T local_var(non_uniform_group.get_local_linear_id() + 1); + + // predicates + // The variable is never 1 for any member of the non-uniform group + auto none_true = [&](T i) { return i == 0; }; + // Exactly one member of the non-uniform group has value 1 (the first) + auto one_true = [&](T i) { return i == 1; }; + // Some (or all, for non-uniform groups of size 1) members of the + // non-uniform group have this value + auto some_true = [&](T i) { return i > size / 2; }; + // The variable is less than or equal to the non-uniform group size + // for all members of the non-uniform group. + auto all_true = [&](T i) { return i <= size; }; + + { + ASSERT_RETURN_TYPE( + bool, + sycl::any_of_group(non_uniform_group, none_true(local_var)), + "Return type of any_of_group(GroupT g, bool pred) is wrong\n"); + res_acc[0] &= + !sycl::any_of_group(non_uniform_group, none_true(local_var)); + res_acc[1] &= + sycl::any_of_group(non_uniform_group, one_true(local_var)); + res_acc[2] &= + sycl::any_of_group(non_uniform_group, some_true(local_var)); + res_acc[3] &= + sycl::any_of_group(non_uniform_group, all_true(local_var)); + + ASSERT_RETURN_TYPE( + bool, + sycl::all_of_group(non_uniform_group, none_true(local_var)), + "Return type of all_of_group(GroupT g, bool pred) is wrong\n"); + res_acc[4] = + !sycl::all_of_group(non_uniform_group, none_true(local_var)); + // Note that 'one_true' returns true for the first item. Thus in the + // case that the non-uniform group size is 1, check that all items + // match; otherwise check that not all items match. + res_acc[5] &= + sycl::all_of_group(non_uniform_group, one_true(local_var)) ^ + (size != 1); + // Note that 'some_true' returns true for the first item if the + // non-uniform group size is 1. In that case, check that all items + // match; otherwise check that not all items match. + res_acc[6] &= + sycl::all_of_group(non_uniform_group, some_true(local_var)) ^ + (size != 1); + res_acc[7] &= + sycl::all_of_group(non_uniform_group, all_true(local_var)); + + ASSERT_RETURN_TYPE( + bool, + sycl::none_of_group(non_uniform_group, none_true(local_var)), + "Return type of none_of_group(GroupT g, bool pred) is " + "wrong\n"); + res_acc[8] &= + sycl::none_of_group(non_uniform_group, none_true(local_var)); + res_acc[9] &= + !sycl::none_of_group(non_uniform_group, one_true(local_var)); + res_acc[10] &= + !sycl::none_of_group(non_uniform_group, some_true(local_var)); + res_acc[11] &= + !sycl::none_of_group(non_uniform_group, all_true(local_var)); + } + }); + }); + } + int index = 0; + for (int i = 0; i < test_matrix; ++i) + for (int j = 0; j < test_cases; ++j) { + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group); + INFO("Value of " << test_names[i] << " with " << test_cases_names[j] + << " predicate is " + << (res[index] ? "right" : "wrong")); + CHECK(res[index++]); + } + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute.cpp new file mode 100644 index 000000000..5b311f183 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_permute.cpp @@ -0,0 +1,36 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_permute.h" + +// hipSYCL does not permute right 8-bit types inside groups +TEMPLATE_LIST_TEST_CASE("Non-uniform-group permute", + "[oneapi_non_uniform_groups][group_func][type_list]", + CustomTypes) { + auto queue = once_per_unit::get_queue(); + + permute_non_uniform_group, + TestType>(queue); + permute_non_uniform_group, + TestType>(queue); + permute_non_uniform_group, + TestType>(queue); + permute_non_uniform_group(queue); +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute.h b/tests/extension/oneapi_non_uniform_groups/group_permute.h new file mode 100644 index 000000000..f67ba37e6 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_permute.h @@ -0,0 +1,110 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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 + +#include "../../group_functions/group_functions_common.h" +#include "non_uniform_group_common.h" + +template +class permute_non_uniform_group_kernel; + +template +void permute_non_uniform_group(sycl::queue& queue) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing permute for " + group_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + const std::string test_name = + "T permute_group_by_xor(GroupT g, T x, GroupT::linear_id_type mask)"; + + sycl::range<1> work_group_range = sycl_cts::util::work_group_range<1>(queue); + size_t work_group_size = work_group_range.size(); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + // array to return results: + std::valarray res(false, work_group_size); + { + sycl::buffer res_sycl(std::begin(res), + sycl::range<1>(work_group_size)); + + queue.submit([&](sycl::handler& cgh) { + auto res_acc = res_sycl.get_access(cgh); + + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + cgh.parallel_for>( + executionRange, [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, they fill their + // elements in the result with true and leave early. + if (!NonUniformGroupHelper::should_participate( + sub_group, test_case)) { + res_acc[item.get_local_linear_id()] = true; + return; + } + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + + using lin_id_type = typename GroupT::linear_id_type; + const lin_id_type llid = non_uniform_group.get_local_linear_id(); + + T local_var(splat_init(llid + 1)); + T permuted_var(splat_init(llid + 1)); + + ASSERT_RETURN_TYPE( + T, + sycl::permute_group_by_xor(non_uniform_group, local_var, 0), + "Return type of permute_group_by_xor(GroupT g, T x, " + "GroupT::linear_id_type mask) is wrong\n"); + + bool res = true; + for (lin_id_type mask = 1u; mask > 0; mask <<= 1) { + permuted_var = sycl::permute_group_by_xor(non_uniform_group, + local_var, mask); + res &= equal(permuted_var, splat_init((llid ^ mask) + 1)) || + ((llid ^ mask) >= + non_uniform_group.get_local_linear_range()); + } + res_acc[item.get_local_linear_id()] = res; + }); + }); + } + bool result = res[0]; + for (size_t j = 1; j < work_group_size; ++j) result &= res[j]; + + std::string work_group = sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group); + INFO("Value of " << test_name << " with T = " << type_name() << " is " + << (result ? "right" : "wrong")); + CHECK(result); + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp new file mode 100644 index 000000000..a54497bfc --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp @@ -0,0 +1,39 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_permute.h" + +TEST_CASE("Non-uniform-group permute", + "[oneapi_non_uniform_groups][group_func][fp16]") { + auto queue = once_per_unit::get_queue(); + + if (queue.get_device().has(sycl::aspect::fp16)) { + permute_non_uniform_group, + sycl::half>(queue); + permute_non_uniform_group, + sycl::half>(queue); + permute_non_uniform_group, + sycl::half>(queue); + permute_non_uniform_group( + queue); + } else { + WARN("Device does not support half precision floating point operations."); + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp new file mode 100644 index 000000000..4d97f760a --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp @@ -0,0 +1,38 @@ +/******************************************************************************* +// +// 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_permute.h" + +TEST_CASE("Non-uniform-group permute", + "[oneapi_non_uniform_groups][group_func][fp64]") { + auto queue = once_per_unit::get_queue(); + + if (queue.get_device().has(sycl::aspect::fp64)) { + permute_non_uniform_group, + double>(queue); + permute_non_uniform_group, + double>(queue); + permute_non_uniform_group, + double>(queue); + permute_non_uniform_group(queue); + } else { + WARN("Device does not support double precision floating point operations."); + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_reduce.h b/tests/extension/oneapi_non_uniform_groups/group_reduce.h new file mode 100644 index 000000000..ee6d28b0a --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_reduce.h @@ -0,0 +1,560 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_functions/group_functions_common.h" +#include "non_uniform_group_common.h" + +constexpr size_t init = 8; +constexpr size_t test_size = 8; + +template +size_t get_reduce_reference(IteratorT first, IteratorT end) { + // Cast `init` to size_t so that guards are introduced in verification + if constexpr (with_init) + return std::accumulate(first, end, size_t(init), OpT()); + else + return std::accumulate(first + 1, end, size_t(*first), OpT()); +} + +template +void result_verifier(const std::vector& v_input, + const std::vector& v_output, + const std::vector& sg_ids, + const std::vector& nug_ids) { + std::map, OutputT> reference_results; + auto op = OpT(); + + for (size_t i = 0; i < sg_ids.size(); ++i) { + uint32_t sg_id = sg_ids[i]; + // Max values indicate items not participating. + if (sg_id == std::numeric_limits::max()) continue; + + uint32_t nug_id = nug_ids[i]; + auto key = std::make_pair(sg_id, nug_id); + InputT input = v_input[i]; + + auto iter = reference_results.find(key); + if (iter == reference_results.end()) { + // First may need to apply init value. + OutputT value = with_init ? op(InputT(init), input) : OutputT(input); + reference_results.emplace(std::make_pair(key, value)); + } else { + iter->second = op(iter->second, input); + } + } + + bool res = false; + for (size_t i = 0; i < sg_ids.size(); ++i) { + uint32_t sg_id = sg_ids[i]; + // Max values indicate items not participating. + if (sg_id == std::numeric_limits::max()) continue; + + uint32_t nug_id = nug_ids[i]; + auto key = std::make_pair(sg_id, nug_id); + + OutputT expected = reference_results[key]; + OutputT actual = v_output[i]; + + if (expected > util::exact_max) continue; + + INFO("Verifying reduction result of element with sub-group ID " + + std::to_string(sg_id) + " and non-uniform group ID " + + std::to_string(nug_id)); + CHECK(expected == actual); + } +} + +template +class joint_reduce_group_kernel; + +/** + * @brief Provides test for joint reduce by group + * @tparam GroupT Non-uniform group type to use for testing + * @tparam T Type for reduced values + * @tparam OpT Type for binary operator + */ +template +void joint_reduce_group(sycl::queue& queue, const std::string& op_name) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing joint_reduce_group for " + group_name + " and " + op_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + const std::string test_name = + "std::iterator_traits::value_type joint_reduce(GroupT g, Ptr first, " + "Ptr last, BinaryOperation binary_op)"; + + sycl::range<1> work_group_range = + sycl_cts::util::work_group_range<1>(queue, test_size); + size_t work_group_size = work_group_range.size(); + + const size_t sizes[3] = {4, work_group_size / 2, 2 * work_group_size}; + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + for (size_t size : sizes) { + std::vector v(size); + std::iota(v.begin(), v.end(), 1); + + // array to return results + std::vector res(work_group_size); + // participation markers + std::vector participating(work_group_size, 0); + { + sycl::buffer v_sycl(v.data(), sycl::range<1>(size)); + sycl::buffer res_sycl(res.data(), + sycl::range<1>(work_group_size)); + sycl::buffer participating_sycl( + participating.data(), sycl::range<1>(work_group_size)); + + queue.submit([&](sycl::handler& cgh) { + auto v_acc = + v_sycl.template get_access(cgh); + auto res_acc = + res_sycl.template get_access(cgh); + auto participating_acc = + participating_sycl.get_access( + cgh); + + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + cgh.parallel_for>( + executionRange, [=](sycl::nd_item<1> item) { + size_t index = item.get_global_linear_id(); + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate( + sub_group, test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + participating_acc[index] = 1; + + T* v_begin = v_acc.get_pointer(); + T* v_end = v_begin + v_acc.size(); + + ASSERT_RETURN_TYPE( + T, + sycl::joint_reduce(non_uniform_group, v_begin, v_end, + OpT()), + "Return type of joint_reduce(GroupT g, Ptr first, Ptr " + "last, BinaryOperation binary_op) is wrong\n"); + + res_acc[index] = sycl::joint_reduce(non_uniform_group, v_begin, + v_end, OpT()); + }); + }); + } + + const auto expected = + get_reduce_reference(v.begin(), v.end()); + + if (expected <= util::exact_max) { + for (size_t i = 0; i < work_group_size; ++i) { + if (!participating[i]) continue; + + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group, size, i); + INFO("Verifying value of " + << test_name << " with " << op_name + << " operation and Ptr = " << type_name() << "*"); + CHECK(res[i] == expected); + } + } + } + } +} + +template +class invoke_joint_reduce_group { + public: + void operator()(sycl::queue& queue, const std::string& op_name) { + if constexpr (type_traits::group_algorithms::is_legal_operator_v< + T, OperatorT>) { + joint_reduce_group(queue, op_name); + } + } +}; + +template +class init_joint_reduce_group_kernel; + +/** + * @brief Provides test for joint reduce by group with init + * @tparam GroupT Non-uniform group type to use for testing + * @tparam T Type for init and result values + * @tparam U Type for reduced values + * @tparam OpT Type for binary operator + */ +template +void init_joint_reduce_group(sycl::queue& queue, const std::string& op_name) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing joint_reduce_group with init for " + group_name + " and " + + op_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + const std::string test_name = + "T joint_reduce(GroupT g, Ptr first, Ptr last, T init, " + "BinaryOperation binary_op)"; + + sycl::range<1> work_group_range = + sycl_cts::util::work_group_range<1>(queue, test_size); + size_t work_group_size = work_group_range.size(); + + const size_t sizes[3] = {4, work_group_size / 2, 2 * work_group_size}; + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + for (size_t size : sizes) { + std::vector v(size); + std::iota(v.begin(), v.end(), 1); + + // array to return results + std::vector res(work_group_size); + // participation markers + std::vector participating(work_group_size, 0); + { + sycl::buffer v_sycl(v.data(), sycl::range<1>(size)); + sycl::buffer res_sycl(res.data(), + sycl::range<1>(work_group_size)); + sycl::buffer participating_sycl( + participating.data(), sycl::range<1>(work_group_size)); + + queue.submit([&](sycl::handler& cgh) { + auto v_acc = + v_sycl.template get_access(cgh); + auto res_acc = + res_sycl.template get_access(cgh); + auto participating_acc = + participating_sycl.get_access( + cgh); + + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + cgh.parallel_for>( + executionRange, [=](sycl::nd_item<1> item) { + size_t index = item.get_global_linear_id(); + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate( + sub_group, test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + participating_acc[index] = 1; + + U* v_begin = v_acc.get_pointer(); + U* v_end = v_begin + v_acc.size(); + + ASSERT_RETURN_TYPE( + T, + sycl::joint_reduce(non_uniform_group, v_begin, v_end, + T(init), OpT()), + "Return type of joint_reduce(GroupT g, Ptr first, Ptr " + "last, T init, BinaryOperation binary_op) is wrong\n"); + + res_acc[index] = sycl::joint_reduce(non_uniform_group, v_begin, + v_end, T(init), OpT()); + }); + }); + } + const auto expected = get_reduce_reference(v.begin(), v.end()); + + if (expected <= util::exact_max) { + for (size_t i = 0; i < work_group_size; ++i) { + if (!participating[i]) continue; + + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group, size, i); + INFO("Verifying value of " + << test_name << " with " << op_name + << " operation and Ptr = " << type_name() << "*"); + CHECK(res[i] == expected); + } + } + } + } +} + +template +class invoke_init_joint_reduce_group { + public: + void operator()(sycl::queue& queue, const std::string& op_name) { + if constexpr (type_traits::group_algorithms::is_legal_operator_v< + RetT, OperatorT>) { + init_joint_reduce_group(queue, + op_name); + } + } +}; + +template +class reduce_over_group_kernel; + +/** + * @brief Provides test for reduce over group values + * @tparam GroupT Non-uniform group type to use for testing + * @tparam T Type for reduced values + * @tparam OpT Type for binary operator + */ +template +void reduce_over_group(sycl::queue& queue, const std::string& op_name) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing reduce_over_group for " + group_name + " and " + op_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + const std::string test_name = + "T reduce_over_group(GroupT g, T x, BinaryOperation binary_op)"; + + sycl::range<1> work_group_range = + sycl_cts::util::work_group_range<1>(queue, test_size); + size_t work_group_size = work_group_range.size(); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + bool res = false; + // array to input data + std::vector v(work_group_size); + std::iota(v.begin(), v.end(), 1); + // array to reduce results + std::vector nug_output(work_group_size, 0); + // Sub-group ID and non-uniform group ID (Max int means the item is not + // participating in a reduction) + std::vector sg_id(work_group_size, + std::numeric_limits::max()); + std::vector nug_id(work_group_size, + std::numeric_limits::max()); + + { + sycl::buffer v_sycl(v.data(), sycl::range<1>(work_group_size)); + sycl::buffer nug_output_sycl(nug_output.data(), + sycl::range<1>(work_group_size)); + sycl::buffer sg_id_sycl(sg_id.data(), + sycl::range<1>(work_group_size)); + sycl::buffer nug_id_sycl(nug_id.data(), + sycl::range<1>(work_group_size)); + + queue.submit([&](sycl::handler& cgh) { + auto v_acc = v_sycl.template get_access(cgh); + auto nug_output_acc = + nug_output_sycl.template get_access( + cgh); + auto sg_id_acc = + sg_id_sycl.template get_access(cgh); + auto nug_id_acc = + nug_id_sycl.template get_access( + cgh); + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + cgh.parallel_for>( + executionRange, [=](sycl::nd_item<1> item) { + size_t index = item.get_global_linear_id(); + + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate(sub_group, + test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + + sg_id_acc[index] = sub_group.get_group_linear_id(); + nug_id_acc[index] = non_uniform_group.get_group_linear_id(); + + ASSERT_RETURN_TYPE(T, + sycl::reduce_over_group(non_uniform_group, + v_acc[index], OpT()), + "Return type of reduce_over_group(GroupT g, " + "T x, BinaryOperation binary_op) is wrong\n"); + nug_output_acc[index] = sycl::reduce_over_group( + non_uniform_group, v_acc[index], OpT()); + }); + }); + } + + // Verify return value for reduce_over_group on GroupT + { + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group); + result_verifier(v, nug_output, sg_id, nug_id); + } + } +} + +template +class invoke_reduce_over_group { + public: + void operator()(sycl::queue& queue, const std::string& op_name) { + if constexpr (type_traits::group_algorithms::is_legal_operator_v< + T, OperatorT>) { + reduce_over_group(queue, op_name); + } + } +}; + +template +class init_reduce_over_group_kernel; + +/** + * @brief Provides test for reduce over group values with init + * @tparam GroupT Non-uniform group type to use for testing + * @tparam T Type for init and result values + * @tparam U Type for group values + * @tparam OpT Type for binary operator + */ +template +void init_reduce_over_group(sycl::queue& queue, const std::string& op_name) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing reduce_over_group with init for " + group_name + " and " + + op_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + const std::string test_name = + "T reduce_over_group(GroupT g, V x, T init, BinaryOperation binary_op)"; + + sycl::range<1> work_group_range = + sycl_cts::util::work_group_range<1>(queue, test_size); + size_t work_group_size = work_group_range.size(); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + bool res = false; + // array to input data + std::vector v(work_group_size); + std::iota(v.begin(), v.end(), 1); + // array to reduce results + std::vector nug_output(work_group_size, 0); + // Sub-group ID and non-uniform group ID (Max int means the item is not + // participating in a reduction) + std::vector sg_id(work_group_size, + std::numeric_limits::max()); + std::vector nug_id(work_group_size, + std::numeric_limits::max()); + + { + sycl::buffer v_sycl(v.data(), sycl::range<1>(work_group_size)); + sycl::buffer nug_output_sycl(nug_output.data(), + sycl::range<1>(work_group_size)); + sycl::buffer sg_id_sycl(sg_id.data(), + sycl::range<1>(work_group_size)); + sycl::buffer nug_id_sycl(nug_id.data(), + sycl::range<1>(work_group_size)); + + queue.submit([&](sycl::handler& cgh) { + auto v_acc = v_sycl.template get_access(cgh); + auto nug_output_acc = + nug_output_sycl.template get_access( + cgh); + auto sg_id_acc = + sg_id_sycl.template get_access(cgh); + auto nug_id_acc = + nug_id_sycl.template get_access( + cgh); + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + cgh.parallel_for>( + executionRange, [=](sycl::nd_item<1> item) { + size_t index = item.get_global_linear_id(); + + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate(sub_group, + test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + + sg_id_acc[index] = sub_group.get_group_linear_id(); + nug_id_acc[index] = non_uniform_group.get_group_linear_id(); + + ASSERT_RETURN_TYPE( + T, + sycl::reduce_over_group(non_uniform_group, v_acc[index], + T(init), OpT()), + "Return type of reduce_over_group(GroupT g, V x, T init, " + "BinaryOperation binary_op) is wrong\n"); + nug_output_acc[index] = sycl::reduce_over_group( + non_uniform_group, v_acc[index], T(init), OpT()); + }); + }); + } + + // Verify return value for reduce_over_group on GroupT + { + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group); + result_verifier(v, nug_output, sg_id, nug_id); + } + } +} + +template +class invoke_init_reduce_over_group { + public: + void operator()(sycl::queue& queue, const std::string& op_name) { + if constexpr (type_traits::group_algorithms::is_legal_operator_v< + RetT, OperatorT>) { + init_reduce_over_group(queue, op_name); + } + } +}; diff --git a/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in new file mode 100644 index 000000000..ef51ca672 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in @@ -0,0 +1,88 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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" + +// clang-format off +#cmakedefine CTS_TYPE @CTS_TYPE@ +#cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") +// clang-format on +using ReduceTypes = Types; + +TEST_CASE(CTS_TYPE_NAME + " non-uniform group reduce functions", + "[oneapi_non_uniform_groups][group_func][type_list]") { + auto queue = once_per_unit::get_queue(); + // Get binary operators from TestType + const auto Operators = get_op_types(); + const auto RetType = unnamed_type_pack(); + const auto GroupTypes = unnamed_type_pack< + oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::tangle_group, + oneapi_ext::opportunistic_group>(); + + 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(GroupTypes, RetType, Operators, + queue); +} + +TEMPLATE_LIST_TEST_CASE(CTS_TYPE_NAME + + " non-uniform group reduce functions with init", + "[oneapi_non_uniform_groups][group_func][type_list]", + ReduceTypes) { + auto queue = once_per_unit::get_queue(); + + // 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 GroupTypes = unnamed_type_pack< + oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::tangle_group, + oneapi_ext::opportunistic_group>(); + + 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( + GroupTypes, RetType, ReducedType, Operators, queue); +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_scan.h b/tests/extension/oneapi_non_uniform_groups/group_scan.h new file mode 100644 index 000000000..98d84534c --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_scan.h @@ -0,0 +1,594 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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 + +#include "../../group_functions/group_functions_common.h" +#include "non_uniform_group_common.h" + +template +class joint_scan_group_kernel; + +// This should never be higher than std::numeric_limits::max() for the +// smallest type tested. Currently, the smallest type tested is +// char/int8_t, so it shouldn't be higher than 127. +constexpr int init = 42; +constexpr size_t test_size = 12; + +template +auto joint_inclusive_scan_helper(Group group, T* v_begin, T* v_end, + U* r_i_begin, OpT op, bool with_init) { + if (with_init) { + return sycl::joint_inclusive_scan(group, v_begin, v_end, r_i_begin, op, + I(init)); + } + assert((std::is_same_v && + "Without init value I and U should be the same type.")); + return (U*)sycl::joint_inclusive_scan(group, v_begin, v_end, (I*)r_i_begin, + op); +} + +template +auto joint_exclusive_scan_helper(Group group, T* v_begin, T* v_end, + U* r_e_begin, OpT op, bool with_init) { + if (with_init) { + return sycl::joint_exclusive_scan(group, v_begin, v_end, r_e_begin, I(init), + op); + } + assert((std::is_same_v && + "Without init value I and U should be the same type.")); + return (U*)sycl::joint_exclusive_scan(group, v_begin, v_end, (I*)r_e_begin, + op); +} + +template +struct JointScanDataStruct { + JointScanDataStruct(size_t range_size, OpT op, bool with_init) + : ref_input(range_size), res(range_size * 2, U(-1)) { + std::iota(ref_input.begin(), ref_input.end(), T(1)); + if constexpr (std::is_same_v> || + std::is_same_v>) { + auto identity = sycl::known_identity_v; + auto acc = with_init ? I{init} : identity; + for (size_t i = 0; i < range_size; ++i) { + I tmp = op(I(acc), I(ref_input[i])); + if (tmp > std::numeric_limits::max()) { + ref_input[i] = identity; + } + acc = op(acc, ref_input[i]); + } + } + } + + void check_results(size_t range_size, OpT op, const std::string& op_name, + const std::string& group_name, bool with_init) { + CHECK(end[0]); + CHECK(end[1]); + CHECK(ret_type[0]); + CHECK(ret_type[1]); + + I init_value = with_init ? I(init) : sycl::known_identity::value; + + std::vector reference_e(range_size, U(-1)); + std::vector reference_i(range_size, U(-1)); + std::exclusive_scan(ref_input.begin(), ref_input.end(), reference_e.begin(), + init_value, op); + std::inclusive_scan(ref_input.begin(), ref_input.end(), reference_i.begin(), + op, init_value); + for (int i = 0; i < range_size; i++) { + { + INFO("Check joint_exclusive_scan on " + group_name + " for element " + + std::to_string(i) + " (Operator: " + op_name + ")"); + INFO("Result: " + std::to_string(res[i])); + INFO("Expected: " + std::to_string(reference_e[i])); + CHECK(res[i] == reference_e[i]); + } + { + INFO("Check joint_inclusive_scan on " + group_name + " for element " + + std::to_string(i) + " (Operator: " + op_name + ")"); + INFO("Result: " + std::to_string(res[i + range_size])); + INFO("Expected: " + std::to_string(reference_i[i])); + CHECK(res[i + range_size] == reference_i[i]); + } + } + } + + sycl::buffer create_ref_input_buffer() { + return {ref_input.data(), ref_input.size()}; + } + + sycl::buffer create_res_buffer() { return {res.data(), res.size()}; } + + sycl::buffer create_end_buffer() { return {end, 2}; } + + sycl::buffer create_ret_type_buffer() { return {ret_type, 2}; } + + std::vector ref_input; + std::vector res; + bool end[2] = {false, false}; + bool ret_type[2] = {false, false}; + std::vector local_id; +}; + +template +void check_scan(sycl::queue& queue, size_t size, + sycl::nd_range<1> executionRange, OpT op, + const std::string& op_name, bool with_init) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + JointScanDataStruct host_data{size, op, with_init}; + { + sycl::buffer ref_input_sycl = host_data.create_ref_input_buffer(); + sycl::buffer res_sycl = host_data.create_res_buffer(); + sycl::buffer end_sycl = host_data.create_end_buffer(); + sycl::buffer ret_type_sycl = host_data.create_ret_type_buffer(); + + queue + .submit([&](sycl::handler& cgh) { + sycl::accessor ref_input_acc(ref_input_sycl, cgh); + sycl::accessor res_acc(res_sycl, cgh); + sycl::accessor end_acc(end_sycl, cgh); + sycl::accessor ret_type_acc(ret_type_sycl, cgh); + + cgh.parallel_for>( + executionRange, [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave + // early. + if (!NonUniformGroupHelper::should_participate( + sub_group, test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, + test_case); + + // Likewise, we only use a single group to compute these + // values. + if (sub_group.get_group_linear_id() != 0 || + non_uniform_group.get_group_linear_id() != + NonUniformGroupHelper:: + preferred_single_worker_group_id(test_case)) + return; + + T* v_begin = ref_input_acc.get_pointer(); + T* v_end = v_begin + ref_input_acc.size(); + + U* r_nug_e_begin = res_acc.get_pointer(); + U* r_nug_i_begin = res_acc.get_pointer() + size; + + auto r_nug_e_end = joint_exclusive_scan_helper( + non_uniform_group, v_begin, v_end, r_nug_e_begin, op, + with_init); + ret_type_acc[0] = std::is_same_v; + + auto r_nug_i_end = joint_inclusive_scan_helper( + non_uniform_group, v_begin, v_end, r_nug_i_begin, op, + with_init); + ret_type_acc[1] = std::is_same_v; + + end_acc[0] = (r_nug_e_begin + size == r_nug_e_end); + end_acc[1] = (r_nug_i_begin + size == r_nug_i_end); + }); + }) + .wait_and_throw(); + } + + host_data.check_results(size, op, op_name, group_name, with_init); + } +} + +/** + * @brief Provides test for joint scans + * @tparam GroupT Group type to test with + * @tparam T Type pointed by InPtr + * @tparam U Type pointed by OutPtr + * @tparam OperatorT Type of binary operation + */ +template +struct joint_scan_group { + void operator()(sycl::queue& queue, const std::string& op_name) { + if constexpr (type_traits::group_algorithms::is_legal_operator_v< + U, OperatorT>) { + INFO(" with types " + type_name() + " and " + type_name()); + + sycl::range<1> work_group_range = + sycl_cts::util::work_group_range<1>(queue, test_size); + + size_t work_group_size = work_group_range.size(); + + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + const size_t sizes[2] = {5, 2}; + for (size_t size : sizes) { + check_scan(queue, size, executionRange, OperatorT(), + op_name, false); + } + } + } +}; + +template +class invoke_joint_scan_group { + public: + void operator()(sycl::queue& queue) { + const auto operators = get_op_types(); + for_all_combinations(operators, queue); + } +}; + +// FIXME: Helper for implementations that cannot handle cases of different types +template +class invoke_joint_scan_group_same_type { + public: + void operator()(sycl::queue& queue) { + const auto operators = get_op_types(); + for_all_combinations(operators, queue); + } +}; + +template +class init_joint_scan_group_kernel; + +/** + * @brief Provides test for joint scans with init + * @tparam GroupT Group type to test with + * @tparam T Type pointed by InPtr + * @tparam U Type pointed by OutPtr + * @tparam I Type used for init value + * @tparam OperatorT Type of binary operation + */ +template +struct init_joint_scan_group { + void operator()(sycl::queue& queue, const std::string& op_name) { + if constexpr (type_traits::group_algorithms::is_legal_operator_v< + I, OperatorT>) { + INFO(" with types " + type_name() + " and " + type_name() + + ", init type " + type_name()); + + sycl::range<1> work_group_range = + sycl_cts::util::work_group_range<1>(queue, test_size); + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + size_t work_group_size = work_group_range.size(); + + const size_t sizes[2] = {5, 2}; + for (size_t size : sizes) { + check_scan(queue, size, executionRange, OperatorT(), + op_name, true); + } + } + } +}; + +template +class invoke_init_joint_scan_group { + public: + void operator()(sycl::queue& queue) { + const auto operators = get_op_types(); + for_all_combinations(operators, + queue); + } +}; + +// FIXME: Helper for implementations that cannot handle cases of different types +template +class invoke_init_joint_scan_group_same_type { + public: + void operator()(sycl::queue& queue) { + const auto operators = get_op_types(); + for_all_combinations(operators, + queue); + } +}; + +template +class scan_over_group_kernel; + +template +auto inclusive_scan_over_group_helper(Group group, U x, OpT op, + bool with_init) { + if (with_init) { + return sycl::inclusive_scan_over_group(group, x, op, T(init)); + } + assert((std::is_same_v && + "Without init value T and U should be the same type.")); + return sycl::inclusive_scan_over_group(group, T(x), op); +} + +template +auto exclusive_scan_over_group_helper(Group group, U x, OpT op, + bool with_init) { + if (with_init) { + return sycl::exclusive_scan_over_group(group, x, T(init), op); + } + assert((std::is_same_v && + "Without init value T and U should be the same type.")); + return sycl::exclusive_scan_over_group(group, T(x), op); +} + +template +struct ScanOverGroupDataStruct { + ScanOverGroupDataStruct(size_t range_size) + : ref_input(range_size), + res(range_size * 2, T(-1)), + local_id(range_size, 0), + sub_group_id(range_size, 0), + non_uniform_group_id(range_size, 0) { + std::iota(ref_input.begin(), ref_input.end(), U(1)); + std::fill(local_id.begin(), local_id.end(), + std::numeric_limits::max()); + std::fill(sub_group_id.begin(), sub_group_id.end(), + std::numeric_limits::max()); + std::fill(non_uniform_group_id.begin(), non_uniform_group_id.end(), + std::numeric_limits::max()); + } + + template + void check_results(size_t range_size, OpT op, const std::string& op_name, + const std::string& group_name, bool with_init) { + CHECK(ret_type[0]); + CHECK(ret_type[1]); + + T init_value = with_init ? T(init) : sycl::known_identity::value; + { + // Mapping from "sub-group id" and "non-uniform group id" to "vector of + // input data (ordered by item linear id within the sub-group)" + std::map, std::vector> ref_input_per_group; + for (int i = 0; i < range_size; i++) { + size_t sgid = sub_group_id[i]; + // If sub-group id is max value it was an item not participating, so + // skip it. + if (sgid == std::numeric_limits::max()) continue; + + size_t nugid = non_uniform_group_id[i]; + auto key = std::make_pair(sgid, nugid); + std::vector& input_vec = ref_input_per_group[key]; + // Extend input vector dynamically. + size_t lid = local_id[i]; + if (input_vec.size() <= lid) input_vec.resize(lid + 1); + // Place the data identified by (sgid, lid). + input_vec[lid] = ref_input[i]; + } + // Compute the reference results and verify. + for (int i = 0; i < range_size; i++) { + size_t sgid = sub_group_id[i]; + // If sub-group id is max value it was an item not participating, so + // skip it. + if (sgid == std::numeric_limits::max()) continue; + + size_t nugid = non_uniform_group_id[i]; + auto key = std::make_pair(sgid, nugid); + const std::vector& input_vec = ref_input_per_group[key]; + // Scan over the first (lid + 1) elements of input_vec to obtain the + // result identified by i. + size_t lid = local_id[i]; + std::vector reference(lid + 1, T(-1)); + std::exclusive_scan(input_vec.begin(), input_vec.begin() + lid + 1, + reference.begin(), init_value, op); + { + int res_i = i; + INFO("Check exclusive_scan_over_group on " + group_name + + " for element " + std::to_string(i) + " (Operator: " + op_name + + ")"); + INFO("Result: " + std::to_string(res[i])); + INFO("Expected: " + std::to_string(reference[lid])); + CHECK(res[i] == reference[lid]); + } + std::inclusive_scan(input_vec.begin(), input_vec.begin() + lid + 1, + reference.begin(), op, init_value); + { + INFO("Check inclusive_scan_over_group on " + group_name + + " for element " + std::to_string(i) + " (Operator: " + op_name + + ")"); + INFO("Result: " + std::to_string(res[range_size + i])); + INFO("Expected: " + std::to_string(reference[lid])); + CHECK(res[range_size + i] == reference[lid]); + } + } + } + } + + sycl::buffer create_ref_input_buffer() { + return {ref_input.data(), ref_input.size()}; + } + + sycl::buffer create_res_buffer() { return {res.data(), res.size()}; } + + sycl::buffer create_ret_type_buffer() { return {ret_type, 4}; } + + sycl::buffer create_local_id_buffer() { + return {local_id.data(), local_id.size()}; + } + + sycl::buffer create_sub_group_id_buffer() { + return {sub_group_id.data(), sub_group_id.size()}; + } + + sycl::buffer create_non_uniform_group_id_buffer() { + return {non_uniform_group_id.data(), non_uniform_group_id.size()}; + } + + std::vector ref_input; + std::vector res; + bool ret_type[2] = {false, false}; + std::vector local_id; + std::vector sub_group_id; + std::vector non_uniform_group_id; +}; + +template +void check_scan_over_group(sycl::queue& queue, sycl::range<1> range, OpT op, + const std::string& op_name, bool with_init) { + const std::string group_name = NonUniformGroupHelper::get_name(); + auto range_size = range.size(); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + ScanOverGroupDataStruct host_data{range_size}; + { + auto ref_input_sycl = host_data.create_ref_input_buffer(); + auto res_sycl = host_data.create_res_buffer(); + auto ret_type_sycl = host_data.create_ret_type_buffer(); + auto local_id_sycl = host_data.create_local_id_buffer(); + auto sg_id_sycl = host_data.create_sub_group_id_buffer(); + auto nug_id_sycl = host_data.create_non_uniform_group_id_buffer(); + + queue + .submit([&](sycl::handler& cgh) { + sycl::accessor ref_input_acc( + ref_input_sycl, cgh); + sycl::accessor res_acc(res_sycl, cgh); + sycl::accessor ret_type_acc(ret_type_sycl, cgh); + sycl::accessor local_id_acc(local_id_sycl, cgh); + sycl::accessor sg_id_acc(sg_id_sycl, cgh); + sycl::accessor nug_id_acc(nug_id_sycl, cgh); + + cgh.parallel_for>( + sycl::nd_range<1>(range, range), [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave + // early. + if (!NonUniformGroupHelper::should_participate( + sub_group, test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, + test_case); + + auto g_index = item.get_global_linear_id(); + + local_id_acc[g_index] = + non_uniform_group.get_local_linear_id(); + sg_id_acc[g_index] = sub_group.get_group_linear_id(); + nug_id_acc[g_index] = non_uniform_group.get_group_linear_id(); + + auto res_nug_e = exclusive_scan_over_group_helper( + non_uniform_group, ref_input_acc[g_index], op, with_init); + res_acc[g_index] = res_nug_e; + ret_type_acc[0] = std::is_same_v; + + auto res_nug_i = inclusive_scan_over_group_helper( + non_uniform_group, ref_input_acc[g_index], op, with_init); + res_acc[range_size + g_index] = res_nug_i; + ret_type_acc[1] = std::is_same_v; + }); + }) + .wait_and_throw(); + } + + host_data.check_results(range_size, op, op_name, group_name, with_init); + } +} + +/** + * @brief Provides test for scans over group values + * @tparam GroupT Group type to test with + * @tparam T Type used for value + * @tparam OperatorT Type of binary operation + */ +template +struct scan_over_group { + void operator()(sycl::queue& queue, const std::string& op_name) { + if constexpr (type_traits::group_algorithms::is_legal_operator_v< + T, OperatorT>) { + INFO(" with type " + type_name()); + + sycl::range<1> work_group_range = + sycl_cts::util::work_group_range<1>(queue, test_size); + size_t work_group_size = work_group_range.size(); + + check_scan_over_group(queue, work_group_range, OperatorT(), + op_name, false); + } + } +}; + +template +class invoke_scan_over_group { + public: + void operator()(sycl::queue& queue) { + const auto operators = get_op_types(); + for_all_combinations(operators, queue); + } +}; + +template +class init_scan_over_group_kernel; + +// many errors with short types for hipSYCL +// it means conversion and calculation patterns are not OK +/** + * @brief Provides test for scans over group with an init value + * @tparam GroupT Group type to test with + * @tparam T Type used for init value and result + * @tparam U Type used for group values + * @tparam OperatorT Type of binary operation + */ +template +struct init_scan_over_group { + void operator()(sycl::queue& queue, const std::string& op_name) { + if constexpr (type_traits::group_algorithms::is_legal_operator_v< + T, OperatorT>) { + INFO(" with types " + type_name() + " and " + type_name()); + + sycl::range<1> work_group_range = + sycl_cts::util::work_group_range<1>(queue, test_size); + + check_scan_over_group(queue, work_group_range, OperatorT(), + op_name, true); + } + } +}; + +template +class invoke_init_scan_over_group { + public: + void operator()(sycl::queue& queue) { + const auto operators = get_op_types(); + for_all_combinations(operators, queue); + } +}; + +// FIXME: Helper for implementations that cannot handle cases of different types +template +class invoke_init_scan_over_group_same_type { + public: + void operator()(sycl::queue& queue) { + const auto operators = get_op_types(); + for_all_combinations(operators, queue); + } +}; diff --git a/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in new file mode 100644 index 000000000..b0afdacc2 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in @@ -0,0 +1,76 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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. +// +*******************************************************************************/ + +// clang-format off +#cmakedefine CTS_TYPE @CTS_TYPE@ +#cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") +// clang-format on + +#include "group_scan.h" + +using TestType = unnamed_type_pack; +using ScanTypes = Types; + +TEST_CASE(CTS_TYPE_NAME + " non-uniform group scan functions", + "[oneapi_non_uniform_groups][group_func][type_list]"){ + auto queue = once_per_unit::get_queue(); + const auto GroupTypes = unnamed_type_pack< + oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::tangle_group, + oneapi_ext::opportunistic_group>(); + + 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(GroupTypes, TestType{}, queue); +}; + +TEST_CASE(CTS_TYPE_NAME + " non-uniform group scan functions with init", + "[oneapi_non_uniform_groups][group_func][type_list]"){ + auto queue = once_per_unit::get_queue(); + const auto GroupTypes = unnamed_type_pack< + oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::tangle_group, + oneapi_ext::opportunistic_group>(); + + 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(GroupTypes, TestType{}, + ScanTypes{}, queue); +}; diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift.cpp b/tests/extension/oneapi_non_uniform_groups/group_shift.cpp new file mode 100644 index 000000000..efa8c2a41 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_shift.cpp @@ -0,0 +1,36 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_shift.h" + +// errors in hipSYCL with bool and 8-bit types - only in group shifts +TEMPLATE_LIST_TEST_CASE("Non-uniform-group shift", + "[oneapi_non_uniform_groups][group_func][type_list]", + CustomTypes) { + auto queue = once_per_unit::get_queue(); + + shift_non_uniform_group, TestType>( + queue); + shift_non_uniform_group, + TestType>(queue); + shift_non_uniform_group, TestType>( + queue); + shift_non_uniform_group(queue); +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift.h b/tests/extension/oneapi_non_uniform_groups/group_shift.h new file mode 100644 index 000000000..db39cd3b3 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_shift.h @@ -0,0 +1,149 @@ +/******************************************************************************* +// +// 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 + +#include "../../group_functions/group_functions_common.h" +#include "non_uniform_group_common.h" + +template +class shift_non_uniform_group_kernel; + +template +void shift_non_uniform_group(sycl::queue& queue) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing permute for " + group_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + // 4 functions + constexpr int test_matrix = 4; + const std::string test_names[test_matrix] = { + "T shift_group_left(GroupT g, T x)", + "T shift_group_left(GroupT g, T x, GroupT::linear_id_type delta)", + "T shift_group_right(GroupT g, T x)", + "T shift_group_right(GroupT g, T x, GroupT::linear_id_type delta)"}; + + sycl::range<1> work_group_range = sycl_cts::util::work_group_range<1>(queue); + size_t work_group_size = work_group_range.size(); + + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + // array to return results: + std::valarray res(false, test_matrix * work_group_size); + { + sycl::buffer res_sycl( + std::begin(res), sycl::range<1>(test_matrix * work_group_size)); + + queue.submit([&](sycl::handler& cgh) { + auto res_acc = res_sycl.get_access(cgh); + + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + cgh.parallel_for>( + executionRange, [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, they fill their + // elements in the result with true and leave early. + if (!NonUniformGroupHelper::should_participate( + sub_group, test_case)) { + res_acc[0 * work_group_size + item.get_local_linear_id()] = + true; + res_acc[1 * work_group_size + item.get_local_linear_id()] = + true; + res_acc[2 * work_group_size + item.get_local_linear_id()] = + true; + res_acc[3 * work_group_size + item.get_local_linear_id()] = + true; + return; + } + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + const typename GroupT::linear_id_type llid = + non_uniform_group.get_local_linear_id(); + + T local_var(splat_init(llid + 1)); + T shifted_var(splat_init(llid + 1)); + + ASSERT_RETURN_TYPE( + T, sycl::shift_group_left(non_uniform_group, local_var), + "Return type of shift_group_left(GroupT g, T x) is wrong\n"); + + shifted_var = + sycl::shift_group_left(non_uniform_group, local_var); + res_acc[0 * work_group_size + item.get_local_linear_id()] = + equal(shifted_var, splat_init(llid + 2)) || + (llid + 1 >= non_uniform_group.get_local_linear_range()); + + ASSERT_RETURN_TYPE( + T, sycl::shift_group_left(non_uniform_group, local_var, 3), + "Return type of shift_group_left(GroupT g, T x, " + "GroupT::linear_id_type delta) is wrong\n"); + + shifted_var = + sycl::shift_group_left(non_uniform_group, local_var, 3); + res_acc[1 * work_group_size + item.get_local_linear_id()] = + equal(shifted_var, splat_init(llid + 4)) || + (llid + 3 >= non_uniform_group.get_local_linear_range()); + + ASSERT_RETURN_TYPE( + T, sycl::shift_group_right(non_uniform_group, local_var), + "Return type of shift_group_right(GroupT g, T x) is wrong\n"); + + shifted_var = + sycl::shift_group_right(non_uniform_group, local_var); + res_acc[2 * work_group_size + item.get_local_linear_id()] = + equal(shifted_var, splat_init(llid)) || (llid < 1); + + ASSERT_RETURN_TYPE( + T, sycl::shift_group_right(non_uniform_group, local_var, 2), + "Return type of shift_group_right(GroupT g, T x, " + "GroupT::linear_id_type delta) is wrong\n"); + + shifted_var = + sycl::shift_group_right(non_uniform_group, local_var, 2); + res_acc[3 * work_group_size + item.get_local_linear_id()] = + equal(shifted_var, splat_init(llid - 1)) || (llid < 2); + }); + }); + } + for (int i = 0; i < test_matrix; ++i) { + bool result = res[i * work_group_size]; + for (size_t j = 1; j < work_group_size; ++j) + result &= res[i * work_group_size + j]; + + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group); + INFO("Value of " << test_names[i] << " with T = " << type_name() + << " is " << (result ? "right" : "wrong")); + CHECK(result); + } + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp new file mode 100644 index 000000000..48c250e17 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp @@ -0,0 +1,38 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_shift.h" + +TEST_CASE("Non-uniform-group shift", + "[oneapi_non_uniform_groups][group_func][fp16]") { + auto queue = once_per_unit::get_queue(); + + if (queue.get_device().has(sycl::aspect::fp16)) { + shift_non_uniform_group, + sycl::half>(queue); + shift_non_uniform_group, + sycl::half>(queue); + shift_non_uniform_group, + sycl::half>(queue); + shift_non_uniform_group(queue); + } else { + WARN("Device does not support half precision floating point operations."); + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp new file mode 100644 index 000000000..bdb72faef --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp @@ -0,0 +1,38 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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_shift.h" + +TEST_CASE("Non-uniform-group shift", + "[oneapi_non_uniform_groups][group_func][fp64]") { + auto queue = once_per_unit::get_queue(); + + if (queue.get_device().has(sycl::aspect::fp64)) { + shift_non_uniform_group, double>( + queue); + shift_non_uniform_group, + double>(queue); + shift_non_uniform_group, double>( + queue); + shift_non_uniform_group(queue); + } else { + WARN("Device does not support double precision floating point operations."); + } +} diff --git a/tests/extension/oneapi_non_uniform_groups/is_fixed_topology_group.cpp b/tests/extension/oneapi_non_uniform_groups/is_fixed_topology_group.cpp new file mode 100644 index 000000000..28d1f94ae --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/is_fixed_topology_group.cpp @@ -0,0 +1,59 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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 non_uniform_groups::tests { + +TEST_CASE("Test for is_fixed_topology_group trait with existing groups.", + "[oneapi_non_uniform_groups]") { +#ifndef SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS + SKIP("SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS is not defined"); +#else + namespace oneapi_ext = sycl::ext::oneapi::experimental; + + STATIC_CHECK(oneapi_ext::is_fixed_topology_group>::value); + STATIC_CHECK(oneapi_ext::is_fixed_topology_group_v>); + STATIC_CHECK(oneapi_ext::is_fixed_topology_group>::value); + STATIC_CHECK(oneapi_ext::is_fixed_topology_group_v>); + STATIC_CHECK(oneapi_ext::is_fixed_topology_group>::value); + STATIC_CHECK(oneapi_ext::is_fixed_topology_group_v>); + + STATIC_CHECK(oneapi_ext::is_fixed_topology_group::value); + STATIC_CHECK(oneapi_ext::is_fixed_topology_group_v); + +#ifdef SYCL_EXT_ONEAPI_ROOT_GROUP + STATIC_CHECK( + oneapi_ext::is_fixed_topology_group>::value); + STATIC_CHECK( + oneapi_ext::is_fixed_topology_group_v>); + STATIC_CHECK( + oneapi_ext::is_fixed_topology_group>::value); + STATIC_CHECK( + oneapi_ext::is_fixed_topology_group_v>); + STATIC_CHECK( + oneapi_ext::is_fixed_topology_group>::value); + STATIC_CHECK( + oneapi_ext::is_fixed_topology_group_v>); +#endif +#endif +} + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h b/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h new file mode 100644 index 000000000..69b819f91 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h @@ -0,0 +1,207 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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 oneapi_ext = sycl::ext::oneapi::experimental; + +// Helper class for working with non-uniform group of type GroupT. If the +// result is empty the work-item does not participate in the execution. +template +struct NonUniformGroupHelper; + +template <> +struct NonUniformGroupHelper> { + static constexpr size_t num_test_cases = 4; + + static bool is_supported(const sycl::device& dev) { + return dev.has(sycl::aspect::ext_oneapi_ballot_group); + } + + static bool should_participate(sycl::sub_group sg, int test_case) { + return true; + } + + static oneapi_ext::ballot_group create(sycl::sub_group sg, + int test_case) { + assert(test_case < num_test_cases); + // Split it so that 1/3rd of the items are in the first "true" group and the + // rest are in "false" group. + switch (test_case) { + case 0: + return oneapi_ext::get_ballot_group( + sg, sg.get_local_linear_id() < sg.get_local_range().size() / 3); + case 1: + return oneapi_ext::get_ballot_group(sg, sg.get_local_linear_id() & 1); + case 2: + return oneapi_ext::get_ballot_group(sg, true); + case 3: + return oneapi_ext::get_ballot_group(sg, false); + } + return oneapi_ext::get_ballot_group(sg, false); + } + + static size_t preferred_single_worker_group_id(int test_case) { + // Some work requires us to pick a single work-group to do work. Generally + // we would pick group 0, but in case 2 it is empty so we pick 1 for that + // instead. + return test_case == 2; + } + + static std::string get_name() { return "ballot_group"; } + + static std::string get_test_case_name(int test_case) { + assert(test_case < num_test_cases); + switch (test_case) { + case 0: + return "predicate is true for first N items."; + case 1: + return "predicate is true for all work-items with odd local id"; + case 2: + return "predicate is true for all work-items"; + case 3: + return "predicate is false for all work-items"; + } + return ""; + } +}; + +template +struct NonUniformGroupHelper< + oneapi_ext::fixed_size_group> { + static constexpr size_t num_test_cases = 1; + + static bool is_supported(const sycl::device& dev) { + return dev.has(sycl::aspect::ext_oneapi_fixed_size_group); + } + + static bool should_participate(sycl::sub_group sg, int test_case) { + return true; + } + + static oneapi_ext::fixed_size_group create( + sycl::sub_group sg, int test_case) { + return oneapi_ext::get_fixed_size_group(sg); + } + + static size_t preferred_single_worker_group_id(int) { + return 0; + } + + static std::string get_name() { + return "fixed_size_group<" + std::to_string(PartitionSize) + + ", sycl::sub_group>"; + } + + static std::string get_test_case_name(int) { + return "testing fixed_size_group"; + } +}; + +template <> +struct NonUniformGroupHelper> { + static constexpr size_t num_test_cases = 3; + + static bool is_supported(const sycl::device& dev) { + return dev.has(sycl::aspect::ext_oneapi_tangle_group); + } + + static bool should_participate(sycl::sub_group sg, int test_case) { + assert(test_case < num_test_cases); + switch (test_case) { + case 0: + return sg.get_local_linear_id() < sg.get_local_range().size() / 3; + case 1: + return sg.get_local_linear_id() & 1; + case 2: + return true; + } + return false; + } + + static oneapi_ext::tangle_group create(sycl::sub_group sg, + int test_case) { + return oneapi_ext::get_tangle_group(sg); + } + + static size_t preferred_single_worker_group_id(int) { + return 0; + } + + static std::string get_name() { return "tangle_group"; } + + static std::string get_test_case_name(int test_case) { + assert(test_case < num_test_cases); + switch (test_case) { + case 0: + return "predicate is true for first N items."; + case 1: + return "predicate is true for all work-items with odd local id"; + case 2: + return "predicate is true for all work-items"; + } + return ""; + } +}; + +template <> +struct NonUniformGroupHelper { + static constexpr size_t num_test_cases = 3; + + static bool is_supported(const sycl::device& dev) { + return dev.has(sycl::aspect::ext_oneapi_opportunistic_group); + } + + static bool should_participate(sycl::sub_group sg, int test_case) { + assert(test_case < num_test_cases); + switch (test_case) { + case 0: + return sg.get_local_linear_id() < sg.get_local_range().size() / 3; + case 1: + return sg.get_local_linear_id() & 1; + case 2: + return true; + } + return false; + } + + static oneapi_ext::opportunistic_group create(sycl::sub_group, int) { + return oneapi_ext::this_kernel::get_opportunistic_group(); + } + + static size_t preferred_single_worker_group_id(int) { + return 0; + } + + static std::string get_name() { return "opportunistic_group"; } + + static std::string get_test_case_name(int test_case) { + assert(test_case < num_test_cases); + switch (test_case) { + case 0: + return "predicate is true for first N items."; + case 1: + return "predicate is true for all work-items with odd local id"; + case 2: + return "predicate is true for all work-items"; + } + return ""; + } +}; diff --git a/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp new file mode 100644 index 000000000..cb686b805 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp @@ -0,0 +1,182 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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 non_uniform_groups::tests { + +struct checks { + enum { + get_group_id, + get_local_id, + get_group_range, + get_local_range, + get_group_linear_id, + get_local_linear_id, + get_group_linear_range, + get_local_linear_range, + leader, + COUNT, + }; +}; + +TEST_CASE("Test for opportunistic_group apis.", "[oneapi_non_uniform_groups]") { +#ifndef SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS + SKIP("SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS is not defined"); +#else + namespace oneapi_ext = sycl::ext::oneapi::experimental; + using opportunistic_group_t = oneapi_ext::opportunistic_group; + using CheckResults = bool[checks::COUNT]; + + constexpr size_t num_items = 64; + + sycl::buffer results_buffer{num_items}; + + auto q = sycl_cts::util::get_cts_object::queue(); + + if (!q.get_device().has(sycl::aspect::ext_oneapi_opportunistic_group)) { + SKIP("Device does not support opportunistic_group."); + } + + q.submit([&](sycl::handler& cgh) { + sycl::accessor acc{results_buffer, cgh, sycl::write_only}; + + sycl::nd_range<1> nd_range{sycl::range<1>{num_items}, + sycl::range<1>{num_items}}; + + cgh.parallel_for( + nd_range, [=](sycl::nd_item<1> it) { + auto& results = acc[it.get_global_id()]; + + size_t sub_group_size = it.get_sub_group().get_local_range().size(); + + auto opportunistic = + oneapi_ext::this_kernel::get_opportunistic_group(); + static_assert( + std::is_same_v); + + results[checks::get_group_id] = opportunistic.get_group_id() == 0; + results[checks::get_local_id] = + opportunistic.get_local_id() < + opportunistic.get_local_range().size(); + results[checks::get_group_range] = + opportunistic.get_group_range().size() == 1; + results[checks::get_local_range] = + opportunistic.get_local_range().size() <= sub_group_size; + results[checks::get_group_linear_id] = + opportunistic.get_group_linear_id() == 0; + results[checks::get_local_linear_id] = + opportunistic.get_local_linear_id() == + opportunistic.get_local_id(); + results[checks::get_group_linear_range] = + opportunistic.get_group_linear_range() == + opportunistic.get_group_range().size(); + results[checks::get_local_linear_range] = + opportunistic.get_local_linear_range() == + opportunistic.get_local_range().size(); + results[checks::leader] = + opportunistic.leader() == (opportunistic.get_local_id() == 0); + }); + }); + q.wait(); + + CheckResults results = {}; + sycl::accessor acc = results_buffer.get_host_access(); + for (size_t check = 0; check < checks::COUNT; check++) + results[check] = std::all_of(acc.cbegin(), acc.cend(), + [=](const auto& it) { return it[check]; }); + + // Group-category traits. + STATIC_CHECK(sycl::is_group::value); + STATIC_CHECK(sycl::is_group_v); + STATIC_CHECK( + oneapi_ext::is_user_constructed_group::value); + STATIC_CHECK(oneapi_ext::is_user_constructed_group_v); + STATIC_CHECK( + !oneapi_ext::is_fixed_topology_group::value); + STATIC_CHECK(!oneapi_ext::is_fixed_topology_group_v); + + // Aliases. + STATIC_CHECK(std::is_same_v>); + STATIC_CHECK( + std::is_same_v>); + STATIC_CHECK(std::is_same_v); + + // Static constexpr members. + STATIC_CHECK(opportunistic_group_t::dimensions == 1); + STATIC_CHECK(opportunistic_group_t::fence_scope == + sycl::sub_group::fence_scope); + + // get_group_id + CHECK(std::is_same_v< + decltype(std::declval().get_group_id()), + opportunistic_group_t::id_type>); + CHECK(results[checks::get_group_id]); + + // get_local_id + CHECK(std::is_same_v< + decltype(std::declval().get_local_id()), + opportunistic_group_t::id_type>); + CHECK(results[checks::get_local_id]); + + // get_group_range + CHECK(std::is_same_v< + decltype(std::declval().get_group_range()), + opportunistic_group_t::range_type>); + CHECK(results[checks::get_group_range]); + + // get_local_range + CHECK(std::is_same_v< + decltype(std::declval().get_local_range()), + opportunistic_group_t::range_type>); + CHECK(results[checks::get_local_range]); + + // get_group_linear_id + CHECK(std::is_same_v< + decltype(std::declval().get_group_linear_id()), + opportunistic_group_t::linear_id_type>); + CHECK(results[checks::get_group_linear_id]); + + // get_local_linear_id + CHECK(std::is_same_v< + decltype(std::declval().get_local_linear_id()), + opportunistic_group_t::linear_id_type>); + CHECK(results[checks::get_local_linear_id]); + + // get_group_linear_range + CHECK(std::is_same_v() + .get_group_linear_range()), + opportunistic_group_t::linear_id_type>); + CHECK(results[checks::get_group_linear_range]); + + // get_local_linear_range + CHECK(std::is_same_v() + .get_local_linear_range()), + opportunistic_group_t::linear_id_type>); + CHECK(results[checks::get_local_linear_range]); + + // leader + CHECK(std::is_same_v().leader()), + bool>); + CHECK(results[checks::leader]); +#endif +} + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp new file mode 100644 index 000000000..dbdb860e9 --- /dev/null +++ b/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp @@ -0,0 +1,182 @@ +/******************************************************************************* +// +// SYCL 2020 Conformance Test Suite +// +// Copyright (c) 2024 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 non_uniform_groups::tests { + +struct checks { + enum { + get_group_id, + get_local_id, + get_group_range, + get_local_range, + get_group_linear_id, + get_local_linear_id, + get_group_linear_range, + get_local_linear_range, + leader, + COUNT, + }; +}; + +TEST_CASE("Test for tangle_group apis.", "[oneapi_non_uniform_groups]") { +#ifndef SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS + SKIP("SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS is not defined"); +#else + namespace oneapi_ext = sycl::ext::oneapi::experimental; + using tangle_group_t = oneapi_ext::tangle_group; + using CheckResults = bool[checks::COUNT]; + + constexpr size_t num_items = 64; + + sycl::buffer results_buffer{num_items}; + + auto q = sycl_cts::util::get_cts_object::queue(); + + if (!q.get_device().has(sycl::aspect::ext_oneapi_tangle_group)) { + SKIP("Device does not support tangle_group."); + } + + q.submit([&](sycl::handler& cgh) { + sycl::accessor acc{results_buffer, cgh, sycl::write_only}; + + sycl::nd_range<1> nd_range{sycl::range<1>{num_items}, + sycl::range<1>{num_items}}; + + cgh.parallel_for(nd_range, [=](sycl::nd_item<1> it) { + auto& results = acc[it.get_global_id()]; + + sycl::sub_group sg = it.get_sub_group(); + size_t sub_group_size = sg.get_local_range().size(); + size_t split = sub_group_size / 2; + + if (sg.get_local_linear_id() < split) { + auto tangle = oneapi_ext::get_tangle_group(sg); + static_assert(std::is_same_v); + + results[checks::get_group_id] = tangle.get_group_id() == 0; + results[checks::get_local_id] = tangle.get_local_id() < split; + results[checks::get_group_range] = tangle.get_group_range().size() == 1; + results[checks::get_local_range] = + tangle.get_local_range().size() == split; + results[checks::get_group_linear_id] = + tangle.get_group_linear_id() == 0; + results[checks::get_local_linear_id] = + tangle.get_local_linear_id() == tangle.get_local_id(); + results[checks::get_group_linear_range] = + tangle.get_group_linear_range() == tangle.get_group_range().size(); + results[checks::get_local_linear_range] = + tangle.get_local_linear_range() == tangle.get_local_range().size(); + results[checks::leader] = + tangle.leader() == (tangle.get_local_id() == 0); + } else { + // All excluded items simply fill with true. + results[checks::get_group_id] = true; + results[checks::get_local_id] = true; + results[checks::get_group_range] = true; + results[checks::get_local_range] = true; + results[checks::get_group_linear_id] = true; + results[checks::get_local_linear_id] = true; + results[checks::get_group_linear_range] = true; + results[checks::get_local_linear_range] = true; + results[checks::leader] = true; + } + }); + }); + q.wait(); + + CheckResults results = {}; + sycl::accessor acc = results_buffer.get_host_access(); + for (size_t check = 0; check < checks::COUNT; check++) + results[check] = std::all_of(acc.cbegin(), acc.cend(), + [=](const auto& it) { return it[check]; }); + + // Group-category traits. + STATIC_CHECK(sycl::is_group::value); + STATIC_CHECK(sycl::is_group_v); + STATIC_CHECK(oneapi_ext::is_user_constructed_group::value); + STATIC_CHECK(oneapi_ext::is_user_constructed_group_v); + STATIC_CHECK(!oneapi_ext::is_fixed_topology_group::value); + STATIC_CHECK(!oneapi_ext::is_fixed_topology_group_v); + + // Aliases. + STATIC_CHECK(std::is_same_v>); + STATIC_CHECK(std::is_same_v>); + STATIC_CHECK(std::is_same_v); + + // Static constexpr members. + STATIC_CHECK(tangle_group_t::dimensions == 1); + STATIC_CHECK(tangle_group_t::fence_scope == sycl::sub_group::fence_scope); + + // get_group_id + CHECK(std::is_same_v().get_group_id()), + tangle_group_t::id_type>); + CHECK(results[checks::get_group_id]); + + // get_local_id + CHECK(std::is_same_v().get_local_id()), + tangle_group_t::id_type>); + CHECK(results[checks::get_local_id]); + + // get_group_range + CHECK( + std::is_same_v().get_group_range()), + tangle_group_t::range_type>); + CHECK(results[checks::get_group_range]); + + // get_local_range + CHECK( + std::is_same_v().get_local_range()), + tangle_group_t::range_type>); + CHECK(results[checks::get_local_range]); + + // get_group_linear_id + CHECK(std::is_same_v< + decltype(std::declval().get_group_linear_id()), + tangle_group_t::linear_id_type>); + CHECK(results[checks::get_group_linear_id]); + + // get_local_linear_id + CHECK(std::is_same_v< + decltype(std::declval().get_local_linear_id()), + tangle_group_t::linear_id_type>); + CHECK(results[checks::get_local_linear_id]); + + // get_group_linear_range + CHECK(std::is_same_v< + decltype(std::declval().get_group_linear_range()), + tangle_group_t::linear_id_type>); + CHECK(results[checks::get_group_linear_range]); + + // get_local_linear_range + CHECK(std::is_same_v< + decltype(std::declval().get_local_linear_range()), + tangle_group_t::linear_id_type>); + CHECK(results[checks::get_local_linear_range]); + + // leader + CHECK( + std::is_same_v().leader()), bool>); + CHECK(results[checks::leader]); +#endif +} + +} // namespace non_uniform_groups::tests From 5383ebc1c4c55040bf708410d1b88c6ceac40518 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 21 Feb 2024 09:39:59 -0800 Subject: [PATCH 02/13] Fix formatting Signed-off-by: Larsen, Steffen --- .../non_uniform_group_common.h | 12 +++--------- 1 file changed, 3 insertions(+), 9 deletions(-) diff --git a/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h b/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h index 69b819f91..d5e723692 100644 --- a/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h +++ b/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h @@ -101,9 +101,7 @@ struct NonUniformGroupHelper< return oneapi_ext::get_fixed_size_group(sg); } - static size_t preferred_single_worker_group_id(int) { - return 0; - } + static size_t preferred_single_worker_group_id(int) { return 0; } static std::string get_name() { return "fixed_size_group<" + std::to_string(PartitionSize) + @@ -141,9 +139,7 @@ struct NonUniformGroupHelper> { return oneapi_ext::get_tangle_group(sg); } - static size_t preferred_single_worker_group_id(int) { - return 0; - } + static size_t preferred_single_worker_group_id(int) { return 0; } static std::string get_name() { return "tangle_group"; } @@ -186,9 +182,7 @@ struct NonUniformGroupHelper { return oneapi_ext::this_kernel::get_opportunistic_group(); } - static size_t preferred_single_worker_group_id(int) { - return 0; - } + static size_t preferred_single_worker_group_id(int) { return 0; } static std::string get_name() { return "opportunistic_group"; } From e303f967fa30ec3220f6874096e9da503101f9eb Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 21 Feb 2024 09:57:24 -0800 Subject: [PATCH 03/13] Add branching tangle_group API tests Signed-off-by: Larsen, Steffen --- .../tangle_group_api.cpp | 28 ++++++++----------- 1 file changed, 12 insertions(+), 16 deletions(-) diff --git a/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp index dbdb860e9..27629c528 100644 --- a/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp +++ b/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp @@ -66,17 +66,14 @@ TEST_CASE("Test for tangle_group apis.", "[oneapi_non_uniform_groups]") { sycl::sub_group sg = it.get_sub_group(); size_t sub_group_size = sg.get_local_range().size(); - size_t split = sub_group_size / 2; - - if (sg.get_local_linear_id() < split) { - auto tangle = oneapi_ext::get_tangle_group(sg); - static_assert(std::is_same_v); + size_t split = sub_group_size / 3; + auto run_checks = [&](tangle_group_t tangle, size_t expected_size) { results[checks::get_group_id] = tangle.get_group_id() == 0; results[checks::get_local_id] = tangle.get_local_id() < split; results[checks::get_group_range] = tangle.get_group_range().size() == 1; results[checks::get_local_range] = - tangle.get_local_range().size() == split; + tangle.get_local_range().size() == expected_size; results[checks::get_group_linear_id] = tangle.get_group_linear_id() == 0; results[checks::get_local_linear_id] = @@ -87,17 +84,16 @@ TEST_CASE("Test for tangle_group apis.", "[oneapi_non_uniform_groups]") { tangle.get_local_linear_range() == tangle.get_local_range().size(); results[checks::leader] = tangle.leader() == (tangle.get_local_id() == 0); + }; + + if (sg.get_local_linear_id() < split) { + auto tangle = oneapi_ext::get_tangle_group(sg); + static_assert(std::is_same_v); + run_checks(tangle, split); } else { - // All excluded items simply fill with true. - results[checks::get_group_id] = true; - results[checks::get_local_id] = true; - results[checks::get_group_range] = true; - results[checks::get_local_range] = true; - results[checks::get_group_linear_id] = true; - results[checks::get_local_linear_id] = true; - results[checks::get_group_linear_range] = true; - results[checks::get_local_linear_range] = true; - results[checks::leader] = true; + auto tangle = oneapi_ext::get_tangle_group(sg); + static_assert(std::is_same_v); + run_checks(tangle, sub_group_size - split); } }); }); From c13fe69263acd655c4df91ee499176537cbbd79c Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 21 Feb 2024 10:10:28 -0800 Subject: [PATCH 04/13] Multiple fixed_size_group sizes Signed-off-by: Larsen, Steffen --- .../oneapi_non_uniform_groups/group_barrier.cpp | 6 ++++++ .../oneapi_non_uniform_groups/group_broadcast.cpp | 6 ++++++ .../group_broadcast_fp16.cpp | 6 ++++++ .../group_broadcast_fp64.cpp | 6 ++++++ .../group_joint_reduce.cpp.in | 6 ++++++ .../group_joint_scan.cpp.in | 6 ++++++ .../extension/oneapi_non_uniform_groups/group_of.cpp | 12 ++++++++++++ .../oneapi_non_uniform_groups/group_permute.cpp | 6 ++++++ .../oneapi_non_uniform_groups/group_permute_fp16.cpp | 6 ++++++ .../oneapi_non_uniform_groups/group_permute_fp64.cpp | 6 ++++++ .../group_reduce_over_group.cpp.in | 6 ++++++ .../group_scan_over_group.cpp.in | 6 ++++++ .../oneapi_non_uniform_groups/group_shift.cpp | 6 ++++++ .../oneapi_non_uniform_groups/group_shift_fp16.cpp | 6 ++++++ .../oneapi_non_uniform_groups/group_shift_fp64.cpp | 6 ++++++ .../non_uniform_group_common.h | 10 +++++++++- 16 files changed, 105 insertions(+), 1 deletion(-) diff --git a/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp b/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp index 6942a3327..7db96edb3 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp @@ -30,8 +30,14 @@ TEST_CASE("Non-uniform-group barriers", auto queue = once_per_unit::get_queue(); non_uniform_group_barrier>(queue); + non_uniform_group_barrier>( + queue); + non_uniform_group_barrier>( + queue); non_uniform_group_barrier>( queue); + non_uniform_group_barrier>( + queue); non_uniform_group_barrier>(queue); non_uniform_group_barrier(queue); } diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp index 630e01694..8ad263b46 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp @@ -28,8 +28,14 @@ TEMPLATE_LIST_TEST_CASE("Non-uniform group broadcast and select", auto queue = once_per_unit::get_queue(); broadcast_non_uniform_group, TestType>(queue); + broadcast_non_uniform_group, + TestType>(queue); + broadcast_non_uniform_group, + TestType>(queue); broadcast_non_uniform_group, TestType>(queue); + broadcast_non_uniform_group, + TestType>(queue); broadcast_non_uniform_group, TestType>(queue); broadcast_non_uniform_group(queue); diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp index c5b7103dd..e2d11ee77 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp @@ -26,8 +26,14 @@ TEST_CASE("Non-uniform group broadcast and select", if (queue.get_device().has(sycl::aspect::fp16)) { broadcast_non_uniform_group, sycl::half>(queue); + broadcast_non_uniform_group< + oneapi_ext::fixed_size_group<1, sycl::sub_group>, sycl::half>(queue); + broadcast_non_uniform_group< + oneapi_ext::fixed_size_group<2, sycl::sub_group>, sycl::half>(queue); broadcast_non_uniform_group< oneapi_ext::fixed_size_group<4, sycl::sub_group>, sycl::half>(queue); + broadcast_non_uniform_group< + oneapi_ext::fixed_size_group<8, sycl::sub_group>, sycl::half>(queue); broadcast_non_uniform_group, sycl::half>(queue); broadcast_non_uniform_group( diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp index adc3a65ea..ff64caa0b 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp @@ -25,8 +25,14 @@ TEST_CASE("Non-uniform group broadcast and select", "[group_func][fp64][dim]") { if (queue.get_device().has(sycl::aspect::fp64)) { broadcast_non_uniform_group, double>(queue); + broadcast_non_uniform_group< + oneapi_ext::fixed_size_group<1, sycl::sub_group>, double>(queue); + broadcast_non_uniform_group< + oneapi_ext::fixed_size_group<2, sycl::sub_group>, double>(queue); broadcast_non_uniform_group< oneapi_ext::fixed_size_group<4, sycl::sub_group>, double>(queue); + broadcast_non_uniform_group< + oneapi_ext::fixed_size_group<8, sycl::sub_group>, double>(queue); broadcast_non_uniform_group, double>(queue); broadcast_non_uniform_group(queue); diff --git a/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in index 22eac4e4b..d07e358a6 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in +++ b/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in @@ -33,7 +33,10 @@ TEST_CASE(CTS_TYPE_NAME + " non-uniform group joint reduce functions", const auto RetType = unnamed_type_pack(); const auto GroupTypes = unnamed_type_pack< oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<1, sycl::sub_group>, + oneapi_ext::fixed_size_group<2, sycl::sub_group>, oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::fixed_size_group<8, sycl::sub_group>, oneapi_ext::tangle_group, oneapi_ext::opportunistic_group>(); @@ -63,7 +66,10 @@ TEMPLATE_LIST_TEST_CASE( const auto ReducedType = unnamed_type_pack(); const auto GroupTypes = unnamed_type_pack< oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<1, sycl::sub_group>, + oneapi_ext::fixed_size_group<2, sycl::sub_group>, oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::fixed_size_group<8, sycl::sub_group>, oneapi_ext::tangle_group, oneapi_ext::opportunistic_group>(); diff --git a/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in index 71af50f63..0449bbd3b 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in +++ b/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in @@ -35,7 +35,10 @@ TEST_CASE(CTS_TYPE_NAME + " non-uniform group joint scan functions", auto queue = once_per_unit::get_queue(); const auto GroupTypes = unnamed_type_pack< oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<1, sycl::sub_group>, + oneapi_ext::fixed_size_group<2, sycl::sub_group>, oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::fixed_size_group<8, sycl::sub_group>, oneapi_ext::tangle_group, oneapi_ext::opportunistic_group>(); @@ -60,7 +63,10 @@ TEST_CASE(CTS_TYPE_NAME + " non-uniform group joint scan functions with init", auto queue = once_per_unit::get_queue(); const auto GroupTypes = unnamed_type_pack< oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<1, sycl::sub_group>, + oneapi_ext::fixed_size_group<2, sycl::sub_group>, oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::fixed_size_group<8, sycl::sub_group>, oneapi_ext::tangle_group, oneapi_ext::opportunistic_group>(); diff --git a/tests/extension/oneapi_non_uniform_groups/group_of.cpp b/tests/extension/oneapi_non_uniform_groups/group_of.cpp index f0532e1e8..a0b9b22ab 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_of.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_of.cpp @@ -29,8 +29,14 @@ TEMPLATE_LIST_TEST_CASE( auto queue = once_per_unit::get_queue(); predicate_function_of_non_uniform_group< oneapi_ext::ballot_group, TestType>(queue); + predicate_function_of_non_uniform_group< + oneapi_ext::fixed_size_group<1, sycl::sub_group>, TestType>(queue); + predicate_function_of_non_uniform_group< + oneapi_ext::fixed_size_group<2, sycl::sub_group>, TestType>(queue); predicate_function_of_non_uniform_group< oneapi_ext::fixed_size_group<4, sycl::sub_group>, TestType>(queue); + predicate_function_of_non_uniform_group< + oneapi_ext::fixed_size_group<8, sycl::sub_group>, TestType>(queue); predicate_function_of_non_uniform_group< oneapi_ext::tangle_group, TestType>(queue); predicate_function_of_non_uniform_group>( queue); + bool_function_of_non_uniform_group< + oneapi_ext::fixed_size_group<1, sycl::sub_group>>(queue); + bool_function_of_non_uniform_group< + oneapi_ext::fixed_size_group<2, sycl::sub_group>>(queue); bool_function_of_non_uniform_group< oneapi_ext::fixed_size_group<4, sycl::sub_group>>(queue); + bool_function_of_non_uniform_group< + oneapi_ext::fixed_size_group<8, sycl::sub_group>>(queue); bool_function_of_non_uniform_group>( queue); bool_function_of_non_uniform_group(queue); diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute.cpp index 5b311f183..5a344e7e7 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_permute.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_permute.cpp @@ -28,8 +28,14 @@ TEMPLATE_LIST_TEST_CASE("Non-uniform-group permute", permute_non_uniform_group, TestType>(queue); + permute_non_uniform_group, + TestType>(queue); + permute_non_uniform_group, + TestType>(queue); permute_non_uniform_group, TestType>(queue); + permute_non_uniform_group, + TestType>(queue); permute_non_uniform_group, TestType>(queue); permute_non_uniform_group(queue); diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp index a54497bfc..efd085cac 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp @@ -27,8 +27,14 @@ TEST_CASE("Non-uniform-group permute", if (queue.get_device().has(sycl::aspect::fp16)) { permute_non_uniform_group, sycl::half>(queue); + permute_non_uniform_group, + sycl::half>(queue); + permute_non_uniform_group, + sycl::half>(queue); permute_non_uniform_group, sycl::half>(queue); + permute_non_uniform_group, + sycl::half>(queue); permute_non_uniform_group, sycl::half>(queue); permute_non_uniform_group( diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp index 4d97f760a..ed279d148 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp @@ -27,8 +27,14 @@ TEST_CASE("Non-uniform-group permute", if (queue.get_device().has(sycl::aspect::fp64)) { permute_non_uniform_group, double>(queue); + permute_non_uniform_group, + double>(queue); + permute_non_uniform_group, + double>(queue); permute_non_uniform_group, double>(queue); + permute_non_uniform_group, + double>(queue); permute_non_uniform_group, double>(queue); permute_non_uniform_group(queue); diff --git a/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in index ef51ca672..f5404a73d 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in +++ b/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in @@ -34,7 +34,10 @@ TEST_CASE(CTS_TYPE_NAME + " non-uniform group reduce functions", const auto RetType = unnamed_type_pack(); const auto GroupTypes = unnamed_type_pack< oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<1, sycl::sub_group>, + oneapi_ext::fixed_size_group<2, sycl::sub_group>, oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::fixed_size_group<8, sycl::sub_group>, oneapi_ext::tangle_group, oneapi_ext::opportunistic_group>(); @@ -66,7 +69,10 @@ TEMPLATE_LIST_TEST_CASE(CTS_TYPE_NAME + const auto ReducedType = unnamed_type_pack(); const auto GroupTypes = unnamed_type_pack< oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<1, sycl::sub_group>, + oneapi_ext::fixed_size_group<2, sycl::sub_group>, oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::fixed_size_group<8, sycl::sub_group>, oneapi_ext::tangle_group, oneapi_ext::opportunistic_group>(); diff --git a/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in index b0afdacc2..aa19e768b 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in +++ b/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in @@ -33,7 +33,10 @@ TEST_CASE(CTS_TYPE_NAME + " non-uniform group scan functions", auto queue = once_per_unit::get_queue(); const auto GroupTypes = unnamed_type_pack< oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<1, sycl::sub_group>, + oneapi_ext::fixed_size_group<2, sycl::sub_group>, oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::fixed_size_group<8, sycl::sub_group>, oneapi_ext::tangle_group, oneapi_ext::opportunistic_group>(); @@ -56,7 +59,10 @@ TEST_CASE(CTS_TYPE_NAME + " non-uniform group scan functions with init", auto queue = once_per_unit::get_queue(); const auto GroupTypes = unnamed_type_pack< oneapi_ext::ballot_group, + oneapi_ext::fixed_size_group<1, sycl::sub_group>, + oneapi_ext::fixed_size_group<2, sycl::sub_group>, oneapi_ext::fixed_size_group<4, sycl::sub_group>, + oneapi_ext::fixed_size_group<8, sycl::sub_group>, oneapi_ext::tangle_group, oneapi_ext::opportunistic_group>(); diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift.cpp b/tests/extension/oneapi_non_uniform_groups/group_shift.cpp index efa8c2a41..9a18b2cb9 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_shift.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_shift.cpp @@ -28,8 +28,14 @@ TEMPLATE_LIST_TEST_CASE("Non-uniform-group shift", shift_non_uniform_group, TestType>( queue); + shift_non_uniform_group, + TestType>(queue); + shift_non_uniform_group, + TestType>(queue); shift_non_uniform_group, TestType>(queue); + shift_non_uniform_group, + TestType>(queue); shift_non_uniform_group, TestType>( queue); shift_non_uniform_group(queue); diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp index 48c250e17..55c503f60 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp @@ -27,8 +27,14 @@ TEST_CASE("Non-uniform-group shift", if (queue.get_device().has(sycl::aspect::fp16)) { shift_non_uniform_group, sycl::half>(queue); + shift_non_uniform_group, + sycl::half>(queue); + shift_non_uniform_group, + sycl::half>(queue); shift_non_uniform_group, sycl::half>(queue); + shift_non_uniform_group, + sycl::half>(queue); shift_non_uniform_group, sycl::half>(queue); shift_non_uniform_group(queue); diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp index bdb72faef..208c4a842 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp @@ -27,8 +27,14 @@ TEST_CASE("Non-uniform-group shift", if (queue.get_device().has(sycl::aspect::fp64)) { shift_non_uniform_group, double>( queue); + shift_non_uniform_group, + double>(queue); + shift_non_uniform_group, + double>(queue); shift_non_uniform_group, double>(queue); + shift_non_uniform_group, + double>(queue); shift_non_uniform_group, double>( queue); shift_non_uniform_group(queue); diff --git a/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h b/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h index d5e723692..a6a18077f 100644 --- a/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h +++ b/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h @@ -86,10 +86,18 @@ struct NonUniformGroupHelper> { template struct NonUniformGroupHelper< oneapi_ext::fixed_size_group> { + // Fixed-size group testing have one case per size we test, but since the size + // is part of the type it will have to be specified further out. static constexpr size_t num_test_cases = 1; static bool is_supported(const sycl::device& dev) { - return dev.has(sycl::aspect::ext_oneapi_fixed_size_group); + // For these tests we consider them unsupported if the smallest supported + // sub-group size is less than the partition size. + std::vector sg_sizes = + dev.get_info(); + size_t min_sg_size = *std::min_element(sg_sizes.cbegin(), sg_sizes.cend()); + return min_sg_size >= PartitionSize && + dev.has(sycl::aspect::ext_oneapi_fixed_size_group); } static bool should_participate(sycl::sub_group sg, int test_case) { From 7878306cc5fdce6dc7d439b91927cc8cfb5db5ef Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 8 Mar 2024 01:06:40 -0800 Subject: [PATCH 05/13] Minor adjustments Signed-off-by: Larsen, Steffen --- .../oneapi_non_uniform_groups/group_broadcast.h | 9 ++++----- .../oneapi_non_uniform_groups/group_broadcast_fp16.cpp | 2 +- .../oneapi_non_uniform_groups/group_broadcast_fp64.cpp | 2 +- .../oneapi_non_uniform_groups/non_uniform_group_common.h | 3 --- 4 files changed, 6 insertions(+), 10 deletions(-) diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast.h b/tests/extension/oneapi_non_uniform_groups/group_broadcast.h index f3be91049..c2f9de166 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast.h +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast.h @@ -95,19 +95,18 @@ void broadcast_non_uniform_group(sycl::queue& queue) { static_cast(item.get_global_linear_id() * 100 + non_uniform_group.get_local_id()))); - T local_var(splat_init(0)); - // To simplify the test, we are only checking the first group in // the first sub-group. + size_t preferred_group_id = NonUniformGroupHelper< + GroupT>::preferred_single_worker_group_id(test_case); if (item.get_sub_group().get_group_id()[0] == 0 && - non_uniform_group.get_group_id()[0] == 0) { + non_uniform_group.get_group_id()[0] == preferred_group_id) { // Find local id of first, last and some third sub-group item in // between. Will be used to check different combinations of // broadcasting and receiving work-items sycl::id<1> first_id = 0; sycl::id<1> mid_id = non_uniform_group.get_local_range() / 2; - sycl::id<1> last_id = non_uniform_group.get_local_range(); - --last_id[0]; + sycl::id<1> last_id = non_uniform_group.get_local_range() - 1; // Broadcast from the first work-item ASSERT_RETURN_TYPE( diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp index e2d11ee77..68d89d76f 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp @@ -21,7 +21,7 @@ #include "group_broadcast.h" TEST_CASE("Non-uniform group broadcast and select", - "[oneapi_non_uniform_groups][group_func][fp16][dim]") { + "[oneapi_non_uniform_groups][group_func][fp16]") { auto queue = once_per_unit::get_queue(); if (queue.get_device().has(sycl::aspect::fp16)) { broadcast_non_uniform_group, diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp index ff64caa0b..dedf65513 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp @@ -20,7 +20,7 @@ #include "group_broadcast.h" -TEST_CASE("Non-uniform group broadcast and select", "[group_func][fp64][dim]") { +TEST_CASE("Non-uniform group broadcast and select", "[group_func][fp64]") { auto queue = once_per_unit::get_queue(); if (queue.get_device().has(sycl::aspect::fp64)) { broadcast_non_uniform_group, diff --git a/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h b/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h index a6a18077f..b011fb563 100644 --- a/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h +++ b/tests/extension/oneapi_non_uniform_groups/non_uniform_group_common.h @@ -41,7 +41,6 @@ struct NonUniformGroupHelper> { static oneapi_ext::ballot_group create(sycl::sub_group sg, int test_case) { - assert(test_case < num_test_cases); // Split it so that 1/3rd of the items are in the first "true" group and the // rest are in "false" group. switch (test_case) { @@ -130,7 +129,6 @@ struct NonUniformGroupHelper> { } static bool should_participate(sycl::sub_group sg, int test_case) { - assert(test_case < num_test_cases); switch (test_case) { case 0: return sg.get_local_linear_id() < sg.get_local_range().size() / 3; @@ -174,7 +172,6 @@ struct NonUniformGroupHelper { } static bool should_participate(sycl::sub_group sg, int test_case) { - assert(test_case < num_test_cases); switch (test_case) { case 0: return sg.get_local_linear_id() < sg.get_local_range().size() / 3; From cd099dc8beb0079fc2b512d85853dd59c8fbcb36 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 16 Apr 2024 06:01:18 -0700 Subject: [PATCH 06/13] Fix opportunistic fence-scope check Signed-off-by: Larsen, Steffen --- .../oneapi_non_uniform_groups/opportunistic_group_api.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp index cb686b805..d63cbd478 100644 --- a/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp +++ b/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp @@ -122,7 +122,7 @@ TEST_CASE("Test for opportunistic_group apis.", "[oneapi_non_uniform_groups]") { // Static constexpr members. STATIC_CHECK(opportunistic_group_t::dimensions == 1); STATIC_CHECK(opportunistic_group_t::fence_scope == - sycl::sub_group::fence_scope); + sycl::memory_scope::sub_group); // get_group_id CHECK(std::is_same_v< From c911a4140184d6d27dac9205f61c46ddad06e012 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 16 Apr 2024 07:02:51 -0700 Subject: [PATCH 07/13] Add missing joint_of tests Signed-off-by: Larsen, Steffen --- .../oneapi_non_uniform_groups/group_of.cpp | 17 +++ .../oneapi_non_uniform_groups/group_of.h | 143 ++++++++++++++++++ 2 files changed, 160 insertions(+) diff --git a/tests/extension/oneapi_non_uniform_groups/group_of.cpp b/tests/extension/oneapi_non_uniform_groups/group_of.cpp index a0b9b22ab..c12440e3f 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_of.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_of.cpp @@ -23,6 +23,23 @@ // use wide types to exclude truncation of init values using WideTypes = std::tuple; +TEMPLATE_LIST_TEST_CASE("Non-uniform group joint of bool functions", + "[oneapi_non_uniform_groups][group_func][type_list]", + WideTypes) { + auto queue = once_per_unit::get_queue(); + joint_of_group, TestType>(queue); + joint_of_group, TestType>( + queue); + joint_of_group, TestType>( + queue); + joint_of_group, TestType>( + queue); + joint_of_group, TestType>( + queue); + joint_of_group, TestType>(queue); + joint_of_group(queue); +} + TEMPLATE_LIST_TEST_CASE( "Non-uniform group of bool functions with predicate functions", "[oneapi_non_uniform_groups][group_func][type_list]", WideTypes) { diff --git a/tests/extension/oneapi_non_uniform_groups/group_of.h b/tests/extension/oneapi_non_uniform_groups/group_of.h index 512fb56bd..18dbb9d99 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_of.h +++ b/tests/extension/oneapi_non_uniform_groups/group_of.h @@ -21,6 +21,149 @@ #include "../../group_functions/group_functions_common.h" #include "non_uniform_group_common.h" +template +class joint_of_group_kernel; + +/** + * @brief Provides test for joint non-uniform group bool of operations with + * predicate functions + * @tparam GroupT Type of the non-uniform group to test with + * @tparam T Type pointed by Ptr + */ +template +void joint_of_group(sycl::queue& queue) { + const std::string group_name = NonUniformGroupHelper::get_name(); + + INFO("Testing group-of predicate function for " + group_name); + if (!NonUniformGroupHelper::is_supported(queue.get_device())) { + SKIP("Device does not support " + group_name); + } + + // 3 functions * 4 predicates + constexpr int test_matrix = 3; + const std::string test_names[test_matrix] = { + "bool joint_any_of(GroupT g, Ptr first, Ptr last, Predicate pred)", + "bool joint_all_of(GroupT g, Ptr first, Ptr last, Predicate pred)", + "bool joint_none_of(GroupT g, Ptr first, Ptr last, Predicate pred)"}; + constexpr int test_cases = 4; + const std::string test_cases_names[test_cases] = {"none true", "one true", + "some true", "all true"}; + + sycl::range<1> work_group_range = sycl_cts::util::work_group_range<1>(queue); + size_t work_group_size = work_group_range.size(); + + const size_t sizes[3] = {5, work_group_size / 2, 3 * work_group_size}; + for (size_t test_case = 0; + test_case < NonUniformGroupHelper::num_test_cases; ++test_case) { + const std::string test_case_name = + NonUniformGroupHelper::get_test_case_name(test_case); + INFO("Running test case (" + std::to_string(test_case) + ") with " + + test_case_name); + + for (size_t size : sizes) { + std::vector v(size); + std::iota(v.begin(), v.end(), 1); + + // array to return results: + bool res[test_matrix * test_cases] = {false}; + { + sycl::buffer v_sycl(v.data(), sycl::range<1>(size)); + + sycl::buffer res_sycl( + res, sycl::range<1>(test_matrix * test_cases)); + + queue.submit([&](sycl::handler& cgh) { + auto v_acc = + v_sycl.template get_access(cgh); + auto res_acc = + res_sycl.get_access(cgh); + + sycl::nd_range<1> executionRange(work_group_range, work_group_range); + + cgh.parallel_for>(executionRange, [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate(sub_group, + test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + + T* v_begin = v_acc.get_pointer(); + T* v_end = v_begin + v_acc.size(); + + // predicates + auto none_true = [&](T i) { return i == 0; }; + auto one_true = [&](T i) { return i == 1; }; + auto some_true = [&](T i) { return i > size / 2; }; + auto all_true = [&](T i) { return i <= size; }; + + ASSERT_RETURN_TYPE( + bool, + sycl::joint_any_of(non_uniform_group, v_begin, v_end, + none_true), + "Return type of joint_any_of(GroupT g, Ptr first, Ptr last, " + "Predicate pred) is wrong\n"); + res_acc[0] = !sycl::joint_any_of(non_uniform_group, v_begin, v_end, + none_true); + res_acc[1] = + sycl::joint_any_of(non_uniform_group, v_begin, v_end, one_true); + res_acc[2] = sycl::joint_any_of(non_uniform_group, v_begin, v_end, + some_true); + res_acc[3] = + sycl::joint_any_of(non_uniform_group, v_begin, v_end, all_true); + + ASSERT_RETURN_TYPE( + bool, + sycl::joint_all_of(non_uniform_group, v_begin, v_end, + none_true), + "Return type of joint_all_of(GroupT g, Ptr first, Ptr last, " + "Predicate pred) is wrong\n"); + res_acc[4] = !sycl::joint_all_of(non_uniform_group, v_begin, v_end, + none_true); + res_acc[5] = !sycl::joint_all_of(non_uniform_group, v_begin, v_end, + one_true); + res_acc[6] = !sycl::joint_all_of(non_uniform_group, v_begin, v_end, + some_true); + res_acc[7] = + sycl::joint_all_of(non_uniform_group, v_begin, v_end, all_true); + + ASSERT_RETURN_TYPE( + bool, + sycl::joint_none_of(non_uniform_group, v_begin, v_end, + none_true), + "Return type of joint_none_of(GroupT g, Ptr first, Ptr last, " + "Predicate pred) is wrong\n"); + res_acc[8] = sycl::joint_none_of(non_uniform_group, v_begin, v_end, + none_true); + res_acc[9] = !sycl::joint_none_of(non_uniform_group, v_begin, v_end, + one_true); + res_acc[10] = !sycl::joint_none_of(non_uniform_group, v_begin, + v_end, some_true); + res_acc[11] = !sycl::joint_none_of(non_uniform_group, v_begin, + v_end, all_true); + }); + }); + } + int index = 0; + for (int i = 0; i < test_matrix; ++i) + for (int j = 0; j < test_cases; ++j) { + std::string work_group = + sycl_cts::util::work_group_print(work_group_range); + CAPTURE(group_name, work_group); + INFO("Value of " << test_names[i] << " with " << test_cases_names[j] + << " predicate" + " is " + << (res[index] ? "right" : "wrong")); + CHECK(res[index++]); + } + } + } +} + template class predicate_function_of_non_uniform_group_kernel; From 3a5917a10ddb45e9ae2956fc7b05bc4074a3350d Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 16 Apr 2024 09:04:49 -0700 Subject: [PATCH 08/13] Fix copyright year Signed-off-by: Larsen, Steffen --- tests/extension/oneapi_non_uniform_groups/group_barrier.cpp | 2 +- .../oneapi_non_uniform_groups/group_joint_reduce.cpp.in | 2 +- .../extension/oneapi_non_uniform_groups/group_permute_fp64.cpp | 2 +- tests/group_functions/group_shift.h | 2 +- 4 files changed, 4 insertions(+), 4 deletions(-) diff --git a/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp b/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp index 7db96edb3..4142d1579 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp @@ -2,7 +2,7 @@ // // SYCL 2020 Conformance Test Suite // -// Copyright (c) 2023 The Khronos Group Inc. +// Copyright (c) 2024 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. diff --git a/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in index d07e358a6..5b953eba1 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in +++ b/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in @@ -2,7 +2,7 @@ // // SYCL 2020 Conformance Test Suite // -// Copyright (c) 2023 The Khronos Group Inc. +// Copyright (c) 2024 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. diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp index ed279d148..570b29a52 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp @@ -2,7 +2,7 @@ // // SYCL 2020 Conformance Test Suite // -// Copyright (c) 2023 The Khronos Group Inc. +// Copyright (c) 2024 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. diff --git a/tests/group_functions/group_shift.h b/tests/group_functions/group_shift.h index aa395d683..d295aff0a 100644 --- a/tests/group_functions/group_shift.h +++ b/tests/group_functions/group_shift.h @@ -2,7 +2,7 @@ // // SYCL 2020 Conformance Test Suite // -// Copyright (c) 2023 The Khronos Group Inc. +// Copyright (c) 2024 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. From bf6e897978819c131ab7dfc5036ed81718d3d816 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 13 May 2024 00:05:17 -0700 Subject: [PATCH 09/13] Address comments Signed-off-by: Larsen, Steffen --- .../oneapi_non_uniform_groups/CMakeLists.txt | 19 +++--- .../ballot_group_api.cpp | 1 - .../fixed_size_group_api.cpp | 1 - .../group_barrier.cpp | 4 ++ .../oneapi_non_uniform_groups/group_barrier.h | 17 ++--- .../group_broadcast.cpp | 4 ++ .../group_broadcast.h | 32 ++++----- .../group_broadcast_fp16.cpp | 4 ++ .../group_broadcast_fp64.cpp | 4 ++ .../group_joint_reduce.cpp.in | 4 ++ .../group_joint_scan.cpp.in | 4 ++ .../oneapi_non_uniform_groups/group_of.cpp | 4 ++ .../oneapi_non_uniform_groups/group_of.h | 66 ++++++++++--------- .../group_permute.cpp | 4 ++ .../oneapi_non_uniform_groups/group_permute.h | 6 +- .../group_permute_fp16.cpp | 4 ++ .../group_permute_fp64.cpp | 4 ++ .../oneapi_non_uniform_groups/group_reduce.h | 34 +++++----- .../group_reduce_over_group.cpp.in | 4 ++ .../group_scan_over_group.cpp.in | 4 ++ .../oneapi_non_uniform_groups/group_shift.cpp | 4 ++ .../oneapi_non_uniform_groups/group_shift.h | 22 ++++--- .../group_shift_fp16.cpp | 4 ++ .../group_shift_fp64.cpp | 4 ++ .../opportunistic_group_api.cpp | 1 - .../tangle_group_api.cpp | 1 - tests/group_functions/group_shift.h | 2 +- 27 files changed, 165 insertions(+), 97 deletions(-) diff --git a/tests/extension/oneapi_non_uniform_groups/CMakeLists.txt b/tests/extension/oneapi_non_uniform_groups/CMakeLists.txt index 79b819177..d5ec87a0b 100644 --- a/tests/extension/oneapi_non_uniform_groups/CMakeLists.txt +++ b/tests/extension/oneapi_non_uniform_groups/CMakeLists.txt @@ -2,10 +2,10 @@ if(SYCL_CTS_ENABLE_EXT_ONEAPI_NON_UNIFORM_GROUPS_TESTS) 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) + 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 @@ -23,16 +23,17 @@ if(SYCL_CTS_ENABLE_EXT_ONEAPI_NON_UNIFORM_GROUPS_TESTS) foreach(TEMP IN LISTS TEMPLATE_LIST) foreach(TY IN LISTS TYPE_LIST) if(TY STREQUAL "bool") - continue() + 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) + TYPE "${TY}" + IN_FILENAME "${TEMP}.cpp.in" + OUT_FILENAME ${OUT_FILE} + TEST_LIST test_cases_list + ) endforeach() endforeach() diff --git a/tests/extension/oneapi_non_uniform_groups/ballot_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/ballot_group_api.cpp index a54fedb66..b6b18db16 100644 --- a/tests/extension/oneapi_non_uniform_groups/ballot_group_api.cpp +++ b/tests/extension/oneapi_non_uniform_groups/ballot_group_api.cpp @@ -94,7 +94,6 @@ TEST_CASE("Test for ballot_group apis.", "[oneapi_non_uniform_groups]") { results[checks::leader] = ballot.leader() == (ballot.get_local_id() == 0); }); }); - q.wait(); CheckResults results = {}; sycl::accessor acc = results_buffer.get_host_access(); diff --git a/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp index 742cefc86..6b9efe1de 100644 --- a/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp +++ b/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp @@ -99,7 +99,6 @@ TEST_CASE("Test for fixed_size_group apis.", "[oneapi_non_uniform_groups]") { fixed_size.leader() == (fixed_size.get_local_id() == 0); }); }); - q.wait(); CheckResults results = {}; sycl::accessor acc = results_buffer.get_host_access(); diff --git a/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp b/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp index 4142d1579..a7e9b357b 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_barrier.cpp @@ -22,6 +22,8 @@ #include "group_barrier.h" +namespace non_uniform_groups::tests { + template class test_fence; @@ -41,3 +43,5 @@ TEST_CASE("Non-uniform-group barriers", non_uniform_group_barrier>(queue); non_uniform_group_barrier(queue); } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_barrier.h b/tests/extension/oneapi_non_uniform_groups/group_barrier.h index 652b9e0a4..3c00a4898 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_barrier.h +++ b/tests/extension/oneapi_non_uniform_groups/group_barrier.h @@ -103,7 +103,8 @@ void non_uniform_group_barrier(sycl::queue& queue) { sycl::range<1>(work_group_size)); sycl::buffer non_uniform_group_barriers_buf( - non_uniform_group_barriers.data(), sycl::range<1>(5)); + non_uniform_group_barriers.data(), + sycl::range<1>(non_uniform_group_barrier_variants)); queue.submit([&](sycl::handler& cgh) { sycl::nd_range<1> executionRange(work_group_range, work_group_range); @@ -132,13 +133,13 @@ void non_uniform_group_barrier(sycl::queue& queue) { size_t llid = non_uniform_group.get_local_linear_id(); size_t max_id = non_uniform_group.get_local_linear_range() - 1; - ASSERT_RETURN_TYPE( - void, sycl::group_barrier(non_uniform_group), - "Return type of group_barrier(GroupT g) is wrong\n"); - ASSERT_RETURN_TYPE( - void, - sycl::group_barrier(non_uniform_group, - non_uniform_group.fence_scope), + static_assert(std::is_same_v, + "Return type of group_barrier(GroupT g) is wrong\n"); + static_assert( + std::is_same_v, "Return type of group_barrier(GroupT g, " "memory_scope fence_scope) is wrong\n"); diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp index 8ad263b46..410578a94 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast.cpp @@ -20,6 +20,8 @@ #include "group_broadcast.h" +namespace non_uniform_groups::tests { + using BroadcastTypes = CustomTypes; TEMPLATE_LIST_TEST_CASE("Non-uniform group broadcast and select", @@ -40,3 +42,5 @@ TEMPLATE_LIST_TEST_CASE("Non-uniform group broadcast and select", TestType>(queue); broadcast_non_uniform_group(queue); } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast.h b/tests/extension/oneapi_non_uniform_groups/group_broadcast.h index c2f9de166..b9a14c0e7 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast.h +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast.h @@ -109,10 +109,10 @@ void broadcast_non_uniform_group(sycl::queue& queue) { sycl::id<1> last_id = non_uniform_group.get_local_range() - 1; // Broadcast from the first work-item - ASSERT_RETURN_TYPE( - T, - sycl::group_broadcast(non_uniform_group, - value_to_broadcast), + static_assert( + std::is_same_v, "Return type of group_broadcast(GroupT g, T x) is wrong\n"); if (non_uniform_group.leader()) { @@ -127,10 +127,10 @@ void broadcast_non_uniform_group(sycl::queue& queue) { broadcasted_values_acc[0] = broadcasted_value; // Broadcast from the last work-item - ASSERT_RETURN_TYPE( - T, - sycl::group_broadcast(non_uniform_group, value_to_broadcast, - last_id), + static_assert( + std::is_same_v, "Return type of group_broadcast(GroupT g, T x, " "GroupT::linear_id_type local_linear_id) is wrong\n"); @@ -148,10 +148,10 @@ void broadcast_non_uniform_group(sycl::queue& queue) { broadcasted_values_acc[1] = broadcasted_value; // Broadcast from a mid work-item - ASSERT_RETURN_TYPE( - T, - sycl::group_broadcast(non_uniform_group, value_to_broadcast, - mid_id), + static_assert( + std::is_same_v, "Return type of group_broadcast(GroupT g, T x, " "GroupT::id_type local_id) is wrong\n"); @@ -167,10 +167,10 @@ void broadcast_non_uniform_group(sycl::queue& queue) { broadcasted_values_acc[2] = broadcasted_value; // Select from the first work-item - ASSERT_RETURN_TYPE( - T, - sycl::select_from_group(non_uniform_group, - value_to_broadcast, first_id), + static_assert( + std::is_same_v, "Return type of select_from_group(GroupT g, T x, " "GroupT::id_type local_id) is wrong\n"); diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp index 68d89d76f..ecf78db75 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp16.cpp @@ -20,6 +20,8 @@ #include "group_broadcast.h" +namespace non_uniform_groups::tests { + TEST_CASE("Non-uniform group broadcast and select", "[oneapi_non_uniform_groups][group_func][fp16]") { auto queue = once_per_unit::get_queue(); @@ -42,3 +44,5 @@ TEST_CASE("Non-uniform group broadcast and select", WARN("Device does not support half precision floating point operations."); } } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp index dedf65513..1a55529e6 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast_fp64.cpp @@ -20,6 +20,8 @@ #include "group_broadcast.h" +namespace non_uniform_groups::tests { + TEST_CASE("Non-uniform group broadcast and select", "[group_func][fp64]") { auto queue = once_per_unit::get_queue(); if (queue.get_device().has(sycl::aspect::fp64)) { @@ -40,3 +42,5 @@ TEST_CASE("Non-uniform group broadcast and select", "[group_func][fp64]") { WARN("Device does not support double precision floating point operations."); } } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in index 5b953eba1..c8ff919dd 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in +++ b/tests/extension/oneapi_non_uniform_groups/group_joint_reduce.cpp.in @@ -20,6 +20,8 @@ #include "group_reduce.h" +namespace non_uniform_groups::tests { + // clang-format off #cmakedefine CTS_TYPE @CTS_TYPE@ #cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") @@ -89,3 +91,5 @@ TEMPLATE_LIST_TEST_CASE( for_all_combinations( GroupTypes, RetType, ReducedType, Operators, queue); } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in index 0449bbd3b..c5351edf8 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in +++ b/tests/extension/oneapi_non_uniform_groups/group_joint_scan.cpp.in @@ -26,6 +26,8 @@ #if !SYCL_CTS_COMPILING_WITH_HIPSYCL #include "group_scan.h" +namespace non_uniform_groups::tests { + using TestType = unnamed_type_pack; using ScanTypes = Types; #endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL @@ -85,3 +87,5 @@ TEST_CASE(CTS_TYPE_NAME + " non-uniform group joint scan functions with init", for_all_combinations( GroupTypes, TestType{}, ScanTypes{}, ScanTypes{}, queue); }; + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_of.cpp b/tests/extension/oneapi_non_uniform_groups/group_of.cpp index c12440e3f..17787b09c 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_of.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_of.cpp @@ -20,6 +20,8 @@ #include "group_of.h" +namespace non_uniform_groups::tests { + // use wide types to exclude truncation of init values using WideTypes = std::tuple; @@ -77,3 +79,5 @@ TEST_CASE("Non-uniform group of bool functions", queue); bool_function_of_non_uniform_group(queue); } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_of.h b/tests/extension/oneapi_non_uniform_groups/group_of.h index 18dbb9d99..08bdeea76 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_of.h +++ b/tests/extension/oneapi_non_uniform_groups/group_of.h @@ -101,10 +101,10 @@ void joint_of_group(sycl::queue& queue) { auto some_true = [&](T i) { return i > size / 2; }; auto all_true = [&](T i) { return i <= size; }; - ASSERT_RETURN_TYPE( - bool, - sycl::joint_any_of(non_uniform_group, v_begin, v_end, - none_true), + static_assert( + std::is_same_v, "Return type of joint_any_of(GroupT g, Ptr first, Ptr last, " "Predicate pred) is wrong\n"); res_acc[0] = !sycl::joint_any_of(non_uniform_group, v_begin, v_end, @@ -116,10 +116,10 @@ void joint_of_group(sycl::queue& queue) { res_acc[3] = sycl::joint_any_of(non_uniform_group, v_begin, v_end, all_true); - ASSERT_RETURN_TYPE( - bool, - sycl::joint_all_of(non_uniform_group, v_begin, v_end, - none_true), + static_assert( + std::is_same_v, "Return type of joint_all_of(GroupT g, Ptr first, Ptr last, " "Predicate pred) is wrong\n"); res_acc[4] = !sycl::joint_all_of(non_uniform_group, v_begin, v_end, @@ -131,10 +131,10 @@ void joint_of_group(sycl::queue& queue) { res_acc[7] = sycl::joint_all_of(non_uniform_group, v_begin, v_end, all_true); - ASSERT_RETURN_TYPE( - bool, - sycl::joint_none_of(non_uniform_group, v_begin, v_end, - none_true), + static_assert( + std::is_same_v, "Return type of joint_none_of(GroupT g, Ptr first, Ptr last, " "Predicate pred) is wrong\n"); res_acc[8] = sycl::joint_none_of(non_uniform_group, v_begin, v_end, @@ -249,9 +249,10 @@ void predicate_function_of_non_uniform_group(sycl::queue& queue) { auto all_true = [&](T i) { return i <= size; }; { - ASSERT_RETURN_TYPE( - bool, - sycl::any_of_group(non_uniform_group, local_var, none_true), + static_assert( + std::is_same_v, "Return type of any_of_group(GroupT g, bool pred) is wrong\n"); res_acc[0] &= !sycl::any_of_group(non_uniform_group, local_var, none_true); @@ -262,9 +263,10 @@ void predicate_function_of_non_uniform_group(sycl::queue& queue) { res_acc[3] &= sycl::any_of_group(non_uniform_group, local_var, all_true); - ASSERT_RETURN_TYPE( - bool, - sycl::all_of_group(non_uniform_group, local_var, none_true), + static_assert( + std::is_same_v, "Return type of all_of_group(GroupT g, bool pred) is wrong\n"); res_acc[4] &= !sycl::all_of_group(non_uniform_group, local_var, none_true); @@ -283,9 +285,10 @@ void predicate_function_of_non_uniform_group(sycl::queue& queue) { res_acc[7] &= sycl::all_of_group(non_uniform_group, local_var, all_true); - ASSERT_RETURN_TYPE( - bool, - sycl::none_of_group(non_uniform_group, local_var, none_true), + static_assert( + std::is_same_v, "Return type of none_of_group(GroupT g, bool pred) is " "wrong\n"); res_acc[8] &= @@ -399,9 +402,10 @@ void bool_function_of_non_uniform_group(sycl::queue& queue) { auto all_true = [&](T i) { return i <= size; }; { - ASSERT_RETURN_TYPE( - bool, - sycl::any_of_group(non_uniform_group, none_true(local_var)), + static_assert( + std::is_same_v, "Return type of any_of_group(GroupT g, bool pred) is wrong\n"); res_acc[0] &= !sycl::any_of_group(non_uniform_group, none_true(local_var)); @@ -412,9 +416,10 @@ void bool_function_of_non_uniform_group(sycl::queue& queue) { res_acc[3] &= sycl::any_of_group(non_uniform_group, all_true(local_var)); - ASSERT_RETURN_TYPE( - bool, - sycl::all_of_group(non_uniform_group, none_true(local_var)), + static_assert( + std::is_same_v, "Return type of all_of_group(GroupT g, bool pred) is wrong\n"); res_acc[4] = !sycl::all_of_group(non_uniform_group, none_true(local_var)); @@ -433,9 +438,10 @@ void bool_function_of_non_uniform_group(sycl::queue& queue) { res_acc[7] &= sycl::all_of_group(non_uniform_group, all_true(local_var)); - ASSERT_RETURN_TYPE( - bool, - sycl::none_of_group(non_uniform_group, none_true(local_var)), + static_assert( + std::is_same_v, "Return type of none_of_group(GroupT g, bool pred) is " "wrong\n"); res_acc[8] &= diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute.cpp index 5a344e7e7..e187671e2 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_permute.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_permute.cpp @@ -20,6 +20,8 @@ #include "group_permute.h" +namespace non_uniform_groups::tests { + // hipSYCL does not permute right 8-bit types inside groups TEMPLATE_LIST_TEST_CASE("Non-uniform-group permute", "[oneapi_non_uniform_groups][group_func][type_list]", @@ -40,3 +42,5 @@ TEMPLATE_LIST_TEST_CASE("Non-uniform-group permute", TestType>(queue); permute_non_uniform_group(queue); } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute.h b/tests/extension/oneapi_non_uniform_groups/group_permute.h index f67ba37e6..eff0fd4ca 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_permute.h +++ b/tests/extension/oneapi_non_uniform_groups/group_permute.h @@ -80,9 +80,9 @@ void permute_non_uniform_group(sycl::queue& queue) { T local_var(splat_init(llid + 1)); T permuted_var(splat_init(llid + 1)); - ASSERT_RETURN_TYPE( - T, - sycl::permute_group_by_xor(non_uniform_group, local_var, 0), + static_assert( + std::is_same_v, "Return type of permute_group_by_xor(GroupT g, T x, " "GroupT::linear_id_type mask) is wrong\n"); diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp index efd085cac..16cbac422 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_permute_fp16.cpp @@ -20,6 +20,8 @@ #include "group_permute.h" +namespace non_uniform_groups::tests { + TEST_CASE("Non-uniform-group permute", "[oneapi_non_uniform_groups][group_func][fp16]") { auto queue = once_per_unit::get_queue(); @@ -43,3 +45,5 @@ TEST_CASE("Non-uniform-group permute", WARN("Device does not support half precision floating point operations."); } } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp index 570b29a52..5b56a070e 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_permute_fp64.cpp @@ -20,6 +20,8 @@ #include "group_permute.h" +namespace non_uniform_groups::tests { + TEST_CASE("Non-uniform-group permute", "[oneapi_non_uniform_groups][group_func][fp64]") { auto queue = once_per_unit::get_queue(); @@ -42,3 +44,5 @@ TEST_CASE("Non-uniform-group permute", WARN("Device does not support double precision floating point operations."); } } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_reduce.h b/tests/extension/oneapi_non_uniform_groups/group_reduce.h index ee6d28b0a..6063c9a3e 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_reduce.h +++ b/tests/extension/oneapi_non_uniform_groups/group_reduce.h @@ -159,10 +159,10 @@ void joint_reduce_group(sycl::queue& queue, const std::string& op_name) { T* v_begin = v_acc.get_pointer(); T* v_end = v_begin + v_acc.size(); - ASSERT_RETURN_TYPE( - T, - sycl::joint_reduce(non_uniform_group, v_begin, v_end, - OpT()), + static_assert( + std::is_same_v, "Return type of joint_reduce(GroupT g, Ptr first, Ptr " "last, BinaryOperation binary_op) is wrong\n"); @@ -283,10 +283,10 @@ void init_joint_reduce_group(sycl::queue& queue, const std::string& op_name) { U* v_begin = v_acc.get_pointer(); U* v_end = v_begin + v_acc.size(); - ASSERT_RETURN_TYPE( - T, - sycl::joint_reduce(non_uniform_group, v_begin, v_end, - T(init), OpT()), + static_assert( + std::is_same_v, "Return type of joint_reduce(GroupT g, Ptr first, Ptr " "last, T init, BinaryOperation binary_op) is wrong\n"); @@ -409,11 +409,11 @@ void reduce_over_group(sycl::queue& queue, const std::string& op_name) { sg_id_acc[index] = sub_group.get_group_linear_id(); nug_id_acc[index] = non_uniform_group.get_group_linear_id(); - ASSERT_RETURN_TYPE(T, - sycl::reduce_over_group(non_uniform_group, - v_acc[index], OpT()), - "Return type of reduce_over_group(GroupT g, " - "T x, BinaryOperation binary_op) is wrong\n"); + static_assert(std::is_same_v, + "Return type of reduce_over_group(GroupT g, " + "T x, BinaryOperation binary_op) is wrong\n"); nug_output_acc[index] = sycl::reduce_over_group( non_uniform_group, v_acc[index], OpT()); }); @@ -526,10 +526,10 @@ void init_reduce_over_group(sycl::queue& queue, const std::string& op_name) { sg_id_acc[index] = sub_group.get_group_linear_id(); nug_id_acc[index] = non_uniform_group.get_group_linear_id(); - ASSERT_RETURN_TYPE( - T, - sycl::reduce_over_group(non_uniform_group, v_acc[index], - T(init), OpT()), + static_assert( + std::is_same_v, "Return type of reduce_over_group(GroupT g, V x, T init, " "BinaryOperation binary_op) is wrong\n"); nug_output_acc[index] = sycl::reduce_over_group( diff --git a/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in index f5404a73d..5b3885deb 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in +++ b/tests/extension/oneapi_non_uniform_groups/group_reduce_over_group.cpp.in @@ -20,6 +20,8 @@ #include "group_reduce.h" +namespace non_uniform_groups::tests { + // clang-format off #cmakedefine CTS_TYPE @CTS_TYPE@ #cmakedefine CTS_TYPE_NAME std::string("@CTS_TYPE_NAME@") @@ -92,3 +94,5 @@ TEMPLATE_LIST_TEST_CASE(CTS_TYPE_NAME + for_all_combinations( GroupTypes, RetType, ReducedType, Operators, queue); } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in b/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in index aa19e768b..c4c1e214a 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in +++ b/tests/extension/oneapi_non_uniform_groups/group_scan_over_group.cpp.in @@ -25,6 +25,8 @@ #include "group_scan.h" +namespace non_uniform_groups::tests { + using TestType = unnamed_type_pack; using ScanTypes = Types; @@ -80,3 +82,5 @@ TEST_CASE(CTS_TYPE_NAME + " non-uniform group scan functions with init", for_all_combinations(GroupTypes, TestType{}, ScanTypes{}, queue); }; + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift.cpp b/tests/extension/oneapi_non_uniform_groups/group_shift.cpp index 9a18b2cb9..8b474e15b 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_shift.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_shift.cpp @@ -20,6 +20,8 @@ #include "group_shift.h" +namespace non_uniform_groups::tests { + // errors in hipSYCL with bool and 8-bit types - only in group shifts TEMPLATE_LIST_TEST_CASE("Non-uniform-group shift", "[oneapi_non_uniform_groups][group_func][type_list]", @@ -40,3 +42,5 @@ TEMPLATE_LIST_TEST_CASE("Non-uniform-group shift", queue); shift_non_uniform_group(queue); } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift.h b/tests/extension/oneapi_non_uniform_groups/group_shift.h index db39cd3b3..45cbbe499 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_shift.h +++ b/tests/extension/oneapi_non_uniform_groups/group_shift.h @@ -2,7 +2,7 @@ // // SYCL 2020 Conformance Test Suite // -// Copyright (c) 2023 The Khronos Group Inc. +// Copyright (c) 2024 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. @@ -91,8 +91,9 @@ void shift_non_uniform_group(sycl::queue& queue) { T local_var(splat_init(llid + 1)); T shifted_var(splat_init(llid + 1)); - ASSERT_RETURN_TYPE( - T, sycl::shift_group_left(non_uniform_group, local_var), + static_assert( + std::is_same_v, "Return type of shift_group_left(GroupT g, T x) is wrong\n"); shifted_var = @@ -101,8 +102,9 @@ void shift_non_uniform_group(sycl::queue& queue) { equal(shifted_var, splat_init(llid + 2)) || (llid + 1 >= non_uniform_group.get_local_linear_range()); - ASSERT_RETURN_TYPE( - T, sycl::shift_group_left(non_uniform_group, local_var, 3), + static_assert( + std::is_same_v, "Return type of shift_group_left(GroupT g, T x, " "GroupT::linear_id_type delta) is wrong\n"); @@ -112,8 +114,9 @@ void shift_non_uniform_group(sycl::queue& queue) { equal(shifted_var, splat_init(llid + 4)) || (llid + 3 >= non_uniform_group.get_local_linear_range()); - ASSERT_RETURN_TYPE( - T, sycl::shift_group_right(non_uniform_group, local_var), + static_assert( + std::is_same_v, "Return type of shift_group_right(GroupT g, T x) is wrong\n"); shifted_var = @@ -121,8 +124,9 @@ void shift_non_uniform_group(sycl::queue& queue) { res_acc[2 * work_group_size + item.get_local_linear_id()] = equal(shifted_var, splat_init(llid)) || (llid < 1); - ASSERT_RETURN_TYPE( - T, sycl::shift_group_right(non_uniform_group, local_var, 2), + static_assert( + std::is_same_v, "Return type of shift_group_right(GroupT g, T x, " "GroupT::linear_id_type delta) is wrong\n"); diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp b/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp index 55c503f60..483c8daeb 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_shift_fp16.cpp @@ -20,6 +20,8 @@ #include "group_shift.h" +namespace non_uniform_groups::tests { + TEST_CASE("Non-uniform-group shift", "[oneapi_non_uniform_groups][group_func][fp16]") { auto queue = once_per_unit::get_queue(); @@ -42,3 +44,5 @@ TEST_CASE("Non-uniform-group shift", WARN("Device does not support half precision floating point operations."); } } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp b/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp index 208c4a842..047f10764 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp +++ b/tests/extension/oneapi_non_uniform_groups/group_shift_fp64.cpp @@ -20,6 +20,8 @@ #include "group_shift.h" +namespace non_uniform_groups::tests { + TEST_CASE("Non-uniform-group shift", "[oneapi_non_uniform_groups][group_func][fp64]") { auto queue = once_per_unit::get_queue(); @@ -42,3 +44,5 @@ TEST_CASE("Non-uniform-group shift", WARN("Device does not support double precision floating point operations."); } } + +} // namespace non_uniform_groups::tests diff --git a/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp index d63cbd478..c2f02078f 100644 --- a/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp +++ b/tests/extension/oneapi_non_uniform_groups/opportunistic_group_api.cpp @@ -95,7 +95,6 @@ TEST_CASE("Test for opportunistic_group apis.", "[oneapi_non_uniform_groups]") { opportunistic.leader() == (opportunistic.get_local_id() == 0); }); }); - q.wait(); CheckResults results = {}; sycl::accessor acc = results_buffer.get_host_access(); diff --git a/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp index 27629c528..30d72520f 100644 --- a/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp +++ b/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp @@ -97,7 +97,6 @@ TEST_CASE("Test for tangle_group apis.", "[oneapi_non_uniform_groups]") { } }); }); - q.wait(); CheckResults results = {}; sycl::accessor acc = results_buffer.get_host_access(); diff --git a/tests/group_functions/group_shift.h b/tests/group_functions/group_shift.h index d295aff0a..aa395d683 100644 --- a/tests/group_functions/group_shift.h +++ b/tests/group_functions/group_shift.h @@ -2,7 +2,7 @@ // // SYCL 2020 Conformance Test Suite // -// Copyright (c) 2024 The Khronos Group Inc. +// 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. From 0f4992c47c05241938904ed73f9e97222bc7995d Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 13 May 2024 00:18:14 -0700 Subject: [PATCH 10/13] Add TODO Signed-off-by: Larsen, Steffen --- .../oneapi_non_uniform_groups/fixed_size_group_api.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp index 6b9efe1de..f92c06c09 100644 --- a/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp +++ b/tests/extension/oneapi_non_uniform_groups/fixed_size_group_api.cpp @@ -37,6 +37,8 @@ struct checks { }; }; +// TODO: Parameterize the tests on this and test edge-cases, like 1 and +// the sub-group size. constexpr size_t partition_size = 4; TEST_CASE("Test for fixed_size_group apis.", "[oneapi_non_uniform_groups]") { From 1a1fa4293062111153692dc2682115bdcb5659ca Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 13 May 2024 00:18:45 -0700 Subject: [PATCH 11/13] Fix formatting Signed-off-by: Larsen, Steffen --- .../group_broadcast.h | 213 +++++++++--------- 1 file changed, 105 insertions(+), 108 deletions(-) diff --git a/tests/extension/oneapi_non_uniform_groups/group_broadcast.h b/tests/extension/oneapi_non_uniform_groups/group_broadcast.h index b9a14c0e7..c68e142ab 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_broadcast.h +++ b/tests/extension/oneapi_non_uniform_groups/group_broadcast.h @@ -78,114 +78,111 @@ void broadcast_non_uniform_group(sycl::queue& queue) { REQUIRE(executionRange.get_global_range().size() < std::numeric_limits::max() / 100); - cgh.parallel_for>( - executionRange, [=](sycl::nd_item<1> item) { - sycl::sub_group sub_group = item.get_sub_group(); - - // If this item is not participating in the group, leave early. - if (!NonUniformGroupHelper::should_participate(sub_group, - test_case)) - return; - - GroupT non_uniform_group = - NonUniformGroupHelper::create(sub_group, test_case); - - // Each work-item computes a unique value - T value_to_broadcast(splat_init( - static_cast(item.get_global_linear_id() * 100 + - non_uniform_group.get_local_id()))); - - // To simplify the test, we are only checking the first group in - // the first sub-group. - size_t preferred_group_id = NonUniformGroupHelper< - GroupT>::preferred_single_worker_group_id(test_case); - if (item.get_sub_group().get_group_id()[0] == 0 && - non_uniform_group.get_group_id()[0] == preferred_group_id) { - // Find local id of first, last and some third sub-group item in - // between. Will be used to check different combinations of - // broadcasting and receiving work-items - sycl::id<1> first_id = 0; - sycl::id<1> mid_id = non_uniform_group.get_local_range() / 2; - sycl::id<1> last_id = non_uniform_group.get_local_range() - 1; - - // Broadcast from the first work-item - static_assert( - std::is_same_v, - "Return type of group_broadcast(GroupT g, T x) is wrong\n"); - - if (non_uniform_group.leader()) { - // Work-item which does the broadcast stores value to - // broadcast to use it later as a reference - origin_values_acc[0] = value_to_broadcast; - } - auto broadcasted_value = sycl::group_broadcast( - non_uniform_group, value_to_broadcast); - // We read broadcasted value in another work-item - if (non_uniform_group.get_local_id() == last_id) - broadcasted_values_acc[0] = broadcasted_value; - - // Broadcast from the last work-item - static_assert( - std::is_same_v, - "Return type of group_broadcast(GroupT g, T x, " - "GroupT::linear_id_type local_linear_id) is wrong\n"); - - if (non_uniform_group.get_local_id() == last_id) { - // Work-item which does the broadcast stores value to - // broadcast to use it later as a reference - origin_values_acc[1] = value_to_broadcast; - } - - broadcasted_value = sycl::group_broadcast( - non_uniform_group, value_to_broadcast, - non_uniform_group.get_local_linear_range() - 1); - // We read broadcasted value in another work-item - if (non_uniform_group.get_local_id() == mid_id) - broadcasted_values_acc[1] = broadcasted_value; - - // Broadcast from a mid work-item - static_assert( - std::is_same_v, - "Return type of group_broadcast(GroupT g, T x, " - "GroupT::id_type local_id) is wrong\n"); - - if (non_uniform_group.get_local_id() == mid_id) { - // Work-item which does the broadcast stores value to - // broadcast to use it later as a reference - origin_values_acc[2] = value_to_broadcast; - } - broadcasted_value = sycl::group_broadcast( - non_uniform_group, value_to_broadcast, mid_id); - // We read broadcasted value in another work-item - if (non_uniform_group.get_local_id() == first_id) - broadcasted_values_acc[2] = broadcasted_value; - - // Select from the first work-item - static_assert( - std::is_same_v, - "Return type of select_from_group(GroupT g, T x, " - "GroupT::id_type local_id) is wrong\n"); - - if (non_uniform_group.get_local_id() == first_id) { - // Work-item which does the broadcast stores value to - // broadcast to use it later as a reference - origin_values_acc[3] = value_to_broadcast; - } - broadcasted_value = sycl::select_from_group( - non_uniform_group, value_to_broadcast, first_id); - // We read broadcasted value in another work-item - if (non_uniform_group.get_local_id() == mid_id) - broadcasted_values_acc[3] = broadcasted_value; - } - }); + cgh.parallel_for>(executionRange, [=](sycl::nd_item<1> item) { + sycl::sub_group sub_group = item.get_sub_group(); + + // If this item is not participating in the group, leave early. + if (!NonUniformGroupHelper::should_participate(sub_group, + test_case)) + return; + + GroupT non_uniform_group = + NonUniformGroupHelper::create(sub_group, test_case); + + // Each work-item computes a unique value + T value_to_broadcast(splat_init( + static_cast(item.get_global_linear_id() * 100 + + non_uniform_group.get_local_id()))); + + // To simplify the test, we are only checking the first group in + // the first sub-group. + size_t preferred_group_id = + NonUniformGroupHelper::preferred_single_worker_group_id( + test_case); + if (item.get_sub_group().get_group_id()[0] == 0 && + non_uniform_group.get_group_id()[0] == preferred_group_id) { + // Find local id of first, last and some third sub-group item in + // between. Will be used to check different combinations of + // broadcasting and receiving work-items + sycl::id<1> first_id = 0; + sycl::id<1> mid_id = non_uniform_group.get_local_range() / 2; + sycl::id<1> last_id = non_uniform_group.get_local_range() - 1; + + // Broadcast from the first work-item + static_assert( + std::is_same_v, + "Return type of group_broadcast(GroupT g, T x) is wrong\n"); + + if (non_uniform_group.leader()) { + // Work-item which does the broadcast stores value to + // broadcast to use it later as a reference + origin_values_acc[0] = value_to_broadcast; + } + auto broadcasted_value = + sycl::group_broadcast(non_uniform_group, value_to_broadcast); + // We read broadcasted value in another work-item + if (non_uniform_group.get_local_id() == last_id) + broadcasted_values_acc[0] = broadcasted_value; + + // Broadcast from the last work-item + static_assert(std::is_same_v, + "Return type of group_broadcast(GroupT g, T x, " + "GroupT::linear_id_type local_linear_id) is wrong\n"); + + if (non_uniform_group.get_local_id() == last_id) { + // Work-item which does the broadcast stores value to + // broadcast to use it later as a reference + origin_values_acc[1] = value_to_broadcast; + } + + broadcasted_value = sycl::group_broadcast( + non_uniform_group, value_to_broadcast, + non_uniform_group.get_local_linear_range() - 1); + // We read broadcasted value in another work-item + if (non_uniform_group.get_local_id() == mid_id) + broadcasted_values_acc[1] = broadcasted_value; + + // Broadcast from a mid work-item + static_assert(std::is_same_v, + "Return type of group_broadcast(GroupT g, T x, " + "GroupT::id_type local_id) is wrong\n"); + + if (non_uniform_group.get_local_id() == mid_id) { + // Work-item which does the broadcast stores value to + // broadcast to use it later as a reference + origin_values_acc[2] = value_to_broadcast; + } + broadcasted_value = sycl::group_broadcast( + non_uniform_group, value_to_broadcast, mid_id); + // We read broadcasted value in another work-item + if (non_uniform_group.get_local_id() == first_id) + broadcasted_values_acc[2] = broadcasted_value; + + // Select from the first work-item + static_assert(std::is_same_v, + "Return type of select_from_group(GroupT g, T x, " + "GroupT::id_type local_id) is wrong\n"); + + if (non_uniform_group.get_local_id() == first_id) { + // Work-item which does the broadcast stores value to + // broadcast to use it later as a reference + origin_values_acc[3] = value_to_broadcast; + } + broadcasted_value = sycl::select_from_group( + non_uniform_group, value_to_broadcast, first_id); + // We read broadcasted value in another work-item + if (non_uniform_group.get_local_id() == mid_id) + broadcasted_values_acc[3] = broadcasted_value; + } + }); }); } for (int i = 0; i < test_matrix; ++i) { From 75ae361d3633f4843ad942656ed4ea7d6f124ab1 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 24 May 2024 06:13:07 -0700 Subject: [PATCH 12/13] Adjust test-size to adhere to fixed_size requirements Signed-off-by: Larsen, Steffen --- tests/extension/oneapi_non_uniform_groups/group_reduce.h | 2 ++ tests/extension/oneapi_non_uniform_groups/group_scan.h | 4 +++- 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/tests/extension/oneapi_non_uniform_groups/group_reduce.h b/tests/extension/oneapi_non_uniform_groups/group_reduce.h index 6063c9a3e..2b91daa11 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_reduce.h +++ b/tests/extension/oneapi_non_uniform_groups/group_reduce.h @@ -22,6 +22,8 @@ #include "non_uniform_group_common.h" constexpr size_t init = 8; + +// Test size must be a multiple of all the used fixed_size_group sizes. constexpr size_t test_size = 8; template diff --git a/tests/extension/oneapi_non_uniform_groups/group_scan.h b/tests/extension/oneapi_non_uniform_groups/group_scan.h index 98d84534c..e454e0ea6 100644 --- a/tests/extension/oneapi_non_uniform_groups/group_scan.h +++ b/tests/extension/oneapi_non_uniform_groups/group_scan.h @@ -30,7 +30,9 @@ class joint_scan_group_kernel; // smallest type tested. Currently, the smallest type tested is // char/int8_t, so it shouldn't be higher than 127. constexpr int init = 42; -constexpr size_t test_size = 12; + +// Test size must be a multiple of all the used fixed_size_group sizes. +constexpr size_t test_size = 8; template auto joint_inclusive_scan_helper(Group group, T* v_begin, T* v_end, From e03d8eb0eb6fae38f293a6b31b641fcb949ac649 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 24 May 2024 06:17:24 -0700 Subject: [PATCH 13/13] Fix tangle_group local_id expectation Signed-off-by: Larsen, Steffen --- tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp b/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp index 30d72520f..05c59565b 100644 --- a/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp +++ b/tests/extension/oneapi_non_uniform_groups/tangle_group_api.cpp @@ -70,7 +70,7 @@ TEST_CASE("Test for tangle_group apis.", "[oneapi_non_uniform_groups]") { auto run_checks = [&](tangle_group_t tangle, size_t expected_size) { results[checks::get_group_id] = tangle.get_group_id() == 0; - results[checks::get_local_id] = tangle.get_local_id() < split; + results[checks::get_local_id] = tangle.get_local_id() < expected_size; results[checks::get_group_range] = tangle.get_group_range().size() == 1; results[checks::get_local_range] = tangle.get_local_range().size() == expected_size;