diff --git a/tests/common/common.h b/tests/common/common.h index e9f04d932..41ad9243c 100644 --- a/tests/common/common.h +++ b/tests/common/common.h @@ -752,4 +752,22 @@ inline bool kernel_supports_wg_size(sycl_cts::util::logger& log, } \ } \ } + +/// Linearizes a multi-dimensional index according to the specification. +template +size_t linearize(sycl::range range, sycl::id id); + +inline size_t linearize(sycl::range<1> range, sycl::id<1> id) { + static_cast(range); + return id[0]; +} + +inline size_t linearize(sycl::range<2> range, sycl::id<2> id) { + return id[1] + id[0] * range[1]; +} + +inline size_t linearize(sycl::range<3> range, sycl::id<3> id) { + return id[2] + id[1] * range[2] + id[0] * range[1] * range[2]; +} + #endif // __SYCLCTS_TESTS_COMMON_COMMON_H diff --git a/tests/group/group_api.cpp b/tests/group/group_api.cpp index 99f7da20b..b4829beae 100644 --- a/tests/group/group_api.cpp +++ b/tests/group/group_api.cpp @@ -19,376 +19,462 @@ // *******************************************************************************/ -#include "../common/common.h" #include "../../util/array.h" +#include "../common/common.h" -#define TEST_NAME group_api +#include +#include +#include -namespace TEST_NAMESPACE { using namespace sycl_cts; -static const size_t GROUP_RANGE_1D = 2; -static const size_t GROUP_RANGE_2D = 4; -static const size_t GROUP_RANGE_3D = 8; -static const size_t DEFAULT_LOCAL_RANGE[3] = {4, 3, 2}; -static const size_t NUM_GROUPS = - GROUP_RANGE_1D * GROUP_RANGE_2D * GROUP_RANGE_3D; -static const size_t NUM_METHODS = 9; - -class getter -{ -public: +constexpr size_t GROUP_RANGE[3] = {2, 4, 8}; +constexpr size_t DEFAULT_LOCAL_RANGE[3] = {4, 3, 2}; +constexpr size_t GROUP_LINEAR_RANGE = + GROUP_RANGE[0] * GROUP_RANGE[1] * GROUP_RANGE[2]; +constexpr size_t DEFAULT_LOCAL_LINEAR_RANGE = + DEFAULT_LOCAL_RANGE[0] * DEFAULT_LOCAL_RANGE[1] * DEFAULT_LOCAL_RANGE[2]; + +class getter { + public: enum class method : size_t { - get = 0, - get_dims = 1, - local_range = 2, - local_range_dims = 3, - global_range = 4, - global_range_dims = 5, - group_range = 6, - group_range_dims = 7, - subscript = 8, + group_id = 0, + group_id_dims, + local_id, + local_id_dims, + local_range, + local_range_dims, + group_range, + group_range_dims, + max_local_range, + subscript, + group_linear_id, + local_linear_id, + group_linear_range, + local_linear_range, + leader, + method_count // defines size, should be last }; - static inline size_t get_index(size_t groupLinearID, - getter::method getterMethod) { - const auto offset = to_integral(getterMethod); - return (groupLinearID * NUM_METHODS) + offset; - } + static constexpr auto method_cnt = to_integral(method::method_count); - static const char *name(getter::method getterMethod) { + static const char* name(getter::method getterMethod) { switch (getterMethod) { - case method::get: - return "get()"; - case method::get_dims: - return "get(int)"; + case method::group_id: + return "get_group_id()"; + case method::group_id_dims: + return "get_group_id(int)"; + case method::local_id: + return "get_local_id()"; + case method::local_id_dims: + return "get_local_id(int)"; case method::local_range: return "get_local_range()"; case method::local_range_dims: return "get_local_range(int)"; - case method::global_range: - return "get_global_range()"; - case method::global_range_dims: - return "get_global_range(int)"; case method::group_range: return "get_group_range()"; case method::group_range_dims: return "get_group_range(int)"; + case method::max_local_range: + return "get_max_local_range"; case method::subscript: return "operator[](int)"; - default: - return "__unknown__"; + case method::group_linear_id: + return "get_group_linear_id()"; + case method::local_linear_id: + return "get_local_linear_id()"; + case method::group_linear_range: + return "get_group_linear_range()"; + case method::local_linear_range: + return "get_local_linear_range()"; + case method::leader: + return "get_leader()"; + case method::method_count: + return "invalid enum value"; } + // no default case to allow for compiler warning + return nullptr; } }; -template +template class test_kernel; template class test_helper { -public: - using range_t = sycl::range; - -private: - static constexpr int NUM_RESULTS = NUM_GROUPS * NUM_METHODS; - struct call_result_t - { + private: + /** Maximum size. */ + static constexpr int NUM_RESULTS = + GROUP_LINEAR_RANGE * DEFAULT_LOCAL_LINEAR_RANGE * getter::method_cnt; + struct call_result_t { bool hasValidType; sycl_cts::util::array values; }; - std::vector m_callResults; - sycl::range m_globalRange; - sycl::range m_localRange; + std::vector m_callResults{}; + sycl::range m_globalRange; + sycl::range m_localRange; + + /** + * Convert a linear group id, a linear local range, a linear local id, + * and a getter method to an index in \p m_callResults. */ + static inline size_t get_index(size_t glid, size_t llrange, size_t llid, + getter::method getterMethod) { + const auto offset = to_integral(getterMethod); + return glid * llrange * getter::method_cnt + llid * getter::method_cnt + + offset; + } -public: + public: test_helper(sycl::range globalRange, - sycl::range localRange): - m_callResults(NUM_RESULTS), - m_globalRange(globalRange), - m_localRange(localRange) { - for (size_t i = 0; i < NUM_RESULTS; i++) { - auto& callResult = m_callResults.data()[i]; - callResult.hasValidType = false; - for (auto& value: callResult.values) - value = 0; + sycl::range localRange) + : m_callResults(NUM_RESULTS), + m_globalRange(globalRange), + m_localRange(localRange) { + for (size_t i = 0; i < NUM_RESULTS; i++) { + auto& callResult = m_callResults[i]; + callResult.hasValidType = false; + for (size_t j = 0; j < dimensions; j++) { + // special value to verify that a value was written + callResult.values[j] = std::numeric_limits::max(); } + } } - void collect_group_indicies(sycl::queue& queue) { - sycl::buffer buf(m_callResults.data(), - sycl::range<1>(NUM_RESULTS)); - - queue.submit([&](sycl::handler &cgh) { - auto a_dev = - buf.template get_access(cgh); - - cgh.parallel_for_work_group>( - m_globalRange, - m_localRange, - [=](sycl::group my_group) { - - const size_t groupLinearID = my_group.get_linear_id(); - - // get() - { - call_result_t& callResult = - a_dev[getter::get_index(groupLinearID, getter::method::get)]; - - auto m_get_group = my_group.get_id(); - for (size_t i = 0; i < dimensions; ++i) - callResult.values[i] = m_get_group.get(i); - - using expected_t = sycl::id; - callResult.hasValidType = - std::is_same::value; - } - - // get(int) - { - call_result_t& callResult = - a_dev[getter::get_index(groupLinearID, getter::method::get_dims)]; - - for (size_t i = 0; i < dimensions; ++i) - callResult.values[i] = my_group.get_id(i); - - using expected_t = size_t; - callResult.hasValidType = - std::is_same::value; - } - - // get_local_range() - { - call_result_t& callResult = a_dev[getter::get_index(groupLinearID, getter::method::local_range)]; - - auto m_get_local_range = my_group.get_local_range(); - for (size_t i = 0; i < dimensions; ++i) - callResult.values[i] = m_get_local_range.get(i); - - using expected_t = sycl::range; - callResult.hasValidType = - std::is_same::value; - } - - // get_local_range(int) - { - call_result_t& callResult = a_dev[getter::get_index(groupLinearID, getter::method::local_range_dims)]; - - for (size_t i = 0; i < dimensions; ++i) - callResult.values[i] = my_group.get_local_range(i); - - using expected_t = size_t; - callResult.hasValidType = - std::is_same::value; - } - - // get_global_range() - { - call_result_t& callResult = a_dev[getter::get_index(groupLinearID, getter::method::global_range)]; - - auto m_get_global_range = my_group.get_global_range(); - for (size_t i = 0; i < dimensions; ++i) - callResult.values[i] = m_get_global_range.get(i); - - using expected_t = sycl::range; - callResult.hasValidType = - std::is_same::value; - } - - // get_global_range(int) - { - call_result_t& callResult = a_dev[getter::get_index(groupLinearID, getter::method::global_range_dims)]; - - for (size_t i = 0; i < dimensions; ++i) - callResult.values[i] = my_group.get_global_range(i); - - using expected_t = size_t; - callResult.hasValidType = - std::is_same::value; - } - - // get_group_range() - { - call_result_t& callResult = a_dev[getter::get_index(groupLinearID, getter::method::group_range)]; - - auto m_get_group_range = my_group.get_group_range(); - for (size_t i = 0; i < dimensions; ++i) - callResult.values[i] = m_get_group_range.get(i); - - using expected_t = sycl::range; - callResult.hasValidType = - std::is_same::value; - } - - // get_group_range(int) - { - call_result_t& callResult = a_dev[getter::get_index(groupLinearID, getter::method::group_range_dims)]; - - for (size_t i = 0; i < dimensions; ++i) - callResult.values[i] = my_group.get_group_range(i); - - using expected_t = size_t; - callResult.hasValidType = - std::is_same::value; - } + /** Set the call result if the return type is a one-dimensional value. */ + template + static void set_val(call_result_t& call_result, data_type val) { + call_result.values[0] = val; + } - // operator[] - { - call_result_t& callResult = a_dev[getter::get_index(groupLinearID, getter::method::subscript)]; + /** Set the call result if the return type is a multi-dimensional id. */ + static void set_id(call_result_t& call_result, sycl::id id) { + for (size_t i = 0; i < dimensions; ++i) { + call_result.values[i] = id.get(i); + } + } - for (size_t i = 0; i < dimensions; ++i) - callResult.values[i] = my_group[i]; + /** + * Set the call result if the return type provides a \p get method + * for each index. */ + template + static void set_for_dim(call_result_t& call_result, function func) { + for (size_t i = 0; i < dimensions; ++i) { + call_result.values[i] = func(i); + } + } - using expected_t = size_t; - callResult.hasValidType = - std::is_same::value; - } - }); + /** + * Launch a kernel to collect the results of calling the group's member + * functions. */ + void collect_results(sycl::queue& queue) { + sycl::buffer buffer(m_callResults.data(), + sycl::range<1>(NUM_RESULTS)); + + const size_t local_linear_range = m_localRange.size(); + queue.submit([&](sycl::handler& cgh) { + auto accessor_device = + buffer.template get_access(cgh); + // use parallel_for(nd_range) to execute a kernel and control the + // number of work-groups and work-items + cgh.parallel_for>( + sycl::nd_range(m_globalRange * m_localRange, + m_localRange), + [=](sycl::nd_item item) { + // obtain indices independently of group api + const size_t glid = item.get_group_linear_id(); + const size_t llrange = local_linear_range; + const size_t llid = item.get_local_linear_id(); + sycl::group group = item.get_group(); + + // helper function to obtain call result struct + const auto get_res = + [=](const getter::method& method) -> call_result_t& { + return accessor_device[get_index(glid, llrange, llid, method)]; + }; + + { // get_group_id() + call_result_t& res = get_res(getter::method::group_id); + set_id(res, group.get_group_id()); + res.hasValidType = std::is_same_v, + decltype(group.get_group_id())>; + } + { // get_group_id(int) + call_result_t& res = get_res(getter::method::group_id_dims); + set_for_dim(res, [=](size_t i) { return group.get_group_id(i); }); + res.hasValidType = + std::is_same_v; + } + { // get_local_id() + call_result_t& res = get_res(getter::method::local_id); + set_id(res, group.get_local_id()); + res.hasValidType = std::is_same_v, + decltype(group.get_local_id())>; + } + { // get_local_id(int) + call_result_t& res = get_res(getter::method::local_id_dims); + set_for_dim(res, [=](size_t i) { return group.get_local_id(i); }); + res.hasValidType = + std::is_same_v; + } + { // get_local_range() + call_result_t& res = get_res(getter::method::local_range); + set_id(res, group.get_local_range()); + res.hasValidType = + std::is_same_v, + decltype(group.get_local_range())>; + } + { // get_local_range(int) + call_result_t& res = get_res(getter::method::local_range_dims); + set_for_dim(res, + [=](size_t i) { return group.get_local_range(i); }); + res.hasValidType = + std::is_same_v; + } + { // get_group_range() + call_result_t& res = get_res(getter::method::group_range); + set_id(res, group.get_group_range()); + res.hasValidType = + std::is_same_v, + decltype(group.get_group_range())>; + } + { // get_group_range(int) + call_result_t& res = get_res(getter::method::group_range_dims); + set_for_dim(res, + [=](size_t i) { return group.get_group_range(i); }); + res.hasValidType = + std::is_same_v; + } + { // get_max_local_range + call_result_t& res = get_res(getter::method::max_local_range); + set_id(res, group.get_max_local_range()); + res.hasValidType = + std::is_same_v, + decltype(group.get_max_local_range())>; + } + { // operator[] + call_result_t& res = get_res(getter::method::subscript); + set_for_dim(res, [=](size_t i) { return group[i]; }); + res.hasValidType = std::is_same_v; + } + { // get_group_linear_id + call_result_t& res = get_res(getter::method::group_linear_id); + set_val(res, group.get_group_linear_id()); + res.hasValidType = + std::is_same_v; + } + { // get_group_linear_range + call_result_t& res = get_res(getter::method::group_linear_range); + set_val(res, group.get_group_linear_range()); + res.hasValidType = + std::is_same_v; + } + { // get_local_linear_id + call_result_t& res = get_res(getter::method::local_linear_id); + set_val(res, group.get_local_linear_id()); + res.hasValidType = + std::is_same_v; + } + { // get_local_linear_range + call_result_t& res = get_res(getter::method::local_linear_range); + set_val(res, group.get_local_linear_range()); + res.hasValidType = + std::is_same_v; + } + { // leader + call_result_t& res = get_res(getter::method::leader); + set_val(res, group.leader()); + res.hasValidType = std::is_same_v; + } + }); }); } - void validate_group_indicies(util::logger &log) const - { - // For each work item - validate_group_indicies(log, std::integral_constant{}); + static std::string format_range(sycl::range range) { + std::ostringstream ss; + ss << "["; + for (size_t i = 0; i < dimensions; i++) { + ss << range[i] << (i + 1 < dimensions ? ", " : "]"); + } + return ss.str(); } -private: - void validate_group_indicies( - util::logger &log, std::integral_constant loopSelector) const - { + /** Validates the results obtained by \p collect_results. */ + void validate_results() const { + INFO("group range: " << format_range(m_globalRange) + << " local range: " << format_range(m_localRange)); + + // for each work-group and each work-item + validate_results(std::integral_constant{}); + } + + private: + void validate_results(std::integral_constant loopSelector) const { static_cast(loopSelector); - for (size_t groupID0 = 0; groupID0 < m_globalRange[0]; ++groupID0) { - validate_group_indicies(log, groupID0, std::array{groupID0}); + for (size_t gid0 = 0; gid0 < m_globalRange[0]; ++gid0) { + size_t glid = linearize({m_globalRange[0]}, {gid0}); + for (size_t lid0 = 0; lid0 < m_localRange[0]; ++lid0) { + size_t llid = linearize({m_localRange[0]}, {lid0}); + validate_results_impl(glid, {gid0}, llid, {lid0}); + } } } - void validate_group_indicies( - util::logger &log, std::integral_constant loopSelector) const - { + void validate_results(std::integral_constant loopSelector) const { static_cast(loopSelector); - for (size_t groupID0 = 0; groupID0 < m_globalRange[0]; ++groupID0) { - for (size_t groupID1 = 0; groupID1 < m_globalRange[1]; ++groupID1) { - const size_t groupLinearID = groupID1 + (groupID0 * GROUP_RANGE_2D); - validate_group_indicies( - log, groupLinearID, std::array{groupID0, groupID1}); + for (size_t gid0 = 0; gid0 < m_globalRange[0]; ++gid0) { + for (size_t gid1 = 0; gid1 < m_globalRange[1]; ++gid1) { + size_t glid = + linearize({m_globalRange[0], m_globalRange[1]}, {gid0, gid1}); + for (size_t lid0 = 0; lid0 < m_localRange[0]; ++lid0) { + for (size_t lid1 = 0; lid1 < m_localRange[1]; ++lid1) { + size_t llid = + linearize({m_localRange[0], m_localRange[1]}, {lid0, lid1}); + validate_results_impl(glid, {gid0, gid1}, llid, {lid0, lid1}); + } + } } } } - void validate_group_indicies( - util::logger &log, std::integral_constant loopSelector) const - { + void validate_results(std::integral_constant loopSelector) const { static_cast(loopSelector); - for (size_t groupID0 = 0; groupID0 < m_globalRange[0]; ++groupID0) { - for (size_t groupID1 = 0; groupID1 < m_globalRange[1]; ++groupID1) { - for (size_t groupID2 = 0; groupID2 < m_globalRange[2]; ++groupID2) { - const size_t groupLinearID = - (groupID2 + (groupID1 * GROUP_RANGE_3D) + - (groupID0 * GROUP_RANGE_3D * GROUP_RANGE_2D)); - validate_group_indicies( - log, - groupLinearID, - std::array{groupID0, groupID1, groupID2}); + for (size_t gid0 = 0; gid0 < m_globalRange[0]; ++gid0) { + for (size_t gid1 = 0; gid1 < m_globalRange[1]; ++gid1) { + for (size_t gid2 = 0; gid2 < m_globalRange[2]; ++gid2) { + size_t glid = + linearize({m_globalRange[0], m_globalRange[1], m_globalRange[2]}, + {gid0, gid1, gid2}); + for (size_t lid0 = 0; lid0 < m_localRange[0]; ++lid0) { + for (size_t lid1 = 0; lid1 < m_localRange[1]; ++lid1) { + for (size_t lid2 = 0; lid2 < m_localRange[2]; ++lid2) { + size_t llid = linearize( + {m_localRange[0], m_localRange[1], m_localRange[2]}, + {lid0, lid1, lid2}); + validate_results_impl(glid, {gid0, gid1, gid2}, llid, + {lid0, lid1, lid2}); + } + } + } } } } } - void validate_group_indicies(util::logger &log, size_t groupLinearId, - std::array groupId) const - { - // get(), get(int), operator[] - { - const auto& expected = groupId; - check_indices(log, groupLinearId, getter::method::get, expected); - check_indices(log, groupLinearId, getter::method::get_dims, expected); - check_indices(log, groupLinearId, getter::method::subscript, expected); + void validate_results_impl(size_t glid, std::array gid, + size_t llid, + std::array lid) const { + { // operator[], get_group_id(), get_group_id(int) + const std::array& expected = gid; + check(glid, llid, getter::method::subscript, expected); + check(glid, llid, getter::method::group_id, expected); + check(glid, llid, getter::method::group_id_dims, expected); } - - // get_local_range(), get_local_range(int) - { - const auto& expected = get_local_range_values(); - check_indices(log, groupLinearId, getter::method::local_range, expected); - check_indices(log, groupLinearId, getter::method::local_range_dims, - expected); + { // get_local_id(), get_local_id(int) + const std::array expected = lid; + check(glid, llid, getter::method::local_id, expected); + check(glid, llid, getter::method::local_id_dims, expected); } - - // get_global_range(), get_global_range(int) - { - const auto& expected = get_global_range_values(); - check_indices(log, groupLinearId, getter::method::global_range, expected); - check_indices(log, groupLinearId, getter::method::global_range_dims, - expected); + { // get_local_range(), get_local_range(int) + const std::array& expected = get_local_range_values(); + check(glid, llid, getter::method::local_range, expected); + check(glid, llid, getter::method::local_range_dims, expected); } - - // get_group_range(), get_group_range(int) - { - const std::array expected {GROUP_RANGE_1D, GROUP_RANGE_2D, - GROUP_RANGE_3D}; - check_indices(log, groupLinearId, getter::method::group_range, expected); - check_indices(log, groupLinearId, getter::method::group_range_dims, - expected); + { // get_group_range(), get_group_range(int) + std::array expected; + for (size_t i = 0; i < dimensions; i++) { + expected[i] = GROUP_RANGE[i]; + } + check(glid, llid, getter::method::group_range, expected); + check(glid, llid, getter::method::group_range_dims, expected); + } + { // get_max_local_range() + size_t expected = 0; + for (size_t i = 0; i < dimensions; i++) { + expected = std::max(expected, m_localRange.get(i)); + } + check(glid, llid, getter::method::max_local_range, expected); + } + { // get_group_linear_id() + const size_t expected = glid; + check(glid, llid, getter::method::group_linear_id, expected); + } + { // get_local_linear_id() + const size_t expected = llid; + check(glid, llid, getter::method::local_linear_id, expected); + } + { // get_group_linear_range() + size_t expected = 1; + for (size_t i = 0; i < dimensions; i++) { + expected *= m_globalRange.get(i); + } + check(glid, llid, getter::method::group_linear_range, expected); + } + { // get_local_linear_range() + size_t expected = 1; + for (size_t i = 0; i < dimensions; i++) { + expected *= m_localRange.get(i); + } + check(glid, llid, getter::method::local_linear_range, expected); + } + { // leader() + const size_t expected = llid == 0; + check(glid, llid, getter::method::leader, expected); } } + /** Checks the result for functions that return a multi-dimensional value. */ template - void check_indices( - util::logger &log, size_t groupLinearId, - getter::method getterMethod, - const std::array& expected) const - { - static_assert(expectedDimensions >= dimensions, - "Invalid call for check_indices"); - - const auto& callResult = m_callResults[getter::get_index(groupLinearId, - getterMethod)]; - for (size_t dim = 0; dim < dimensions; ++dim) - { - if (!CHECK_VALUE(log, callResult.values[dim], expected[dim], - static_cast(dim))) { - log.note(" -> group %d: %s", groupLinearId, - getter::name(getterMethod)); - }; - } - if (!callResult.hasValidType) { - FAIL(log, "Invalid return value type"); - log.note(" -> group %d: %s", groupLinearId, - getter::name(getterMethod)); + void check(size_t glid, size_t llid, getter::method getterMethod, + const std::array& expected) const { + static_assert(expectedDimensions >= dimensions, "Invalid call for check"); + INFO("linear group id: " << glid << ", linear local id: " << llid); + INFO("" << getter::name(getterMethod)); + + const size_t llrange = m_localRange.size(); + const call_result_t& callResult = + m_callResults[get_index(glid, llrange, llid, getterMethod)]; + for (size_t dim = 0; dim < dimensions; ++dim) { + INFO("dim: " << dim); + INFO("actual: " << callResult.values[dim] + << " expected: " << expected[dim]); + CHECK((callResult.values[dim] == expected[dim])); } + CHECK(callResult.hasValidType); } - std::array get_local_range_values() const - { - std::array result{}; - int i = 0; - std::generate(result.begin(), result.end(), [this, &i] () mutable { - return m_localRange.get(i++); - }); - return result; + /** Checks the result for functions that return a one-dimensional value. */ + void check(size_t glid, size_t llid, getter::method getterMethod, + size_t expected) const { + INFO("linear group id: " << glid << ", linear local id: " << llid); + INFO("" << getter::name(getterMethod)); + + const size_t llrange = m_localRange.size(); + const call_result_t& callResult = + m_callResults[get_index(glid, llrange, llid, getterMethod)]; + + INFO("actual: " << callResult.values[0] << " expected: " << expected); + CHECK((callResult.values[0] == expected)); + CHECK(callResult.hasValidType); } - std::array get_global_range_values() const - { + + std::array get_local_range_values() const { std::array result{}; int i = 0; - std::generate(result.begin(), result.end(), [this, &i] () mutable { - const auto index = i++; - return m_globalRange.get(index) * m_localRange.get(index); - }); + std::generate(result.begin(), result.end(), + [this, &i]() mutable { return m_localRange.get(i++); }); return result; } }; +/** + * Checks a work-group size against the maximum as defined by the device + * associated with \p queue. */ template -bool reduce_size(sycl::queue& queue) { - bool res = true; +bool wg_size_too_large(sycl::queue& queue, + sycl::range local_range) { using k_name = test_kernel; auto ctx = queue.get_context(); auto kb = @@ -399,69 +485,49 @@ bool reduce_size(sycl::queue& queue) { auto work_group_size_limit = device.template get_info(); - size_t default_wg_size = 1; + size_t wg_size = 1; for (size_t dim = 0; dim < dimensions; ++dim) { - default_wg_size *= DEFAULT_LOCAL_RANGE[dim]; + wg_size *= local_range[dim]; } - res = default_wg_size > work_group_size_limit; - return res; + return wg_size > work_group_size_limit; } -class TEST_NAME : public util::test_base { - public: - /** return information about this test - */ - void get_info(test_base::info &out) const override { - set_test_info(out, TOSTRING(TEST_NAME), TEST_FILE); - } +TEST_CASE("group api", "[group]") { + auto queue = util::get_cts_object::queue(); - /** execute the test - */ - void run(util::logger &log) override { - { - auto queue = util::get_cts_object::queue(); - - // Validate for each dimension possible - { - bool reduce_wg_size = reduce_size<1>(queue); - size_t LOCAL_RANGE_1D = reduce_wg_size ? 1 : DEFAULT_LOCAL_RANGE[0]; - auto validator = test_helper<1>( - sycl::range<1>(GROUP_RANGE_1D), - sycl::range<1>(LOCAL_RANGE_1D)); - - validator.collect_group_indicies(queue); - validator.validate_group_indicies(log); - } - { - // Adjust work-group size - bool reduce_wg_size = reduce_size<2>(queue); - size_t LOCAL_RANGE_1D = reduce_wg_size ? 1 : DEFAULT_LOCAL_RANGE[0]; - size_t LOCAL_RANGE_2D = reduce_wg_size ? 1 : DEFAULT_LOCAL_RANGE[1]; - auto validator = test_helper<2>( - sycl::range<2>(GROUP_RANGE_1D, GROUP_RANGE_2D), - sycl::range<2>(LOCAL_RANGE_1D, LOCAL_RANGE_2D)); - - validator.collect_group_indicies(queue); - validator.validate_group_indicies(log); - } - { - // Adjust work-group size - bool reduce_wg_size = reduce_size<3>(queue); - size_t LOCAL_RANGE_1D = reduce_wg_size ? 1 : DEFAULT_LOCAL_RANGE[0]; - size_t LOCAL_RANGE_2D = reduce_wg_size ? 1 : DEFAULT_LOCAL_RANGE[1]; - size_t LOCAL_RANGE_3D = reduce_wg_size ? 1 : DEFAULT_LOCAL_RANGE[2]; - auto validator = test_helper<3>( - sycl::range<3>(GROUP_RANGE_1D, GROUP_RANGE_2D, GROUP_RANGE_3D), - sycl::range<3>(LOCAL_RANGE_1D, LOCAL_RANGE_2D, LOCAL_RANGE_3D)); - - validator.collect_group_indicies(queue); - validator.validate_group_indicies(log); - } + // validate for dimensions 1, 2, and 3 + { + sycl::range<1> local_range(DEFAULT_LOCAL_RANGE[0]); + if (wg_size_too_large(queue, local_range)) { + WARN("cannot run with default local range, running with range [1]"); + local_range = sycl::range<1>(1); } + auto helper = test_helper(sycl::range<1>(GROUP_RANGE[0]), local_range); + helper.collect_results(queue); + helper.validate_results(); } -}; - -// construction of this proxy will register the above test -util::test_proxy proxy; - -} // namespace TEST_NAMESPACE + { + sycl::range<2> local_range(DEFAULT_LOCAL_RANGE[0], DEFAULT_LOCAL_RANGE[1]); + if (wg_size_too_large(queue, local_range)) { + WARN("cannot run with default local range, running with range [1, 1]"); + local_range = sycl::range<2>(1, 1); + } + auto helper = test_helper(sycl::range<2>(GROUP_RANGE[0], GROUP_RANGE[1]), + local_range); + helper.collect_results(queue); + helper.validate_results(); + } + { + sycl::range<3> local_range(DEFAULT_LOCAL_RANGE[0], DEFAULT_LOCAL_RANGE[1], + DEFAULT_LOCAL_RANGE[2]); + if (wg_size_too_large(queue, local_range)) { + WARN("cannot run with default local range, running with range [1, 1, 1]"); + local_range = sycl::range<3>(1, 1, 1); + } + auto helper = test_helper( + sycl::range<3>(GROUP_RANGE[0], GROUP_RANGE[1], GROUP_RANGE[2]), + local_range); + helper.collect_results(queue); + helper.validate_results(); + } +} diff --git a/tests/group/group_combined_mem_fence.cpp b/tests/group/group_combined_mem_fence.cpp deleted file mode 100644 index df465ee38..000000000 --- a/tests/group/group_combined_mem_fence.cpp +++ /dev/null @@ -1,106 +0,0 @@ -/******************************************************************************* -// -// SYCL 2020 Conformance Test Suite -// -// Provide coverage for global_and_local mem_fence -// -*******************************************************************************/ - -#include "../common/common.h" -#include "../group/group_mem_fence_common.h" - -#define TEST_NAME group_combined_mem_fence - -namespace group_combined_mem_fence__ { -using namespace sycl_cts; - -template -class combined_mem_fence_kernel_local; - -template -class combined_mem_fence_kernel_global; - -/** - * @brief Test mem_fence works for global_and_local fence space - * @param accessGroup Fence group to use for tests - * @param dim Dimension to use - * @param log Logger to use - * @param queue Queue to use - */ -template -void test_mem_fence(util::logger &log, sycl::queue &queue) { - const auto fenceSpace = sycl::access::fence_space::global_and_local; - const auto testName = test_name::get(fenceSpace); - - using localKernelT = combined_mem_fence_kernel_local; - using globalKernelT = combined_mem_fence_kernel_global; - - const auto fenceCallFactory = make_fence_call_factory( - [=](sycl::group item) { - item.mem_fence(fenceSpace); - }, - [=](sycl::group item) { - item.template mem_fence(fenceSpace); - }, - [=](sycl::group item) { - item.template mem_fence(fenceSpace); - }, - [=](sycl::group item) { - item.template mem_fence(fenceSpace); - }); - const auto access = std::integral_constant{}; - - const auto readFenceCall = fenceCallFactory.get_read(access); - const auto writeFenceCall = fenceCallFactory.get_write(access); - - // Verify global_and_local mem_fence works for local address space - { - const bool passed = test_rw_mem_fence_local_space( - log, queue, readFenceCall, writeFenceCall); - - if (!passed) { - FAIL(log, testName + "failed for local address space"); - } - } - // Verify global_and_local mem_fence works for global address space - { - const bool passed = test_rw_mem_fence_global_space( - log, queue, readFenceCall, writeFenceCall); - - if (!passed) { - FAIL(log, testName + "failed for global address space"); - } - } -} - -/** test sycl::group mem_fence functions -*/ -class TEST_NAME : public util::test_base { - public: - /** return information about this test - * @param info, test_base::info structure as output - */ - void get_info(test_base::info &out) const override { - set_test_info(out, TOSTRING(TEST_NAME), TEST_FILE); - } - - /** execute the test - * @param log, test transcript logging class - */ - void run(util::logger &log) override { - { - auto queue = util::get_cts_object::queue(); - - test_mem_fence(log, queue); - test_mem_fence(log, queue); - test_mem_fence(log, queue); - - queue.wait_and_throw(); - } - } -}; - -// construction of this proxy will register the above test -util::test_proxy proxy; - -} // namespace group_combined_mem_fence__ diff --git a/tests/group/group_default_mem_fence.cpp b/tests/group/group_default_mem_fence.cpp deleted file mode 100644 index 0e471876d..000000000 --- a/tests/group/group_default_mem_fence.cpp +++ /dev/null @@ -1,105 +0,0 @@ -/******************************************************************************* -// -// SYCL 2020 Conformance Test Suite -// -// Provide coverage for default mem_fence -// -*******************************************************************************/ - -#include "../common/common.h" -#include "../group/group_mem_fence_common.h" - -#define TEST_NAME group_default_mem_fence - -namespace group_default_mem_fence__ { -using namespace sycl_cts; - -template -class default_mem_fence_kernel_local; - -template -class default_mem_fence_kernel_global; - -/** - * @brief Test mem_fence works for default fence space - * @param accessGroup Fence group to use for tests - * @param dim Dimension to use - * @param log Logger to use - * @param queue Queue to use - */ -template -void test_mem_fence(util::logger &log, sycl::queue &queue) { - const auto testName = test_name::get(); - - using localKernelT = default_mem_fence_kernel_local; - using globalKernelT = default_mem_fence_kernel_global; - - const auto fenceCallFactory = make_fence_call_factory( - [=](sycl::group item) { - item.mem_fence(); - }, - [=](sycl::group item) { - item.template mem_fence(); - }, - [=](sycl::group item) { - item.template mem_fence(); - }, - [=](sycl::group item) { - item.template mem_fence(); - }); - const auto access = std::integral_constant{}; - - const auto readFenceCall = fenceCallFactory.get_read(access); - const auto writeFenceCall = fenceCallFactory.get_write(access); - - // Verify default mem_fence works for local address space - { - const bool passed = test_rw_mem_fence_local_space( - log, queue, readFenceCall, writeFenceCall); - - if (!passed) { - FAIL(log, testName + "failed for local address space"); - } - } - // Verify default mem_fence works for global address space - { - const bool passed = test_rw_mem_fence_global_space( - log, queue, readFenceCall, writeFenceCall); - - if (!passed) { - FAIL(log, testName + "failed for global address space"); - } - } -} - -/** test sycl::group mem_fence functions -*/ -class TEST_NAME : public util::test_base { - public: - /** return information about this test - * @param info, test_base::info structure as output - */ - void get_info(test_base::info &out) const override { - set_test_info(out, TOSTRING(TEST_NAME), TEST_FILE); - } - - /** execute the test - * @param log, test transcript logging class - */ - void run(util::logger &log) override { - { - auto queue = util::get_cts_object::queue(); - - test_mem_fence(log, queue); - test_mem_fence(log, queue); - test_mem_fence(log, queue); - - queue.wait_and_throw(); - } - } -}; - -// construction of this proxy will register the above test -util::test_proxy proxy; - -} // namespace group_default_mem_fence__ diff --git a/tests/group/group_global_mem_fence.cpp b/tests/group/group_global_mem_fence.cpp deleted file mode 100644 index 5caaba6c4..000000000 --- a/tests/group/group_global_mem_fence.cpp +++ /dev/null @@ -1,93 +0,0 @@ -/******************************************************************************* -// -// SYCL 2020 Conformance Test Suite -// -// Provide coverage for global mem_fence -// -*******************************************************************************/ - -#include "../common/common.h" -#include "../group/group_mem_fence_common.h" - -#define TEST_NAME group_global_mem_fence - -namespace group_global_mem_fence__ { -using namespace sycl_cts; - -template -class global_mem_fence_kernel_global; - -/** - * @brief Test mem_fence works for global fence space - * @param accessGroup Fence group to use for tests - * @param dim Dimension to use - * @param log Logger to use - * @param queue Queue to use - */ -template -void test_mem_fence(util::logger &log, sycl::queue &queue) { - const auto fenceSpace = sycl::access::fence_space::global_space; - const auto testName = test_name::get(fenceSpace); - - using globalKernelT = global_mem_fence_kernel_global; - - const auto fenceCallFactory = make_fence_call_factory( - [=](sycl::group item) { - item.mem_fence(fenceSpace); - }, - [=](sycl::group item) { - item.template mem_fence(fenceSpace); - }, - [=](sycl::group item) { - item.template mem_fence(fenceSpace); - }, - [=](sycl::group item) { - item.template mem_fence(fenceSpace); - }); - const auto access = std::integral_constant{}; - - const auto readFenceCall = fenceCallFactory.get_read(access); - const auto writeFenceCall = fenceCallFactory.get_write(access); - - // Verify global mem_fence works for global address space - { - const bool passed = test_rw_mem_fence_global_space( - log, queue, readFenceCall, writeFenceCall); - - if (!passed) { - FAIL(log, testName + "failed for global address space"); - } - } -} - -/** test sycl::group mem_fence functions -*/ -class TEST_NAME : public util::test_base { - public: - /** return information about this test - * @param info, test_base::info structure as output - */ - void get_info(test_base::info &out) const override { - set_test_info(out, TOSTRING(TEST_NAME), TEST_FILE); - } - - /** execute the test - * @param log, test transcript logging class - */ - void run(util::logger &log) override { - { - auto queue = util::get_cts_object::queue(); - - test_mem_fence(log, queue); - test_mem_fence(log, queue); - test_mem_fence(log, queue); - - queue.wait_and_throw(); - } - } -}; - -// construction of this proxy will register the above test -util::test_proxy proxy; - -} // namespace group_global_mem_fence__ diff --git a/tests/group/group_local_mem_fence.cpp b/tests/group/group_local_mem_fence.cpp deleted file mode 100644 index 2f14850bf..000000000 --- a/tests/group/group_local_mem_fence.cpp +++ /dev/null @@ -1,93 +0,0 @@ -/******************************************************************************* -// -// SYCL 2020 Conformance Test Suite -// -// Provide coverage for local mem_fence -// -*******************************************************************************/ - -#include "../common/common.h" -#include "../group/group_mem_fence_common.h" - -#define TEST_NAME group_local_mem_fence - -namespace group_local_mem_fence__ { -using namespace sycl_cts; - -template -class local_mem_fence_kernel_local; - -/** - * @brief Test mem_fence works for local fence space - * @param accessGroup Fence group to use for tests - * @param dim Dimension to use - * @param log Logger to use - * @param queue Queue to use - */ -template -void test_mem_fence(util::logger &log, sycl::queue &queue) { - const auto fenceSpace = sycl::access::fence_space::local_space; - const auto testName = test_name::get(fenceSpace); - - using localKernelT = local_mem_fence_kernel_local; - - const auto fenceCallFactory = make_fence_call_factory( - [=](sycl::group item) { - item.mem_fence(fenceSpace); - }, - [=](sycl::group item) { - item.template mem_fence(fenceSpace); - }, - [=](sycl::group item) { - item.template mem_fence(fenceSpace); - }, - [=](sycl::group item) { - item.template mem_fence(fenceSpace); - }); - const auto access = std::integral_constant{}; - - const auto readFenceCall = fenceCallFactory.get_read(access); - const auto writeFenceCall = fenceCallFactory.get_write(access); - - // Verify local mem_fence works for local address space - { - const bool passed = test_rw_mem_fence_local_space( - log, queue, readFenceCall, writeFenceCall); - - if (!passed) { - FAIL(log, testName + "failed for local address space"); - } - } -} - -/** test sycl::group mem_fence functions -*/ -class TEST_NAME : public util::test_base { - public: - /** return information about this test - * @param info, test_base::info structure as output - */ - void get_info(test_base::info &out) const override { - set_test_info(out, TOSTRING(TEST_NAME), TEST_FILE); - } - - /** execute the test - * @param log, test transcript logging class - */ - void run(util::logger &log) override { - { - auto queue = util::get_cts_object::queue(); - - test_mem_fence(log, queue); - test_mem_fence(log, queue); - test_mem_fence(log, queue); - - queue.wait_and_throw(); - } - } -}; - -// construction of this proxy will register the above test -util::test_proxy proxy; - -} // namespace group_local_mem_fence__ diff --git a/tests/group/group_mem_fence_common.h b/tests/group/group_mem_fence_common.h deleted file mode 100644 index 467ed4118..000000000 --- a/tests/group/group_mem_fence_common.h +++ /dev/null @@ -1,371 +0,0 @@ -/******************************************************************************* -// -// SYCL 2020 Conformance Test Suite -// -// Provides common methods for group mem_fence tests -// -*******************************************************************************/ - -#ifndef __SYCLCTS_TESTS_GROUP_GROUP_MEM_FENCE_COMMON_H -#define __SYCLCTS_TESTS_GROUP_GROUP_MEM_FENCE_COMMON_H - -#include "../common/common.h" -#include -#include -#include - -namespace { - -/** - * @brief Test memory fence works for global address space - * @tparam kernelT Kernel to run onto - * @tparam dim Group dimension to use; TBD - * @tparam readFenceCallT Type of read fence call. Deduced. - * @tparam writeFenceCallT Type of write fence call. Deduced. - * @param log Logger to use - * @param queue Queue to use - * @param readMemFence Lambda with the read (load) fence call within - * @param writeMemFence Lambda with the write (store) fence call within - */ -template -bool test_rw_mem_fence_global_space(sycl_cts::util::logger& log, - sycl::queue &queue, - const readFenceCallT& readMemFence, - const writeFenceCallT& writeMemFence) -{ - static_assert(dim == 1, - "Multidimensional groups are not supported currently"); - // Set workspace size and cut-off read iterations limit - const size_t globalSize = 64; - const size_t localSize = 2; - - // Check work-group size limits; skip test in case it cannot be run - if (!device_supports_wg_size(log, queue, localSize) || - !kernel_supports_wg_size(log, queue, localSize)) - return true; - - bool passed = true; - - // Init ranges - sycl::range<1> globalRange(globalSize); - sycl::range<1> workGroupRange(globalSize / localSize); - sycl::range<1> localRange(localSize); - - // Run kernel to verify memory ordering works for adjacent work-items - { - sycl::buffer data(globalRange); - sycl::buffer passedBuf(&passed, sycl::range<1>(1)); - - // Initialize data state - { - auto ptr = data.get_access(); - for (size_t i = 0; i < ptr.size(); ++i) - ptr[i] = -1; - } - queue.submit([&](sycl::handler &cgh) { - auto ptr = data.get_access(cgh); - auto pass = passedBuf.get_access(cgh); - - cgh.parallel_for_work_group( - workGroupRange, localRange, [=](sycl::group<1> group) { - group.parallel_for_work_item( - [&](sycl::h_item<1> item) { - const size_t current = item.get_global().get_linear_id(); - const size_t other = current ^ 1U; - - const int nWrites = 100; - const int nReads = 100000; - const int writesPerIteration = - static_cast(2U * (current + 1U)); - const int readsPerIteration = - nReads / nWrites * writesPerIteration; - - int previousValue = ptr[other].load(); - int currentValue = 0; - - for (int i = 0; i < nWrites; i+= writesPerIteration) { - // Run sequence of writes - for (int j = 0; j < writesPerIteration; ++j) { - ptr[current].store(i + j); - writeMemFence(group); - } - // Run sequence of reads - for (int k = 0; k < readsPerIteration; ++k) { - // Memory fence on read; should be either read_write or read - readMemFence(group); - currentValue = ptr[other].load(); - - // Verify memory order from other work-item - if (currentValue < previousValue) { - pass[0] = false; - } - previousValue = currentValue; - } - } - }); - }); - }); - } - - return passed; -} - -/** - * @brief Test memory fence works for local address space - * @tparam kernelT Kernel to run onto - * @tparam dim Group dimension to use - * @tparam readFenceCallT Type of read fence call. Deduced. - * @tparam writeFenceCallT Type of write fence call. Deduced. - * @param log Logger to use - * @param queue Queue to use - * @param readMemFence Lambda with the read (load) fence call within - * @param writeMemFence Lambda with the write (store) fence call within - */ -template -bool test_rw_mem_fence_local_space(sycl_cts::util::logger& log, - sycl::queue &queue, - const readFenceCallT& readMemFence, - const writeFenceCallT& writeMemFence) -{ - static_assert(dim == 1, - "Multidimensional groups are not supported currently"); - // Set workspace size and cut-off read iterations limit - const size_t globalSize = 64; - const size_t localSize = 2; - - // Check work-group size limits; skip test in case it cannot be run - if (!device_supports_wg_size(log, queue, localSize) || - !kernel_supports_wg_size(log, queue, localSize)) - return true; - - bool passed = true; - - // Init ranges - sycl::range<1> globalRange(globalSize); - sycl::range<1> workGroupRange(globalSize / localSize); - sycl::range<1> localRange(localSize); - - // Run kernel to verify memory ordering works for adjacent work-items - { - sycl::buffer passedBuf(&passed, sycl::range<1>(1)); - - queue.submit([&](sycl::handler &cgh) { - auto pass = passedBuf.get_access(cgh); - sycl::accessor - ptr(globalRange, cgh); - - cgh.parallel_for_work_group( - workGroupRange, localRange, [=](sycl::group<1> group) { - // Initialize data state - group.parallel_for_work_item( - [&](sycl::h_item<1> item) { - const size_t idx = item.get_global().get_linear_id(); - - ptr[idx].store(-1); - }); - // Implicit barrier - group.parallel_for_work_item( - [&](sycl::h_item<1> item) { - const size_t current = item.get_global().get_linear_id(); - const size_t other = current ^ 1U; - - const int nWrites = 100; - const int nReads = 100000; - const int writesPerIteration = - static_cast(2U * (current + 1U)); - const int readsPerIteration = - nReads / nWrites * writesPerIteration; - - int previousValue = ptr[other].load(); - int currentValue = 0; - - for (int i = 0; i < nWrites; i+= writesPerIteration) { - // Run sequence of writes - for (int j = 0; j < writesPerIteration; ++j) { - ptr[current].store(i + j); - writeMemFence(group); - } - // Run sequence of reads - for (int k = 0; k < readsPerIteration; ++k) { - // Memory fence on read; should be either read_write or read - readMemFence(group); - currentValue = ptr[other].load(); - - // Verify memory order from other work-item - if (currentValue < previousValue) { - pass[0] = false; - } - previousValue = currentValue; - } - } - }); - }); - }); - } - - return passed; -} - -/** Memory fence access group, by access_mode used for read and write access - */ -enum class access_group: int { - useDefault = 0, - useCombined, - useSeparate -}; - -/** - * @brief Test name factory, to use for logs - * @param accessGroup Fence access group to use - * @param dim Group dimension - */ -template -class test_name -{ -public: - /** - * @brief Retrieve test name for explicit fence space usage - * @param fenceSpace Fence space value - */ - static std::string get(sycl::access::fence_space fenceSpace) - { - switch (fenceSpace) { - case sycl::access::fence_space::global_and_local: - return "global_and_local space " + mem_fence_name(); - case sycl::access::fence_space::local_space: - return "local space " + mem_fence_name(); - case sycl::access::fence_space::global_space: - return "global space " + mem_fence_name(); - default: - return "__unknown__"; - }; - } - - /** - * @brief Retrieve test name for default fence space usage - */ - static std::string get() - { - return "default space " + mem_fence_name(); - } -private: - static std::string mem_fence_name() { - const auto dimensions = std::to_string(dim); - switch (accessGroup) { - case access_group::useDefault: - return "default memory fence(" + dimensions + ")"; - case access_group::useCombined: - return "read_write memory fence(" + dimensions + ")"; - case access_group::useSeparate: - return "read and write memory fences(" + dimensions + ")"; - default: - return "__unknown__"; - }; - } -}; - -/** - * @brief Provides access to stored lambdas using the access_group mode - * as selector. - * @param defaultFenceCall Lambda for mem_fence call without access specified - * @param readWriteFenceCall Lambda for mem_fence call with read_write access - * @param readFenceCall Lambda for mem_fence call with read access - * @param writeFenceCall Lambda for mem_fence call with write access - */ -template -class fence_call_factory -{ - defaultFenceCallT defaultFenceCall; - readWriteFenceCallT readWriteFenceCall; - readFenceCallT readFenceCall; - writeFenceCallT writeFenceCall; -public: - fence_call_factory(defaultFenceCallT defaultCall, - readWriteFenceCallT readWriteCall, - readFenceCallT readCall, - writeFenceCallT writeCall): - defaultFenceCall(defaultCall), - readWriteFenceCall(readWriteCall), - readFenceCall(readCall), - writeFenceCall(writeCall) { - } - - using defaultAccessT = - std::integral_constant; - using combinedAccessT = - std::integral_constant; - using separateAccessT = - std::integral_constant; - - /** Retrieve read fence call for 'default' access group - */ - const defaultFenceCallT& get_read(defaultAccessT modeSelector) const { - static_cast(modeSelector); - return defaultFenceCall; - } - /** Retrieve read fence call for 'combined' access group - */ - const readWriteFenceCallT& get_read(combinedAccessT modeSelector) const { - static_cast(modeSelector); - return readWriteFenceCall; - } - /** Retrieve read fence call for 'separate' access group - */ - const readFenceCallT& get_read(separateAccessT modeSelector) const { - static_cast(modeSelector); - return readFenceCall; - } - /** Retrieve write fence call for 'default' access group - */ - const defaultFenceCallT& get_write(defaultAccessT modeSelector) const { - static_cast(modeSelector); - return defaultFenceCall; - } - /** Retrieve write fence call for 'combined' access group - */ - const readWriteFenceCallT& get_write(combinedAccessT modeSelector) const { - static_cast(modeSelector); - return readWriteFenceCall; - } - /** Retrieve write fence call for 'separate' access group - */ - const writeFenceCallT& get_write(separateAccessT modeSelector) const { - static_cast(modeSelector); - return writeFenceCall; - } -}; - -/** - * @brief Object generator to store memory fence lambdas for different usage - * modes. Hides actual lambda type. - * @param defaultFenceCall Lambda for mem_fence call without access specified - * @param readWriteFenceCall Lambda for mem_fence call with read_write access - * @param readFenceCall Lambda for mem_fence call with read access - * @param writeFenceCall Lambda for mem_fence call with write access - */ -template -fence_call_factory - make_fence_call_factory(defaultFenceCallT defaultFenceCall, - readWriteFenceCallT readWriteFenceCall, - readFenceCallT readFenceCall, - writeFenceCallT writeFenceCall) { - return fence_call_factory( - defaultFenceCall, readWriteFenceCall, readFenceCall, writeFenceCall); -} - -} //namespace - -#endif // __SYCLCTS_TESTS_GROUP_GROUP_MEM_FENCE_COMMON_H