[SYCL][ESIMD][EMU] ESIMD test updates for esimd_emulator backend (#869)

* [SYCL][ESIMD][EMU] ESIMD test updates for esimd_emulator backend

- Reducing number of iterations / Limiting matrix size in order to
avoid timeout failure
- Removing 'XFAIL' from passing tests
- Required memory intrinsic implementations are supported in
intel/llvm/sycl
- Recovering XFAIL for 'SYCL/ESIMD/api/simd_subscript_operator.cpp'
  For unimplemented 'half' type

Author: dongkyunahn-intel

SYCL/ESIMD/BitonicSortK.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
@@ -602,6 +600,12 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
   double kernel_times = 0;
   unsigned num_iters = 10;
 
+  // Reducing number of iterations for esimd_emulator backend in order
+  // to avoid timeout failure
+  if (pQueue_->get_backend() == cl::sycl::backend::ext_intel_esimd_emulator) {
+    num_iters = 2;
+  }
+
   // num_iters + 1, iteration#0 is for warmup
   for (int iter = 0; iter <= num_iters; ++iter) {
     try {

SYCL/ESIMD/BitonicSortKv2.cpp

@@ -521,6 +521,12 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
   double kernel_times = 0;
   unsigned num_iters = 10;
 
+  // Reducing number of iterations for esimd_emulator backend in order
+  // to avoid timeout failure
+  if (pQueue_->get_backend() == cl::sycl::backend::ext_intel_esimd_emulator) {
+    num_iters = 2;
+  }
+
   // num_iters + 1, iteration#0 is for warmup
   for (int iter = 0; iter <= num_iters; ++iter) {
     // enqueue sort265 kernel

SYCL/ESIMD/accessor.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl -D_CRT_SECURE_NO_WARNINGS=1 %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/aot_mixed.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
+// UNSUPPORTED: esimd_emulator
 // RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device gen9" -o %t.sycl.out -DENABLE_SYCL=0 %s
 // RUN: %GPU_RUN_PLACEHOLDER %t.sycl.out
 // RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device gen9" -o %t.out %s

SYCL/ESIMD/api/ballot.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to outdated memory intrinsic
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/api/simd_binop_integer_promotion.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to outdated memory intrinsic
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/api/simd_negation_operator.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/api/simd_subscript_operator.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to outdated memory intrinsic
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/api/simd_view_negation_operator.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/api/simd_view_subscript_operator.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/api/slm_gather_scatter.cpp

@@ -1,7 +1,5 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/api/slm_gather_scatter_rgba.cpp

@@ -1,7 +1,5 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/fp_in_phi.cpp

@@ -14,8 +14,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 // UNSUPPORTED: ze_debug-1,ze_debug4
-// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled
-// XFAIL: esimd_emulator
 //
 // The test checks that ESIMD kernels correctly handle function pointers as
 // arguments of LLVM's PHI function.

SYCL/ESIMD/histogram.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to outdated __esimd_media_ld
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/histogram_256_slm.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/histogram_2d.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to outdated __esimd_media_ld
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/histogram_raw_send.cpp

@@ -9,8 +9,6 @@
 // REQUIRES: gpu
 // UNSUPPORTED: gpu-intel-dg1,cuda,hip
 // UNSUPPORTED: ze_debug-1,ze_debug4
-// TODO: esimd_emulator fails due to outdated __esimd_media_ld
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/linear/linear.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to outdated __esimd_media_ld
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -I%S/.. -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %S/linear_in.bmp %S/linear_gold_hw.bmp
 

SYCL/ESIMD/matrix_transpose2.cpp

@@ -25,6 +25,8 @@ using namespace cl::sycl;
 using namespace std;
 using namespace sycl::ext::intel::experimental::esimd;
 
+const unsigned int ESIMD_EMULATOR_SIZE_LIMIT = 1U << 10;
+
 void initMatrix(int *M, unsigned N) {
   assert(N >= 8 && (((N - 1) & N) == 0) &&
          "only power of 2 (>= 16) is supported");
@@ -278,6 +280,16 @@ bool runTest(unsigned MZ, unsigned block_size, unsigned num_iters,
   cerr << "\nTranspose square matrix of size " << MZ << "\n";
   // printMatrix("Initial matrix:", M, MZ);
 
+  if ((q.get_backend() == cl::sycl::backend::ext_intel_esimd_emulator) &&
+      (MZ > ESIMD_EMULATOR_SIZE_LIMIT)) {
+    cerr << "Matrix Size larger than " << ESIMD_EMULATOR_SIZE_LIMIT
+         << " is skipped"
+         << "\n";
+    cerr << "for esimd_emulator backend due to timeout"
+         << "\n";
+    return true;
+  }
+
   // Each C-for-Metal thread works on one or two blocks of size 8 x 8.
   int thread_width = MZ / block_size;
   int thread_height = MZ / block_size;

SYCL/ESIMD/preemption.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu && linux
 // UNSUPPORTED: cuda || hip
+// UNSUPPORTED: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER IGC_DumpToCustomDir=%t.dump IGC_ShaderDumpEnable=1 %t.out
 // RUN: grep enablePreemption %t.dump/*.asm

SYCL/ESIMD/regression/operator_decrement.cpp

@@ -9,8 +9,6 @@
 //
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/regression/unused_load.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/regression/variable_gather_mask.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
+// TODO: esimd_emulator fails due to unimplemented 'single_task()' method
+// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/slm_barrier.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled4
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/slm_split_barrier.cpp

@@ -9,8 +9,6 @@
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_scatter4_scaled
-// XFAIL: esimd_emulator
 
 #include "esimd_test_utils.hpp"
 

SYCL/ESIMD/vadd_1d.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/vadd_2d.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to outdated __esimd_media_ld
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/vadd_2d_acc.cpp

@@ -10,8 +10,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned
-// XFAIL: esimd_emulator
 
 // The test checks that 2D workitem addressing works correctly with SIMD
 // kernels.