[OpenACC] Add warning for VLAs in a private/firstprivate clause

private/firstprivate typically do copy operations, however copying a VLA
isn't really possible.  This patch introduces a warning to alert the
person that this copy isn't happening correctly.

As a future direction, we MIGHT consider doing additional work to make
sure they are initialized/copied/deleted/etc correctly.

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13531,6 +13531,11 @@ def err_acc_invalid_default_type
 def err_acc_device_type_multiple_archs
     : Error<"OpenACC 'device_type' clause on a 'set' construct only permits "
             "one architecture">;
+def warn_acc_var_referenced_non_const_array
+    : Warning<"variable of array type %0 referenced in OpenACC '%1' clause "
+              "does not have constant bounds; initialization will happen after "
+              "decay to pointer">,
+      InGroup<DiagGroup<"openacc-var-non-const-array">>;
 def warn_acc_var_referenced_lacks_op
     : Warning<"variable of type %0 referenced in OpenACC '%1' clause does not "
               "have a %enum_select<AccVarReferencedReason>{%DefCtor{default "

clang/lib/Sema/SemaOpenACC.cpp

@@ -646,8 +646,17 @@ ExprResult CheckVarType(SemaOpenACC &S, OpenACCClauseKind CK, Expr *VarExpr,
   if (auto *RefTy = InnerTy->getAs<ReferenceType>())
     InnerTy = RefTy->getPointeeType();
 
-  if (auto *ArrTy = InnerTy->getAsArrayTypeUnsafe())
+  if (auto *ArrTy = InnerTy->getAsArrayTypeUnsafe()) {
+    // Non constant arrays decay to 'pointer', so warn and return that we're
+    // successful.
+    if (!ArrTy->isConstantArrayType()) {
+      S.Diag(InnerLoc, clang::diag::warn_acc_var_referenced_non_const_array)
+          << InnerTy << CK;
+      return VarExpr;
+    }
+
     return CheckVarType(S, CK, VarExpr, InnerLoc, ArrTy->getElementType());
+  }
 
   auto *RD = InnerTy->getAsCXXRecordDecl();
 

clang/test/SemaOpenACC/private_firstprivate_reduction_required_ops.cpp

@@ -275,3 +275,20 @@ void firstprivate_arrays() {
 #pragma acc parallel firstprivate(UDCopyArr)
   ;
 }
+
+template<unsigned I>
+void non_const_array_templ() {
+  int CArr[I];
+
+#pragma acc parallel firstprivate(CArr)
+  ;
+}
+
+void non_const_arrays(int I) {
+  non_const_array_templ<5>();
+
+  int NCArr[I];
+  // expected-warning@+1{{variable of array type 'int[I]' referenced in OpenACC 'firstprivate' clause does not have constant bounds; initialization will happen after decay to pointer}}
+#pragma acc parallel firstprivate(NCArr)
+  ;
+}

clang/test/SemaOpenACC/sub-array.cpp

@@ -61,13 +61,15 @@ void Func(int i, int j) {
 #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}}
+  // expected-error@+2{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is an array of unknown bound}}
+  // expected-warning@+1{{variable of array type 'int[i]' referenced in OpenACC 'private' clause does not have constant bounds; initialization will happen after decay to pointer}}
 #pragma acc parallel private(VLA[3:])
   while (true);
 
 #pragma acc parallel private(ptr[:3])
   while (true);
 
+  // expected-warning@+1{{variable of array type 'int[i]' referenced in OpenACC 'private' clause does not have constant bounds; initialization will happen after decay to pointer}}
 #pragma acc parallel private(VLA[:3])
   while (true);
 
@@ -159,11 +161,13 @@ void Templ(int i){
   // 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}}
+  // expected-error@+2{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is an array of unknown bound}}
+  // expected-warning@+1{{variable of array type 'int[i]' referenced in OpenACC 'private' clause does not have constant bounds; initialization will happen after decay to pointer}}
 #pragma acc parallel private(VLA[Conv:])
   while (true);
 #pragma acc parallel private(ptr[:Conv])
   while (true);
+  // expected-warning@+1{{variable of array type 'int[i]' referenced in OpenACC 'private' clause does not have constant bounds; initialization will happen after decay to pointer}}
 #pragma acc parallel private(VLA[:Conv])
   while (true);