[OpenMP] Add support for multidimensional array sections in map clause SEMA.

[sfantao]

Summary

In some cases it can be proved statically that multidimensional array section refer to contiguous storage and can therefore be allowed in a map clause. This patch adds support for those cases in SEMA.

[ABataev]

include/clang/Basic/DiagnosticSemaKinds.td

7787–7789

If you say 'can be' incompatible, then this must be a warning, not an error

lib/Sema/SemaOpenMP.cpp

9016–9023

I can't agree with that. For example:

const int n = 0;
arr[n:]

It won't work with your solution, though we shall support it

9236–9248

OMPArraySectionExpr has static function getBaseOriginalType()

[sfantao]

Use better disgnostic message and OMPArraySectionExpr::getBaseOriginalType.

[sfantao]

Hi Alexey,

Thanks for the review!

include/clang/Basic/DiagnosticSemaKinds.td

7787–7789

I agree the diagnostic description is not the best. I changed it to: "can't prove employed array section specifies contiguous storage".

This has to be an error because there is no way to express non-contiguous in the offload runtime interface. So, we would have to codegen traps. The user would end up getting runtime crashes instead of a diagnostic that would be far more useful to direct the user to take the right actions.

lib/Sema/SemaOpenMP.cpp

9016–9023

Hi Alexey,

That example works fine. Note that CheckArrayExpressionReferToWholeSize only results in a error if AllowWholeSizeArraySection = false. It is initially set to true, and will be set to false as components of the expression prove to be incompatible with that. I added a few regression test for when the bounds come from variables.

So:

struct S1 {
  int a;
  int b;
}
struct S2 {
  S1 a[10];
  int b;
}

foo (int arg) {
  int a[5][6];
  const int n = 0;
  S2 s;

  // valid - the array expression is in the right most component.
  #pragma omp target map(a[arg:])
  #pragma omp target map(a[:arg])
  #pragma omp target map(a[n:])
  // valid - this is valid only if n is zero and the compiler can prove that.
  #pragma omp target map(a[:][n:])

  // invalid - is only valid if arg is zero and that cannot be proved.
  #pragma omp target map(a[:][arg:])

  // invalid - it is contiguous storage but the OpenMP 4.5 spec explicitly say array sections only allowed in the rightmost expression if struct fields are involved.
  #pragma omp target map(s.a[n:1].b)
}

9236–9248

Oh, great, I hadn't noticed. I am using that now. Thanks!

[ABataev]

include/clang/Basic/DiagnosticSemaKinds.td

7787–7789

'can't prove' again is not good for an error. Still think this must be a warning or you don't need to diagnose anything if you can't make a decision. In this case, you must consider specified array section as a contiguous.

lib/Sema/SemaOpenMP.cpp

9016–9023

I agree with all your examples except for this one

#pragma omp target map(a[:][arg:])

For me, this is valid if you can't prove it is non-contiguous. You should consider this as a possibly contiguous. If it is not contiguous, it is user's problem. But if he needs to use such form of expression, I don't see why we should not allow him to do this.

[ABataev]

lib/Sema/SemaOpenMP.cpp

9016–9023

I don't think this code

const int n = 0;
arr[n:]

will work in C. Try to add a test for C, not C++. In C 'const' does not mean you can evaluate the value of the variable during compilation.

[sfantao]

• Allow array expressions whose memory contiguous checks are inconclusive.
• Add C test to check the constant expressions no longer can prove bounds in array sections.

include/clang/Basic/DiagnosticSemaKinds.td

7787–7789

Now I am only emitting an error for the cases that can be proved wrong. For those that are inconclusive I don't issue any warning. The message also changed to reflect that.

lib/Sema/SemaOpenMP.cpp

9016–9023

In the new diff I am allowing this given that the ranges check is inconclusive. At runtime the behavior will be the same as map(a[:]), so the runtime map will contain the holes. This falls into those cases that is users' fault.

9016–9023

So if you have:

const int n = 0;
arr[:][n:]

This is not a probably anymore because the check in the second dimension will be inconclusive, so it will be allowed even if n can't be proved to be zero.

I'm adding a test anyways.

[ABataev]

lib/Sema/SemaOpenMP.cpp

9013

Don't use 'int' as a result of the functions, add enumeric and use it.

9042–9044

Remove extra braces

9069

Again, use enum as a return result

9228–9230

Extra braces

[sfantao]

• Remove extra braces and revert the sense of the whole/unity array section check.

lib/Sema/SemaOpenMP.cpp

9013

I changed this back to bool and reverted the sense given that I am only taking any action if something is proved NOT to be whole or unity.

9069

Using bool instead with reverted sense.

[ABataev]

lib/Sema/SemaOpenMP.cpp

9066

This must return 'bool', not 'int'

[sfantao]

• Change return value from int to bool.

[sfantao]

lib/Sema/SemaOpenMP.cpp

9066

Sorry, my bad. It is fixed in the new diff.

[ABataev]

LG

include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_expected_var_name_member_expr : Error<
   "expected variable name%select{| or data member of current class}0">;
 def err_omp_expected_var_name_member_expr_or_array_item : Error<
   "expected variable name%select{|, data member of current class}0, array element or array section">;
 def err_omp_expected_named_var_member_or_array_expression: Error<
   "expected expression containing only member accesses and/or array sections based on named variables">;
 def err_omp_bit_fields_forbidden_in_map_clause : Error<
   "bit fields cannot be used to specify storage in a map clause">;
-def err_omp_array_section_in_rightmost_expression : Error<
-  "array section can only be associated with the rightmost variable in a map clause expression">;
+def err_cannot_prove_omp_array_section_specifies_contiguous_storage : Error<
+  "can't prove employed array section specifies contiguous storage">;
 def err_omp_union_type_not_allowed : Error<
   "mapped storage cannot be derived from a union">;
 def err_omp_expected_access_to_data_field : Error<
   "expected access to data field">;
 def err_omp_multiple_array_items_in_map_clause : Error<
   "multiple array elements associated with the same variable are not allowed in map clauses of the same construct">;
 def err_omp_pointer_mapped_along_with_derived_section : Error<
   "pointer cannot be mapped along with a section derived from itself">;
 def err_omp_original_storage_is_shared_and_does_not_contain : Error<

lib/Sema/SemaOpenMP.cpp

static bool CheckTypeMappable(SourceLocation SL, SourceRange SR, Sema &SemaRef,
   } else if (CXXRecordDecl *RD = dyn_cast_or_null<CXXRecordDecl>(ND)) {
     if (!RD->isInvalidDecl() &&
         !IsCXXRecordForMappable(SemaRef, SL, Stack, RD))
       return false;
   }
   return true;
 }
 
+/// \brief Return true if it can be proven that the provided array expression
+/// (array section or array subscript) specify the whole size of the array whose
+/// base type is \a BaseQTy.
+static bool CheckArrayExpressionReferToWholeSize(Sema &SemaRef, const Expr *E,
+                                                 QualType BaseQTy) {
+  auto *OASE = dyn_cast<OMPArraySectionExpr>(E);
+
+  // If this is an array subscript, it refers to the whole size if the size of
+  // the dimension is constant and equals 1. Also, an array section assumes the
+  // format of an array subscript if no colon is used.
+  if (isa<ArraySubscriptExpr>(E) || (OASE && OASE->getColonLoc().isInvalid())) {
+    if (auto *ATy = dyn_cast<ConstantArrayType>(BaseQTy.getTypePtr()))
+      return ATy->getSize().getSExtValue() == 1;
+    return false;
+  }
+
+  assert(OASE && "Expecting array section if not an array subscript.");
+  auto *LowerBound = OASE->getLowerBound();
+  auto *Length = OASE->getLength();
+
+  // If there is a lower bound that does not evaluates to zero, we are not
+  // convering the whole dimension.
+  if (LowerBound) {
+    llvm::APSInt ConstLowerBound;
+    if (!LowerBound->EvaluateAsInt(ConstLowerBound, SemaRef.getASTContext()))
+      return false; // Can't get the integer value as a constant.
+    if (ConstLowerBound.getSExtValue())
+      return false;
+  }
+
+  // If we don't have a length we covering the whole dimension.
+  if (!Length) {
+    return true;
+  }
+
+  // If the base is a pointer, we don't have a way to get the size of the
+  // pointee.
+  if (BaseQTy->isPointerType())
+    return false;
+
+  // We can only check if the length is the same as the size of the dimension
+  // if we have a constant array.
+  auto *CATy = dyn_cast<ConstantArrayType>(BaseQTy.getTypePtr());
+  if (!CATy)
+    return false;
+
+  llvm::APSInt ConstLength;
+  if (!Length->EvaluateAsInt(ConstLength, SemaRef.getASTContext()))
+    return false; // Can't get the integer value as a constant.
+
+  return CATy->getSize().getSExtValue() == ConstLength.getSExtValue();
+}
+
+// Return true if it can be proven that the provided array expression (array
+// section or array subscript) specify a single element of the array whose base
+// type is \a BaseQTy.
+static bool CheckArrayExpressionReferToUnitySize(Sema &SemaRef, const Expr *E,
+                                                 QualType BaseQTy) {
+  auto *OASE = dyn_cast<OMPArraySectionExpr>(E);
+
+  // An array subscript always refer to a single element. Also, an array section
+  // assumes the format of an array subscript if no colon is used.
+  if (isa<ArraySubscriptExpr>(E) || (OASE && OASE->getColonLoc().isInvalid()))
+    return true;
+
+  assert(OASE && "Expecting array section if not an array subscript.");
+  auto *Length = OASE->getLength();
+
+  // If we don't have a length we have to check if the array has unitary size
+  // for this dimension. Also, we should always expect a length if the base type
+  // is pointer.
+  if (!Length) {
+    if (auto *ATy = dyn_cast<ConstantArrayType>(BaseQTy.getTypePtr()))
+      return ATy->getSize().getSExtValue() == 1;
+    return false;
+  }
+
+  // Check if the length evaluates to 1.
+  llvm::APSInt ConstLength;
+  if (!Length->EvaluateAsInt(ConstLength, SemaRef.getASTContext()))
+    return false; // Can't get the integer value as a constant.
+
+  return ConstLength.getSExtValue() == 1;
+}
+
 // Return the expression of the base of the map clause or null if it cannot
 // be determined and do all the necessary checks to see if the expression is
 // valid as a standalone map clause expression.
 static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E) {
   SourceLocation ELoc = E->getExprLoc();
   SourceRange ERange = E->getSourceRange();
 
   // The base of elements of list in a map clause have to be either:
static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E) {
   //     #pragma omp target map (S.Arr[:12]);
   //   }
   // }
   //
   // We want to retrieve the member expression 'this->S';
 
   Expr *RelevantExpr = nullptr;
 
-  // Flags to help capture some memory
-
   // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.2]
   //  If a list item is an array section, it must specify contiguous storage.
   //
   // For this restriction it is sufficient that we make sure only references
   // to variables or fields and array expressions, and that no array sections
-  // exist except in the rightmost expression. E.g. these would be invalid:
+  // exist except in the rightmost expression (unless they cover the whole
+  // dimension of the array). E.g. these would be invalid:
   //
   //   r.ArrS[3:5].Arr[6:7]
   //
   //   r.ArrS[3:5].x
   //
   // but these would be valid:
   //   r.ArrS[3].Arr[6:7]
   //
   //   r.ArrS[3].x
 
-  bool IsRightMostExpression = true;
-
-  while (!RelevantExpr) {
-    auto AllowArraySection = IsRightMostExpression;
-    IsRightMostExpression = false;
+  bool AllowUnitySizeArraySection = true;
+  bool AllowWholeSizeArraySection = true;
 
+  for (bool IsRightMostExpression = true; !RelevantExpr;
+       IsRightMostExpression = false) {
     E = E->IgnoreParenImpCasts();
 
     if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
       if (!isa<VarDecl>(CurE->getDecl()))
         break;
 
       RelevantExpr = CurE;
+
+      // If we got a reference to a declaration, we should not expect any array
+      // section before that.
+      AllowUnitySizeArraySection = false;
+      AllowWholeSizeArraySection = false;
       continue;
     }
 
     if (auto *CurE = dyn_cast<MemberExpr>(E)) {
       auto *BaseE = CurE->getBase()->IgnoreParenImpCasts();
 
       if (isa<CXXThisExpr>(BaseE))
         // We found a base expression: this->Val.
if (auto *CurE = dyn_cast<MemberExpr>(E)) {
       //
       if (auto *RT = CurType->getAs<RecordType>())
         if (RT->isUnionType()) {
           SemaRef.Diag(ELoc, diag::err_omp_union_type_not_allowed)
               << CurE->getSourceRange();
           break;
         }
 
+      // If we got a member expression, we should not expect any array section
+      // before that:
+      //
+      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.7]
+      //  If a list item is an element of a structure, only the rightmost symbol
+      //  of the variable reference can be an array section.
+      //
+      AllowUnitySizeArraySection = false;
+      AllowWholeSizeArraySection = false;
       continue;
     }
 
     if (auto *CurE = dyn_cast<ArraySubscriptExpr>(E)) {
       E = CurE->getBase()->IgnoreParenImpCasts();
 
       if (!E->getType()->isAnyPointerType() && !E->getType()->isArrayType()) {
         SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name)
             << 0 << CurE->getSourceRange();
         break;
       }
+
+      // If we got an array subscript that express the whole dimension we
+      // can have any array expressions before. If it only expressing part of
+      // the dimension, we can only have unitary-size array expressions.
+      if (!CheckArrayExpressionReferToWholeSize(SemaRef, CurE, E->getType())) {
+        AllowWholeSizeArraySection = false;
+      }
       continue;
     }
 
     if (auto *CurE = dyn_cast<OMPArraySectionExpr>(E)) {
-      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.7]
-      //  If a list item is an element of a structure, only the rightmost symbol
-      //  of the variable reference can be an array section.
-      //
-      if (!AllowArraySection) {
-        SemaRef.Diag(ELoc, diag::err_omp_array_section_in_rightmost_expression)
-            << CurE->getSourceRange();
-        break;
-      }
-
       E = CurE->getBase()->IgnoreParenImpCasts();
 
+      auto CurType =
+          OMPArraySectionExpr::getBaseOriginalType(E).getCanonicalType();
+
       // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
       //  If the type of a list item is a reference to a type T then the type
       //  will be considered to be T for all purposes of this clause.
-      QualType CurType = E->getType();
       if (CurType->isReferenceType())
         CurType = CurType->getPointeeType();
 
-      if (!CurType->isAnyPointerType() && !CurType->isArrayType()) {
+      bool IsPointer = CurType->isAnyPointerType();
+
+      if (!IsPointer && !CurType->isArrayType()) {
         SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name)
             << 0 << CurE->getSourceRange();
         break;
       }
 
+      bool IsWhole =
+          CheckArrayExpressionReferToWholeSize(SemaRef, CurE, CurType);
+      bool IsUnity =
+          CheckArrayExpressionReferToUnitySize(SemaRef, CurE, CurType);
+
+      if (AllowWholeSizeArraySection && AllowUnitySizeArraySection) {
+        // Any array section is currently allowed.
+        //
+        // If this array section refers to the whole dimension we can still
+        // accept other array sections before this one, except if the base is a
+        // pointer. Otherwise, only unitary sections are accepted.
+        if (!IsWhole || IsPointer)
+          AllowWholeSizeArraySection = false;
+      } else if ((AllowUnitySizeArraySection && !IsUnity) ||
+                 (AllowWholeSizeArraySection && !IsWhole)) {
+        // A unity or whole array section is not allowed and that is not
+        // compatible with the properties of the current array section.
+        SemaRef.Diag(
+            ELoc,
+            diag::
+                err_cannot_prove_omp_array_section_specifies_contiguous_storage)
+            << CurE->getSourceRange();
+        break;
+      }
       continue;
     }
 
     // If nothing else worked, this is not a valid map clause expression.
     SemaRef.Diag(ELoc,
                  diag::err_omp_expected_named_var_member_or_array_expression)
         << ERange;
     break;

test/OpenMP/target_map_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 200 %s
 template <typename T, int I>
 struct SA {
   static int ss;
   #pragma omp threadprivate(ss) // expected-note {{defined as threadprivate or thread local}}
   float a;
   int b[12];
   float *c;
   T d;

 union SD {
   unsigned A;
   float B;
 };
 
 void SAclient(int arg) {
   SA<int,123> s;
   s.func(arg); // expected-note {{in instantiation of member function}}
+  double marr[10][10][10];
+  double marr2[5][10][1];
+  double mvla[5][arg][10];
+  double ***mptr;
+  const int n = 0;
+  const int m = 1;
 
   SB *p;
 
   SD u;
   SC r(p),t(p);
   #pragma omp target map(r)
   {}
+  #pragma omp target map(marr[2][0:2][0:2]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[:][0:2][0:2]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[2][3][0:2])
+  {}
+  #pragma omp target map(marr[:][:][:])
+  {}
+  #pragma omp target map(marr[:2][:][:])
+  {}
+  #pragma omp target map(marr[arg:][:][:])
+  {}
+  #pragma omp target map(marr[arg:])
+  {}
+  #pragma omp target map(marr[arg:][:arg][:]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[:arg][:])
+  {}
+  #pragma omp target map(marr[:arg][n:])
+  {}
+  #pragma omp target map(marr[:][:arg][n:]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[n:m][:arg][n:])
+  {}
+  #pragma omp target map(marr[:2][:1][:]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[:2][1:][:]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[:2][:][:1]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[:2][:][1:]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[:1][:2][:])
+  {}
+  #pragma omp target map(marr[:1][0][:])
+  {}
+  #pragma omp target map(marr[:arg][:2][:]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[:1][3:1][:2])
+  {}
+  #pragma omp target map(marr[:1][3:arg][:2]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[:1][3:2][:2]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr[:2][:10][:])
+  {}
+  #pragma omp target map(marr[:2][:][:5+5])
+  {}
+  #pragma omp target map(marr[:2][2+2-4:][0:5+5])
+  {}
+
+  #pragma omp target map(marr[:1][:2][0]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(marr2[:1][:2][0])
+  {}
+
+  #pragma omp target map(mvla[:1][:][0]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(mvla[:2][:arg][:]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(mvla[1][2:arg][:])
+  {}
+  #pragma omp target map(mvla[:1][:][:])
+  {}
+
+  #pragma omp target map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(mptr[:1][:2-1][2:4-3])
+  {}
+  #pragma omp target map(mptr[:1][:arg][2:4-3]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(mptr[:1][:2-1][0:2])
+  {}
+  #pragma omp target map(mptr[:1][:2][0:2]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+  #pragma omp target map(mptr[:1][:][0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+  {}
+  #pragma omp target map(mptr[:2][:1][0:2]) // expected-error {{can't prove employed array section specifies contiguous storage}}
+  {}
+
   #pragma omp target map(r.ArrS[0].B)
   {}
+  #pragma omp target map(r.ArrS[:1].B) // expected-error {{OpenMP array section is not allowed here}}
+  {}
+  #pragma omp target map(r.ArrS[:arg].B) // expected-error {{OpenMP array section is not allowed here}}
+  {}
   #pragma omp target map(r.ArrS[0].Arr[1:23])
   {}
+  #pragma omp target map(r.ArrS[0].Arr[1:arg])
+  {}
+  #pragma omp target map(r.ArrS[0].Arr[arg:23])
+  {}
   #pragma omp target map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}}
   {}
   #pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}}
   {}
   #pragma omp target map(r.ArrS[0].A, t.ArrS[1].A)
   {}
   #pragma omp target map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer derreferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
   {}