Skip to content

Commit

Permalink
Merge pull request #740 from AmirIpma/sub_group_semantics
Browse files Browse the repository at this point in the history
sycl::sub_group by-value semantics tests
  • Loading branch information
bader committed Jul 17, 2023
2 parents 2902e86 + 0e6d7fc commit acc4407
Show file tree
Hide file tree
Showing 3 changed files with 193 additions and 36 deletions.
81 changes: 46 additions & 35 deletions tests/common/invoke.h
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@ namespace {
* @brief Functor to invoke kernels with nd_item in use
* @tparam dim Dimension to use
* @tparam kernelT Type to use as the kernel name
*/
*/
template <int dim, typename kernelT>
struct invoke_nd_item {
static constexpr int dimensions = dim;
Expand All @@ -33,26 +33,23 @@ struct invoke_nd_item {
* @param kernelBody Kernel body to call
*/
template <typename kernelBodyT>
void operator()(sycl::handler &cgh,
sycl::range<dim> numWorkItems,
sycl::range<dim> workGroupSize,
kernelBodyT kernelBody) {

cgh.parallel_for<kernelT>(
sycl::nd_range<dim>(numWorkItems, workGroupSize),
[=](sycl::nd_item<dim> ndItem) {
const size_t index = ndItem.get_global_linear_id();

kernelBody(ndItem, index);
});
void operator()(sycl::handler& cgh, sycl::range<dim> numWorkItems,
sycl::range<dim> workGroupSize, kernelBodyT kernelBody) {
cgh.parallel_for<kernelT>(sycl::nd_range<dim>(numWorkItems, workGroupSize),
[=](sycl::nd_item<dim> ndItem) {
const size_t index =
ndItem.get_global_linear_id();

kernelBody(ndItem, index);
});
}
};

/**
* @brief Functor to invoke kernels with group in use
* @tparam dim Dimension to use
* @tparam kernelT Type to use as the kernel name
*/
*/
template <int dim, typename kernelT>
struct invoke_group {
static constexpr int dimensions = dim;
Expand All @@ -67,31 +64,49 @@ struct invoke_group {
* @param kernelBody Kernel body to call
*/
template <typename kernelBodyT>
void operator()(sycl::handler &cgh,
sycl::range<dim> numWorkItems,
sycl::range<dim> workGroupSize,
kernelBodyT kernelBody) {
void operator()(sycl::handler& cgh, sycl::range<dim> numWorkItems,
sycl::range<dim> workGroupSize, kernelBodyT kernelBody) {
sycl::range<dim> numWorkGroups = numWorkItems / workGroupSize;

cgh.parallel_for_work_group<kernelT>(
numWorkGroups, workGroupSize,
[=](sycl::group<dim> group) {
const size_t index = group.get_linear_id();
numWorkGroups, workGroupSize, [=](sycl::group<dim> group) {
const size_t index = group.get_linear_id();

kernelBody(group, index);
});
kernelBody(group, index);
});
}
};

template <int dim, typename kernelT>
struct invoke_sub_group {
static constexpr int dimensions = dim;
using parameterT = sycl::sub_group;

template <typename kernelBodyT>
void operator()(sycl::handler& cgh, sycl::range<dim> numWorkItems,
sycl::range<dim> workGroupSize, kernelBodyT kernelBody) {
sycl::range<dim> numWorkGroups = numWorkItems / workGroupSize;

cgh.parallel_for<kernelT>(
sycl::nd_range<dim>(numWorkItems, workGroupSize),
[=](sycl::nd_item<3> item) {
const size_t index = item.get_global_linear_id();
sycl::sub_group sub_group = item.get_sub_group();

kernelBody(sub_group, index);
});
}
};

/**
* @brief Generate and store the given number of nd_item/group/h_item instances
* @brief Generate and store the given number of nd_item/group/h_item/sub_group
* instances
* @retval Array of instances
* @tparam numItems Number of instances to store
* @tparam kernelInvokeT Invocation functor to use
*/
template <size_t numItems, class kernelInvokeT>
std::array<typename kernelInvokeT::parameterT, numItems> store_instances()
{
std::array<typename kernelInvokeT::parameterT, numItems> store_instances() {
constexpr auto numDims = kernelInvokeT::dimensions;
using item_t = typename kernelInvokeT::parameterT;
using item_array_t = std::array<item_t, numItems>;
Expand All @@ -104,25 +119,21 @@ std::array<typename kernelInvokeT::parameterT, numItems> store_instances()
sycl_cts::util::get_cts_object::range<numDims>::get(numItems, 1, 1);

{
sycl::buffer<item_t> itemBuf(items.data(),
sycl::range<1>(items.size()));
sycl::buffer<item_t> itemBuf(items.data(), sycl::range<1>(items.size()));

auto queue = sycl_cts::util::get_cts_object::queue();
queue.submit([&](sycl::handler& cgh) {
auto itemAcc =
itemBuf.template get_access<sycl::access_mode::write>(cgh);
auto itemAcc = itemBuf.template get_access<sycl::access_mode::write>(cgh);

kernelInvokeT{}(
cgh, itemRange, oneElemRange,
[=](item_t& item, const size_t index) {
itemAcc[index] = item;
});
[=](item_t& item, const size_t index) { itemAcc[index] = item; });
});
queue.wait_and_throw();
}
return items;
}

} // namespace
} // namespace

#endif // __SYCLCTS_TESTS_COMMON_INVOKE_H
#endif // __SYCLCTS_TESTS_COMMON_INVOKE_H
2 changes: 1 addition & 1 deletion tests/common/semantics_by_value.h
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,7 @@ namespace common_by_value_semantics {
* @brief Provides a safe index for checking an operation
*/
enum class current_check : size_t {
reflexivity_equal_self,
reflexivity_equal_self = 0,
reflexivity_not_equal_self,
equal_copy,
equal_copy_symmetry,
Expand Down
146 changes: 146 additions & 0 deletions tests/sub_group/sub_group_semantics.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,146 @@
/*******************************************************************************
//
// SYCL 2020 Conformance Test Suite
//
// Copyright (c) 2023 The Khronos Group Inc.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
*******************************************************************************/

#include "../common/common.h"
#include "../common/invoke.h"
#include "../common/once_per_unit.h"
#include "../common/semantics_by_value.h"

namespace sub_group_semantics {

enum class op_codes : size_t {
ctor_copy = 0,
ctor_move,
assign_copy,
assign_move,
code_count
};

constexpr size_t error_count = to_integral(op_codes::code_count);

static constexpr size_t sizes[] = {16, 32, 64};

static const std::array<std::string, error_count> error_strings{
"sub_group with sub_group was not constructed correctly",
"sub_group with sub_group was not move constructed correctly",
"sub_group with sub_group was not copy assigned correctly",
"sub_group with sub_group was not move assigned correctly",
};

template <op_codes Code, typename ResultArray>
void set_success_operation(ResultArray& result, bool success) {
int index = to_integral(Code);
result[index] = success;
}

std::string get_error_string(int code) { return error_strings[code]; }

struct sub_group_semantics_kernel;
struct sub_group_equality_kernel;
struct setup_kernel;

bool check_equality_by_id(const sycl::sub_group& actual,
sycl::id<1>* expected_ids) {
return actual.get_group_id() == expected_ids[0] &&
actual.get_local_id() == expected_ids[1];
}

template <typename ResultArray>
void check_by_value_semantics(sycl::sub_group& sub_group, ResultArray& result) {
sycl::id<1> expected_ids[] = {sub_group.get_group_id(),
sub_group.get_local_id()};
// Check copy constructor
sycl::sub_group copied(sub_group);
set_success_operation<op_codes::ctor_copy>(
result, check_equality_by_id(copied, expected_ids));

// Check copy assignment
sycl::sub_group copy_assigned(sub_group);
copy_assigned = sub_group;
set_success_operation<op_codes::assign_copy>(
result, check_equality_by_id(copy_assigned, expected_ids));

// Check move constructor; invalidates sub_group
sycl::sub_group moved(std::move(sub_group));
set_success_operation<op_codes::ctor_move>(
result, check_equality_by_id(moved, expected_ids));

// Check move assignment
sycl::sub_group move_assigned(copy_assigned);
move_assigned = std::move(copy_assigned);
set_success_operation<op_codes::assign_move>(
result, check_equality_by_id(move_assigned, expected_ids));
}

TEST_CASE("sub_group by-value semantics", "[sub_group]") {
bool result[error_count];
std::fill(result, result + error_count, false);
{
sycl::buffer<bool, 1> res_buf(result, sycl::range(error_count));

sycl::queue queue = once_per_unit::get_queue();
const sycl::range<3> r{1, 1, 1};
sycl::nd_range<3> nd_range(r, r);
queue
.submit([&](sycl::handler& cgh) {
auto res_acc = res_buf.get_access<sycl::access_mode::read_write>(cgh);
cgh.parallel_for<sub_group_semantics_kernel>(
nd_range, [=](sycl::nd_item<3> nd_item) {
sycl::sub_group sub_group = nd_item.get_sub_group();
check_by_value_semantics(sub_group, res_acc);
});
})
.wait_and_throw();
}
for (int i = 0; i < error_count; ++i) {
INFO(get_error_string(i));
CHECK(result[i]);
}
}

// FIXME: re-enable when == and != operators is implemented
#if !SYCL_CTS_COMPILING_WITH_HIPSYCL && !SYCL_CTS_COMPILING_WITH_COMPUTECPP
TEST_CASE("Check sycl::sub_group equality", "[sub_group]") {
size_t code_count =
to_integral(common_by_value_semantics::current_check::size);
bool result[code_count];
std::fill(result, result + code_count, false);
auto items = store_instances<2, invoke_sub_group<3, setup_kernel>>();
{
sycl::buffer<bool, 1> res_buf(result, sycl::range(code_count));
auto queue = once_per_unit::get_queue();
queue
.submit([&](sycl::handler& cgh) {
auto res_acc = res_buf.get_access(cgh);
cgh.single_task<sub_group_equality_kernel>([=] {
common_by_value_semantics::check_equality(items[0], items[1],
res_acc);
});
})
.wait_and_throw();
}
for (int i = 0; i < code_count; ++i) {
INFO(common_by_value_semantics::get_error_string(i));
CHECK(result[i]);
}
}
#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL &&
// !SYCL_CTS_COMPILING_WITH_COMPUTECPP
} // namespace sub_group_semantics

0 comments on commit acc4407

Please sign in to comment.