Added SLM gather/scatter for 1- and 2-byte block sizes

Author: DenisBakhvalov

sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp

@@ -649,25 +649,73 @@ inline ESIMD_NODEBUG void esimd_sbarrier(split_barrier_action flag) {
 /// Declare per-work-group slm size.
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void slm_init(uint32_t size);
 
-/// SLM gather.
-///
-/// Only allow simd-16 and simd-32.
-template <typename T, int n>
+/// SLM gather (version for 4-byte block size).
+/// \tparam T element type of the input vector.
+/// \tparam N size of the \p offsets , \p pred and returned vectors. Must be 16
+/// or 32.
+/// @param offsets byte-offsets within the SLM.
+/// @param pred predication control used for masking lanes.
+/// @return vector of read values of type \p T.
+/// \ingroup sycl_esimd
+template <typename T, int N>
 ESIMD_INLINE ESIMD_NODEBUG
-    typename sycl::detail::enable_if_t<(n == 16 || n == 32) && (sizeof(T) <= 4),
-                                       simd<T, n>>
-    slm_load(simd<uint32_t, n> offsets, simd<uint16_t, n> pred = 1) {
-  return __esimd_slm_read<T, n>(offsets.data(), pred.data());
+    typename sycl::detail::enable_if_t<(N == 16 || N == 32) && (sizeof(T) == 4),
+                                       simd<T, N>>
+    slm_load(simd<uint32_t, N> offsets, simd<uint16_t, N> pred = 1) {
+  return __esimd_slm_read<T, N>(offsets.data(), pred.data());
 }
 
-/// SLM scatter.
-template <typename T, int n>
+/// SLM gather (version for 1- and 2-byte block size).
+/// \tparam T element type of the input vector.
+/// \tparam N size of the \p offsets , \p pred and returned vectors. Must be 16
+/// or 32.
+/// @param offsets byte-offsets within the SLM.
+/// @param pred predication control used for masking lanes.
+/// @return vector of read values of type \p T.
+/// \ingroup sycl_esimd
+template <typename T, int N>
+ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t<
+    (N == 16 || N == 32) && (sizeof(T) == 1 || sizeof(T) == 2), simd<T, N>>
+slm_load(simd<uint32_t, N> offsets, simd<uint16_t, N> pred = 1) {
+  typedef typename detail::dword_type<T>::type T1;
+  simd<T1, N> temp = __esimd_slm_read<T1, N>(offsets.data(), pred.data());
+  simd<T, N> res = temp;
+  return res;
+}
+
+/// SLM scatter (version for 4-byte block size).
+/// \tparam T element type of the input vector.
+/// \tparam N size of the \p offsets , \p pred and \p vals vectors. Must be 16
+/// or 32.
+/// @param vals values to be written.
+/// @param offsets byte-offsets within the SLM.
+/// @param pred predication control used for masking lanes.
+/// \ingroup sycl_esimd
+template <typename T, int N>
 ESIMD_INLINE ESIMD_NODEBUG
-    typename sycl::detail::enable_if_t<(n == 16 || n == 32) && (sizeof(T) <= 4),
+    typename sycl::detail::enable_if_t<(N == 16 || N == 32) && (sizeof(T) == 4),
                                        void>
-    slm_store(simd<T, n> vals, simd<uint32_t, n> offsets,
-              simd<uint16_t, n> pred = 1) {
-  __esimd_slm_write<T, n>(offsets.data(), vals.data(), pred.data());
+    slm_store(simd<T, N> vals, simd<uint32_t, N> offsets,
+              simd<uint16_t, N> pred = 1) {
+  __esimd_slm_write<T, N>(offsets.data(), vals.data(), pred.data());
+}
+
+/// SLM scatter (version for 1- and 2-byte block size).
+/// \tparam T element type of the input vector.
+/// \tparam N size of the \p offsets , \p pred and \p vals vectors. Must be 16
+/// or 32.
+/// @param vals values to be written.
+/// @param offsets byte-offsets within the SLM.
+/// @param pred predication control used for masking lanes.
+/// \ingroup sycl_esimd
+template <typename T, int N>
+ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t<
+    (N == 16 || N == 32) && (sizeof(T) == 1 || sizeof(T) == 2), void>
+slm_store(simd<T, N> vals, simd<uint32_t, N> offsets,
+          simd<uint16_t, N> pred = 1) {
+  typedef typename detail::dword_type<T>::type T1;
+  simd<T1, N> temp = vals;
+  __esimd_slm_write<T1, N>(offsets.data(), temp.data(), pred.data());
 }
 
 /// SLM gather4.

sycl/test/esimd/slm_load.cpp

@@ -29,7 +29,7 @@ void slm_supported_types() __attribute__((sycl_device)) {
   auto v3 = slm_load<int, 32>(offsets);
   auto v4 = slm_load<float, 32>(offsets);
   // expected-error@+2 {{no matching function for call to 'slm_load'}}
-  // expected-note@sycl/ext/intel/experimental/esimd/memory.hpp:* {{candidate template ignored}}
+  // expected-note@sycl/ext/intel/experimental/esimd/memory.hpp:* 2 {{candidate template ignored}}
   auto v5 = slm_load<double, 32>(offsets);
 
   slm_store<char, 32>(v1, offsets);
@@ -38,6 +38,6 @@ void slm_supported_types() __attribute__((sycl_device)) {
   slm_store<float, 32>(v4, offsets);
   simd<double, 32> v6(0, 1);
   // expected-error@+2 {{no matching function for call to 'slm_store'}}
-  // expected-note@sycl/ext/intel/experimental/esimd/memory.hpp:* {{candidate template ignored}}
+  // expected-note@sycl/ext/intel/experimental/esimd/memory.hpp:* 2 {{candidate template ignored}}
   slm_store<double, 32>(v6, offsets);
 }