[SYCL][NFC] Update FE tests to have a common infrastructure

clang/test/SemaSYCL/accessors-targets-image.cpp

@@ -1,66 +1,81 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump -sycl-std=2020 %s | FileCheck %s
 
-// This test checks that compiler generates correct kernel wrapper arguments for
+// This test checks if the compiler generates correct kernel wrapper arguments for
 // image accessors targets.
 
 #include "Inputs/sycl.hpp"
 
-using namespace cl::sycl;
-
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
-  kernelFunc();
-}
+cl::sycl::queue q;
 
 int main() {
 
-  accessor<int, 1, access::mode::read,
-           access::target::image, access::placeholder::false_t>
+  // 1-dimensional accessor with Read-only access
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read,
+                     cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t>
       image_acc1d_read;
-  kernel<class use_image1d_r>(
-      [=]() {
-        image_acc1d_read.use();
-      });
 
-  accessor<int, 2, access::mode::read,
-           access::target::image, access::placeholder::false_t>
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class use_image1d_r>(
+        [=] {
+          image_acc1d_read.use();
+        });
+  });
+
+  // 2-dimensional accessor with Read-only access
+  cl::sycl::accessor<int, 2, cl::sycl::access::mode::read,
+                     cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t>
       image_acc2d_read;
-  kernel<class use_image2d_r>(
-      [=]() {
-        image_acc2d_read.use();
-      });
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class use_image2d_r>(
+        [=] {
+          image_acc2d_read.use();
+        });
+  });
 
-  accessor<int, 3, access::mode::read,
-           access::target::image, access::placeholder::false_t>
+  // 3-dimensional accessor with Read-only access
+  cl::sycl::accessor<int, 3, cl::sycl::access::mode::read,
+                     cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t>
       image_acc3d_read;
-  kernel<class use_image3d_r>(
-      [=]() {
-        image_acc3d_read.use();
-      });
 
-  accessor<int, 1, access::mode::write,
-           access::target::image, access::placeholder::false_t>
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class use_image3d_r>(
+        [=] {
+          image_acc3d_read.use();
+        });
+  });
+
+  // 1-dimensional accessor with Write-only access
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::write,
+                     cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t>
       image_acc1d_write;
-  kernel<class use_image1d_w>(
-      [=]() {
-        image_acc1d_write.use();
-      });
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class use_image1d_w>(
+        [=] {
+          image_acc1d_write.use();
+        });
+  });
 
-  accessor<int, 2, access::mode::write,
-           access::target::image, access::placeholder::false_t>
+  // 2-dimensional accessor with Write-only access
+  cl::sycl::accessor<int, 2, cl::sycl::access::mode::write,
+                     cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t>
       image_acc2d_write;
-  kernel<class use_image2d_w>(
-      [=]() {
-        image_acc2d_write.use();
-      });
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class use_image2d_w>(
+        [=] {
+          image_acc2d_write.use();
+        });
+  });
 
-  accessor<int, 3, access::mode::write,
-           access::target::image, access::placeholder::false_t>
+  // 3-dimensional accessor with Write-only access
+  cl::sycl::accessor<int, 3, cl::sycl::access::mode::write,
+                     cl::sycl::access::target::image, cl::sycl::access::placeholder::false_t>
       image_acc3d_write;
-  kernel<class use_image3d_w>(
-      [=]() {
-        image_acc3d_write.use();
-      });
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class use_image3d_w>(
+        [=] {
+          image_acc3d_write.use();
+        });
+  });
 }
 
 // CHECK: {{.*}}use_image1d_r 'void (__read_only image1d_t)'

clang/test/SemaSYCL/accessors-targets.cpp

@@ -1,40 +1,46 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump -sycl-std=2020 %s | FileCheck %s
 
-// This test checks that compiler generates correct kernel wrapper arguments for
+// This test checks that the compiler generates correct kernel wrapper arguments for
 // different accessors targets.
 
 #include "Inputs/sycl.hpp"
 
-using namespace cl::sycl;
-
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
-  kernelFunc();
-}
+cl::sycl::queue q;
 
 int main() {
-
-  accessor<int, 1, access::mode::read_write,
-           access::target::local>
+  // Access work-group local memory with read and write access.
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                     cl::sycl::access::target::local>
       local_acc;
-  accessor<int, 1, access::mode::read_write,
-           access::target::global_buffer>
+  // Access buffer via global memory with read and write access.
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                     cl::sycl::access::target::global_buffer>
       global_acc;
-  accessor<int, 1, access::mode::read_write,
-           access::target::constant_buffer>
+  // Access buffer via constant memory with read and write access.
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                     cl::sycl::access::target::constant_buffer>
       constant_acc;
-  kernel<class use_local>(
-      [=]() {
-        local_acc.use();
-      });
-  kernel<class use_global>(
-      [=]() {
-        global_acc.use();
-      });
-  kernel<class use_constant>(
-      [=]() {
-        constant_acc.use();
-      });
+
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class use_local>(
+        [=] {
+          local_acc.use();
+        });
+  });
+
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class use_global>(
+        [=] {
+          global_acc.use();
+        });
+  });
+
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class use_constant>(
+        [=] {
+          constant_acc.use();
+        });
+  });
 }
 // CHECK: {{.*}}use_local{{.*}} 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
 // CHECK: {{.*}}use_global{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'

clang/test/SemaSYCL/allow-constexpr-recursion.cpp

@@ -1,9 +1,10 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-sycl-2017-compat -verify -fsyntax-only -std=c++20 -Werror=vla %s
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  kernelFunc();
-}
+// This test verifies that a SYCL kernel executed on a device, cannot call a recursive function.
+
+#include "Inputs/sycl.hpp"
+
+cl::sycl::queue q;
 
 // expected-note@+1{{function implemented using recursion declared here}}
 constexpr int constexpr_recurse1(int n);
@@ -71,6 +72,11 @@ void constexpr_recurse_test_err() {
 }
 
 int main() {
-  kernel_single_task<class fake_kernel>([]() { constexpr_recurse_test(); });
-  kernel_single_task<class fake_kernel>([]() { constexpr_recurse_test_err(); });
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class fake_kernel>([]() { constexpr_recurse_test(); });
+  });
+
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class fake_kernel>([]() { constexpr_recurse_test_err(); });
+  });
 }

clang/test/SemaSYCL/array-kernel-param-neg.cpp

@@ -1,43 +1,49 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-sycl-2017-compat -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -sycl-std=2020 -verify -fsyntax-only %s
 
 // This test checks if compiler reports compilation error on an attempt to pass
 // an array of non-trivially copyable structs as SYCL kernel parameter or
 // a non-constant size array.
 
-struct B {
+#include "Inputs/sycl.hpp"
+
+cl::sycl::queue q;
+
+struct NonTrivialCopyStruct {
   int i;
-  B(int _i) : i(_i) {}
-  B(const B &x) : i(x.i) {}
+  NonTrivialCopyStruct(int _i) : i(_i) {}
+  NonTrivialCopyStruct(const NonTrivialCopyStruct &x) : i(x.i) {}
 };
 
-struct D {
+struct NonTrivialDestructorStruct {
   int i;
-  ~D();
+  ~NonTrivialDestructorStruct();
 };
 
-class E {
+class Array {
   // expected-error@+1 {{kernel parameter is not a constant size array}}
-  int i[];
+  int NonConstantSizeArray[];
 
 public:
-  int operator()() const { return i[0]; }
+  int operator()() const { return NonConstantSizeArray[0]; }
 };
 
-template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  kernelFunc();
-}
-
 void test() {
-  B nsl1[4] = {1, 2, 3, 4};
-  D nsl2[5];
-  E es;
-  kernel_single_task<class kernel_capture_refs>([=] {
-    // expected-error@+1 {{kernel parameter has non-trivially copy constructible class/struct type}}
-    int b = nsl1[2].i;
-    // expected-error@+1 {{kernel parameter has non-trivially destructible class/struct type}}
-    int d = nsl2[4].i;
+  NonTrivialCopyStruct NTCSObject[4] = {1, 2, 3, 4};
+  NonTrivialDestructorStruct NTDSObject[5];
+  // expected-note@+1 {{'UnknownSizeArrayObj' declared here}}
+  Array UnknownSizeArrayObj;
+
+  q.submit([&](cl::sycl::handler &h) {
+    h.single_task<class kernel_capture_refs>([=] {
+      // expected-error@+1 {{kernel parameter has non-trivially copy constructible class/struct type}}
+      int b = NTCSObject[2].i;
+      // expected-error@+1 {{kernel parameter has non-trivially destructible class/struct type}}
+      int d = NTDSObject[4].i;
+    });
   });
 
-  kernel_single_task<class kernel_bad_array>(es);
+  q.submit([&](cl::sycl::handler &h) {
+    // expected-error@+1 {{variable 'UnknownSizeArrayObj' with flexible array member cannot be captured in a lambda expression}}
+    h.single_task<class kernel_bad_array>(UnknownSizeArrayObj);
+  });
 }