[ESIMD] Add more user-friendly error msg when write beyond simd obj (#…

…12059)

It is still not prohibited to create 'simd_view' that views beyond the
bounds of viewed 'simd' object. And it is not yet prohibited to read
using such 'simd_view' object, but the commit
17f8939 prohibited writing beyong simd
bounds. The error message was difficult to read and did not give any
guidance on the proposed fix.

The old error message only said that __esimd_wrregion call could not be
matched. The newly added static assert says that the viewed object in
LHS does not fit the value in RHS.

Example of problem situation:
  simd<float, 16> vec; // vec is only 16-elements
  auto vec_too_bit_view = vec.select<128, 1>();// 128-elems - too many
  vec_too_bit_view = simd<float, 128>{0}; //error: write 128-elem to
16-elem storage

Signed-off-by: Klochkov, Vyacheslav N <vyacheslav.n.klochkov@intel.com>

sycl/include/sycl/ext/intel/esimd/detail/intrin.hpp

@@ -113,8 +113,7 @@ __esimd_rdindirect(__ESIMD_DNS::vector_type_t<T, N> Input,
 // for (int i = 0; i < NumRows; ++i) {
 //   for (int j = 0; j < Width; ++j) {
 //       if (Mask[Index])
-//           Result[i * VStride +  j * Stride + EltOffset] =
-//           NewVal[Index];
+//           Result[i * VStride +  j * Stride + EltOffset] = NewVal[Index];
 //       ++Index;
 //   }
 // }
@@ -143,9 +142,9 @@ template <class T> using __st = __raw_t<T>;
 
 /// read from a basic region of a vector, return a vector
 template <typename BT, int BN, typename RTy>
-__ESIMD_DNS::vector_type_t<__st<typename RTy::element_type>, RTy::length>
-    ESIMD_INLINE readRegion(
-        const __ESIMD_DNS::vector_type_t<__st<BT>, BN> &Base, RTy Region) {
+__ESIMD_DNS::vector_type_t<__st<typename RTy::element_type>,
+                           RTy::length> ESIMD_INLINE
+readRegion(const __ESIMD_DNS::vector_type_t<__st<BT>, BN> &Base, RTy Region) {
   using ElemTy = __st<typename RTy::element_type>;
   auto Base1 = bitcast<ElemTy, __st<BT>, BN>(Base);
   constexpr int Bytes = BN * sizeof(BT);

sycl/include/sycl/ext/intel/esimd/detail/simd_obj_impl.hpp

@@ -650,6 +650,13 @@ class [[__sycl_detail__::__uses_aspects__(
       constexpr int M = RTy::Size_x;
       constexpr int Stride = RTy::Stride_x;
       uint16_t Offset = Region.M_offset_x * sizeof(ElemTy);
+      static_assert(M > 0, "Malformed RHS region.");
+      static_assert(M <= BN, "Attempt to write beyond viewed area: The viewed "
+                             "object in LHS does not fit RHS.");
+      // (M > BN) condition is added below to not duplicate the above assert
+      // for big values of M. The assert below is for 'Stride'.
+      static_assert((M > BN) || (M - 1) * Stride < BN,
+                    "Malformed RHS region - too big stride.");
 
       // Merge and update.
       auto Merged = __esimd_wrregion<ElemTy, BN, M,
@@ -685,6 +692,11 @@ class [[__sycl_detail__::__uses_aspects__(
         constexpr int Stride = TR::Stride_x;
         uint16_t Offset = Region.first.M_offset_x * sizeof(ElemTy);
 
+        static_assert(M <= BN1, "Attempt to write beyond viewed area: The "
+                                "viewed object in LHS does not fit RHS.");
+        static_assert(M > 0, "Malformed RHS region.");
+        static_assert((M - 1) * Stride < BN,
+                      "Malformed RHS region - too big stride.");
         // Merge and update.
         Base1 = __esimd_wrregion<ElemTy, BN1, M,
                                  /*VS*/ 0, M, Stride>(Base1, Val, Offset);
@@ -702,6 +714,12 @@ class [[__sycl_detail__::__uses_aspects__(
             (Region.first.M_offset_y * PaTy::Size_x + Region.first.M_offset_x) *
             sizeof(ElemTy));
 
+        static_assert(M <= BN1, "Attempt to write beyond viewed area: The "
+                                "viewed object in LHS does not fit RHS.");
+        static_assert(M > 0 && W > 0 && M % W == 0, "Malformed RHS region.");
+        static_assert(W == 0 || ((M / W) - 1) * VS + (W - 1) * HS < BN1,
+                      "Malformed RHS region - too big vertical and/or "
+                      "horizontal stride.");
         // Merge and update.
         Base1 = __esimd_wrregion<ElemTy, BN1, M, VS, W, HS, ParentWidth>(
             Base1, Val, Offset);

sycl/test/esimd/wrregion.cpp

@@ -9,10 +9,21 @@ using namespace sycl::ext::intel::esimd;
 // test wrregion size checks.
 SYCL_ESIMD_FUNCTION void test_wrregion_size_check() {
   simd<int, 16> v16 = 0;
-  v16.template select<64, 1>(0) = slm_block_load<int, 64>(0);
-  // expected-error@* {{no matching function for call to '__esimd_wrregion'}}
+  simd<int, 64> v64;
+  v16.template select<64, 1>(0) = v64;
+  // expected-error@* {{static assertion failed due to requirement 'M <= BN'}}
   // expected-note@sycl/ext/intel/esimd/detail/simd_view_impl.hpp:* {{in instantiation of function template specialization}}
   // expected-note@sycl/ext/intel/esimd/detail/simd_view_impl.hpp:* {{in instantiation of member function}}
-  // expected-note@* {{operator=' requested here}}
+  // expected-note@wrregion.cpp:* {{in instantiation of member function}}
+  // expected-note@sycl/ext/intel/esimd/detail/simd_obj_impl.hpp:* {{expression evaluates to '64 <= 16'}}
+
+  // expected-error@* {{no matching function for call to '__esimd_wrregion'}}
   // expected-note@sycl/ext/intel/esimd/detail/intrin.hpp:* {{candidate template ignored: requirement '64 <= 16' was not satisfied}}
+
+  simd<int, 2> v2;
+  v16.template select<2, 64>() = v2;
+  // expected-error@* {{static assertion failed due to requirement '(M > BN) || (M - 1) * Stride < BN'}}
+  // expected-note@sycl/ext/intel/esimd/detail/simd_view_impl.hpp:* {{in instantiation of function template specialization}}
+  // expected-note@sycl/ext/intel/esimd/detail/simd_view_impl.hpp:* {{in instantiation of member function}}
+  // expected-note@wrregion.cpp:* {{in instantiation of member function}}
 }