[ESIMD] Only dword types are supported by slm_load/store

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

@@ -650,22 +650,36 @@ inline ESIMD_NODEBUG void esimd_sbarrier(split_barrier_action flag) {
 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>
+/// \tparam T element type of the input vector, must be 4-byte type.
+/// \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), 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>
+/// \tparam T element type of the input vector, must be 4-byte type.
+/// \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), 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());
+    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 gather4.

sycl/test/esimd/slm_load.cpp

@@ -1,5 +1,4 @@
 // RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
-// expected-no-diagnostics
 
 #include <CL/sycl.hpp>
 #include <sycl/ext/intel/experimental/esimd.hpp>
@@ -14,6 +13,10 @@ void kernel() __attribute__((sycl_device)) {
   simd<int, 32> v1(0, 1);
 
   auto v0 = slm_load<int, 32>(offsets);
+  auto v2 = 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}}
+  auto v3 = slm_load<double, 32>(offsets);
 
   esimd_fence(3);
   esimd_barrier();