[SYCL][ESIMD][EMU] Marking ESIMD kernels for esimd_emulator backend

SYCL/ESIMD/BitonicSortK.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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/accessor.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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/accessor_gather_scatter.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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/accessor_load_store.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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/ballot.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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/bin_and_cmp_ops_heavy.cpp

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

SYCL/ESIMD/api/esimd_bit_ops.cpp

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

SYCL/ESIMD/api/simd_binop_integer_promotion.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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_copy_to_from.cpp

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

SYCL/ESIMD/api/simd_memory_access.cpp

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

SYCL/ESIMD/api/simd_negation_operator.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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_copy_move_assign.cpp

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

SYCL/ESIMD/api/simd_view_negation_operator.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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,5 +1,7 @@
 // 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_heavy.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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/slm_gather_scatter_rgba.cpp

@@ -1,5 +1,7 @@
 // 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/unary_ops_heavy.cpp

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

SYCL/ESIMD/esimd_check_vc_codegen.cpp

@@ -0,0 +1,42 @@
+//==-------- esimd_check_vc_codegen.cpp - DPC++ ESIMD on-device test -------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// UNSUPPORTED: cuda || hip
+// esimd_emulator does not support online-compiler that invokes 'piProgramBuild'
+// UNSUPPORTED: esimd_emulator
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: env SYCL_PI_TRACE=-1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER
+
+#include "esimd_test_utils.hpp"
+
+#include <CL/sycl.hpp>
+#include <iostream>
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+using namespace cl::sycl;
+
+int main(void) {
+  try {
+    queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
+
+    auto dev = q.get_device();
+    std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+
+    auto e =
+        q.submit([&](handler &cgh) { cgh.single_task<class Test>([] {}); });
+    e.wait();
+  } catch (sycl::exception const &e) {
+    std::cout << "SYCL exception caught: " << e.what() << '\n';
+    return 1;
+  }
+  return 0;
+}
+
+// CHECK: ---> piProgramBuild(
+// CHECK: <const char *>: {{.*}}-vc-codegen
+// CHECK: ) ---> pi_result : PI_SUCCESS

SYCL/ESIMD/ext_math.cpp

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

SYCL/ESIMD/fp_in_phi.cpp

@@ -14,6 +14,8 @@
 // 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,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out

SYCL/ESIMD/histogram_256_slm.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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_256_slm_spec_2020.cpp

@@ -1,6 +1,7 @@
 // TODO enable on Windows
 // REQUIRES: linux && gpu
 // UNSUPPORTED: cuda || hip
+// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out 16
 

SYCL/ESIMD/histogram_2d.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out

SYCL/ESIMD/histogram_raw_send.cpp

@@ -9,6 +9,8 @@
 // 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: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out

SYCL/ESIMD/linear/linear.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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: %HOST_RUN_PLACEHOLDER %t.out %S/linear_in.bmp %S/linear_gold_hw.bmp
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %S/linear_in.bmp %S/linear_gold_hw.bmp

SYCL/ESIMD/mandelbrot/mandelbrot.cpp

@@ -8,6 +8,8 @@
 
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
+// TODO: esimd_emulator fails due to outdated __esimd_media_st
+// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -I%S/.. -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %T/output.ppm %S/golden_hw.ppm
 

SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp

@@ -9,6 +9,7 @@
 // TODO enable on Windows
 // REQUIRES: linux && gpu
 // UNSUPPORTED: cuda || hip
+// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -I%S/.. -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %T/output_spec.ppm %S/golden_hw.ppm 512 -2.09798 -1.19798 0.004 4.0
 

SYCL/ESIMD/matrix_transpose.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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/matrix_transpose2.cpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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/matrix_transpose_glb.cpp

@@ -20,6 +20,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");
@@ -246,6 +248,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/printf.cpp

@@ -9,6 +9,8 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
 // CUDA and HIP don't support printf.
+// TODO: esimd_emulator fails due to unimplemented 'single_task()' method
+// UNSUPPORTED: esimd_emulator
 //
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER

SYCL/ESIMD/regression/big_const_initializer.cpp

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

SYCL/ESIMD/regression/complex-lib-lin.cpp

@@ -7,6 +7,7 @@
 //
 // REQUIRES: linux,gpu
 // UNSUPPORTED: cuda || hip
+// XFAIL: esimd_emulator
 //
 // RUN: %clangxx -fsycl -fPIC -O3 %S/Inputs/complex-lib-sycl.cpp -c -o %t-lib-sycl.o
 // RUN: %clangxx -fsycl -fPIC -O3 %S/Inputs/complex-lib-esimd.cpp -c -o %t-lib-esimd.o

SYCL/ESIMD/regression/globals.cpp

@@ -1,5 +1,7 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
+// Sub-group is not supported
+// UNSUPPORTED: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 

SYCL/ESIMD/regression/operator_decrement.cpp

@@ -9,6 +9,8 @@
 //
 // 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,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 // 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