Fixed sycl/test/esimd/*

Author: DenisBakhvalov

sycl/test/esimd/glob.cpp

@@ -1,5 +1,6 @@
-// RUN: %clangxx -fsycl -fsycl-explicit-simd -c -fsycl-device-only -Xclang -emit-llvm %s -o - | \
-// RUN:  FileCheck %s
+// RUN: %clangxx -fsycl -fsycl-explicit-simd -c -fsycl-device-only -Xclang -emit-llvm %s -o %t
+// RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %t -o %t.table
+// RUN: FileCheck %s -input-file=%t_esimd_0.ll
 
 // This test checks that globals with register attribute are allowed in ESIMD
 // mode, can be accessed in functions and correct LLVM IR is generated
@@ -9,7 +10,6 @@
 #include <CL/sycl/INTEL/esimd.hpp>
 #include <iostream>
 
-using namespace cl::sycl;
 using namespace sycl::INTEL::gpu;
 
 constexpr unsigned VL = 16;
@@ -22,7 +22,16 @@ ESIMD_PRIVATE ESIMD_REGISTER(17 + VL) simd<int, VL> vc1;
 // CHECK-DAG: @vc1 = {{.+}} <16 x i32> zeroinitializer, align 64 #1
 // CHECK-DAG: attributes #1 = { {{.*}}"VCByteOffset"="33" "VCGlobalVariable" "VCVolatile"{{.*}} }
 
-SYCL_EXTERNAL ESIMD_NOINLINE void init_vc(int x) {
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+  kernelFunc();
+}
+
+SYCL_ESIMD_FUNCTION SYCL_EXTERNAL ESIMD_NOINLINE void init_vc(int x) {
   vc1 = vc + 1;
   vc = x;
 }
+
+void caller(int x) {
+  kernel<class kernel_esimd>([=]() SYCL_ESIMD_KERNEL { init_vc(x); });
+}

sycl/test/esimd/intrins_trans.cpp

@@ -1,5 +1,6 @@
-// RUN: %clangxx -O0 -fsycl -fsycl-explicit-simd -fsycl-device-only -Xclang -emit-llvm %s -o - | \
-// RUN:   FileCheck %s
+// RUN: %clangxx -O0 -fsycl -fsycl-explicit-simd -fsycl-device-only -Xclang -emit-llvm %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
 
 // Checks ESIMD intrinsic translation.
 // NOTE: must be run in -O0, as optimizer optimizes away some of the code
@@ -109,7 +110,7 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<float, 16> foo() {
   // CHECK: call void @llvm.genx.media.st.v32i32(i32 0, i32 %[[SI2]], i32 0, i32 32, i32 %{{[0-9a-zA-Z_.]+}}, i32 %{{[0-9a-zA-Z_.]+}}, <32 x i32> %{{[0-9a-zA-Z_.]+}})
 
   auto ee = __esimd_vload<int, 16>((vector_type_t<int, 16> *)(&vg));
-  // CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.vload.v16i32.p4v16i32(<16 x i32> addrspace(4)* {{.*}})
+  // CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.vload.v16i32.p0v16i32(<16 x i32>* {{.*}})
   __esimd_vstore<int, 32>(&vc, va.data());
   // CHECK: store <32 x i32>  %{{[0-9a-zA-Z_.]+}}, <32 x i32> addrspace(4)* {{.*}}
 

sycl/test/esimd/spirv_intrins_trans.cpp

@@ -1,6 +1,9 @@
-// RUN: %clangxx -fsycl -fsycl-explicit-simd -fsycl-device-only -O0 -S -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clangxx -fsycl -fsycl-explicit-simd -fsycl-device-only -O0 -S -emit-llvm -x c++ %s -o %t
+// RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %t -o %t.table
+// RUN: FileCheck %s -input-file=%t_esimd_0.ll
+
 // This test checks that all SPIR-V intrinsics are correctly
-// translated into GenX counterparts (implemented in LowerCM.cpp)
+// translated into GenX counterparts (implemented in LowerESIMD.cpp)
 
 #include <CL/sycl.hpp>
 #include <CL/sycl/INTEL/esimd.hpp>