Skip to content

Commit

Permalink
Merge branch 'SYCL-2020' into root_group_test_option
Browse files Browse the repository at this point in the history 
Signed-off-by: Michael Aziz <[email protected]>
  • Loading branch information
@0x12CC
0x12CC committed Jun 10, 2024
2 parents 0ead589 + 7e1d2a2 commit 78d2a27
Show file tree
Hide file tree
Showing 36 changed files with 4,353 additions and 78 deletions.
4 changes: 4 additions & 0 deletions CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -91,6 +91,10 @@ add_cts_option(SYCL_CTS_ENABLE_EXT_ONEAPI_COMPOSITE_DEVICE_TESTS
"Enable extension oneAPI composite_device 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})

add_cts_option(SYCL_CTS_ENABLE_EXT_ONEAPI_ROOT_GROUP_TESTS
"Enable extension oneAPI root group tests" OFF
FORCE_ON ${SYCL_CTS_ENABLE_EXT_ONEAPI_TESTS})
Expand Down
6 changes: 5 additions & 1 deletion tests/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -198,12 +198,16 @@ function(add_cts_test)
# To make check that any .cpp files are passed
# List created because, direct check on "${ARGN}" gives false-positive result
set(tests_list "${ARGN}")
get_filename_component(test_dir ${CMAKE_CURRENT_SOURCE_DIR} NAME)
if (tests_list)
get_filename_component(test_dir ${CMAKE_CURRENT_SOURCE_DIR} NAME)
set(test_exe_name ${test_dir})
set(test_cases_list "${ARGN}")

add_cts_test_helper(${test_exe_name} "${test_cases_list}")
else()
if(${test_dir} IN_LIST exclude_categories)
message(STATUS "Skipping excluded test: " test_${test_dir})
endif()
endif()
endfunction()

Expand Down
2 changes: 1 addition & 1 deletion tests/buffer/buffer_storage_common.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -70,7 +70,7 @@ class buffer_storage_test {

// Case 5 - Vector data
std::vector<T> data_vector;
data_vector.reserve(size);
data_vector.resize(size);
auto data_final5 = data_vector.begin();

check_write_back(log, r, data_final1.get());
Expand Down
41 changes: 41 additions & 0 deletions tests/extension/oneapi_non_uniform_groups/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,41 @@
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()
174 changes: 174 additions & 0 deletions tests/extension/oneapi_non_uniform_groups/ballot_group_api.cpp
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,174 @@
/*******************************************************************************
//
// 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<sycl::sub_group>;
using CheckResults = bool[checks::COUNT];

constexpr size_t num_items = 64;

sycl::buffer<CheckResults, 1> 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<struct BallotGroupAPI>(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<decltype(ballot), ballot_group_t>);

// 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);
});
});

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<ballot_group_t>::value);
STATIC_CHECK(sycl::is_group_v<ballot_group_t>);
STATIC_CHECK(oneapi_ext::is_user_constructed_group<ballot_group_t>::value);
STATIC_CHECK(oneapi_ext::is_user_constructed_group_v<ballot_group_t>);
STATIC_CHECK(!oneapi_ext::is_fixed_topology_group<ballot_group_t>::value);
STATIC_CHECK(!oneapi_ext::is_fixed_topology_group_v<ballot_group_t>);

// Aliases.
STATIC_CHECK(std::is_same_v<ballot_group_t::id_type, sycl::id<1>>);
STATIC_CHECK(std::is_same_v<ballot_group_t::range_type, sycl::range<1>>);
STATIC_CHECK(std::is_same_v<ballot_group_t::linear_id_type, uint32_t>);

// 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<decltype(std::declval<ballot_group_t>().get_group_id()),
ballot_group_t::id_type>);
CHECK(results[checks::get_group_id]);

// get_local_id
CHECK(std::is_same_v<decltype(std::declval<ballot_group_t>().get_local_id()),
ballot_group_t::id_type>);
CHECK(results[checks::get_local_id]);

// get_group_range
CHECK(
std::is_same_v<decltype(std::declval<ballot_group_t>().get_group_range()),
ballot_group_t::range_type>);
CHECK(results[checks::get_group_range]);

// get_local_range
CHECK(
std::is_same_v<decltype(std::declval<ballot_group_t>().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<ballot_group_t>().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<ballot_group_t>().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<ballot_group_t>().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<ballot_group_t>().get_local_linear_range()),
ballot_group_t::linear_id_type>);
CHECK(results[checks::get_local_linear_range]);

// leader
CHECK(
std::is_same_v<decltype(std::declval<ballot_group_t>().leader()), bool>);
CHECK(results[checks::leader]);
#endif
}

} // namespace non_uniform_groups::tests
Loading

0 comments on commit 78d2a27

Please sign in to comment.