diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 4b074b853bfe6..fbb107679d94b 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12291,8 +12291,8 @@ def warn_acc_if_self_conflict
               "evaluates to true">,
       InGroup<DiagGroup<"openacc-self-if-potential-conflict">>;
 def err_acc_int_expr_requires_integer
-    : Error<"OpenACC %select{clause|directive}0 '%1' requires expression of "
-            "integer type (%2 invalid)">;
+    : Error<"OpenACC %select{clause '%1'|directive '%2'|sub-array bound}0 "
+            "requires expression of integer type (%3 invalid)">;
 def err_acc_int_expr_incomplete_class_type
     : Error<"OpenACC integer expression has incomplete class type %0">;
 def err_acc_int_expr_explicit_conversion
@@ -12310,4 +12310,24 @@ def err_acc_num_gangs_num_args
 def err_acc_not_a_var_ref
     : Error<"OpenACC variable is not a valid variable name, sub-array, array "
             "element, or composite variable member">;
+def err_acc_typecheck_subarray_value
+    : Error<"OpenACC sub-array subscripted value is not an array or pointer">;
+def err_acc_subarray_function_type
+    : Error<"OpenACC sub-array cannot be of function type %0">;
+def err_acc_subarray_incomplete_type
+    : Error<"OpenACC sub-array base is of incomplete type %0">;
+def err_acc_subarray_no_length
+    : Error<"OpenACC sub-array length is unspecified and cannot be inferred "
+            "because the subscripted value is %select{not an array|an array of "
+            "unknown bound}0">;
+def err_acc_subarray_negative
+    : Error<"OpenACC sub-array %select{lower bound|length}0 evaluated to "
+            "negative value %1">;
+def err_acc_subarray_out_of_range
+    : Error<"OpenACC sub-array %select{lower bound|length}0 evaluated to a "
+            "value (%1) that would be out of the range of the subscripted "
+            "array size of %2">;
+def err_acc_subarray_base_plus_length_out_of_range
+    : Error<"OpenACC sub-array specified range [%0:%1] would be out of the "
+            "range of the subscripted array size of %2">;
 } // end of sema component.
diff --git a/clang/lib/Parse/ParseExpr.cpp b/clang/lib/Parse/ParseExpr.cpp
index 7d6febb04a82c..5f5f9a79c8c40 100644
--- a/clang/lib/Parse/ParseExpr.cpp
+++ b/clang/lib/Parse/ParseExpr.cpp
@@ -2039,7 +2039,8 @@ Parser::ParsePostfixExpressionSuffix(ExprResult LHS) {
         if (Tok.is(tok::colon)) {
           // Consume ':'
           ColonLocFirst = ConsumeToken();
-          Length = Actions.CorrectDelayedTyposInExpr(ParseExpression());
+          if (Tok.isNot(tok::r_square))
+            Length = Actions.CorrectDelayedTyposInExpr(ParseExpression());
         }
       } else if (ArgExprs.size() <= 1 && getLangOpts().OpenMP) {
         ColonProtectionRAIIObject RAII(*this);
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 3ea81e0497c20..e07c5b2317f27 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -16,6 +16,7 @@
 #include "clang/Basic/DiagnosticSema.h"
 #include "clang/Basic/OpenACCKinds.h"
 #include "clang/Sema/Sema.h"
+#include "llvm/ADT/StringExtras.h"
 #include "llvm/Support/Casting.h"
 
 using namespace clang;
@@ -367,7 +368,9 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK,
   assert(((DK != OpenACCDirectiveKind::Invalid &&
            CK == OpenACCClauseKind::Invalid) ||
           (DK == OpenACCDirectiveKind::Invalid &&
-           CK != OpenACCClauseKind::Invalid)) &&
+           CK != OpenACCClauseKind::Invalid) ||
+          (DK == OpenACCDirectiveKind::Invalid &&
+           CK == OpenACCClauseKind::Invalid)) &&
          "Only one of directive or clause kind should be provided");
 
   class IntExprConverter : public Sema::ICEConvertDiagnoser {
@@ -375,6 +378,16 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK,
     OpenACCClauseKind ClauseKind;
     Expr *IntExpr;
 
+    // gets the index into the diagnostics so we can use this for clauses,
+    // directives, and sub array.s
+    unsigned getDiagKind() const {
+      if (ClauseKind != OpenACCClauseKind::Invalid)
+        return 0;
+      if (DirectiveKind != OpenACCDirectiveKind::Invalid)
+        return 1;
+      return 2;
+    }
+
   public:
     IntExprConverter(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
                      Expr *IntExpr)
@@ -390,12 +403,8 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK,
     }
     SemaBase::SemaDiagnosticBuilder diagnoseNotInt(Sema &S, SourceLocation Loc,
                                                    QualType T) override {
-      if (ClauseKind != OpenACCClauseKind::Invalid)
-        return S.Diag(Loc, diag::err_acc_int_expr_requires_integer) <<
-               /*Clause=*/0 << ClauseKind << T;
-
-      return S.Diag(Loc, diag::err_acc_int_expr_requires_integer) <<
-             /*Directive=*/1 << DirectiveKind << T;
+      return S.Diag(Loc, diag::err_acc_int_expr_requires_integer)
+             << getDiagKind() << ClauseKind << DirectiveKind << T;
     }
 
     SemaBase::SemaDiagnosticBuilder
@@ -503,12 +512,211 @@ ExprResult SemaOpenACC::ActOnArraySectionExpr(Expr *Base, SourceLocation LBLoc,
                                               SourceLocation RBLoc) {
   ASTContext &Context = getASTContext();
 
-  // TODO OpenACC: We likely have to reproduce a lot of the same logic from the
-  // OMP version of this, but at the moment we don't have a good way to test it,
-  // so for now we'll just create the node.
+  // Handle placeholders.
+  if (Base->hasPlaceholderType() &&
+      !Base->hasPlaceholderType(BuiltinType::ArraySection)) {
+    ExprResult Result = SemaRef.CheckPlaceholderExpr(Base);
+    if (Result.isInvalid())
+      return ExprError();
+    Base = Result.get();
+  }
+  if (LowerBound && LowerBound->getType()->isNonOverloadPlaceholderType()) {
+    ExprResult Result = SemaRef.CheckPlaceholderExpr(LowerBound);
+    if (Result.isInvalid())
+      return ExprError();
+    Result = SemaRef.DefaultLvalueConversion(Result.get());
+    if (Result.isInvalid())
+      return ExprError();
+    LowerBound = Result.get();
+  }
+  if (Length && Length->getType()->isNonOverloadPlaceholderType()) {
+    ExprResult Result = SemaRef.CheckPlaceholderExpr(Length);
+    if (Result.isInvalid())
+      return ExprError();
+    Result = SemaRef.DefaultLvalueConversion(Result.get());
+    if (Result.isInvalid())
+      return ExprError();
+    Length = Result.get();
+  }
+
+  // Check the 'base' value, it must be an array or pointer type, and not to/of
+  // a function type.
+  QualType OriginalBaseTy = ArraySectionExpr::getBaseOriginalType(Base);
+  QualType ResultTy;
+  if (!Base->isTypeDependent()) {
+    if (OriginalBaseTy->isAnyPointerType()) {
+      ResultTy = OriginalBaseTy->getPointeeType();
+    } else if (OriginalBaseTy->isArrayType()) {
+      ResultTy = OriginalBaseTy->getAsArrayTypeUnsafe()->getElementType();
+    } else {
+      return ExprError(
+          Diag(Base->getExprLoc(), diag::err_acc_typecheck_subarray_value)
+          << Base->getSourceRange());
+    }
+
+    if (ResultTy->isFunctionType()) {
+      Diag(Base->getExprLoc(), diag::err_acc_subarray_function_type)
+          << ResultTy << Base->getSourceRange();
+      return ExprError();
+    }
+
+    if (SemaRef.RequireCompleteType(Base->getExprLoc(), ResultTy,
+                                    diag::err_acc_subarray_incomplete_type,
+                                    Base))
+      return ExprError();
+
+    if (!Base->hasPlaceholderType(BuiltinType::ArraySection)) {
+      ExprResult Result = SemaRef.DefaultFunctionArrayLvalueConversion(Base);
+      if (Result.isInvalid())
+        return ExprError();
+      Base = Result.get();
+    }
+  }
+
+  auto GetRecovery = [&](Expr *E, QualType Ty) {
+    ExprResult Recovery =
+        SemaRef.CreateRecoveryExpr(E->getBeginLoc(), E->getEndLoc(), E, Ty);
+    return Recovery.isUsable() ? Recovery.get() : nullptr;
+  };
+
+  // Ensure both of the expressions are int-exprs.
+  if (LowerBound && !LowerBound->isTypeDependent()) {
+    ExprResult LBRes =
+        ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Invalid,
+                     LowerBound->getExprLoc(), LowerBound);
+
+    if (LBRes.isUsable())
+      LBRes = SemaRef.DefaultLvalueConversion(LBRes.get());
+    LowerBound =
+        LBRes.isUsable() ? LBRes.get() : GetRecovery(LowerBound, Context.IntTy);
+  }
+
+  if (Length && !Length->isTypeDependent()) {
+    ExprResult LenRes =
+        ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Invalid,
+                     Length->getExprLoc(), Length);
+
+    if (LenRes.isUsable())
+      LenRes = SemaRef.DefaultLvalueConversion(LenRes.get());
+    Length =
+        LenRes.isUsable() ? LenRes.get() : GetRecovery(Length, Context.IntTy);
+  }
+
+  // Length is required if the base type is not an array of known bounds.
+  if (!Length && (OriginalBaseTy.isNull() ||
+                  (!OriginalBaseTy->isDependentType() &&
+                   !OriginalBaseTy->isConstantArrayType() &&
+                   !OriginalBaseTy->isDependentSizedArrayType()))) {
+    bool IsArray = !OriginalBaseTy.isNull() && OriginalBaseTy->isArrayType();
+    Diag(ColonLoc, diag::err_acc_subarray_no_length) << IsArray;
+    // Fill in a dummy 'length' so that when we instantiate this we don't
+    // double-diagnose here.
+    ExprResult Recovery = SemaRef.CreateRecoveryExpr(
+        ColonLoc, SourceLocation(), ArrayRef<Expr *>{std::nullopt},
+        Context.IntTy);
+    Length = Recovery.isUsable() ? Recovery.get() : nullptr;
+  }
+
+  // Check the values of each of the arguments, they cannot be negative(we
+  // assume), and if the array bound is known, must be within range. As we do
+  // so, do our best to continue with evaluation, we can set the
+  // value/expression to nullptr/nullopt if they are invalid, and treat them as
+  // not present for the rest of evaluation.
+
+  // We don't have to check for dependence, because the dependent size is
+  // represented as a different AST node.
+  std::optional<llvm::APSInt> BaseSize;
+  if (!OriginalBaseTy.isNull() && OriginalBaseTy->isConstantArrayType()) {
+    const auto *ArrayTy = Context.getAsConstantArrayType(OriginalBaseTy);
+    BaseSize = ArrayTy->getSize();
+  }
+
+  auto GetBoundValue = [&](Expr *E) -> std::optional<llvm::APSInt> {
+    if (!E || E->isInstantiationDependent())
+      return std::nullopt;
+
+    Expr::EvalResult Res;
+    if (!E->EvaluateAsInt(Res, Context))
+      return std::nullopt;
+    return Res.Val.getInt();
+  };
+
+  std::optional<llvm::APSInt> LowerBoundValue = GetBoundValue(LowerBound);
+  std::optional<llvm::APSInt> LengthValue = GetBoundValue(Length);
+
+  // Check lower bound for negative or out of range.
+  if (LowerBoundValue.has_value()) {
+    if (LowerBoundValue->isNegative()) {
+      Diag(LowerBound->getExprLoc(), diag::err_acc_subarray_negative)
+          << /*LowerBound=*/0 << toString(*LowerBoundValue, /*Radix=*/10);
+      LowerBoundValue.reset();
+      LowerBound = GetRecovery(LowerBound, LowerBound->getType());
+    } else if (BaseSize.has_value() &&
+               llvm::APSInt::compareValues(*LowerBoundValue, *BaseSize) >= 0) {
+      // Lower bound (start index) must be less than the size of the array.
+      Diag(LowerBound->getExprLoc(), diag::err_acc_subarray_out_of_range)
+          << /*LowerBound=*/0 << toString(*LowerBoundValue, /*Radix=*/10)
+          << toString(*BaseSize, /*Radix=*/10);
+      LowerBoundValue.reset();
+      LowerBound = GetRecovery(LowerBound, LowerBound->getType());
+    }
+  }
+
+  // Check length for negative or out of range.
+  if (LengthValue.has_value()) {
+    if (LengthValue->isNegative()) {
+      Diag(Length->getExprLoc(), diag::err_acc_subarray_negative)
+          << /*Length=*/1 << toString(*LengthValue, /*Radix=*/10);
+      LengthValue.reset();
+      Length = GetRecovery(Length, Length->getType());
+    } else if (BaseSize.has_value() &&
+               llvm::APSInt::compareValues(*LengthValue, *BaseSize) > 0) {
+      // Length must be lessthan or EQUAL to the size of the array.
+      Diag(Length->getExprLoc(), diag::err_acc_subarray_out_of_range)
+          << /*Length=*/1 << toString(*LengthValue, /*Radix=*/10)
+          << toString(*BaseSize, /*Radix=*/10);
+      LengthValue.reset();
+      Length = GetRecovery(Length, Length->getType());
+    }
+  }
+
+  // Adding two APSInts requires matching sign, so extract that here.
+  auto AddAPSInt = [](llvm::APSInt LHS, llvm::APSInt RHS) -> llvm::APSInt {
+    if (LHS.isSigned() == RHS.isSigned())
+      return LHS + RHS;
+
+    unsigned Width = std::max(LHS.getBitWidth(), RHS.getBitWidth()) + 1;
+    return llvm::APSInt(LHS.sext(Width) + RHS.sext(Width), /*Signed=*/true);
+  };
+
+  // If we know all 3 values, we can diagnose that the total value would be out
+  // of range.
+  if (BaseSize.has_value() && LowerBoundValue.has_value() &&
+      LengthValue.has_value() &&
+      llvm::APSInt::compareValues(AddAPSInt(*LowerBoundValue, *LengthValue),
+                                  *BaseSize) > 0) {
+    Diag(Base->getExprLoc(),
+         diag::err_acc_subarray_base_plus_length_out_of_range)
+        << toString(*LowerBoundValue, /*Radix=*/10)
+        << toString(*LengthValue, /*Radix=*/10)
+        << toString(*BaseSize, /*Radix=*/10);
+
+    LowerBoundValue.reset();
+    LowerBound = GetRecovery(LowerBound, LowerBound->getType());
+    LengthValue.reset();
+    Length = GetRecovery(Length, Length->getType());
+  }
+
+  // If any part of the expression is dependent, return a dependent sub-array.
+  QualType ArrayExprTy = Context.ArraySectionTy;
+  if (Base->isTypeDependent() ||
+      (LowerBound && LowerBound->isInstantiationDependent()) ||
+      (Length && Length->isInstantiationDependent()))
+    ArrayExprTy = Context.DependentTy;
+
   return new (Context)
-      ArraySectionExpr(Base, LowerBound, Length, Context.ArraySectionTy,
-                       VK_LValue, OK_Ordinary, ColonLoc, RBLoc);
+      ArraySectionExpr(Base, LowerBound, Length, ArrayExprTy, VK_LValue,
+                       OK_Ordinary, ColonLoc, RBLoc);
 }
 
 bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
diff --git a/clang/test/ParserOpenACC/parse-cache-construct.c b/clang/test/ParserOpenACC/parse-cache-construct.c
index de26fc2b277a6..8937aa095d5ea 100644
--- a/clang/test/ParserOpenACC/parse-cache-construct.c
+++ b/clang/test/ParserOpenACC/parse-cache-construct.c
@@ -134,9 +134,8 @@ void func() {
   }
 
   for (int i = 0; i < 10; ++i) {
-    // expected-error@+2{{expected expression}}
     // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
-    #pragma acc cache(readonly:ArrayPtr[5:])
+    #pragma acc cache(readonly:ArrayPtr[5:1])
   }
 
   for (int i = 0; i < 10; ++i) {
diff --git a/clang/test/ParserOpenACC/parse-cache-construct.cpp b/clang/test/ParserOpenACC/parse-cache-construct.cpp
index f1c71e8b58478..374fe2697b63f 100644
--- a/clang/test/ParserOpenACC/parse-cache-construct.cpp
+++ b/clang/test/ParserOpenACC/parse-cache-construct.cpp
@@ -74,12 +74,12 @@ void use() {
   for (int i = 0; i < 10; ++i) {
     // expected-error@+2{{OpenACC sub-array is not allowed here}}
     // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
-    #pragma acc cache(Arrs.MemArr[3:4].array[1:4])
+    #pragma acc cache(Arrs.MemArr[2:1].array[1:4])
   }
   for (int i = 0; i < 10; ++i) {
     // expected-error@+2{{OpenACC sub-array is not allowed here}}
     // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
-    #pragma acc cache(Arrs.MemArr[3:4].array[4])
+    #pragma acc cache(Arrs.MemArr[2:1].array[4])
   }
   for (int i = 0; i < 10; ++i) {
     // expected-error@+3{{expected ']'}}
@@ -88,7 +88,7 @@ void use() {
     #pragma acc cache(Arrs.MemArr[3:4:].array[4])
   }
   for (int i = 0; i < 10; ++i) {
-    // expected-error@+2{{expected expression}}
+    // expected-error@+2{{OpenACC sub-array is not allowed here}}
     // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
     #pragma acc cache(Arrs.MemArr[:].array[4])
   }
@@ -105,7 +105,7 @@ void use() {
     #pragma acc cache(Arrs.MemArr[: :].array[4])
   }
   for (int i = 0; i < 10; ++i) {
-    // expected-error@+2{{expected expression}}
+    // expected-error@+2{{OpenACC sub-array is not allowed here}}
     // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
     #pragma acc cache(Arrs.MemArr[3:].array[4])
   }
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 8a439a5ccd4bd..fa43f42585fc2 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -499,7 +499,6 @@ void VarListClauses() {
 #pragma acc serial copy(HasMem.MemArr[1:3].array[1:2]), seq
   for(;;){}
 
-  // expected-error@+3{{expected expression}}
   // expected-warning@+2{{OpenACC clause 'copy' not yet implemented, clause ignored}}
   // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
 #pragma acc serial copy(HasMem.MemArr[:]), seq
@@ -519,7 +518,6 @@ void VarListClauses() {
 #pragma acc serial copy(HasMem.MemArr[: :]), seq
   for(;;){}
 
-  // expected-error@+3{{expected expression}}
   // expected-warning@+2{{OpenACC clause 'copy' not yet implemented, clause ignored}}
   // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
 #pragma acc serial copy(HasMem.MemArr[3:]), seq
diff --git a/clang/test/ParserOpenACC/parse-sub-array.cpp b/clang/test/ParserOpenACC/parse-sub-array.cpp
new file mode 100644
index 0000000000000..c0d3f89159e89
--- /dev/null
+++ b/clang/test/ParserOpenACC/parse-sub-array.cpp
@@ -0,0 +1,89 @@
+// RUN: %clang_cc1 %s -verify -fopenacc
+
+void Func(int i, int j) {
+  int array[5];
+#pragma acc parallel private(array[:])
+  while (true);
+#pragma acc parallel private(array[i:])
+  while (true);
+#pragma acc parallel private(array[:j])
+  while (true);
+#pragma acc parallel private(array[i:j])
+  while (true);
+#pragma acc parallel private(array[1:2])
+  while (true);
+
+  // expected-error@+1{{expected unqualified-id}}
+#pragma acc parallel private(array[::])
+  while (true);
+  // expected-error@+2{{expected ']'}}
+  // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[1::])
+  while (true);
+  // expected-error@+2{{expected ']'}}
+  // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[:2:])
+  while (true);
+  // expected-error@+3{{expected unqualified-id}}
+  // expected-error@+2{{expected ']'}}
+  // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[::3])
+  while (true);
+  // expected-error@+2{{expected ']'}}
+  // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[1:2:3])
+  while (true);
+}
+
+template<typename T, unsigned I, auto &IPtr>// #IPTR
+void TemplFunc() {
+  T array[I];
+  T array2[2*I];
+  T t; // #tDecl
+#pragma acc parallel private(array[:])
+  while (true);
+#pragma acc parallel private(array[t:])
+  while (true);
+#pragma acc parallel private(array[I-1:])
+  while (true);
+#pragma acc parallel private(array[IPtr:])
+  while (true);
+#pragma acc parallel private(array[:t])
+  while (true);
+#pragma acc parallel private(array[:I])
+  while (true);
+#pragma acc parallel private(array[:IPtr])
+  while (true);
+#pragma acc parallel private(array[t:t])
+  while (true);
+#pragma acc parallel private(array2[I:I])
+  while (true);
+#pragma acc parallel private(array[IPtr:IPtr])
+  while (true);
+
+  // expected-error@+1{{expected unqualified-id}}
+#pragma acc parallel private(array[::])
+  while (true);
+  // expected-error@+3{{'t' is not a class, namespace, or enumeration}}
+  // expected-note@#tDecl{{'t' declared here}}
+  // expected-error@+1{{expected unqualified-id}}
+#pragma acc parallel private(array[t::])
+  while (true);
+  // expected-error@+2{{expected ']'}}
+  // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[:I:])
+  while (true);
+  // expected-error@+2{{no member named 'IPtr' in the global namespace}}
+  // expected-note@#IPTR{{'IPtr' declared here}}
+#pragma acc parallel private(array[::IPtr])
+  while (true);
+  // expected-error@+2{{expected ']'}}
+  // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[IPtr:I:t])
+  while (true);
+}
+
+void use() {
+  static constexpr int SomeI = 1;
+  TemplFunc<int, 5, SomeI>();
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-private-clause.c b/clang/test/SemaOpenACC/compute-construct-private-clause.c
index 15775279fc869..d2615c384cdb1 100644
--- a/clang/test/SemaOpenACC/compute-construct-private-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-private-clause.c
@@ -12,12 +12,12 @@ typedef struct IsComplete {
 
 int GlobalInt;
 float GlobalArray[5];
-void *GlobalPointer;
+short *GlobalPointer;
 Complete GlobalComposite;
 
-void uses(int IntParam, void *PointerParam, float ArrayParam[5], Complete CompositeParam) {
+void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete CompositeParam) {
   int LocalInt;
-  void *LocalPointer;
+  short *LocalPointer;
   float LocalArray[5];
   Complete LocalComposite;
 
@@ -35,17 +35,13 @@ void uses(int IntParam, void *PointerParam, float ArrayParam[5], Complete Compos
   while(1);
 #pragma acc parallel private(LocalArray)
   while(1);
-  // TODO OpenACC: Fix array sections, this should be allowed.
-  // expected-error@+1{{expected expression}}
 #pragma acc parallel private(LocalArray[:])
   while(1);
 #pragma acc parallel private(LocalArray[:5])
   while(1);
-  // TODO OpenACC: Fix array sections, this should be allowed.
-  // expected-error@+1{{expected expression}}
 #pragma acc parallel private(LocalArray[2:])
   while(1);
-#pragma acc parallel private(LocalArray[2:5])
+#pragma acc parallel private(LocalArray[2:1])
   while(1);
 #pragma acc parallel private(LocalArray[2])
   while(1);
@@ -103,40 +99,36 @@ void uses(int IntParam, void *PointerParam, float ArrayParam[5], Complete Compos
 #pragma acc parallel private(+GlobalInt)
   while(1);
 
-  // TODO OpenACC: Fix array sections, this should be allowed.
-  // expected-error@+1{{expected expression}}
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
 #pragma acc parallel private(PointerParam[:])
   while(1);
 #pragma acc parallel private(PointerParam[:5])
   while(1);
 #pragma acc parallel private(PointerParam[:IntParam])
   while(1);
-  // TODO OpenACC: Fix array sections, this should be allowed.
-  // expected-error@+1{{expected expression}}
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
 #pragma acc parallel private(PointerParam[2:])
   while(1);
 #pragma acc parallel private(PointerParam[2:5])
   while(1);
 #pragma acc parallel private(PointerParam[2])
   while(1);
-  // TODO OpenACC: Fix array sections, this should be allowed.
-  // expected-error@+1{{expected expression}}
 #pragma acc parallel private(ArrayParam[:])
   while(1);
 #pragma acc parallel private(ArrayParam[:5])
   while(1);
 #pragma acc parallel private(ArrayParam[:IntParam])
   while(1);
-  // TODO OpenACC: Fix array sections, this should be allowed.
-  // expected-error@+1{{expected expression}}
 #pragma acc parallel private(ArrayParam[2:])
   while(1);
+  // expected-error@+1{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
 #pragma acc parallel private(ArrayParam[2:5])
   while(1);
 #pragma acc parallel private(ArrayParam[2])
   while(1);
 
-  // expected-error@+1{{OpenACC sub-array is not allowed here}}
+  // expected-error@+2{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
 #pragma acc parallel private((float*)ArrayParam[2:5])
   while(1);
   // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
diff --git a/clang/test/SemaOpenACC/compute-construct-private-clause.cpp b/clang/test/SemaOpenACC/compute-construct-private-clause.cpp
index 4dd4e0d8029d6..a776b16f0feb2 100644
--- a/clang/test/SemaOpenACC/compute-construct-private-clause.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-private-clause.cpp
@@ -112,8 +112,7 @@ void TemplUses(T t, T (&arrayT)[I], V TemplComp) {
   while(true);
 #pragma acc parallel private(Pointer[:t])
   while(true);
-  // TODO OpenACC: When fixing sub-arrays, this should be permitted}}
-  // expected-error@+1{{expected expression}}
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
 #pragma acc parallel private(Pointer[1:])
   while(true);
 }
diff --git a/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp b/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
index 341be3c58ebd2..aad4b9963c70d 100644
--- a/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
@@ -56,8 +56,11 @@ void NormalUses(float *PointerParam) {
   // CHECK-NEXT: private clause
   // CHECK-NEXT: DeclRefExpr{{.*}}'short[5]' lvalue Var{{.*}}'GlobalArray' 'short[5]'
   // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float *' <LValueToRValue>
   // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
   // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var{{.*}}'Global' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
   // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var{{.*}}'Global' 'int'
   // CHECK-NEXT: WhileStmt
   // CHECK-NEXT: CXXBoolLiteralExpr
@@ -241,8 +244,10 @@ void TemplUses(T t, U u, T*PointerParam) {
   // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
   // CHECK-NEXT: private clause
   // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <LValueToRValue>
   // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'u' 'int *'
   // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
   // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 't' 'int'
   // CHECK-NEXT: WhileStmt
   // CHECK-NEXT: CXXBoolLiteralExpr
@@ -290,6 +295,7 @@ struct S {
   // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
   // CHECK-NEXT: private clause
   // CHECK-NEXT: ArraySectionExpr{{.*}}
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
   // CHECK-NEXT: MemberExpr{{.*}} 'int[5]' lvalue ->ThisMemberArray
   // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' implicit this
   // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
@@ -331,6 +337,7 @@ struct S {
   // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
   // CHECK-NEXT: private clause
   // CHECK-NEXT: ArraySectionExpr{{.*}}
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
   // CHECK-NEXT: MemberExpr{{.*}} 'int[5]' lvalue ->ThisMemberArray
   // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' implicit this
   // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
@@ -372,6 +379,7 @@ void S::foo() {
   // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
   // CHECK-NEXT: private clause
   // CHECK-NEXT: ArraySectionExpr{{.*}}
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
   // CHECK-NEXT: MemberExpr{{.*}} 'int[5]' lvalue ->ThisMemberArray
   // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' implicit this
   // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
@@ -522,6 +530,7 @@ struct STempl {
   // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
   // CHECK-NEXT: private clause
   // CHECK-NEXT: ArraySectionExpr{{.*}}
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
   // CHECK-NEXT: MemberExpr{{.*}} 'int[5]' lvalue ->ThisMemberArray
   // CHECK-NEXT: CXXThisExpr{{.*}} 'STempl<int> *' implicit this
   // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
diff --git a/clang/test/SemaOpenACC/sub-array-ast.cpp b/clang/test/SemaOpenACC/sub-array-ast.cpp
new file mode 100644
index 0000000000000..094976e164275
--- /dev/null
+++ b/clang/test/SemaOpenACC/sub-array-ast.cpp
@@ -0,0 +1,566 @@
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+constexpr int returns_3() { return 3; }
+
+void Func(int i, int j) {
+  // CHECK: FunctionDecl{{.*}}Func
+  // CHECK-NEXT: ParmVarDecl{{.*}} i 'int'
+  // CHECK-NEXT: ParmVarDecl{{.*}} j 'int'
+  // CHECK-NEXT: CompoundStmt
+  int array[5];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} array 'int[5]'
+  int VLA[i];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} VLA 'int[i]'
+  int *ptr;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} ptr 'int *'
+
+#pragma acc parallel private(array[returns_3():])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[5]' lvalue Var{{.*}} 'array' 'int[5]'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(array[:1])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[5]' lvalue Var{{.*}} 'array' 'int[5]'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(array[returns_3():1])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[5]' lvalue Var{{.*}} 'array' 'int[5]'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(array[i:j])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[5]' lvalue Var{{.*}} 'array' 'int[5]'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'j' 'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(VLA[:1])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[i]' lvalue Var{{.*}} 'VLA' 'int[i]'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(VLA[returns_3():1])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[i]' lvalue Var{{.*}} 'VLA' 'int[i]'
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(VLA[i:j])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[i]' lvalue Var{{.*}} 'VLA' 'int[i]'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'j' 'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(ptr[:1])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue Var{{.*}} 'ptr' 'int *'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(ptr[returns_3():1])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue Var{{.*}} 'ptr' 'int *'
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(ptr[i:j])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue Var{{.*}} 'ptr' 'int *'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'j' 'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+}
+
+template<typename T, unsigned I, auto &CEArray>
+void Templ(int i){
+  // CHECK-NEXT: FunctionTemplateDecl{{.*}}Templ
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}} typename depth 0 index 0 T
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 I
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'auto &' depth 0 index 2 CEArray
+  // CHECK-NEXT: FunctionDecl{{.*}}Templ 'void (int)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} i 'int'
+  // CHECK-NEXT: CompoundStmt
+  T array[I+2];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} array 'T[I + 2]'
+  T VLA[i];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} VLA 'T[i]'
+  T *ptr;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} ptr 'T *'
+
+#pragma acc parallel private(array[returns_3():])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T[I + 2]' lvalue Var{{.*}} 'array' 'T[I + 2]'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}}'returns_3' 'int ()'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(array[:I])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T[I + 2]' lvalue Var{{.*}} 'array' 'T[I + 2]'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'I' 'unsigned int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(array[returns_3()-2:I])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T[I + 2]' lvalue Var{{.*}} 'array' 'T[I + 2]'
+  // CHECK-NEXT: BinaryOperator{{.*}} 'int' '-'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 2
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'I' 'unsigned int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(array[i:i])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T[I + 2]' lvalue Var{{.*}} 'array' 'T[I + 2]'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(VLA[:I])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T[i]' lvalue Var{{.*}} 'VLA' 'T[i]'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'I' 'unsigned int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(VLA[returns_3():I])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T[i]' lvalue Var{{.*}} 'VLA' 'T[i]'
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'I' 'unsigned int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(VLA[i:i])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T[i]' lvalue Var{{.*}} 'VLA' 'T[i]'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(ptr[:I])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T *' lvalue Var{{.*}} 'ptr' 'T *'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'I' 'unsigned int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(ptr[returns_3():I])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T *' lvalue Var{{.*}} 'ptr' 'T *'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'I' 'unsigned int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(ptr[i:i])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T *' lvalue Var{{.*}} 'ptr' 'T *'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(CEArray[returns_3() - 2: I])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'auto' lvalue NonTypeTemplateParm{{.*}} 'CEArray' 'auto &'
+  // CHECK-NEXT: BinaryOperator{{.*}} 'int' '-'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 2
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'I' 'unsigned int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel private(CEArray[: I])
+  while (true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'auto' lvalue NonTypeTemplateParm{{.*}} 'CEArray' 'auto &'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'I' 'unsigned int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+
+  // Instantiation:
+  // CHECK-NEXT: FunctionDecl{{.*}} Templ 'void (int)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument{{.*}} 'int'
+  // CHECK-NEXT: BuiltinType{{.*}} 'int'
+  // CHECK-NEXT: TemplateArgument integral 3
+  // CHECK-NEXT: TemplateArgument decl
+  // CHECK-NEXT: Var{{.*}} 'CEArray' 'const int[5]'
+  // CHECK-NEXT: ParmVarDecl{{.*}} i 'int'
+  // CHECK-NEXT: CompoundStmt
+
+  // T array[I+2];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} array 'int[5]'
+  // T VLA[i];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} VLA 'int[i]'
+  // T *ptr;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} ptr 'int *'
+
+//#pragma acc parallel private(array[returns_3():])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[5]' lvalue Var{{.*}} 'array' 'int[5]'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}}'returns_3' 'int ()'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(array[:I])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[5]' lvalue Var{{.*}} 'array' 'int[5]'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(array[returns_3()-2:I])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[5]' lvalue Var{{.*}} 'array' 'int[5]'
+  // CHECK-NEXT: BinaryOperator{{.*}} 'int' '-'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 2
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(array[i:i])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[5]' lvalue Var{{.*}} 'array' 'int[5]'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(VLA[:I])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[i]' lvalue Var{{.*}} 'VLA' 'int[i]'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(VLA[returns_3():I])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[i]' lvalue Var{{.*}} 'VLA' 'int[i]'
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(VLA[i:i])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int[i]' lvalue Var{{.*}} 'VLA' 'int[i]'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(ptr[:I])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue Var{{.*}} 'ptr' 'int *'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(ptr[returns_3():I])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue Var{{.*}} 'ptr' 'int *'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(ptr[i:i])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue Var{{.*}} 'ptr' 'int *'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(CEArray[returns_3() - 2: I])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'const int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'const int[5]' lvalue
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'auto &' depth 0 index 2 CEArray
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const int[5]' lvalue Var{{.*}}'CEArray' 'const int[5]'
+  // CHECK-NEXT: BinaryOperator{{.*}} 'int' '-'
+  // CHECK-NEXT: CallExpr{{.*}} 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int ()' lvalue Function{{.*}} 'returns_3' 'int ()'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 2
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+//#pragma acc parallel private(CEArray[: I])
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+  // CHECK-NEXT: private clause
+  // CHECK-NEXT: ArraySectionExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'const int *' <ArrayToPointerDecay>
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'const int[5]' lvalue
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'auto &' depth 0 index 2 CEArray
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const int[5]' lvalue Var{{.*}}'CEArray' 'const int[5]'
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+}
+
+// CHECK-NEXT: FunctionDecl{{.*}}inst
+void inst() {
+  static constexpr int CEArray[5]={1,2,3,4,5};
+  Templ<int, 3, CEArray>(5);
+}
+#endif
diff --git a/clang/test/SemaOpenACC/sub-array.cpp b/clang/test/SemaOpenACC/sub-array.cpp
new file mode 100644
index 0000000000000..355ac5ef1d3ce
--- /dev/null
+++ b/clang/test/SemaOpenACC/sub-array.cpp
@@ -0,0 +1,208 @@
+// RUN: %clang_cc1 %s -verify -fopenacc
+
+struct Incomplete; // #INCOMPLETE
+struct NotConvertible{} NC;
+
+struct CorrectConvert {
+  operator int();
+} Convert;
+
+constexpr int returns_3() { return 3; }
+
+using FuncPtrTy = void (*)();
+FuncPtrTy FuncPtrTyArray[2];
+
+void Func(int i, int j) {
+  int array[5];
+  int VLA[i];
+  int *ptr;
+  void *void_ptr;
+
+  // Follows int-expr rules, so only convertible to int.
+  // expected-error@+1{{OpenACC sub-array bound requires expression of integer type ('struct NotConvertible' invalid}}
+#pragma acc parallel private(array[NC:])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array bound requires expression of integer type ('struct NotConvertible' invalid}}
+#pragma acc parallel private(array[:NC])
+  while (true);
+
+  // expected-error@+2{{OpenACC sub-array bound requires expression of integer type ('struct NotConvertible' invalid}}
+  // expected-error@+1{{OpenACC sub-array bound requires expression of integer type ('struct NotConvertible' invalid}}
+#pragma acc parallel private(array[NC:NC])
+  while (true);
+
+  // expected-error@+2{{OpenACC sub-array bound requires expression of integer type ('struct NotConvertible' invalid}}
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel private(ptr[NC:])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array bound requires expression of integer type ('struct NotConvertible' invalid}}
+#pragma acc parallel private(ptr[:NC])
+  while (true);
+
+  // expected-error@+2{{OpenACC sub-array bound requires expression of integer type ('struct NotConvertible' invalid}}
+  // expected-error@+1{{OpenACC sub-array bound requires expression of integer type ('struct NotConvertible' invalid}}
+#pragma acc parallel private(ptr[NC:NC])
+  while (true);
+
+  // These are convertible, so they work.
+#pragma acc parallel private(array[Convert:Convert])
+  while (true);
+
+#pragma acc parallel private(ptr[Convert:Convert])
+  while (true);
+
+
+  // The length for "dynamically" allocated dimensions of an array must be
+  // explicitly specified.
+
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel private(ptr[3:])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is an array of unknown bound}}
+#pragma acc parallel private(VLA[3:])
+  while (true);
+
+#pragma acc parallel private(ptr[:3])
+  while (true);
+
+#pragma acc parallel private(VLA[:3])
+  while (true);
+
+  // Error if the length of the array + the initializer is bigger the the array
+  // with known bounds.
+
+  // expected-error@+1{{OpenACC sub-array length evaluated to a value (6) that would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel private(array[i:returns_3() + 3])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array length evaluated to a value (6) that would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel private(array[:returns_3() + 3])
+  while (true);
+
+#pragma acc parallel private(array[:returns_3()])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array specified range [3:3] would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel private(array[returns_3():returns_3()])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array lower bound evaluated to a value (6) that would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel private(array[returns_3() + 3:])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array lower bound evaluated to a value (6) that would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel private(array[returns_3() + 3:1])
+  while (true);
+
+  // Standard doesn't specify this, but negative values are likely not
+  // permitted, so disallow them here until we come up with a good reason to do
+  // otherwise.
+
+  // expected-error@+1{{OpenACC sub-array lower bound evaluated to negative value -1}}
+#pragma acc parallel private(array[returns_3() - 4 : ])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array length evaluated to negative value -1}}
+#pragma acc parallel private(array[: -1])
+  while (true);
+
+  Incomplete *IncompletePtr;
+  // expected-error@+2{{OpenACC sub-array base is of incomplete type 'Incomplete'}}
+  // expected-note@#INCOMPLETE{{forward declaration of 'Incomplete'}}
+#pragma acc parallel private(IncompletePtr[0 :1])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array base is of incomplete type 'void'}}
+#pragma acc parallel private(void_ptr[0:1])
+  while (true);
+
+  // OK: these are function pointers.
+#pragma acc parallel private(FuncPtrTyArray[0 :1])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array cannot be of function type 'void ()'}}
+#pragma acc parallel private(FuncPtrTyArray[0][0 :1])
+  while (true);
+
+
+  // expected-error@+1{{OpenACC sub-array subscripted value is not an array or pointer}}
+#pragma acc parallel private(i[0:1])
+  while (true);
+}
+
+template<typename T, typename U, typename V, unsigned I, auto &CEArray>
+void Templ(int i){
+  T array[I];
+  T VLA[i];
+  T *ptr;
+  U NC;
+  V Conv;
+
+  // Convertible:
+  // expected-error@+2{{OpenACC sub-array bound requires expression of integer type ('NotConvertible' invalid}}
+  // expected-note@#INST{{in instantiation of function template specialization}}
+#pragma acc parallel private(array[NC:])
+  while (true);
+  // expected-error@+1{{OpenACC sub-array bound requires expression of integer type ('NotConvertible' invalid}}
+#pragma acc parallel private(array[:NC])
+  while (true);
+
+#pragma acc parallel private(array[Conv:])
+  while (true);
+#pragma acc parallel private(array[:Conv])
+  while (true);
+
+  // Need a length for unknown size.
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel private(ptr[Conv:])
+  while (true);
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is an array of unknown bound}}
+#pragma acc parallel private(VLA[Conv:])
+  while (true);
+#pragma acc parallel private(ptr[:Conv])
+  while (true);
+#pragma acc parallel private(VLA[:Conv])
+  while (true);
+
+  // Out of bounds.
+  // expected-error@+1{{OpenACC sub-array lower bound evaluated to a value (2) that would be out of the range of the subscripted array size of 2}}
+#pragma acc parallel private(array[I:])
+  while (true);
+
+  // OK, don't know the value.
+#pragma acc parallel private(array[i:])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array length evaluated to a value (3) that would be out of the range of the subscripted array size of 2}}
+#pragma acc parallel private(array[:I + 1])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array lower bound evaluated to a value (5) that would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel private(CEArray[5:])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array length evaluated to a value (6) that would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel private(CEArray[:2 + I + I])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array length evaluated to a value (4294967295) that would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel private(CEArray[:1 - I])
+  while (true);
+
+  // expected-error@+1{{OpenACC sub-array lower bound evaluated to a value (4294967295) that would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel private(CEArray[1 - I:])
+  while (true);
+
+  T not_ptr;
+  // expected-error@+1{{OpenACC sub-array subscripted value is not an array or pointer}}
+#pragma acc parallel private(not_ptr[0:1])
+  while (true);
+}
+
+void inst() {
+  static constexpr int CEArray[5]={1,2,3,4,5};
+  Templ<int, NotConvertible, CorrectConvert, 2, CEArray>(5); // #INST
+}