Allowed 1- and 2-byte types in slm_load/store

Author: DenisBakhvalov

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

@@ -654,7 +654,7 @@ SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void slm_init(uint32_t size);
 /// Only allow simd-16 and simd-32.
 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),
                                        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());
@@ -663,7 +663,7 @@ ESIMD_INLINE ESIMD_NODEBUG
 /// SLM scatter.
 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) {

sycl/test/esimd/slm_load.cpp

@@ -13,10 +13,6 @@ 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();
@@ -25,3 +21,23 @@ void kernel() __attribute__((sycl_device)) {
 
   slm_store<int, 32>(v0, offsets);
 }
+
+void slm_supported_types() __attribute__((sycl_device)) {
+  simd<uint32_t, 32> offsets(0, 1);
+  auto v1 = slm_load<char, 32>(offsets);
+  auto v2 = slm_load<short, 32>(offsets);
+  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}}
+  auto v5 = slm_load<double, 32>(offsets);
+
+  slm_store<char, 32>(v1, offsets);
+  slm_store<short, 32>(v2, offsets);
+  slm_store<int, 32>(v3, offsets);
+  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}}
+  slm_store<double, 32>(v6, offsets);
+}