From f0b2f8a4c3108ca2f5105360d5f68a07d5f60c59 Mon Sep 17 00:00:00 2001 From: vladimirkhashev Date: Thu, 25 May 2023 15:23:16 -0700 Subject: [PATCH 1/5] Covered thr gaps for copy() tests Added checking of copy(acc, acc) operation with target constant_buffer for source accessor Added checking of copy(acc, acc) exception in case of dst_range < src_range --- tests/handler/handler_copy_common.h | 172 +++++++++++++++++++++++++--- tests/handler/handler_copy_core.cpp | 43 +++++++ 2 files changed, 198 insertions(+), 17 deletions(-) diff --git a/tests/handler/handler_copy_common.h b/tests/handler/handler_copy_common.h index a57e777f7..7801d2a7d 100644 --- a/tests/handler/handler_copy_common.h +++ b/tests/handler/handler_copy_common.h @@ -26,6 +26,7 @@ #include #include +#include "../../util/sycl_exceptions.h" #include "../common/common.h" namespace handler_copy_common { @@ -296,6 +297,27 @@ using range_helper = range_id_helper; template using id_helper = range_id_helper; +template +sycl::range<3> default_large_range() { + return range_helper<3>::cast(range_helper::make(5, 7, 9)); +} + +template +sycl::range<3> transform_large_range_into_small(sycl::range<3> largeBufRange) { + sycl::range<3> smallBufRange = sycl::range<3>(1, 1, 1); + + // Condense large range into small range so that both + // have the same size (= same number of items). + for (int d = 0; d < dim_large; ++d) { + if (transposed_copy) { + smallBufRange[std::min(d, dim_small - 1)] *= + largeBufRange[dim_large - d - 1]; + } else { + smallBufRange[std::min(d, dim_small - 1)] *= largeBufRange[d]; + } + } + return smallBufRange; +} /** * @brief The copy_test_context encapsulates all host and device data required * for testing, and provides utility functions for verifying the result @@ -622,23 +644,14 @@ class copy_test_context { * transposed_copy is set, in which case the destination will be transposed. */ void setup_ranges() { - constexpr auto dim_large = dim_src > dim_dst ? dim_src : dim_dst; - constexpr auto dim_small = dim_src <= dim_dst ? dim_src : dim_dst; - - auto largeBufRange = - range_helper<3>::cast(range_helper::make(5, 7, 9)); - auto smallBufRange = sycl::range<3>(1, 1, 1); - - // Condense large range into small range so that both - // have the same size (= same number of items). - for (int d = 0; d < dim_large; ++d) { - if (transposed_copy) { - smallBufRange[std::min(d, dim_small - 1)] *= - largeBufRange[dim_large - d - 1]; - } else { - smallBufRange[std::min(d, dim_small - 1)] *= largeBufRange[d]; - } - } + constexpr auto dim_large = std::max(dim_src, dim_dst); + constexpr auto dim_small = std::min(dim_src, dim_dst); + + auto largeBufRange = default_large_range(); + auto smallBufRange = + transform_large_range_into_small( + largeBufRange); + assert(smallBufRange.size() == largeBufRange.size()); auto largeCopyRange = largeBufRange; @@ -818,6 +831,26 @@ static void test_write_acc_copy_functions(log_helper lh, "copy(accessor<$dataT, $dim_src, $mode_src, $target>, " "accessor<$dataT, $dim_dst, $mode_dst, $target>)")); } + { + if constexpr (mode_src == mode_t::read) { + // Check copy(accessor, accessor) with constant_buffer target + copy_test_context ctx( + queue); + ctx.verify_d2d_copy( + [&](sycl::handler& cgh) { + auto r = + ctx.getSrcBuf() + .template get_access( + cgh, ctx.getSrcCopyRange(), ctx.getSrcCopyOffset()); + auto w = ctx.getDstBuf().template get_access( + cgh, ctx.getDstCopyRange(), ctx.getDstCopyOffset()); + cgh.copy(r, w); + }, + lh.set_line(__LINE__).set_op( + "copy(accessor<$dataT, $dim_src, $mode_src, constant_buffer>, " + "accessor<$dataT, $dim_dst, $mode_dst, $target>)")); + }; + } { // Check fill(accessor, dataT) const auto pattern = type_helper::make(117); @@ -934,5 +967,110 @@ static void test_all_variants(log_helper lh, sycl::queue& queue) { test_all_dimensions(lh, queue); } +/** + * @brief Class provides a test that checks if exception is thrown on explicit + * memory operation copy(acc, acc) in case of destination accessor range less + * than source accessor range + */ +template +class CheckCopyAccToAccException { + static constexpr int dim_src = DimSrcT::value; + static constexpr int dim_dst = DimDstT::value; + static constexpr mode_t mode_src = ModeSrcT::value; + static constexpr mode_t mode_dst = ModeDstT::value; + + sycl::range src_copy_range = range_helper::make(0, 0, 0); + sycl::range dst_copy_range = range_helper::make(0, 0, 0); + + void make_ranges() { + constexpr auto dim_large = std::max(dim_src, dim_dst); + constexpr auto dim_small = std::min(dim_src, dim_dst); + + auto large_range = default_large_range(); + auto small_range = + transform_large_range_into_small(large_range); + + if (dim_src > dim_dst) { + src_copy_range = range_helper::cast(large_range); + // Creating destination range less than source range to force exception on + // explicit memory operation copy(acc, acc) + dst_copy_range = + range_helper::cast(small_range - sycl::range<3>(1, 1, 1)); + } else { + src_copy_range = range_helper::cast(small_range); + // Creating destination range less than source range to force exception on + // explicit memory operation copy(acc, acc) + dst_copy_range = + range_helper::cast(large_range - sycl::range<3>(1, 1, 1)); + } + } + + std::string description(const std::string& type_name, + const std::string& mode_src_name, + const std::string& mode_dst_name, + std::string&& src_target_name) { + std::stringstream ss; + ss << "Check that exception with error code \"errc::invalid\" is thrown on " + "explicit memory operation copy(src_acc, dst_acc) in case of dst_acc " + "with incorrect range size (T: " + << type_name << " dim src: " << dim_src << "dim dist: " << dim_dst + << " acc mode src: " << mode_src_name + << "acc mode dst: " << mode_dst_name + << " source target: " << src_target_name << ")"; + return ss.str(); + } + + public: + CheckCopyAccToAccException() { make_ranges(); } + void operator()(sycl::queue& q, const std::string& type_name, + const std::string&, const std::string&, + const std::string& mode_src_name, + const std::string& mode_dst_name) { + std::shared_ptr src_buf_mem(new DataT[src_copy_range.size()], + std::default_delete()); + std::shared_ptr dst_buf_mem(new DataT[dst_copy_range.size()], + std::default_delete()); + + sycl::buffer src_buf(src_buf_mem, src_copy_range); + sycl::buffer dst_buf(dst_buf_mem, dst_copy_range); + + { + auto check_exception_with_invalid_dst_range = [&] { + q.submit([&](sycl::handler& cgh) { + auto src_acc = + src_buf.template get_access(cgh); + auto dst_acc = + dst_buf.template get_access(cgh); + cgh.copy(src_acc, dst_acc); + }); + }; + INFO(description(type_name, mode_src_name, mode_dst_name, "device")); + CHECK_THROWS_MATCHES( + check_exception_with_invalid_dst_range(), sycl::exception, + sycl_cts::util::equals_exception(sycl::errc::invalid)); + } + + if constexpr (mode_src == mode_t::read) { + auto check_exception_with_invalid_dst_range_constant_buffer = [&] { + q.submit([&](sycl::handler& cgh) { + auto src_acc = + src_buf.template get_access( + cgh); + auto dst_acc = + dst_buf.template get_access(cgh); + cgh.copy(src_acc, dst_acc); + }); + }; + INFO(description(type_name, mode_src_name, mode_dst_name, + "constant_buffer")); + CHECK_THROWS_MATCHES( + check_exception_with_invalid_dst_range_constant_buffer(), + sycl::exception, + sycl_cts::util::equals_exception(sycl::errc::invalid)); + } + } +}; + } // namespace handler_copy_common #endif // __SYCLCTS_TESTS_HANDLER_COPY_COMMON_H diff --git a/tests/handler/handler_copy_core.cpp b/tests/handler/handler_copy_core.cpp index 18290b9fc..ba6f83849 100644 --- a/tests/handler/handler_copy_core.cpp +++ b/tests/handler/handler_copy_core.cpp @@ -10,6 +10,10 @@ #include "catch2/catch_test_macros.hpp" +#include "../common/type_coverage.h" + +#include "../common/string_makers.h" + namespace handler_copy_core { using namespace handler_copy_common; @@ -34,4 +38,43 @@ TEST_CASE("Tests the API for sycl::handler::copy", "[handler]") { #endif } +TEST_CASE( + "Check exception on copy(accessor, accessor) in case of invalid " + "destination accessor size", + "[handler]") { + auto queue = util::get_cts_object::queue(); + + const auto types = + named_type_pack::generate("int" +#if SYCL_CTS_ENABLE_FULL_CONFORMANCE + , + "char", "short", "long", "float", + "sycl::char2", "sycl::short3", "sycl::int4", + "sycl::long8", "sycl::float8" +#endif + ); + + const auto dims = value_pack::generate_named( + "one dim range", "two dim range", "three dim range"); + + const auto src_modes = + value_pack::generate_named(); + + const auto dst_modes = + value_pack::generate_named(); + + for_all_combinations(types, dims, dims, src_modes, + dst_modes, queue); +} + } // namespace handler_copy_core From 7d8c0b417626a29da1047c8a9a6664f6f2c76b49 Mon Sep 17 00:00:00 2001 From: vladimirkhashev Date: Fri, 26 May 2023 00:39:04 -0700 Subject: [PATCH 2/5] Applying comments --- tests/handler/handler_copy_common.h | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tests/handler/handler_copy_common.h b/tests/handler/handler_copy_common.h index 7801d2a7d..7e5df01b3 100644 --- a/tests/handler/handler_copy_common.h +++ b/tests/handler/handler_copy_common.h @@ -831,6 +831,7 @@ static void test_write_acc_copy_functions(log_helper lh, "copy(accessor<$dataT, $dim_src, $mode_src, $target>, " "accessor<$dataT, $dim_dst, $mode_dst, $target>)")); } +#if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS { if constexpr (mode_src == mode_t::read) { // Check copy(accessor, accessor) with constant_buffer target @@ -851,6 +852,7 @@ static void test_write_acc_copy_functions(log_helper lh, "accessor<$dataT, $dim_dst, $mode_dst, $target>)")); }; } +#endif { // Check fill(accessor, dataT) const auto pattern = type_helper::make(117); @@ -1051,6 +1053,7 @@ class CheckCopyAccToAccException { sycl_cts::util::equals_exception(sycl::errc::invalid)); } +#if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS if constexpr (mode_src == mode_t::read) { auto check_exception_with_invalid_dst_range_constant_buffer = [&] { q.submit([&](sycl::handler& cgh) { @@ -1069,6 +1072,7 @@ class CheckCopyAccToAccException { sycl::exception, sycl_cts::util::equals_exception(sycl::errc::invalid)); } +#endif } }; From 748bf430d40234b41871e5e10244db847e752d99 Mon Sep 17 00:00:00 2001 From: vladimirkhashev Date: Wed, 28 Jun 2023 02:56:27 -0700 Subject: [PATCH 3/5] Fix format --- tests/handler/handler_copy_common.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/handler/handler_copy_common.h b/tests/handler/handler_copy_common.h index 7e5df01b3..82ce17120 100644 --- a/tests/handler/handler_copy_common.h +++ b/tests/handler/handler_copy_common.h @@ -831,7 +831,7 @@ static void test_write_acc_copy_functions(log_helper lh, "copy(accessor<$dataT, $dim_src, $mode_src, $target>, " "accessor<$dataT, $dim_dst, $mode_dst, $target>)")); } -#if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS +#if SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS { if constexpr (mode_src == mode_t::read) { // Check copy(accessor, accessor) with constant_buffer target From fc2673c98cc90207cf46fb84461a1e886162b7f9 Mon Sep 17 00:00:00 2001 From: vladimirkhashev Date: Wed, 28 Jun 2023 07:18:58 -0700 Subject: [PATCH 4/5] Fix CI --- tests/handler/handler_copy_core.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/tests/handler/handler_copy_core.cpp b/tests/handler/handler_copy_core.cpp index ba6f83849..9b350d277 100644 --- a/tests/handler/handler_copy_core.cpp +++ b/tests/handler/handler_copy_core.cpp @@ -14,6 +14,8 @@ #include "../common/string_makers.h" +#include "../common/disabled_for_test_case.h" + namespace handler_copy_core { using namespace handler_copy_common; @@ -38,10 +40,11 @@ TEST_CASE("Tests the API for sycl::handler::copy", "[handler]") { #endif } -TEST_CASE( - "Check exception on copy(accessor, accessor) in case of invalid " - "destination accessor size", - "[handler]") { +// FIXME: re-enable when sycl::errc is implemented in computecpp +DISABLED_FOR_TEST_CASE(ComputeCpp) +("Check exception on copy(accessor, accessor) in case of invalid " + "destination accessor size", + "[handler]")({ auto queue = util::get_cts_object::queue(); const auto types = @@ -75,6 +78,6 @@ TEST_CASE( for_all_combinations(types, dims, dims, src_modes, dst_modes, queue); -} +}); } // namespace handler_copy_core From 98aac76447b47b8463788a31baa3d2c6e4382b37 Mon Sep 17 00:00:00 2001 From: vladimirkhashev Date: Thu, 29 Jun 2023 04:33:24 -0700 Subject: [PATCH 5/5] Fix CI --- tests/handler/handler_copy_common.h | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/tests/handler/handler_copy_common.h b/tests/handler/handler_copy_common.h index 82ce17120..ca735aa95 100644 --- a/tests/handler/handler_copy_common.h +++ b/tests/handler/handler_copy_common.h @@ -969,6 +969,8 @@ static void test_all_variants(log_helper lh, sycl::queue& queue) { test_all_dimensions(lh, queue); } +// FIXME: re-enable when sycl::errc is implemented in computecpp +#ifndef SYCL_CTS_COMPILING_WITH_COMPUTECPP /** * @brief Class provides a test that checks if exception is thrown on explicit * memory operation copy(acc, acc) in case of destination accessor range less @@ -1072,9 +1074,9 @@ class CheckCopyAccToAccException { sycl::exception, sycl_cts::util::equals_exception(sycl::errc::invalid)); } -#endif +#endif // SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS } }; - +#endif // SYCL_CTS_COMPILING_WITH_COMPUTECPP } // namespace handler_copy_common #endif // __SYCLCTS_TESTS_HANDLER_COPY_COMMON_H