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

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -3,7 +3,7 @@
 
 // Shared code for SYCL tests
 
-namespace cl {
+inline namespace cl {
 namespace sycl {
 namespace access {
 

clang/test/SemaSYCL/accessor_inheritance.cpp

@@ -1,65 +1,75 @@
-// 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 inheritance support for struct types with accessors
+// passed as kernel arguments, which are decomposed to individual fields.
+
 #include "Inputs/sycl.hpp"
 
-struct Base {
+sycl::queue myQueue;
+
+struct AccessorBase {
   int A, B;
-  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField;
+  sycl::accessor<char, 1, sycl::access::mode::read> AccField;
 };
 
-struct Captured : Base,
-                  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> {
+struct AccessorDerived : AccessorBase,
+                         sycl::accessor<char, 1, sycl::access::mode::read> {
   int C;
 };
 
 int main() {
-  Captured Obj;
-  cl::sycl::kernel_single_task<class kernel>(
-      [=]() {
-        Obj.use();
-      });
+  AccessorDerived DerivedObject;
+  myQueue.submit([&](sycl::handler &h) {
+    h.single_task<class kernel>(
+        [=] {
+          DerivedObject.use();
+        });
+  });
+
+  return 0;
 }
 
 // Check kernel parameters
-// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (int, int, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int)'
+// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (int, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int)'
 // CHECK: ParmVarDecl{{.*}} used _arg_A 'int'
 // CHECK: ParmVarDecl{{.*}} used _arg_B 'int'
 // CHECK: ParmVarDecl{{.*}} used _arg_AccField '__global char *'
-// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'cl::sycl::range<1>'
-// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'cl::sycl::range<1>'
-// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'cl::sycl::id<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'sycl::range<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'sycl::range<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'sycl::id<1>'
 // CHECK: ParmVarDecl{{.*}} used _arg__base '__global char *'
-// CHECK: ParmVarDecl{{.*}} used _arg__base 'cl::sycl::range<1>'
-// CHECK: ParmVarDecl{{.*}} used _arg__base 'cl::sycl::range<1>'
-// CHECK: ParmVarDecl{{.*}} used _arg__base 'cl::sycl::id<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg__base 'sycl::range<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg__base 'sycl::range<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg__base 'sycl::id<1>'
 // CHECK: ParmVarDecl{{.*}} used _arg_C 'int'
 
 // Check lambda initialization
 // CHECK: VarDecl {{.*}} used '(lambda at {{.*}}accessor_inheritance.cpp
 // CHECK-NEXT: InitListExpr {{.*}}
-// CHECK-NEXT: InitListExpr {{.*}} 'Captured'
-// CHECK-NEXT: InitListExpr {{.*}} 'Base'
+// CHECK-NEXT: InitListExpr {{.*}} 'AccessorDerived'
+// CHECK-NEXT: InitListExpr {{.*}} 'AccessorBase'
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_A' 'int'
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_B' 'int'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read>':'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read>':'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_C' 'int'
 
 // Check __init calls
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} .__init
 // CHECK-NEXT: MemberExpr {{.*}} .AccField
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'Base' lvalue <DerivedToBase (Base)>
-// CHECK-NEXT: MemberExpr {{.*}} 'Captured' lvalue .
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'AccessorBase' lvalue <DerivedToBase (AccessorBase)>
+// CHECK-NEXT: MemberExpr {{.*}} 'AccessorDerived' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}}'(lambda at {{.*}}accessor_inheritance.cpp
 // CHECK-NEXT: ImplicitCastExpr {{.*}} '__global char *' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg_AccField' '__global char *'
 
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr{{.*}} lvalue .__init
-// CHECK-NEXT: MemberExpr{{.*}}'Captured' lvalue .
+// CHECK-NEXT: MemberExpr{{.*}}'AccessorDerived' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}accessor_inheritance.cpp
 // CHECK-NEXT: ImplicitCastExpr {{.*}} '__global char *' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg__base' '__global char *'

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();
-}
+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
+  sycl::accessor<int, 1, sycl::access::mode::read,
+                 sycl::access::target::image, 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([&](sycl::handler &h) {
+    h.single_task<class use_image1d_r>(
+        [=] {
+          image_acc1d_read.use();
+        });
+  });
+
+  // 2-dimensional accessor with Read-only access
+  sycl::accessor<int, 2, sycl::access::mode::read,
+                 sycl::access::target::image, sycl::access::placeholder::false_t>
       image_acc2d_read;
-  kernel<class use_image2d_r>(
-      [=]() {
-        image_acc2d_read.use();
-      });
+  q.submit([&](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
+  sycl::accessor<int, 3, sycl::access::mode::read,
+                 sycl::access::target::image, 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([&](sycl::handler &h) {
+    h.single_task<class use_image3d_r>(
+        [=] {
+          image_acc3d_read.use();
+        });
+  });
+
+  // 1-dimensional accessor with Write-only access
+  sycl::accessor<int, 1, sycl::access::mode::write,
+                 sycl::access::target::image, sycl::access::placeholder::false_t>
       image_acc1d_write;
-  kernel<class use_image1d_w>(
-      [=]() {
-        image_acc1d_write.use();
-      });
+  q.submit([&](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
+  sycl::accessor<int, 2, sycl::access::mode::write,
+                 sycl::access::target::image, sycl::access::placeholder::false_t>
       image_acc2d_write;
-  kernel<class use_image2d_w>(
-      [=]() {
-        image_acc2d_write.use();
-      });
+  q.submit([&](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
+  sycl::accessor<int, 3, sycl::access::mode::write,
+                 sycl::access::target::image, sycl::access::placeholder::false_t>
       image_acc3d_write;
-  kernel<class use_image3d_w>(
-      [=]() {
-        image_acc3d_write.use();
-      });
+  q.submit([&](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,41 +1,47 @@
-// 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();
-}
+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.
+  sycl::accessor<int, 1, sycl::access::mode::read_write,
+                 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.
+  sycl::accessor<int, 1, sycl::access::mode::read_write,
+                 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.
+  sycl::accessor<int, 1, sycl::access::mode::read_write,
+                 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([&](sycl::handler &h) {
+    h.single_task<class use_local>(
+        [=] {
+          local_acc.use();
+        });
+  });
+
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class use_global>(
+        [=] {
+          global_acc.use();
+        });
+  });
+
+  q.submit([&](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>)'
-// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
+// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
+// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, sycl::range<1>, sycl::range<1>, 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"
+
+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([&](sycl::handler &h) {
+    h.single_task<class fake_kernel>([]() { constexpr_recurse_test(); });
+  });
+
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class fake_kernel>([]() { constexpr_recurse_test_err(); });
+  });
 }