[SYCL] Correction to recognition of constant-size array. (#2082)

Signed-off-by: rdeodhar <[email protected]>

Author: rdeodhar

clang/lib/Sema/SemaSYCL.cpp

@@ -802,7 +802,9 @@ class KernelObjVisitor {
   template <typename... Handlers>
   void VisitArrayElements(FieldDecl *FD, QualType FieldTy,
                           Handlers &... handlers) {
-    const ConstantArrayType *CAT = cast<ConstantArrayType>(FieldTy);
+    const ConstantArrayType *CAT =
+        SemaRef.getASTContext().getAsConstantArrayType(FieldTy);
+    assert(CAT && "Should only be called on constant-size array.");
     QualType ET = CAT->getElementType();
     int64_t ElemCount = CAT->getSize().getSExtValue();
     std::initializer_list<int>{(handlers.enterArray(), 0)...};
@@ -1014,7 +1016,8 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   // Return true if not copyable, false if copyable.
   bool checkNotCopyableToKernel(const FieldDecl *FD, const QualType &FieldTy) {
     if (FieldTy->isArrayType()) {
-      if (const auto *CAT = dyn_cast<ConstantArrayType>(FieldTy)) {
+      if (const auto *CAT =
+              SemaRef.getASTContext().getAsConstantArrayType(FieldTy)) {
         QualType ET = CAT->getElementType();
         return checkNotCopyableToKernel(FD, ET);
       }

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

@@ -12,6 +12,11 @@ __attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
   kernelFunc();
 }
 
+template <typename T>
+struct S {
+  T a[3];
+};
+
 int main() {
 
   using Accessor =
@@ -22,6 +27,7 @@ int main() {
   struct struct_acc_t {
     Accessor member_acc[2];
   } struct_acc;
+  S<int> s;
 
   struct foo_inner {
     int foo_inner_x;
@@ -56,6 +62,11 @@ int main() {
       [=]() {
         foo local = struct_array[1];
       });
+
+  a_kernel<class kernel_E>(
+      [=]() {
+        int local = s.a[2];
+      });
 }
 
 // Check kernel_A parameters
@@ -100,8 +111,8 @@ int main() {
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
-// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
+// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp:57:7)' cinit
+// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:57:7)'
 // CHECK-NEXT: InitListExpr {{.*}} 'struct_acc_t'
 // CHECK-NEXT: InitListExpr {{.*}} 'Accessor [2]'
 // CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor [2]'
@@ -201,3 +212,21 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
 // CHECK-NEXT: ImplicitCastExpr
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'
+
+// Check kernel_E parameters
+// CHECK: FunctionDecl {{.*}}kernel_E{{.*}} 'void (int, int, int)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp:67:7)' cinit
+// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:67:7)'
+// CHECK-NEXT: InitListExpr {{.*}} 'S<int>'
+// CHECK-NEXT: InitListExpr {{.*}} 'int [3]'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int'
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int'
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int'
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int'

sycl/test/array_param/array-kernel-param-nested-run.cpp

@@ -11,11 +11,11 @@
 
 using namespace cl::sycl;
 
-constexpr size_t c_num_items = 100;
+constexpr size_t c_num_items = 10;
 range<1> num_items{c_num_items}; // range<1>(num_items)
 
 // Change if tests are added/removed
-static int testCount = 1;
+static int testCount = 2;
 static int passCount;
 
 template <typename T>
@@ -84,6 +84,36 @@ bool test_accessor_array_in_struct(queue &myQueue) {
   return verify_1D("Accessor array in struct", c_num_items, output, ref);
 }
 
+template <typename T>
+struct S {
+  T a[c_num_items];
+};
+bool test_templated_array_in_struct(queue &myQueue) {
+  std::array<int, c_num_items> output;
+  std::array<int, c_num_items> ref;
+  init(ref, 3, 3);
+
+  auto out_buffer = buffer<int, 1>(output.data(), num_items);
+
+  S<int> sint;
+  S<long long> sll;
+  init(sint.a, 1, 1);
+  init(sll.a, 2, 2);
+
+  myQueue.submit([&](handler &cgh) {
+    using Accessor =
+        accessor<int, 1, access::mode::read_write, access::target::global_buffer>;
+    auto output_accessor = out_buffer.get_access<access::mode::write>(cgh);
+
+    cgh.parallel_for<class templated_array_in_struct>(num_items, [=](cl::sycl::id<1> index) {
+      output_accessor[index] = sint.a[index] + sll.a[index];
+    });
+  });
+  const auto HostAccessor = out_buffer.get_access<cl::sycl::access::mode::read>();
+
+  return verify_1D("Templated array in struct", c_num_items, output, ref);
+}
+
 bool run_tests() {
   queue Q([](exception_list L) {
     for (auto ep : L) {
@@ -103,6 +133,9 @@ bool run_tests() {
   if (test_accessor_array_in_struct(Q)) {
     ++passCount;
   }
+  if (test_templated_array_in_struct(Q)) {
+    ++passCount;
+  }
 
   auto D = Q.get_device();
   const char *devType = D.is_host() ? "Host" : D.is_cpu() ? "CPU" : "GPU";