[SYCL][Test] Do not include sycl.hpp in clang LIT test (#900)

Signed-off-by: Andrew Savonichev <[email protected]>

Author: Andrew Savonichev

clang/test/CodeGenSYCL/intel-restrict.cpp

@@ -1,68 +1,28 @@
 // RUN: %clang %s -S -emit-llvm -fsycl-device-only -o - | FileCheck %s
 
-#include "CL/sycl.hpp"
-
-constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
-constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;
-
-template <typename Acc1Ty, typename Acc2Ty>
-struct foostr {
-  Acc1Ty A;
-  Acc2Ty B;
-  foostr(Acc1Ty A, Acc2Ty B): A(A), B(B) {}
-  [[intel::kernel_args_restrict]]
-  void operator()() {
-    A[0] = B[0];
-  }
-};
-
-int foo(int X) {
-  int A[] = { 42 };
-  int B[] = { 0 };
-  {
-    cl::sycl::queue Q;
-    cl::sycl::buffer<int, 1> BufA(A, 1);
-    cl::sycl::buffer<int, 1> BufB(B, 1);
-
-    // CHECK: define {{.*}} spir_kernel {{.*}}kernel_norestrict{{.*}}(i32 addrspace(1)* %{{.*}} i32 addrspace(1)* %{{.*}}
-
-    Q.submit([&](cl::sycl::handler& cgh) {
-      auto AccA = BufA.get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      auto AccB = BufB.get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      cgh.single_task<class kernel_norestrict>(
-          [=]() {
-            AccB[0] = AccA[0];
-          });
-    });
-
-    // CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}}
-    Q.submit([&](cl::sycl::handler& cgh) {
-      auto AccA = BufA.get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      auto AccB = BufB.get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      cgh.single_task<class kernel_restrict>(
-          [=]() [[intel::kernel_args_restrict]] {
-            AccB[0] = AccA[0];
-          });
-    });
-
-    // CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict_struct{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}}
-    Q.submit([&](cl::sycl::handler& cgh) {
-      auto AccA = BufA.get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      auto AccB = BufB.get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      foostr<decltype(AccA), decltype(AccB)> f(AccA, AccB);
-      cgh.single_task<class kernel_restrict_struct>(f);
-    });
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+  kernelFunc();
+}
 
-    // CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict_other_params{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}}, i32 %{{[^,]*}})
-    int num = 42;
-    Q.submit([&](cl::sycl::handler& cgh) {
-      auto AccA = BufA.get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      auto AccB = BufB.get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      cgh.single_task<class kernel_restrict_other_params>(
-          [=]() [[intel::kernel_args_restrict]] {
-            AccB[0] = AccA[0] = num;
-          });
-    });
-  }
-  return B[0];
+int main() {
+  int *a;
+  int *b;
+  int *c;
+  kernel<class kernel_restrict>(
+      [a,b,c]() [[intel::kernel_args_restrict]] { c[0] = a[0] + b[0];});
+// CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}})
+
+  int *d;
+  int *e;
+  int *f;
+
+  kernel<class kernel_norestrict>(
+      [d,e,f]() { f[0] = d[0] + e[0];});
+// CHECK: define {{.*}} spir_kernel {{.*}}kernel_norestrict{{.*}}(i32 addrspace(1)* %{{.*}} i32 addrspace(1)* %{{.*}}, i32 addrspace(1)* %{{.*}})
+
+  int g = 42;
+  kernel<class kernel_restrict_other_types>(
+      [a,b,c,g]() [[intel::kernel_args_restrict]] { c[0] = a[0] + b[0] + g;});
+// CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict_other_types{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}}, i32 %{{.*}})
 }