Decomposed array elements, and changed manner of array element initialization.

Author: rdeodhar

clang/lib/Sema/SemaSYCL.cpp

@@ -791,18 +791,23 @@ static void VisitField(CXXRecordDecl *Owner, RangeTy &&Item, QualType ItemTy,
   else if (ItemTy->isStructureOrClassType())
     VisitAccessorWrapper(Owner, Item, ItemTy->getAsCXXRecordDecl(),
                          handlers...);
-#if 0
-  // FIXME Enable this when structs are replaced by their fields
+    // FIXME Enable this when structs are replaced by their fields
+#define STRUCTS_DECOMPOSED 0
+#if STRUCTS_DECOMPOSED
   else if (ItemTy->isArrayType())
     VisitArrayElements(Item, ItemTy, handlers...);
-#endif
+  else if (ItemTy->isScalarType())
+    KF_FOR_EACH(handleScalarType, Item, ItemTy);
+}
+#else
 }
 
 template <typename RangeTy, typename... Handlers>
 static void VisitScalarField(CXXRecordDecl *Owner, RangeTy &&Item,
                              QualType ItemTy, Handlers &... handlers) {
   KF_FOR_EACH(handleScalarType, Item, ItemTy);
 }
+#endif
 
 template <typename RangeTy, typename... Handlers>
 static void VisitArrayElements(RangeTy Item, QualType FieldTy,
@@ -812,13 +817,18 @@ static void VisitArrayElements(RangeTy Item, QualType FieldTy,
   int64_t ElemCount = CAT->getSize().getSExtValue();
   std::initializer_list<int>{(handlers.enterArray(), 0)...};
   for (int64_t Count = 0; Count < ElemCount; Count++) {
+#if STRUCTS_DECOMPOSED
+    VisitField(nullptr, Item, ET, handlers...);
+#else
     if (ET->isScalarType())
       VisitScalarField(nullptr, Item, ET, handlers...);
     else
       VisitField(nullptr, Item, ET, handlers...);
+#endif
     (void)std::initializer_list<int>{(handlers.nextElement(ET), 0)...};
   }
-  (void)std::initializer_list<int>{(handlers.leaveArray(ET, ElemCount), 0)...};
+  (void)std::initializer_list<int>{
+      (handlers.leaveArray(Item, ET, ElemCount), 0)...};
 }
 
 template <typename RangeTy, typename... Handlers>
@@ -932,20 +942,31 @@ template <typename Derived> class SyclKernelFieldHandler {
   // class/field graph. Int Headers use this to calculate offset, most others
   // don't have a need for these.
 
-  virtual void enterStruct(const CXXRecordDecl *, FieldDecl *) {}
-  virtual void leaveStruct(const CXXRecordDecl *, FieldDecl *) {}
-  virtual void enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {}
-  virtual void leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {}
+  virtual bool enterStruct(const CXXRecordDecl *, FieldDecl *) { return true; }
+  virtual bool leaveStruct(const CXXRecordDecl *, FieldDecl *) { return true; }
+  virtual bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {
+    return true;
+  }
+  virtual bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {
+    return true;
+  }
 
   // The following are used for stepping through array elements.
 
-  virtual void enterField(const CXXRecordDecl *, const CXXBaseSpecifier &) {}
-  virtual void leaveField(const CXXRecordDecl *, const CXXBaseSpecifier &) {}
-  virtual void enterField(const CXXRecordDecl *, FieldDecl *) {}
-  virtual void leaveField(const CXXRecordDecl *, FieldDecl *) {}
-  virtual void enterArray() {}
-  virtual void nextElement(QualType) {}
-  virtual void leaveArray(QualType, int64_t) {}
+  virtual bool enterField(const CXXRecordDecl *, const CXXBaseSpecifier &) {
+    return true;
+  }
+  virtual bool leaveField(const CXXRecordDecl *, const CXXBaseSpecifier &) {
+    return true;
+  }
+  virtual bool enterField(const CXXRecordDecl *, FieldDecl *) { return true; }
+  virtual bool leaveField(const CXXRecordDecl *, FieldDecl *) { return true; }
+  virtual bool enterArray() { return true; }
+  virtual bool nextElement(QualType) { return true; }
+  virtual bool leaveArray(const CXXBaseSpecifier &, QualType, int64_t) {
+    return true;
+  }
+  virtual bool leaveArray(FieldDecl *, QualType, int64_t) { return true; }
 };
 
 // A type to check the validity of all of the argument types.
@@ -1242,6 +1263,7 @@ class SyclKernelBodyCreator
   InitializedEntity VarEntity;
   CXXRecordDecl *KernelObj;
   llvm::SmallVector<Expr *, 16> MemberExprBases;
+  uint64_t ArrayIndex;
   FunctionDecl *KernelCallerFunc;
 
   // Using the statements/init expressions that we've created, this generates
@@ -1340,30 +1362,27 @@ class SyclKernelBodyCreator
     InitExprs.push_back(MemberInit.get());
   }
 
-  void createExprForScalarElement(FieldDecl *FD, QualType FieldTy) {
+  void createExprForScalarElement(FieldDecl *FD) {
     InitializedEntity ArrayEntity =
         InitializedEntity::InitializeMember(FD, &VarEntity);
     InitializationKind InitKind =
         InitializationKind::CreateCopy(SourceLocation(), SourceLocation());
     Expr *DRE = createInitExpr(FD);
-    Expr *Idx = dyn_cast<ArraySubscriptExpr>(MemberExprBases.back())->getIdx();
-    llvm::APSInt Result;
-    SemaRef.VerifyIntegerConstantExpression(Idx, &Result);
-    uint64_t IntIdx = Result.getZExtValue();
     InitializedEntity Entity = InitializedEntity::InitializeElement(
-        SemaRef.getASTContext(), IntIdx, ArrayEntity);
+        SemaRef.getASTContext(), ArrayIndex, ArrayEntity);
+    ArrayIndex++;
     InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE);
     ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, DRE);
+    InitExprs.push_back(MemberInit.get());
+  }
+
+  void addArrayInit(FieldDecl *FD, int64_t Count) {
     llvm::SmallVector<Expr *, 16> ArrayInitExprs;
-    if (IntIdx > 0) {
-      // Continue with the current InitList
-      InitListExpr *ILE = cast<InitListExpr>(InitExprs.back());
+    for (int64_t I = 0; I < Count; I++) {
+      ArrayInitExprs.push_back(InitExprs.back());
       InitExprs.pop_back();
-      llvm::ArrayRef<Expr *> L = ILE->inits();
-      for (size_t I = 0; I < L.size(); I++)
-        ArrayInitExprs.push_back(L[I]);
     }
-    ArrayInitExprs.push_back(MemberInit.get());
+    std::reverse(ArrayInitExprs.begin(), ArrayInitExprs.end());
     Expr *ILE = new (SemaRef.getASTContext())
         InitListExpr(SemaRef.getASTContext(), SourceLocation(), ArrayInitExprs,
                      SourceLocation());
@@ -1421,8 +1440,10 @@ class SyclKernelBodyCreator
 
   bool handleSpecialType(FieldDecl *FD, QualType Ty) {
     const auto *RecordDecl = Ty->getAsCXXRecordDecl();
-    // Perform initialization only if it is field of kernel object
-    if (MemberExprBases.size() == 2) {
+    ArraySubscriptExpr *ArrayRef =
+        dyn_cast<ArraySubscriptExpr>(MemberExprBases.back());
+    // Perform initialization only if decomposed from array
+    if (ArrayRef || MemberExprBases.size() == 2) {
       InitializedEntity Entity =
           InitializedEntity::InitializeMember(FD, &VarEntity);
       // Initialize with the default constructor.
@@ -1507,31 +1528,37 @@ class SyclKernelBodyCreator
 
   bool handleScalarType(FieldDecl *FD, QualType FieldTy) final {
     if (dyn_cast<ArraySubscriptExpr>(MemberExprBases.back()))
-      createExprForScalarElement(FD, FieldTy);
+      createExprForScalarElement(FD);
     else
       createExprForStructOrScalar(FD);
     return true;
   }
 
-  void enterField(const CXXRecordDecl *RD, FieldDecl *FD) final {
+  bool enterField(const CXXRecordDecl *RD, FieldDecl *FD) final {
     if (!FD->getType()->isReferenceType())
       MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD));
+    return true;
   }
 
-  void leaveField(const CXXRecordDecl *, FieldDecl *FD) final {
+  bool leaveField(const CXXRecordDecl *, FieldDecl *FD) final {
     if (!FD->getType()->isReferenceType())
       MemberExprBases.pop_back();
+    return true;
   }
 
-  void enterArray() final {
+  bool enterArray() final {
     Expr *ArrayBase = MemberExprBases.back();
     ExprResult IndexExpr = SemaRef.ActOnIntegerConstant(SourceLocation(), 0);
     ExprResult ElementBase = SemaRef.CreateBuiltinArraySubscriptExpr(
         ArrayBase, SourceLocation(), IndexExpr.get(), SourceLocation());
     MemberExprBases.push_back(ElementBase.get());
+    ArrayIndex = 0;
+    return true;
   }
 
-  void nextElement(QualType) final {
+  bool nextElement(QualType ET) final {
+    if (ET->isScalarType())
+      return true;
     ArraySubscriptExpr *LastArrayRef =
         dyn_cast<ArraySubscriptExpr>(MemberExprBases.back());
     MemberExprBases.pop_back();
@@ -1544,14 +1571,20 @@ class SyclKernelBodyCreator
     ExprResult ElementBase = SemaRef.CreateBuiltinArraySubscriptExpr(
         ArrayBase, SourceLocation(), IndexExpr.get(), SourceLocation());
     MemberExprBases.push_back(ElementBase.get());
+    return true;
   }
 
-  void leaveArray(QualType, int64_t) final { MemberExprBases.pop_back(); }
+  bool leaveArray(FieldDecl *FD, QualType, int64_t Count) final {
+    addArrayInit(FD, Count);
+    MemberExprBases.pop_back();
+    return true;
+  }
 
   using SyclKernelFieldHandler::enterArray;
   using SyclKernelFieldHandler::enterField;
   using SyclKernelFieldHandler::handleScalarType;
   using SyclKernelFieldHandler::handleSyclSamplerType;
+  using SyclKernelFieldHandler::leaveArray;
   using SyclKernelFieldHandler::leaveField;
 };
 
@@ -1670,43 +1703,50 @@ class SyclKernelIntHeaderCreator
     return true;
   }
 
-  void enterField(const CXXRecordDecl *RD, FieldDecl *FD) final {
+  bool enterField(const CXXRecordDecl *RD, FieldDecl *FD) final {
     CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8;
+    return true;
   }
 
-  void leaveField(const CXXRecordDecl *, FieldDecl *FD) final {
+  bool leaveField(const CXXRecordDecl *, FieldDecl *FD) final {
     CurOffset -= SemaRef.getASTContext().getFieldOffset(FD) / 8;
+    return true;
   }
 
-  void enterField(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final {
+  bool enterField(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final {
     const ASTRecordLayout &Layout =
         SemaRef.getASTContext().getASTRecordLayout(RD);
     CurOffset += Layout.getBaseClassOffset(BS.getType()->getAsCXXRecordDecl())
                      .getQuantity();
+    return true;
   }
 
-  void leaveField(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final {
+  bool leaveField(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final {
     const ASTRecordLayout &Layout =
         SemaRef.getASTContext().getASTRecordLayout(RD);
     CurOffset -= Layout.getBaseClassOffset(BS.getType()->getAsCXXRecordDecl())
                      .getQuantity();
+    return true;
   }
 
-  void nextElement(QualType ET) final {
+  bool nextElement(QualType ET) final {
     CurOffset += SemaRef.getASTContext().getTypeSizeInChars(ET).getQuantity();
+    return true;
   }
 
-  void leaveArray(QualType ET, int64_t Count) final {
+  bool leaveArray(FieldDecl *, QualType ET, int64_t Count) final {
     int64_t ArraySize =
         SemaRef.getASTContext().getTypeSizeInChars(ET).getQuantity();
     if (!ET->isArrayType()) {
       ArraySize *= Count;
     }
     CurOffset -= ArraySize;
+    return true;
   }
 
   using SyclKernelFieldHandler::handleScalarType;
   using SyclKernelFieldHandler::handleSyclSamplerType;
+  using SyclKernelFieldHandler::leaveArray;
 };
 } // namespace
 

clang/test/CodeGenSYCL/kernel-param-acc-array.cpp

@@ -49,26 +49,30 @@ int main() {
 // CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
 // CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
 
-// CHECK accessor array GEP for acc[0]
+// CHECK accessor array default inits
 // CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
-// CHECK: [[Z0:%[a-zA-Z0-9_]*]] = getelementptr inbounds{{.*}}[[ACCESSOR_ARRAY1]], i64 0, i64 0
+// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0
+// CHECK: [[END:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR:.*]], [[ACCESSOR]]* [[BEGIN]], i64 2
+// CHECK: [[NEXT0:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
+// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
+// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 2
+// CHECK: [[NEXT1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
+
+// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[INDEX:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY2]], i64 0, i64 0
 
 // CHECK load from kernel pointer argument alloca
 // CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]]
 
-// CHECK acc[0] __init method call
-// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z0]] to %"class{{.*}}accessor" addrspace(4)*
+// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)*
 
+// CHECK acc[0] __init method call
 // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
 
-// CHECK accessor array GEP for acc[1]
-// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
-// CHECK: [[Z1:%[a-zA-Z0-9_]*]] = getelementptr inbounds{{.*}}[[ACCESSOR_ARRAY2]], i64 0, i64 1
-
 // CHECK load from kernel pointer argument alloca
 // CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG2]]
 
-// CHECK acc[1] __init method call
-// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)*
+// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)*
 
+// CHECK acc[1] __init method call
 // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])

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

@@ -4,21 +4,12 @@
 // an array of non-trivially copyable structs as SYCL kernel parameter or
 // a non-constant size array.
 
-struct A {
-  int i;
-};
-
 struct B {
   int i;
   B(int _i) : i(_i) {}
   B(const B &x) : i(x.i) {}
 };
 
-struct C : A {
-  const A C2;
-  C() : A{0}, C2{2} {}
-};
-
 struct D {
   int i;
   ~D();
@@ -38,16 +29,12 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
 }
 
 void test() {
-  A cs1[10];
   B nsl1[4] = {1, 2, 3, 4};
-  C cs2[6];
   D nsl2[5];
   E es;
   kernel_single_task<class kernel_capture_refs>([=] {
-    int a = cs1[6].i;
     // expected-error@+1 {{kernel parameter has non-trivially copy constructible class/struct type}}
     int b = nsl1[2].i;
-    int c = cs2[0].i;
     // expected-error@+1 {{kernel parameter has non-trivially destructible class/struct type}}
     int d = nsl2[4].i;
   });

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

@@ -58,8 +58,18 @@ int main() {
 // CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (int, int)'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
+// Check kernel_B inits
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl {{.*}} cinit
+// CHECK-NEXT: InitListExpr
+// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
+// CHECK: ImplicitCastExpr
+// CHECK: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+// CHECK: ImplicitCastExpr
+// CHECK: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
 
-// Correct and enable after struct mebers are extracted into separate parameters
+// Correct and enable after struct members are extracted into separate parameters
 // C HECK kernel_C parameters
 // C HECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (struct {{.*}}, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
 // C HECK-NEXT: ParmVarDecl {{.*}} 'struct {{.*}}'