Fixes ArrayIndex calculation. This information is now extracted

from MemberExprBases. Also fixes a mistake with offset calculation
for integration headers generated. Adds lit tests.

Patch by: Rajiv Deodhar and Elizabeth Andrews

Signed-off-by: Elizabeth Andrews <[email protected]>

Author: elizabethandrews

clang/lib/Sema/SemaSYCL.cpp

@@ -851,7 +851,7 @@ class KernelObjVisitor {
     std::initializer_list<int>{(handlers.enterArray(), 0)...};
     for (int64_t Count = 0; Count < ElemCount; Count++) {
       VisitElement(nullptr, FD, ET, handlers...);
-      (void)std::initializer_list<int>{(handlers.nextElement(ET), 0)...};
+      (void)std::initializer_list<int>{(handlers.nextElement(ET, Count), 0)...};
     }
     (void)std::initializer_list<int>{
         (handlers.leaveArray(FD, ET, ElemCount), 0)...};
@@ -1108,7 +1108,7 @@ class SyclKernelFieldHandler {
   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 nextElement(QualType, int64_t) { return true; }
   virtual bool leaveArray(FieldDecl *, QualType, int64_t) { return true; }
 
   virtual ~SyclKernelFieldHandler() = default;
@@ -1626,7 +1626,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   InitializedEntity VarEntity;
   const CXXRecordDecl *KernelObj;
   llvm::SmallVector<Expr *, 16> MemberExprBases;
-  uint64_t ArrayIndex;
   FunctionDecl *KernelCallerFunc;
 
   // Using the statements/init expressions that we've created, this generates
@@ -1761,27 +1760,29 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   void createExprForScalarElement(FieldDecl *FD) {
     llvm::SmallVector<InitializedEntity, 4> InitEntities;
 
-    InitializedEntity MemberEntity =
-        InitializedEntity::InitializeMember(FD, &VarEntity);
-    InitializedEntity Entity = InitializedEntity::InitializeElement(
-        SemaRef.getASTContext(), ArrayIndex, MemberEntity);
-    InitEntities.push_back(Entity);
     // For multi-dimensional arrays, an initialized entity needs to be
     // generated for each 'dimension'. For example, the initialized entity
     // for s.array[x][y][z] is constructed using initialized entities for
     // s.array[x][y], s.array[x] and s.array. InitEntities is used to maintain
-    // this. MemberExprBases is used to get the array index for 'current
-    // dimension'.
-    for (int i = MemberExprBases.size() - 2; i >= 0; --i) {
-      if (!isa<ArraySubscriptExpr>(MemberExprBases[i]))
-        break;
+    // this.
+    InitializedEntity Entity =
+        InitializedEntity::InitializeMember(FD, &VarEntity);
+    InitEntities.push_back(Entity);
+
+    // Calculate dimension using ArraySubscriptExpressions in MemberExprBases.
+    // Each dimension has an ArraySubscriptExpression (maintains index)
+    // in MemberExprBases. For example, if we are currently handling element
+    // a[0][0][1], the top of stack entries are ArraySubscriptExpressions for
+    // indices 0,0 and 1, with 1 on top.
+    int Dims = getDims();
+    int NIndex = MemberExprBases.size() - 1 - (Dims - 1);
+    for (int i = 0; i < Dims; ++i) {
       InitializedEntity NewEntity = InitializedEntity::InitializeElement(
-          SemaRef.getASTContext(), getArrayIndex(i), InitEntities.back());
+          SemaRef.getASTContext(), getArrayIndex(NIndex), InitEntities.back());
       InitEntities.push_back(NewEntity);
+      ++NIndex;
     }
 
-    ArrayIndex++;
-
     InitializationKind InitKind =
         InitializationKind::CreateCopy(SourceLocation(), SourceLocation());
     Expr *DRE = createInitExpr(FD);
@@ -2067,14 +2068,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     ExprResult ElementBase = SemaRef.CreateBuiltinArraySubscriptExpr(
         ArrayBase, SourceLocation(), IndexExpr.get(), SourceLocation());
     MemberExprBases.push_back(ElementBase.get());
-    ArrayIndex = 0;
     return true;
   }
 
-  bool nextElement(QualType ET) final {
+  bool nextElement(QualType ET, int64_t) final {
     // Top of MemberExprBases holds ArraySubscriptExpression of element
-    // we just finished processing.
-    int64_t nextIndex = getArrayIndex((MemberExprBases.size() - 1)) + 1;
+    // we just handled, or the Array base for the dimension we are
+    // currently visiting.
+    int64_t nextIndex = getArrayIndex(MemberExprBases.size() - 1) + 1;
     MemberExprBases.pop_back();
     Expr *ArrayBase = MemberExprBases.back();
     ExprResult IndexExpr =
@@ -2103,6 +2104,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
   SYCLIntegrationHeader &Header;
   int64_t CurOffset = 0;
+  llvm::SmallVector<size_t, 16> ArrayBases;
   int StructDepth = 0;
 
   void addParam(const FieldDecl *FD, QualType ArgTy,
@@ -2249,18 +2251,20 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool nextElement(QualType ET) final {
-    CurOffset += SemaRef.getASTContext().getTypeSizeInChars(ET).getQuantity();
+  bool enterArray() final {
+    ArrayBases.push_back(CurOffset);
     return true;
   }
 
-  bool leaveArray(FieldDecl *, QualType ET, int64_t Count) final {
-    int64_t ArraySize =
-        SemaRef.getASTContext().getTypeSizeInChars(ET).getQuantity();
-    if (!ET->isArrayType()) {
-      ArraySize *= Count;
-    }
-    CurOffset -= ArraySize;
+  bool nextElement(QualType ET, int64_t Index) final {
+    int64_t Size = SemaRef.getASTContext().getTypeSizeInChars(ET).getQuantity();
+    CurOffset = ArrayBases.back() + Size * (Index + 1);
+    return true;
+  }
+
+  bool leaveArray(FieldDecl *, QualType ET, int64_t) final {
+    CurOffset = ArrayBases.back();
+    ArrayBases.pop_back();
     return true;
   }
   using SyclKernelFieldHandler::enterStruct;

clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp

@@ -13,7 +13,9 @@
 
 // CHECK: static constexpr
 // CHECK-NEXT: const char* const kernel_names[] = {
-// CHECK-NEXT:   "_ZTSZ4mainE8kernel_B"
+// CHECK-NEXT:   "_ZTSZ4mainE8kernel_B",
+// CHECK-NEXT:   "_ZTSZ4mainE8kernel_C",
+// CHECK-NEXT:   "_ZTSZ4mainE8kernel_D"
 // CHECK-NEXT: };
 
 // CHECK: static constexpr
@@ -25,14 +27,40 @@
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 16 },
 // CHECK-EMPTY:
+// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 4 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 16 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 20 },
+// CHECK-EMPTY:
+// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_D
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 4 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 16 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 20 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 24 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 28 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 32 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 36 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 40 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 44 },
+// CHECK-EMPTY:
 // CHECK-NEXT: };
 
 // CHECK: static constexpr
 // CHECK-NEXT: const unsigned kernel_signature_start[] = {
-// CHECK-NEXT:  0 // _ZTSZ4mainE8kernel_B
+// CHECK-NEXT:  0, // _ZTSZ4mainE8kernel_B
+// CHECK-NEXT:  6, // _ZTSZ4mainE8kernel_C
+// CHECK-NEXT:  13 // _ZTSZ4mainE8kernel_D
 // CHECK-NEXT: };
 
 // CHECK: template <> struct KernelInfo<class kernel_B> {
+// CHECK: template <> struct KernelInfo<class kernel_C> {
+// CHECK: template <> struct KernelInfo<class kernel_D> {
 
 #include <sycl.hpp>
 
@@ -46,9 +74,21 @@ __attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
 int main() {
 
   int a[5];
+  int b[2][3];
+  int c[2][3][2];
 
   a_kernel<class kernel_B>(
       [=]() {
         int local = a[3];
       });
+
+  a_kernel<class kernel_C>(
+      [=]() {
+        int local = b[0][1];
+      });
+
+  a_kernel<class kernel_D>(
+      [=]() {
+        int local = c[0][1][1];
+      });
 }

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

@@ -11,7 +11,7 @@
 
 using namespace cl::sycl;
 
-constexpr size_t c_num_items = 100;
+constexpr size_t c_num_items = 4;
 range<1> num_items{c_num_items}; // range<1>(num_items)
 
 // Change if tests are added/removed
@@ -48,20 +48,33 @@ void init(T &A, int value, int increment) {
 }
 
 bool test_one_array(queue &myQueue) {
-  int input1[c_num_items];
+  int input1[c_num_items][c_num_items];
+  int input2[c_num_items][c_num_items][c_num_items];
   int output[c_num_items];
   int ref[c_num_items];
-  init(input1, 1, 1);
-  init(output, 51, 1);
-  init(ref, 2, 1);
+  int value1 = 0;
+  int value2 = 0;
+  int increment = 1;
+  for (int i = 0; i < c_num_items; i++) {
+    for (int j = 0; j < c_num_items; j++) {
+      for (int k = 0; k < c_num_items; k++) {
+        input2[i][j][k] = value1;
+        value1 += increment;
+      }
+      input1[i][j] = value2;
+      value2 += increment;
+    }
+  }
+  init(output, 511, 1);
+  init(ref, 37, 2);
 
   auto out_buffer = buffer<int, 1>(&output[0], num_items);
 
   myQueue.submit([&](handler &cgh) {
     auto output_accessor = out_buffer.get_access<access::mode::write>(cgh);
 
     cgh.parallel_for<class one_array>(num_items, [=](cl::sycl::id<1> index) {
-      output_accessor[index] = input1[index] + 1;
+      output_accessor[index] = input1[0][index] + input2[2][1][index] + 1;
     });
   });
   const auto HostAccessor =