[SYCL][ESIMD] Move/Refactor ESIMD device code checks not in check_device_code (#14218)

In the same vein as #14174, this PR:
- Moves ESIMD device code checks that are not in the `check_device_code`
folder into the `check_device_code/esimd` folder
- Removed host code from tests that check only for device code using
`SYCL_EXTERNAL` functions (as per #14173)

ESIMD folks, I recognize this shuffles a large number of tests around:
Please let me know if there's anything problematic about this change,
and I will do my best to remedy the situation should it arise. Thanks in
advance!

Author: ianayl

sycl/test/esimd/NBarrierAttr.cpp renamed to sycl/test/check_device_code/esimd/NBarrierAttr.cpp

@@ -56,14 +56,14 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void xmx_func() {
   // CHECK-LABEL: define dso_local spir_func void @_Z8xmx_funcv()
 
   { // ======= DPAS BF16 =======================================================
-    simd<bfloat16, M_one *N_pvc> R_bf = 0;
-    simd<float, M_one *N_pvc> R_f = 0;
+    simd<bfloat16, M_one * N_pvc> R_bf = 0;
+    simd<float, M_one * N_pvc> R_f = 0;
 
-    simd<bfloat16, M_one *N_pvc> C_bf = 0;
-    simd<float, M_one *N_pvc> C_f = 0;
+    simd<bfloat16, M_one * N_pvc> C_bf = 0;
+    simd<float, M_one * N_pvc> C_f = 0;
 
-    simd<bfloat16, K_bf16 *N_pvc> B_bf = 0;
-    simd<bfloat16, M_one *K_bf16> A_bf = 0;
+    simd<bfloat16, K_bf16 * N_pvc> B_bf = 0;
+    simd<bfloat16, M_one * K_bf16> A_bf = 0;
 
     R_f = xmx::dpas<8, 1, float>(C_f, B_bf, A_bf);
     zoo(R_f);
@@ -91,14 +91,14 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void xmx_func() {
   }
 
   { // ======= DPAS FP16 =======================================================
-    simd<half, M_one *N_pvc> R_hf = 0;
-    simd<float, M_one *N_pvc> R_f = 0;
+    simd<half, M_one * N_pvc> R_hf = 0;
+    simd<float, M_one * N_pvc> R_f = 0;
 
-    simd<half, M_one *N_pvc> C_hf = 0;
-    simd<float, M_one *N_pvc> C_f = 0;
+    simd<half, M_one * N_pvc> C_hf = 0;
+    simd<float, M_one * N_pvc> C_f = 0;
 
-    simd<half, K_half *N_pvc> B_hf = 0;
-    simd<half, M_one *K_half> A_hf = 0;
+    simd<half, K_half * N_pvc> B_hf = 0;
+    simd<half, M_one * K_half> A_hf = 0;
 
     // ------------------- FP16: WITH ACC OPERAND -----------------------
     R_f = xmx::dpas<8, 1, float>(C_f, B_hf, A_hf);
@@ -128,10 +128,10 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void xmx_func() {
   }
 
   { // ======= DPAS 8-BIT x 2-BIT INT ==========================================
-    simd<int, M_one *N_pvc> R_d = 0;
-    simd<int, M_one *N_pvc> C_d = 0;
-    simd<int, K_int8x2 *N_pvc / 16> B_int2 = 0; // 16 2-bit integers per int32
-    simd<signed char, M_one *K_int8x2> A_int8 = 0;
+    simd<int, M_one * N_pvc> R_d = 0;
+    simd<int, M_one * N_pvc> C_d = 0;
+    simd<int, K_int8x2 * N_pvc / 16> B_int2 = 0; // 16 2-bit integers per int32
+    simd<signed char, M_one * K_int8x2> A_int8 = 0;
 
     // ------------ DPAS s8 x s2: WITH THE ACCUMULATOR OPERAND -----------------
     R_d = xmx::dpas<8, 1, int, int, int, signed char, s2, s8>(C_d, B_int2,
@@ -146,11 +146,11 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void xmx_func() {
   }
 
   { // ======= DPASW BF16 ======================================================
-    simd<float, M_one *N_dg2> R_f = 0;
-    simd<float, M_one *N_dg2> C_f = 0;
+    simd<float, M_one * N_dg2> R_f = 0;
+    simd<float, M_one * N_dg2> C_f = 0;
 
-    simd<bfloat16, K_bf16 *N_dg2> B_bf = 0;
-    simd<bfloat16, M_one *K_bf16 / 2> A_bf = 0;
+    simd<bfloat16, K_bf16 * N_dg2> B_bf = 0;
+    simd<bfloat16, M_one * K_bf16 / 2> A_bf = 0;
 
     // ------------ DPASW BF16: WITH THE ACCUMULATOR OPERAND -------------------
     R_f = xmx::dpasw<8, 1, float>(C_f, B_bf, A_bf);
@@ -164,8 +164,8 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void xmx_func() {
   }
 
   { // ======= DPASW FP16 ======================================================
-    simd<float, M_one *N_dg2> R_f = 0;
-    simd<float, M_one *N_dg2> C_f = 0;
+    simd<float, M_one * N_dg2> R_f = 0;
+    simd<float, M_one * N_dg2> C_f = 0;
 
     simd<half, K_half * N_dg2> B_hf = 0;
     simd<half, M_one * K_half / 2> A_hf = 0;
@@ -182,12 +182,12 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void xmx_func() {
   }
 
   { // ======= DPAS TFLOAT32 ===================================================
-    simd<float, M_one *N_pvc> R_f = 0;
-    simd<float, M_one *N_pvc> C_f = 0;
+    simd<float, M_one * N_pvc> R_f = 0;
+    simd<float, M_one * N_pvc> C_f = 0;
 
-    simd<sycl::ext::intel::experimental::esimd::tfloat32, K_tf32 *N_pvc> B_tf =
+    simd<sycl::ext::intel::experimental::esimd::tfloat32, K_tf32 * N_pvc> B_tf =
         0;
-    simd<sycl::ext::intel::experimental::esimd::tfloat32, M_one *K_tf32> A_tf =
+    simd<sycl::ext::intel::experimental::esimd::tfloat32, M_one * K_tf32> A_tf =
         0;
 
     // ------------------- TFLOAT32: WITH ACC OPERAND --------------------------

sycl/test/esimd/dae.cpp renamed to sycl/test/check_device_code/esimd/dae.cpp

@@ -12,8 +12,7 @@ using namespace sycl::ext::intel::esimd;
 // Wrapper for designating a scalar region of code that will be
 // vectorized by the backend compiler.
 #define SIMT_BEGIN(N, lane)                                                    \
-  [&]() SYCL_ESIMD_FUNCTION ESIMD_NOINLINE                                     \
-      [[intel::sycl_esimd_vectorize(N)]] {                                     \
+  [&]() SYCL_ESIMD_FUNCTION ESIMD_NOINLINE [[intel::sycl_esimd_vectorize(N)]] {                                     \
     int lane = __esimd_lane_id();
 #define SIMT_END                                                               \
   }                                                                            \

sycl/test/esimd/dpas.cpp renamed to sycl/test/check_device_code/esimd/dpas.cpp

@@ -0,0 +1,3 @@
+import platform
+
+config.substitutions.append(("%clang_O0", "-O0 -mllvm -esimd-allow-optnone-noinline"))

sycl/test/esimd/fp16_converts.cpp renamed to sycl/test/check_device_code/esimd/fp16_converts.cpp

@@ -0,0 +1,15 @@
+// RUN: %clangxx -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o - 2>&1 | FileCheck %s
+
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+using namespace sycl::ext::intel::esimd;
+using namespace sycl::ext::intel::experimental::esimd;
+
+SYCL_ESIMD_KERNEL SYCL_EXTERNAL void kernel_esimd() {
+  __ESIMD_NS::named_barrier_init<7>();
+  __ESIMD_NS::named_barrier_wait(2);
+  // CHECK: call spir_func void @_Z13__esimd_fenceh(i8 noundef zeroext 33)
+  // CHECK-NEXT: call spir_func void @_Z23__esimd_nbarrier_arrive{{.*}}
+  __ESIMD_NS::named_barrier_signal(0, 0, 4, 4);
+}

sycl/test/esimd/genx_func_attr.cpp renamed to sycl/test/check_device_code/esimd/genx_func_attr.cpp

@@ -0,0 +1,41 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm -x c++ %s -o %t
+// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table
+// RUN: FileCheck %s -input-file=%t_esimd_0.ll
+
+// This test checks that all LLVM-IR instructions that work with SPIR-V builtins
+// are correctly translated into GenX counterparts (implemented in
+// LowerESIMD.cpp)
+
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+SYCL_ESIMD_KERNEL SYCL_EXTERNAL void
+kernel_SubgroupLocalInvocationId(size_t *DoNotOptimize,
+                                 uint32_t *DoNotOptimize32) {
+  DoNotOptimize[0] = __spirv_SubgroupLocalInvocationId();
+  DoNotOptimize32[0] = __spirv_SubgroupLocalInvocationId() + 3;
+  // CHECK-LABEL: @{{.*}}kernel_SubgroupLocalInvocationId
+  // CHECK: [[ZEXT0:%.*]] = zext i32 0 to i64
+  // CHECK: store i64 [[ZEXT0]]
+  // CHECK: add i32 0, 3
+}
+
+SYCL_ESIMD_KERNEL SYCL_EXTERNAL void
+kernel_SubgroupSize(size_t *DoNotOptimize, uint32_t *DoNotOptimize32) {
+  DoNotOptimize[0] = __spirv_SubgroupSize();
+  DoNotOptimize32[0] = __spirv_SubgroupSize() + 7;
+  // CHECK-LABEL: @{{.*}}kernel_SubgroupSize
+  // CHECK: [[ZEXT0:%.*]] = zext i32 1 to i64
+  // CHECK: store i64 [[ZEXT0]]
+  // CHECK: add i32 1, 7
+}
+
+SYCL_ESIMD_KERNEL SYCL_EXTERNAL void
+kernel_SubgroupMaxSize(size_t *DoNotOptimize, uint32_t *DoNotOptimize32) {
+  DoNotOptimize[0] = __spirv_SubgroupMaxSize();
+  DoNotOptimize32[0] = __spirv_SubgroupMaxSize() + 9;
+  // CHECK-LABEL: @{{.*}}kernel_SubgroupMaxSize
+  // CHECK: [[ZEXT0:%.*]] = zext i32 1 to i64
+  // CHECK: store i64 [[ZEXT0]]
+  // CHECK: add i32 1, 9
+}