[SYCL][ESIMD][EMU] Marking ESIMD kernels for esimd_emulator backend (intel#751)

- 'vadd_1d' and 'sycl_esimd_mix' are duplicated to
'esimd_check_vc_codegen.cpp' and 'sycl_esimd_mix_check_build_opts.cpp'
so that original kernels can run with esimd_emulator without 'CHECK'
commands and duplicated kernels can run with opencl/level_zero
backends with 'CHECK'
- Add XFAIL/UNSUPPORTED markup for
    - Unimplemented features
    - Outdated memory intrinsic implementations
- Reduce workset in matrix_transpose_glb.cpp when running on esimd_emulator
- Removing HOST_RUN_PLACEHOLDER for ESIMD Kernels

Author: dongkyunahn-intel

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/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/replicate_smoke.cpp

@@ -8,10 +8,15 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
 <<<<<<< HEAD
+<<<<<<< HEAD
 // TODO: esimd_emulator fails due to unimplemented 'half' type
 // XFAIL: esimd_emulator
 =======
 >>>>>>> 22f86b532 ([ESIMD] Add smoke test for simd::replicate_vs_w_hs. (#799))
+=======
+// TODO: esimd_emulator fails due to unimplemented 'half' type
+// XFAIL: esimd_emulator
+>>>>>>> d98407d06 ([SYCL][ESIMD][EMU] Marking ESIMD kernels for esimd_emulator backend (#751))
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/api/saturation_smoke.cpp

@@ -8,10 +8,15 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
 <<<<<<< HEAD
+<<<<<<< HEAD
 // TODO: esimd_emulator fails due to unimplemented 'half' type
 // XFAIL: esimd_emulator
 =======
 >>>>>>> 2206708db ([ESIMD] Add "smoke" test for esimd::saturate (#791))
+=======
+// TODO: esimd_emulator fails due to unimplemented 'half' type
+// XFAIL: esimd_emulator
+>>>>>>> d98407d06 ([SYCL][ESIMD][EMU] Marking ESIMD kernels for esimd_emulator backend (#751))
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

SYCL/ESIMD/api/simd_any_all.cpp

@@ -8,10 +8,15 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
 <<<<<<< HEAD
+<<<<<<< HEAD
 // TODO: esimd_emulator fails due to unimplemented 'half' type
 // XFAIL: esimd_emulator
 =======
 >>>>>>> 38486e7dc ([ESIMD] Add smoke test for simd_obj_impl::any/all. (#802))
+=======
+// TODO: esimd_emulator fails due to unimplemented 'half' type
+// XFAIL: esimd_emulator
+>>>>>>> d98407d06 ([SYCL][ESIMD][EMU] Marking ESIMD kernels for esimd_emulator backend (#751))
 // 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_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_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_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/esimd_check_vc_codegen.cpp

@@ -16,7 +16,11 @@
 
 #include <CL/sycl.hpp>
 #include <iostream>
+<<<<<<< HEAD
 #include <sycl/ext/intel/esimd.hpp>
+=======
+#include <sycl/ext/intel/experimental/esimd.hpp>
+>>>>>>> d98407d06 ([SYCL][ESIMD][EMU] Marking ESIMD kernels for esimd_emulator backend (#751))
 
 using namespace cl::sycl;
 
@@ -55,7 +59,11 @@ int main(void) {
       auto PC = bufc.get_access<access::mode::write>(cgh);
       cgh.parallel_for<class Test>(
           GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
+<<<<<<< HEAD
             using namespace sycl::ext::intel::esimd;
+=======
+            using namespace sycl::ext::intel::experimental::esimd;
+>>>>>>> d98407d06 ([SYCL][ESIMD][EMU] Marking ESIMD kernels for esimd_emulator backend (#751))
             unsigned int offset = i * VL * sizeof(float);
             simd<float, VL> va;
             va.copy_from(PA, offset);

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: %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_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: %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: %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: %GPU_RUN_PLACEHOLDER %t.out %S/linear_in.bmp %S/linear_gold_hw.bmp
 

SYCL/ESIMD/matrix_transpose_glb.cpp

@@ -22,6 +22,8 @@ using namespace sycl::ext::intel::esimd;
 
 const unsigned int ESIMD_EMULATOR_SIZE_LIMIT = 1U << 10;
 
+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");

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

@@ -9,8 +9,11 @@
 // UNSUPPORTED: cuda || hip
 // TODO/DEBUG Segmentation fault occurs with esimd_emulator backend
 // XFAIL: esimd_emulator
+<<<<<<< HEAD
 // The test hangs on 22.05.22297 GPU RT on Linux
 // UNSUPPORTED: linux && (opencl || level_zero) && gpu
+=======
+>>>>>>> d98407d06 ([SYCL][ESIMD][EMU] Marking ESIMD kernels for esimd_emulator backend (#751))
 //
 // 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/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
 

SYCL/ESIMD/slm_barrier.cpp

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

SYCL/ESIMD/vadd_2d_acc.cpp

@@ -10,6 +10,8 @@
 // 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.

SYCL/ESIMD/vadd_raw_send.cpp

@@ -9,10 +9,15 @@
 // REQUIRES: gpu
 // UNSUPPORTED: gpu-intel-dg1,cuda,hip
 <<<<<<< HEAD
+<<<<<<< HEAD
 // TODO: esimd_emulator fails due to unimplemented 'raw_send' intrinsic
 // XFAIL: esimd_emulator
 =======
 >>>>>>> 3a9a0d3c0 ([ESIMD] Exclude esimd_vadd_raw_send and histogram_raw_send from DG1 testing (#728))
+=======
+// TODO: esimd_emulator fails due to unimplemented 'raw_send' intrinsic
+// XFAIL: esimd_emulator
+>>>>>>> d98407d06 ([SYCL][ESIMD][EMU] Marking ESIMD kernels for esimd_emulator backend (#751))
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out