[ESIMD] Split the test esimd::wait() to e2e and compile-only (#8821)

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Author: v-klochkov

sycl/test-e2e/ESIMD/wait.cpp

@@ -8,10 +8,6 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
 
-// RUN: %clangxx -fsycl -fsycl-device-only -Xclang -emit-llvm -o %t.comp.ll %s
-// RUN: sycl-post-link -ir-output-only -lower-esimd -S %t.comp.ll -o %t.out.ll
-// RUN: FileCheck --input-file=%t.out.ll %s
-
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
@@ -37,8 +33,6 @@ bool test(sycl::queue Q, int IArg = 128) {
        simd<int, 16> A = IArg;
        simd<int, 16> B = A * A;
        iesimd::wait(B);
-       // CHECK: mul <16 x i32>
-       // CHECK: llvm.genx.dummy.mov
      }).wait();
 
     // Test case 2: check wait() with esimd::simd_view argument.
@@ -48,9 +42,6 @@ bool test(sycl::queue Q, int IArg = 128) {
        auto BView = B.select<8, 2>(0);
        BView += 2;
        iesimd::wait(BView);
-       // CHECK: mul <16 x i32>
-       // CHECK: add <8 x i32>
-       // CHECK: llvm.genx.dummy.mov
      }).wait();
 
     // Test case 3: check wait() that prevesrves one simd and lets
@@ -60,10 +51,6 @@ bool test(sycl::queue Q, int IArg = 128) {
        auto B = A * 17;
        iesimd::wait(B);
        auto C = B * 17;
-       // CHECK: mul <8 x i64>
-       // CHECK-NOT: add <8 x i64>
-       // CHECK: llvm.genx.dummy.mov
-       // CHECK-NEXT: ret void
      }).wait();
   } catch (sycl::exception const &e) {
     std::cout << "SYCL exception caught: " << e.what() << '\n';

sycl/test/esimd/wait.cpp

@@ -0,0 +1,51 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -Xclang -emit-llvm -o %t.comp.ll %s
+// RUN: sycl-post-link -ir-output-only -lower-esimd -S %t.comp.ll -o %t.out.ll
+// RUN: FileCheck --input-file=%t.out.ll %s
+
+// This test verifies that calls of ext::intel::experimental::esimd::wait()
+// are lowered properly, not deleted as unused and they do not let delete
+// the argument of wait() function.
+
+#include <iostream>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::intel::esimd;
+namespace iesimd = sycl::ext::intel::experimental::esimd;
+
+SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void func(int IArg) {
+  // Test case 1: check wait() with esimd::simd argument.
+  {
+    simd<int, 16> A = IArg;
+    simd<int, 16> B = A * A;
+    iesimd::wait(B);
+    // CHECK: mul <16 x i32>
+    // CHECK: llvm.genx.dummy.mov
+  }
+
+  // Test case 2: check wait() with esimd::simd_view argument.
+  {
+    simd<int, 16> A = IArg;
+    simd<int, 16> B = A * 17;
+    auto BView = B.select<8, 2>(0);
+    BView += 2;
+    iesimd::wait(BView);
+    // CHECK: mul <16 x i32>
+    // CHECK: add <8 x i32>
+    // CHECK: llvm.genx.dummy.mov
+  }
+
+  // Test case 3: check wait() that prevesrves one simd and lets
+  // optimize away the other/unused one.
+  {
+    simd<uint64_t, 8> A = IArg;
+    auto B = A * 17;
+    iesimd::wait(B);
+    auto C = B * 17;
+    // CHECK: mul <8 x i64>
+    // CHECK-NOT: add <8 x i64>
+    // CHECK: llvm.genx.dummy.mov
+    // CHECK-NEXT: ret void
+  }
+}