From 197c33a2b17b6c2ce424c9cf53b7a3ea1ddbeca1 Mon Sep 17 00:00:00 2001 From: Vyacheslav Klochkov Date: Mon, 4 Dec 2023 14:22:43 -0600 Subject: [PATCH] [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 17f893972bed74c1c6e147291edefa78ba16181c 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 vec; // vec is only 16-elements auto vec_too_bit_view = vec.select<128, 1>();// 128-elems - too many vec_too_bit_view = simd{0}; //error: write 128-elem to 16-elem storage Signed-off-by: Klochkov, Vyacheslav N --- .../sycl/ext/intel/esimd/detail/intrin.hpp | 9 ++++----- .../ext/intel/esimd/detail/simd_obj_impl.hpp | 18 ++++++++++++++++++ sycl/test/esimd/wrregion.cpp | 17 ++++++++++++++--- 3 files changed, 36 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/ext/intel/esimd/detail/intrin.hpp b/sycl/include/sycl/ext/intel/esimd/detail/intrin.hpp index 17211847314fe..2cfa6547fbc2a 100644 --- a/sycl/include/sycl/ext/intel/esimd/detail/intrin.hpp +++ b/sycl/include/sycl/ext/intel/esimd/detail/intrin.hpp @@ -113,8 +113,7 @@ __esimd_rdindirect(__ESIMD_DNS::vector_type_t 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 using __st = __raw_t; /// read from a basic region of a vector, return a vector template -__ESIMD_DNS::vector_type_t<__st, RTy::length> - ESIMD_INLINE readRegion( - const __ESIMD_DNS::vector_type_t<__st, BN> &Base, RTy Region) { +__ESIMD_DNS::vector_type_t<__st, + RTy::length> ESIMD_INLINE +readRegion(const __ESIMD_DNS::vector_type_t<__st, BN> &Base, RTy Region) { using ElemTy = __st; auto Base1 = bitcast, BN>(Base); constexpr int Bytes = BN * sizeof(BT); diff --git a/sycl/include/sycl/ext/intel/esimd/detail/simd_obj_impl.hpp b/sycl/include/sycl/ext/intel/esimd/detail/simd_obj_impl.hpp index c5dc21860df52..ae6137c01fe12 100644 --- a/sycl/include/sycl/ext/intel/esimd/detail/simd_obj_impl.hpp +++ b/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 0, "Malformed RHS region."); + static_assert((M - 1) * Stride < BN, + "Malformed RHS region - too big stride."); // Merge and update. Base1 = __esimd_wrregion(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( Base1, Val, Offset); diff --git a/sycl/test/esimd/wrregion.cpp b/sycl/test/esimd/wrregion.cpp index 211c192589912..7e9a62347825f 100644 --- a/sycl/test/esimd/wrregion.cpp +++ b/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 v16 = 0; - v16.template select<64, 1>(0) = slm_block_load(0); - // expected-error@* {{no matching function for call to '__esimd_wrregion'}} + simd 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 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}} }