Cover gaps for explicit memory operation copy(acc, acc)

tests/handler/handler_copy_common.h

@@ -26,6 +26,7 @@
 #include <regex>
 #include <sstream>
 
+#include "../../util/sycl_exceptions.h"
 #include "../common/common.h"
 
 namespace handler_copy_common {
@@ -296,6 +297,27 @@ using range_helper = range_id_helper<sycl::range, dims, 1>;
 template <int dims>
 using id_helper = range_id_helper<sycl::id, dims, 0>;
 
+template <int dim>
+sycl::range<3> default_large_range() {
+  return range_helper<3>::cast(range_helper<dim>::make(5, 7, 9));
+}
+
+template <int dim_large, int dim_small, bool transposed_copy = false>
+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<dim_large>::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<dim_large>();
+    auto smallBufRange =
+        transform_large_range_into_small<dim_large, dim_small, transposed_copy>(
+            largeBufRange);
+
     assert(smallBufRange.size() == largeBufRange.size());
 
     auto largeCopyRange = largeBufRange;
@@ -818,6 +831,28 @@ 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
+      copy_test_context<dataT, dim_src, dim_dst, strided, transposed> ctx(
+          queue);
+      ctx.verify_d2d_copy(
+          [&](sycl::handler& cgh) {
+            auto r =
+                ctx.getSrcBuf()
+                    .template get_access<mode_src, target_t::constant_buffer>(
+                        cgh, ctx.getSrcCopyRange(), ctx.getSrcCopyOffset());
+            auto w = ctx.getDstBuf().template get_access<mode_dst, target>(
+                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>)"));
+    };
+  }
+#endif
   {
     // Check fill(accessor, dataT)
     const auto pattern = type_helper<dataT>::make(117);
@@ -934,5 +969,114 @@ static void test_all_variants(log_helper lh, sycl::queue& queue) {
   test_all_dimensions<dataT, true, true>(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
+ * than source accessor range
+ */
+template <typename DataT, typename DimSrcT, typename DimDstT, typename ModeSrcT,
+          typename ModeDstT>
+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<dim_src> src_copy_range = range_helper<dim_src>::make(0, 0, 0);
+  sycl::range<dim_dst> dst_copy_range = range_helper<dim_dst>::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<dim_large>();
+    auto small_range =
+        transform_large_range_into_small<dim_large, dim_small>(large_range);
+
+    if (dim_src > dim_dst) {
+      src_copy_range = range_helper<dim_src>::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<dim_dst>::cast(small_range - sycl::range<3>(1, 1, 1));
+    } else {
+      src_copy_range = range_helper<dim_src>::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<dim_dst>::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<DataT> src_buf_mem(new DataT[src_copy_range.size()],
+                                       std::default_delete<DataT[]>());
+    std::shared_ptr<DataT> dst_buf_mem(new DataT[dst_copy_range.size()],
+                                       std::default_delete<DataT[]>());
+
+    sycl::buffer<DataT, dim_src> src_buf(src_buf_mem, src_copy_range);
+    sycl::buffer<DataT, dim_dst> 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<mode_src, target_t::device>(cgh);
+          auto dst_acc =
+              dst_buf.template get_access<mode_dst, target_t::device>(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 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) {
+          auto src_acc =
+              src_buf.template get_access<mode_src, target_t::constant_buffer>(
+                  cgh);
+          auto dst_acc =
+              dst_buf.template get_access<mode_dst, target_t::device>(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));
+    }
+#endif  // SYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS
+  }
+};
+#endif  // SYCL_CTS_COMPILING_WITH_COMPUTECPP
 }  // namespace handler_copy_common
 #endif  // __SYCLCTS_TESTS_HANDLER_COPY_COMMON_H

tests/handler/handler_copy_core.cpp

@@ -10,6 +10,12 @@
 
 #include "catch2/catch_test_macros.hpp"
 
+#include "../common/type_coverage.h"
+
+#include "../common/string_makers.h"
+
+#include "../common/disabled_for_test_case.h"
+
 namespace handler_copy_core {
 using namespace handler_copy_common;
 
@@ -34,4 +40,44 @@ TEST_CASE("Tests the API for sycl::handler::copy", "[handler]") {
 #endif
 }
 
+// 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 =
+      named_type_pack<int
+#if SYCL_CTS_ENABLE_FULL_CONFORMANCE
+                      ,
+                      char, short, long, float, sycl::char2, sycl::short3,
+                      sycl::int4, sycl::long8, sycl::float8
+#endif
+                      >::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<int, 1, 2, 3>::generate_named(
+      "one dim range", "two dim range", "three dim range");
+
+  const auto src_modes =
+      value_pack<sycl::access_mode, sycl::access_mode::read,
+                 sycl::access_mode::read_write>::generate_named();
+
+  const auto dst_modes =
+      value_pack<sycl::access_mode, sycl::access_mode::write,
+                 sycl::access_mode::read_write,
+                 sycl::access_mode::discard_write,
+                 sycl::access_mode::discard_read_write>::generate_named();
+
+  for_all_combinations<CheckCopyAccToAccException>(types, dims, dims, src_modes,
+                                                   dst_modes, queue);
+});
+
 }  // namespace handler_copy_core