[SYCL][ESIMD] Introduce compile time checks for wrregion API (#10837)

oneapi-src · Sep 6, 2023 · 17f8939 · 17f8939
1 parent 445e6fc
commit 17f8939
Show file tree

Hide file tree

Showing 3 changed files with 29 additions and 7 deletions.
diff --git a/sycl/include/sycl/ext/intel/esimd/detail/intrin.hpp b/sycl/include/sycl/ext/intel/esimd/detail/intrin.hpp
@@ -64,7 +64,8 @@
 //
 template <typename T, int N, int M, int VStride, int Width, int Stride,
           int ParentWidth = 0>
-__ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, M>
+__ESIMD_INTRIN __ESIMD_INTRIN std::enable_if_t<(Width > 0) && M % Width == 0,
+                                               __ESIMD_DNS::vector_type_t<T, M>>
 __esimd_rdregion(__ESIMD_DNS::vector_type_t<T, N> Input, uint16_t Offset);
 
 template <typename T, int N, int M, int ParentWidth = 0>
@@ -121,13 +122,14 @@ __esimd_rdindirect(__ESIMD_DNS::vector_type_t<T, N> Input,
 //
 template <typename T, int N, int M, int VStride, int Width, int Stride,
           int ParentWidth = 0>
-__ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, N>
+__ESIMD_INTRIN std::enable_if_t<M <= N && (Width > 0) && M % Width == 0,
+                                __ESIMD_DNS::vector_type_t<T, N>>
 __esimd_wrregion(__ESIMD_DNS::vector_type_t<T, N> OldVal,
                  __ESIMD_DNS::vector_type_t<T, M> NewVal, uint16_t Offset,
                  __ESIMD_DNS::simd_mask_storage_t<M> Mask = 1);
 
 template <typename T, int N, int M, int ParentWidth = 0>
-__ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, N>
+__ESIMD_INTRIN std::enable_if_t<M <= N, __ESIMD_DNS::vector_type_t<T, N>>
 __esimd_wrindirect(__ESIMD_DNS::vector_type_t<T, N> OldVal,
                    __ESIMD_DNS::vector_type_t<T, M> NewVal,
                    __ESIMD_DNS::vector_type_t<uint16_t, M> Offset,
@@ -262,7 +264,8 @@ __ESIMD_INTRIN uint16_t __esimd_all(__ESIMD_DNS::vector_type_t<T, N> src)
 // Implementations of ESIMD intrinsics for the SYCL host device
 template <typename T, int N, int M, int VStride, int Width, int Stride,
           int ParentWidth>
-__ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, M>
+__ESIMD_INTRIN __ESIMD_INTRIN std::enable_if_t<(Width > 0) && M % Width == 0,
+                                               __ESIMD_DNS::vector_type_t<T, M>>
 __esimd_rdregion(__ESIMD_DNS::vector_type_t<T, N> Input, uint16_t Offset) {
   uint16_t EltOffset = Offset / sizeof(T);
   assert(Offset % sizeof(T) == 0);
@@ -296,7 +299,8 @@ __esimd_rdindirect(__ESIMD_DNS::vector_type_t<T, N> Input,
 
 template <typename T, int N, int M, int VStride, int Width, int Stride,
           int ParentWidth>
-__ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, N>
+__ESIMD_INTRIN std::enable_if_t<M <= N && (Width > 0) && M % Width == 0,
+                                __ESIMD_DNS::vector_type_t<T, N>>
 __esimd_wrregion(__ESIMD_DNS::vector_type_t<T, N> OldVal,
                  __ESIMD_DNS::vector_type_t<T, M> NewVal, uint16_t Offset,
                  __ESIMD_DNS::simd_mask_storage_t<M> Mask) {
@@ -319,7 +323,7 @@ __esimd_wrregion(__ESIMD_DNS::vector_type_t<T, N> OldVal,
 }
 
 template <typename T, int N, int M, int ParentWidth>
-__ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, N>
+__ESIMD_INTRIN std::enable_if_t<M <= N, __ESIMD_DNS::vector_type_t<T, N>>
 __esimd_wrindirect(__ESIMD_DNS::vector_type_t<T, N> OldVal,
                    __ESIMD_DNS::vector_type_t<T, M> NewVal,
                    __ESIMD_DNS::vector_type_t<uint16_t, M> Offset,

diff --git a/sycl/test/esimd/simd.cpp b/sycl/test/esimd/simd.cpp
@@ -190,7 +190,7 @@ template bool test_format_1d_read<int, short>() SYCL_ESIMD_FUNCTION;
 template bool test_format_1d_read<sycl::half, uint8_t>() SYCL_ESIMD_FUNCTION;
 
 template <class T1, class T2> bool test_format_1d_write() SYCL_ESIMD_FUNCTION {
-  simd<T1, 8> r;
+  simd<T1, 32> r;
   auto rl = r.template bit_cast_view<T2>();
   auto rl2 = rl.template select<8, 2>(0);
   auto rh = r.template bit_cast_view<T2>();

diff --git a/sycl/test/esimd/wrregion.cpp b/sycl/test/esimd/wrregion.cpp
@@ -0,0 +1,18 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
+
+#include <limits>
+#include <sycl/ext/intel/esimd.hpp>
+#include <utility>
+
+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'}}
+  // 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@sycl/ext/intel/esimd/detail/intrin.hpp:* {{candidate template ignored: requirement '64 <= 16' was not satisfied}}
+}