[SYCL][ESIMD] reduce restriction on some esimd API based upon user feedback

- allow 256-byte block load for slm
- allow vector-length of 1/2/4/8/16/32 for gather and scatter

Signed-off-by: Gang Y Chen <[email protected]>

Author: Gang Y Chen

sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp

@@ -84,7 +84,7 @@ constexpr unsigned int ElemsPerAddrEncoding() {
 template <typename T, int n, int ElemsPerAddr = 1,
           CacheHint L1H = CacheHint::None, CacheHint L3H = CacheHint::None>
 ESIMD_INLINE ESIMD_NODEBUG
-    typename std::enable_if<((n == 8 || n == 16 || n == 32) &&
+    typename std::enable_if<(__esimd::isPowerOf2(n, 32) &&
                              (ElemsPerAddr == 1 || ElemsPerAddr == 2 ||
                               ElemsPerAddr == 4)),
                             simd<T, n * ElemsPerAddr>>::type
@@ -120,7 +120,7 @@ ESIMD_INLINE ESIMD_NODEBUG
 template <typename T, int n, int ElemsPerAddr = 1,
           CacheHint L1H = CacheHint::None, CacheHint L3H = CacheHint::None>
 ESIMD_INLINE ESIMD_NODEBUG
-    typename std::enable_if<((n == 8 || n == 16 || n == 32) &&
+    typename std::enable_if<(__esimd::isPowerOf2(n, 32) &&
                              (ElemsPerAddr == 1 || ElemsPerAddr == 2 ||
                               ElemsPerAddr == 4)),
                             void>::type
@@ -487,9 +487,9 @@ ESIMD_INLINE ESIMD_NODEBUG simd<T, n> slm_block_load(uint32_t offset) {
   static_assert(Sz % __esimd::OWORD == 0,
                 "block size must be whole number of owords");
   static_assert(__esimd::isPowerOf2(Sz / __esimd::OWORD),
-                "block must be 1, 2, 4 or 8 owords long");
-  static_assert(Sz <= 8 * __esimd::OWORD,
-                "block size must be at most 8 owords");
+                "block must be 1, 2, 4, 8, 16 owords long");
+  static_assert(Sz <= 16 * __esimd::OWORD,
+                "block size must be at most 16 owords");
 
   return __esimd_slm_block_read<T, n>(offset);
 }
@@ -503,9 +503,9 @@ ESIMD_INLINE ESIMD_NODEBUG void slm_block_store(uint32_t offset,
   static_assert(Sz % __esimd::OWORD == 0,
                 "block size must be whole number of owords");
   static_assert(__esimd::isPowerOf2(Sz / __esimd::OWORD),
-                "block must be 1, 2, 4 or 8 owords long");
-  static_assert(Sz <= 8 * __esimd::OWORD,
-                "block size must be at most 8 owords");
+                "block must be 1, 2, 4, 8, or 16 owords long");
+  static_assert(Sz <= 16 * __esimd::OWORD,
+                "block size must be at most 16 owords");
 
   // offset in genx.oword.st is in owords
   __esimd_slm_block_write<T, n>(offset >> 4, vals.data());

sycl/test/basic_tests/esimd/gather_scatter.cpp

@@ -9,7 +9,22 @@
 using namespace sycl::INTEL::gpu;
 using namespace cl::sycl;
 
-void kernel(accessor<int, 1, access::mode::read_write, access::target::global_buffer> &buf) __attribute__((sycl_device)) {
+void kernel0(accessor<int, 1, access::mode::read_write,
+                      access::target::global_buffer> &buf)
+    __attribute__((sycl_device)) {
+  simd<uint32_t, 2> offsets(0, 1);
+  simd<int, 2> v1(0, 1);
+
+  auto v0 = gather<int, 2>(buf.get_pointer(), offsets);
+
+  v0 = v0 + v1;
+
+  scatter<int, 2>(buf.get_pointer(), v0, offsets);
+}
+
+void kernel(accessor<int, 1, access::mode::read_write,
+                     access::target::global_buffer> &buf)
+    __attribute__((sycl_device)) {
   simd<uint32_t, 32> offsets(0, 1);
   simd<int, 32> v1(0, 1);
 
@@ -20,7 +35,9 @@ void kernel(accessor<int, 1, access::mode::read_write, access::target::global_bu
   scatter<int, 32>(buf.get_pointer(), v0, offsets);
 }
 
-void kernel(accessor<uint8_t, 1, access::mode::read_write, access::target::global_buffer> &buf) __attribute__((sycl_device)) {
+void kernel(accessor<uint8_t, 1, access::mode::read_write,
+                     access::target::global_buffer> &buf)
+    __attribute__((sycl_device)) {
   simd<uint32_t, 32> offsets(0, 1);
   simd<uint8_t, 32> v1(0, 1);
 
@@ -33,7 +50,9 @@ void kernel(accessor<uint8_t, 1, access::mode::read_write, access::target::globa
   scatter<uint8_t, 32>(buf.get_pointer(), v0, offsets);
 }
 
-void kernel(accessor<uint16_t, 1, access::mode::read_write, access::target::global_buffer> &buf) __attribute__((sycl_device)) {
+void kernel(accessor<uint16_t, 1, access::mode::read_write,
+                     access::target::global_buffer> &buf)
+    __attribute__((sycl_device)) {
   simd<uint32_t, 32> offsets(0, 1);
   simd<uint16_t, 32> v1(0, 1);
 

sycl/test/basic_tests/esimd/slm_block.cpp

@@ -18,3 +18,13 @@ void kernel() __attribute__((sycl_device)) {
 
   slm_block_store<int, 32>(0, v0);
 }
+
+void kernel2() __attribute__((sycl_device)) {
+  simd<int, 64> v1(0, 1);
+
+  auto v0 = slm_block_load<int, 64>(0);
+
+  v0 = v0 + v1;
+
+  slm_block_store<int, 64>(0, v0);
+}