[SYCL][ESIMD][EMU] Handle intrinsic operations promoted to 4-byte element type (#5727)

* [SYCL][ESIMD][EMU] Handling intrinsic operations promoted to 4-byte element type

- scatter_impl() and gather_impl() in memory.hpp promote
argument/return vector to 4-byte elements

Author: dongkyunahn-intel

sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp

@@ -698,7 +698,7 @@ __esimd_pack_mask(__ESIMD_DNS::vector_type_t<uint16_t, N> src0) {
   // wrapper code (which does the checks already)
   uint32_t retv = 0;
   for (int i = 0; i < N; i++) {
-    if (src0[i] & 0x1) {
+    if (src0[i] != 0) {
       retv |= 0x1 << i;
     }
   }
@@ -709,12 +709,10 @@ __esimd_pack_mask(__ESIMD_DNS::vector_type_t<uint16_t, N> src0) {
 template <int N>
 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<uint16_t, N>
 __esimd_unpack_mask(uint32_t src0) {
-  __ESIMD_DNS::vector_type_t<uint16_t, N> retv;
+  __ESIMD_DNS::vector_type_t<uint16_t, N> retv = 0;
   for (int i = 0; i < N; i++) {
     if ((src0 >> i) & 0x1) {
       retv[i] = 1;
-    } else {
-      retv[i] = 0;
     }
   }
   return retv;

sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp

@@ -42,12 +42,15 @@ namespace __ESIMD_DNS {
 // Provides access to sycl accessor class' private members.
 class AccessorPrivateProxy {
 public:
-#ifdef __SYCL_DEVICE_ONLY__
   template <typename AccessorTy>
   static auto getNativeImageObj(const AccessorTy &Acc) {
+#ifdef __SYCL_DEVICE_ONLY__
     return Acc.getNativeImageObj();
-  }
 #else  // __SYCL_DEVICE_ONLY__
+    return Acc;
+#endif // __SYCL_DEVICE_ONLY__
+  }
+#ifndef __SYCL_DEVICE_ONLY__
   static void *getPtr(const sycl::detail::AccessorBaseHost &Acc) {
     return Acc.getPtr();
   }
@@ -421,18 +424,32 @@ __esimd_scatter_scaled(__ESIMD_DNS::simd_mask_storage_t<N> pred,
   static_assert(TySizeLog2 <= 2);
   static_assert(std::is_integral<Ty>::value || TySizeLog2 == 2);
 
+  // determine the original element's type size (as __esimd_scatter_scaled
+  // requires vals to be a vector of 4-byte integers)
+  constexpr size_t OrigSize = __ESIMD_DNS::ElemsPerAddrDecoding(TySizeLog2);
+  using RestoredTy = __ESIMD_DNS::uint_type_t<OrigSize>;
+
   sycl::detail::ESIMDDeviceInterface *I =
       sycl::detail::getESIMDDeviceInterface();
 
+  __ESIMD_DNS::vector_type_t<RestoredTy, N> TypeAdjustedVals;
+  if constexpr (OrigSize == 4) {
+    TypeAdjustedVals = __ESIMD_DNS::bitcast<RestoredTy, Ty, N>(vals);
+  } else {
+    static_assert(OrigSize == 1 || OrigSize == 2);
+    TypeAdjustedVals = __ESIMD_DNS::convert_vector<RestoredTy, Ty, N>(vals);
+  }
+
   if (surf_ind == __ESIMD_NS::detail::SLM_BTI) {
     // Scattered-store for Shared Local Memory
     // __ESIMD_NS::detail::SLM_BTI is special binding table index for SLM
     assert(global_offset == 0);
     char *SlmBase = I->__cm_emu_get_slm_ptr();
     for (int i = 0; i < N; ++i) {
       if (pred[i]) {
-        Ty *addr = reinterpret_cast<Ty *>(elem_offsets[i] + SlmBase);
-        *addr = vals[i];
+        RestoredTy *addr =
+            reinterpret_cast<RestoredTy *>(elem_offsets[i] + SlmBase);
+        *addr = TypeAdjustedVals[i];
       }
     }
   } else {
@@ -449,8 +466,9 @@ __esimd_scatter_scaled(__ESIMD_DNS::simd_mask_storage_t<N> pred,
 
     for (int idx = 0; idx < N; idx++) {
       if (pred[idx]) {
-        Ty *addr = reinterpret_cast<Ty *>(elem_offsets[idx] + writeBase);
-        *addr = vals[idx];
+        RestoredTy *addr =
+            reinterpret_cast<RestoredTy *>(elem_offsets[idx] + writeBase);
+        *addr = TypeAdjustedVals[idx];
       }
     }
 
@@ -629,7 +647,12 @@ __esimd_gather_masked_scaled2(SurfIndAliasTy surf_ind, uint32_t global_offset,
 {
   static_assert(Scale == 0);
 
-  __ESIMD_DNS::vector_type_t<Ty, N> retv = 0;
+  // determine the original element's type size (as __esimd_scatter_scaled
+  // requires vals to be a vector of 4-byte integers)
+  constexpr size_t OrigSize = __ESIMD_DNS::ElemsPerAddrDecoding(TySizeLog2);
+  using RestoredTy = __ESIMD_DNS::uint_type_t<OrigSize>;
+
+  __ESIMD_DNS::vector_type_t<RestoredTy, N> retv = 0;
   sycl::detail::ESIMDDeviceInterface *I =
       sycl::detail::getESIMDDeviceInterface();
 
@@ -639,7 +662,8 @@ __esimd_gather_masked_scaled2(SurfIndAliasTy surf_ind, uint32_t global_offset,
     char *SlmBase = I->__cm_emu_get_slm_ptr();
     for (int idx = 0; idx < N; ++idx) {
       if (pred[idx]) {
-        Ty *addr = reinterpret_cast<Ty *>(offsets[idx] + SlmBase);
+        RestoredTy *addr =
+            reinterpret_cast<RestoredTy *>(offsets[idx] + SlmBase);
         retv[idx] = *addr;
       }
     }
@@ -655,15 +679,21 @@ __esimd_gather_masked_scaled2(SurfIndAliasTy surf_ind, uint32_t global_offset,
     std::unique_lock<std::mutex> lock(*mutexLock);
     for (int idx = 0; idx < N; idx++) {
       if (pred[idx]) {
-        Ty *addr = reinterpret_cast<Ty *>(offsets[idx] + readBase);
+        RestoredTy *addr =
+            reinterpret_cast<RestoredTy *>(offsets[idx] + readBase);
         retv[idx] = *addr;
       }
     }
 
     // TODO : Optimize
     I->cm_fence_ptr();
   }
-  return retv;
+
+  if constexpr (OrigSize == 4) {
+    return __ESIMD_DNS::bitcast<Ty, RestoredTy, N>(retv);
+  } else {
+    return __ESIMD_DNS::convert_vector<Ty, RestoredTy, N>(retv);
+  }
 }
 #endif // __SYCL_DEVICE_ONLY__
 

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

@@ -61,12 +61,8 @@ __ESIMD_API SurfaceIndex get_surface_index(AccessorTy acc) {
   if constexpr (std::is_same_v<detail::LocalAccessorMarker, AccessorTy>) {
     return detail::SLM_BTI;
   } else {
-#ifdef __SYCL_DEVICE_ONLY__
-    const auto mem_obj = detail::AccessorPrivateProxy::getNativeImageObj(acc);
-    return __esimd_get_surface_index(mem_obj);
-#else  // __SYCL_DEVICE_ONLY__
-    return __esimd_get_surface_index(acc);
-#endif // __SYCL_DEVICE_ONLY__
+    return __esimd_get_surface_index(
+        detail::AccessorPrivateProxy::getNativeImageObj(acc));
   }
 }
 
@@ -253,12 +249,8 @@ __ESIMD_API simd<Tx, N> block_load(AccessorTy acc, uint32_t offset,
   static_assert(Sz <= 8 * detail::OperandSize::OWORD,
                 "block size must be at most 8 owords");
 
-#if defined(__SYCL_DEVICE_ONLY__)
   auto surf_ind = __esimd_get_surface_index(
       detail::AccessorPrivateProxy::getNativeImageObj(acc));
-#else  // __SYCL_DEVICE_ONLY__
-  auto surf_ind = __esimd_get_surface_index(acc);
-#endif // __SYCL_DEVICE_ONLY__
 
   if constexpr (Flags::template alignment<simd<T, N>> >=
                 detail::OperandSize::OWORD) {
@@ -317,12 +309,8 @@ __ESIMD_API void block_store(AccessorTy acc, uint32_t offset,
   static_assert(Sz <= 8 * detail::OperandSize::OWORD,
                 "block size must be at most 8 owords");
 
-#if defined(__SYCL_DEVICE_ONLY__)
   auto surf_ind = __esimd_get_surface_index(
       detail::AccessorPrivateProxy::getNativeImageObj(acc));
-#else //
-  auto surf_ind = __esimd_get_surface_index(acc);
-#endif
   __esimd_oword_st<T, N>(surf_ind, offset >> 4, vals.data());
 }