[SYCL][ESIMD] Update tests after API changes (#974)

* [SYCL][ESIMD] Update tests after API changes

Some arguments have been changes to template parameters to
force them to be compile-time constants which is required
by low-level intrinsics.

Signed-off-by: Sergey Dmitriev <[email protected]>
Co-authored-by: kbobrovs <[email protected]>

Author: sndmitriev

SYCL/ESIMD/api/slm_gather_scatter.cpp

@@ -21,9 +21,9 @@ template <typename T, unsigned VL, unsigned STRIDE> struct Kernel {
     using namespace sycl::ext::intel::esimd;
 
     // In this test, we have a single workitem. No barriers required.
-    slm_init(
-        VL * STRIDE *
-        sizeof(typename sycl::ext::intel::esimd::detail::dword_type<T>::type));
+    slm_init<VL * STRIDE *
+             sizeof(typename sycl::ext::intel::esimd::detail::dword_type<
+                    T>::type)>();
 
     simd<T, VL> valsIn;
     valsIn.copy_from(buf);

SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp

@@ -111,7 +111,7 @@ struct GatherKernel : KernelBase<T, VL, STRIDE> {
   static const char *get_name() { return "slm_gather"; }
 
   void operator()(nd_item<1> i) const SYCL_ESIMD_KERNEL {
-    slm_init(B::SLM_CHUNK_SIZE);
+    slm_init<B::SLM_CHUNK_SIZE>();
 
     // first, read data w/o shuffling into SLM
     simd<T, VL> val;
@@ -174,7 +174,7 @@ struct ScatterKernel : KernelBase<T, VL, STRIDE> {
   static const char *get_name() { return "slm_scatter"; }
 
   ESIMD_INLINE void operator()(nd_item<1> i) const SYCL_ESIMD_KERNEL {
-    slm_init(B::SLM_CHUNK_SIZE);
+    slm_init<B::SLM_CHUNK_SIZE>();
 
     // first, read data from memory into registers w/o shuffling
     simd<T, VL> val;
@@ -242,7 +242,7 @@ struct GatherKernel<T, 1, STRIDE, TEST_VECTOR_NO_MASK>
   static const char *get_name() { return "slm_gather_vl1"; }
 
   void operator()(nd_item<1> i) const SYCL_ESIMD_KERNEL {
-    slm_init(B::SLM_CHUNK_SIZE);
+    slm_init<B::SLM_CHUNK_SIZE>();
 
     // first, read data into SLM
     T val = scalar_load<T>(B::acc_in, B::get_wi_offset(i) * sizeof(T));
@@ -268,7 +268,7 @@ struct ScatterKernel<T, 1, STRIDE, TEST_VECTOR_NO_MASK>
   static const char *get_name() { return "slm_scatter_vl1"; }
 
   ESIMD_INLINE void operator()(nd_item<1> i) const SYCL_ESIMD_KERNEL {
-    slm_init(B::SLM_CHUNK_SIZE);
+    slm_init<B::SLM_CHUNK_SIZE>();
 
     // first, read data from memory into registers
     simd<T, 1> val;

SYCL/ESIMD/api/slm_gather_scatter_rgba.cpp

@@ -32,7 +32,7 @@ template <typename T, unsigned VL, auto CH_MASK> struct Kernel {
     // In this test, each group consist of one workitem. No barriers required.
     // Each workitem accesses contiguous block of VL elements, where
     // each element consists of RGBA channels.
-    slm_init(VL * NUM_RGBA_CHANNELS * sizeof(T));
+    slm_init<VL * NUM_RGBA_CHANNELS * sizeof(T)>();
 
     // Prepare initial values in SLM:
     // 0, -1, -2, -3, -4 ...

SYCL/ESIMD/histogram_256_slm.cpp

@@ -30,7 +30,7 @@ ESIMD_INLINE void histogram_atomic(const uint32_t *input_ptr, uint32_t *output,
                                    uint32_t gid, uint32_t lid,
                                    uint32_t local_size) {
   // Declare and initialize SLM
-  slm_init(SLM_SIZE);
+  slm_init<SLM_SIZE>();
   uint linear_id = gid * local_size + lid;
 
   simd<uint, 16> slm_offset(0, 1);

SYCL/ESIMD/histogram_256_slm_spec_2020.cpp

@@ -28,7 +28,7 @@ ESIMD_INLINE void histogram_atomic(const uint32_t *input_ptr, uint32_t *output,
                                    uint32_t gid, uint32_t lid,
                                    uint32_t local_size, uint32_t num_blocks) {
   // Declare and initialize SLM
-  slm_init(SLM_SIZE);
+  slm_init<SLM_SIZE>();
   uint linear_id = gid * local_size + lid;
 
   simd<uint, 16> slm_offset(0, 1);

SYCL/ESIMD/lsc/lsc_slm.cpp

@@ -60,7 +60,7 @@ int main() {
             auto compare = simd<int, SIMDSize>(id * SIMDSize, 1);
             auto swap = compare * 2;
 
-            slm_init(4096);
+            slm_init<4096>();
             lsc_slm_block_store<int, SIMDSize>(offset, data * 2);
             auto data_0 = lsc_slm_block_load<int, SIMDSize>(offset);
             lsc_block_store<int, SIMDSize>(access_0, offset, data_0);

SYCL/ESIMD/slm_barrier.cpp

@@ -112,7 +112,7 @@ int main(void) {
             uint globalID = ndi.get_global_id(0);
             uint groupID = ndi.get_group(0);
 
-            slm_init(1024);
+            slm_init<1024>();
 
             int grpMemOffset = groupID * groupSize * VL * 4;
 

SYCL/ESIMD/slm_split_barrier.cpp

@@ -19,6 +19,7 @@
 using namespace cl::sycl;
 using namespace sycl::ext::intel;
 using namespace sycl::ext::intel::esimd;
+using namespace sycl::ext::intel::experimental::esimd;
 
 #define LOCAL_SIZE 4
 #define GLOBAL_SIZE 6
@@ -68,11 +69,10 @@ void load_to_slm(uint grpSize, uint localId, uint slmOffset, char *addr,
     vOffsets += (grpSize * 256);
   }
 
-  esimd::fence(fence_mask::global_coherent_fence);
-  experimental::esimd::split_barrier(
-      experimental::esimd::split_barrier_action::signal);
-  experimental::esimd::split_barrier(
-      experimental::esimd::split_barrier_action::wait);
+  // add memory fence and split barriers
+  fence<fence_mask::global_coherent_fence>();
+  split_barrier<split_barrier_action::signal>();
+  split_barrier<split_barrier_action::wait>();
 }
 
 int main(void) {
@@ -115,7 +115,7 @@ int main(void) {
             uint globalID = ndi.get_global_id(0);
             uint groupID = ndi.get_group(0);
 
-            slm_init(1024);
+            slm_init<1024>();
 
             int grpMemOffset = groupID * groupSize * VL * 4;