Author: erichkeane
Date: 2025-07-03T07:13:30-07:00
New Revision: 438863a09e2d907c36c4114228ef1eeee130e16b

URL: 
https://github.com/llvm/llvm-project/commit/438863a09e2d907c36c4114228ef1eeee130e16b
DIFF: 
https://github.com/llvm/llvm-project/commit/438863a09e2d907c36c4114228ef1eeee130e16b.diff

LOG: [OpenACC][Sema] Implement warning for 'cache' invalid var ref

The 'cache' construct is lowered as marking the acc.loop in ACC MLIR.
This results in any variable references that are not inside of the
acc.loop being invalid.  This patch adds a warning to that effect, and
ensures that the variable references won't be added to the AST during
parsing so we don't try to lower them.

This results in loss of instantiation-diagnostics for these, however
that seems like an acceptable consequence to ignoring it.

Added: 
    clang/test/SemaOpenACC/cache-warn-invalid-varloc.cpp

Modified: 
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/include/clang/Sema/Scope.h
    clang/include/clang/Sema/SemaOpenACC.h
    clang/lib/Parse/ParseOpenACC.cpp
    clang/lib/Sema/Scope.cpp
    clang/lib/Sema/SemaExpr.cpp
    clang/lib/Sema/SemaOpenACC.cpp
    clang/test/AST/ast-print-openacc-cache-construct.cpp
    clang/test/ParserOpenACC/parse-cache-construct.c
    clang/test/ParserOpenACC/parse-cache-construct.cpp
    clang/test/SemaOpenACC/cache-construct-ast.cpp
    clang/test/SemaOpenACC/cache-construct.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td 
b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index fac63a500c98f..36f279fb663bb 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13217,8 +13217,12 @@ def err_acc_not_a_var_ref_use_device_declare
     : Error<"OpenACC variable %select{in 'use_device' clause|on 'declare' "
             "construct}0 is not a valid variable name or array name">;
 def err_acc_not_a_var_ref_cache
-    : Error<"OpenACC variable in cache directive is not a valid sub-array or "
+    : Error<"OpenACC variable in 'cache' directive is not a valid sub-array or 
"
             "array element">;
+def warn_acc_cache_var_not_outside_loop
+    : Warning<"OpenACC variable in 'cache' directive was not declared outside "
+              "of the associated 'loop' directive; directive has no effect">,
+      InGroup<DiagGroup<"openacc-cache-var-inside-loop">>;
 def err_acc_typecheck_subarray_value
     : Error<"OpenACC sub-array subscripted value is not an array or pointer">;
 def err_acc_subarray_function_type

diff  --git a/clang/include/clang/Sema/Scope.h 
b/clang/include/clang/Sema/Scope.h
index 07b9e1bc10f5a..95f671c8387c2 100644
--- a/clang/include/clang/Sema/Scope.h
+++ b/clang/include/clang/Sema/Scope.h
@@ -158,11 +158,15 @@ class Scope {
     /// constructs, since they have the same behavior.
     OpenACCComputeConstructScope = 0x10000000,
 
+    /// This is the scope of an OpenACC Loop/Combined construct, which is used
+    /// to determine whether a 'cache' construct variable reference is legal.
+    OpenACCLoopConstructScope = 0x20000000,
+
     /// This is a scope of type alias declaration.
-    TypeAliasScope = 0x20000000,
+    TypeAliasScope = 0x40000000,
 
     /// This is a scope of friend declaration.
-    FriendScope = 0x40000000,
+    FriendScope = 0x80000000,
   };
 
 private:
@@ -549,6 +553,10 @@ class Scope {
     return getFlags() & Scope::OpenACCComputeConstructScope;
   }
 
+  bool isOpenACCLoopConstructScope() const {
+    return getFlags() & Scope::OpenACCLoopConstructScope;
+  }
+
   /// Determine if this scope (or its parents) are a compute construct. If the
   /// argument is provided, the search will stop at any of the specified 
scopes.
   /// Otherwise, it will stop only at the normal 'no longer search' scopes.

diff  --git a/clang/include/clang/Sema/SemaOpenACC.h 
b/clang/include/clang/Sema/SemaOpenACC.h
index c4a986a3eb6f9..b7e7f5d97bcef 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -117,6 +117,15 @@ class SemaOpenACC : public SemaBase {
     OpenACCDirectiveKind DirectiveKind = OpenACCDirectiveKind::Invalid;
   } TileInfo;
 
+  /// The 'cache' var-list requires some additional work to track variable
+  /// references to make sure they are on the 'other' side of a `loop`. This
+  /// structure is used during parse time to track vardecl use while parsing a
+  /// cache var list.
+  struct CacheParseInfo {
+    bool ParsingCacheVarList = false;
+    bool IsInvalidCacheRef = false;
+  } CacheInfo;
+
   /// A list of the active reduction clauses, which allows us to check that all
   /// vars on nested constructs for the same reduction var have the same
   /// reduction operator. Currently this is enforced against all constructs
@@ -861,6 +870,12 @@ class SemaOpenACC : public SemaBase {
   ExprResult ActOnIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
                           SourceLocation Loc, Expr *IntExpr);
 
+  /// Called right before a 'var' is parsed, so we can set the state for 
parsing
+  /// a 'cache' var.
+  void ActOnStartParseVar(OpenACCDirectiveKind DK, OpenACCClauseKind CK);
+  /// Called only if the parse of a 'var' was invalid, else 'ActOnVar' should 
be
+  /// called.
+  void ActOnInvalidParseVar();
   /// Called when encountering a 'var' for OpenACC, ensures it is actually a
   /// declaration reference to a variable of the correct type.
   ExprResult ActOnVar(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
@@ -913,6 +928,10 @@ class SemaOpenACC : public SemaBase {
                            OpenACCDirectiveKind DK, OpenACCGangKind GK,
                            Expr *E);
 
+  // Called when a declaration is referenced, so that we can make sure certain
+  // clauses don't do the 'wrong' thing/have incorrect references.
+  void CheckDeclReference(SourceLocation Loc, Expr *E, Decl *D);
+
   // Does the checking for a 'gang' clause that needs to be done in dependent
   // and not dependent cases.
   OpenACCClause *

diff  --git a/clang/lib/Parse/ParseOpenACC.cpp 
b/clang/lib/Parse/ParseOpenACC.cpp
index f2849c4eac7cc..b79a956d51505 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -606,13 +606,20 @@ unsigned getOpenACCScopeFlags(OpenACCDirectiveKind 
DirKind) {
   case OpenACCDirectiveKind::Parallel:
   case OpenACCDirectiveKind::Serial:
   case OpenACCDirectiveKind::Kernels:
+    // Mark this as a BreakScope/ContinueScope as well as a compute construct
+    // so that we can diagnose trying to 'break'/'continue' inside of one.
+    return Scope::BreakScope | Scope::ContinueScope |
+           Scope::OpenACCComputeConstructScope;
   case OpenACCDirectiveKind::ParallelLoop:
   case OpenACCDirectiveKind::SerialLoop:
   case OpenACCDirectiveKind::KernelsLoop:
     // Mark this as a BreakScope/ContinueScope as well as a compute construct
     // so that we can diagnose trying to 'break'/'continue' inside of one.
     return Scope::BreakScope | Scope::ContinueScope |
-           Scope::OpenACCComputeConstructScope;
+           Scope::OpenACCComputeConstructScope |
+           Scope::OpenACCLoopConstructScope;
+  case OpenACCDirectiveKind::Loop:
+    return Scope::OpenACCLoopConstructScope;
   case OpenACCDirectiveKind::Data:
   case OpenACCDirectiveKind::EnterData:
   case OpenACCDirectiveKind::ExitData:
@@ -621,7 +628,6 @@ unsigned getOpenACCScopeFlags(OpenACCDirectiveKind DirKind) 
{
   case OpenACCDirectiveKind::Init:
   case OpenACCDirectiveKind::Shutdown:
   case OpenACCDirectiveKind::Cache:
-  case OpenACCDirectiveKind::Loop:
   case OpenACCDirectiveKind::Atomic:
   case OpenACCDirectiveKind::Declare:
   case OpenACCDirectiveKind::Routine:
@@ -1416,12 +1422,15 @@ Parser::OpenACCVarParseResult 
Parser::ParseOpenACCVar(OpenACCDirectiveKind DK,
                                                       OpenACCClauseKind CK) {
   OpenACCArraySectionRAII ArraySections(*this);
 
+  getActions().OpenACC().ActOnStartParseVar(DK, CK);
   ExprResult Res = ParseAssignmentExpression();
-  if (!Res.isUsable())
+
+  if (!Res.isUsable()) {
+    getActions().OpenACC().ActOnInvalidParseVar();
     return {Res, OpenACCParseCanContinue::Cannot};
+  }
 
   Res = getActions().OpenACC().ActOnVar(DK, CK, Res.get());
-
   return {Res, OpenACCParseCanContinue::Can};
 }
 

diff  --git a/clang/lib/Sema/Scope.cpp b/clang/lib/Sema/Scope.cpp
index 5bc7e79a68186..ab04fe554be82 100644
--- a/clang/lib/Sema/Scope.cpp
+++ b/clang/lib/Sema/Scope.cpp
@@ -234,7 +234,8 @@ void Scope::dumpImpl(raw_ostream &OS) const {
       {OpenACCComputeConstructScope, "OpenACCComputeConstructScope"},
       {TypeAliasScope, "TypeAliasScope"},
       {FriendScope, "FriendScope"},
-  };
+      {OpenACCComputeConstructScope, "OpenACCComputeConstructScope"},
+      {OpenACCLoopConstructScope, "OpenACCLoopConstructScope"}};
 
   for (auto Info : FlagInfo) {
     if (Flags & Info.first) {

diff  --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 437df742d572b..03dca6f7cb878 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -20169,6 +20169,9 @@ MarkExprReferenced(Sema &SemaRef, SourceLocation Loc, 
Decl *D, Expr *E,
   if (SemaRef.OpenMP().isInOpenMPDeclareTargetContext())
     SemaRef.OpenMP().checkDeclIsAllowedInOpenMPTarget(E, D);
 
+  if (SemaRef.getLangOpts().OpenACC)
+    SemaRef.OpenACC().CheckDeclReference(Loc, E, D);
+
   if (VarDecl *Var = dyn_cast<VarDecl>(D)) {
     DoMarkVarDeclReferenced(SemaRef, Loc, Var, E, RefsMinusAssignments);
     if (SemaRef.getLangOpts().CPlusPlus)

diff  --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index e63d75c934f21..46aa7dd0dcc21 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -519,8 +519,31 @@ bool SemaOpenACC::CheckVarIsPointerType(OpenACCClauseKind 
ClauseKind,
   return false;
 }
 
+void SemaOpenACC::ActOnStartParseVar(OpenACCDirectiveKind DK,
+                                     OpenACCClauseKind CK) {
+  if (DK == OpenACCDirectiveKind::Cache) {
+    CacheInfo.ParsingCacheVarList = true;
+    CacheInfo.IsInvalidCacheRef = false;
+  }
+}
+
+void SemaOpenACC::ActOnInvalidParseVar() {
+  CacheInfo.ParsingCacheVarList = false;
+  CacheInfo.IsInvalidCacheRef = false;
+}
+
 ExprResult SemaOpenACC::ActOnCacheVar(Expr *VarExpr) {
   Expr *CurVarExpr = VarExpr->IgnoreParenImpCasts();
+  // Clear this here, so we can do the returns based on the invalid cache ref
+  // here.  Note all return statements in this function must return ExprError 
if
+  // IsInvalidCacheRef. However, instead of doing an 'early return' in that
+  // case, we can let the rest of the diagnostics happen, as the invalid decl
+  // ref is a warning.
+  bool WasParsingInvalidCacheRef =
+      CacheInfo.ParsingCacheVarList && CacheInfo.IsInvalidCacheRef;
+  CacheInfo.ParsingCacheVarList = false;
+  CacheInfo.IsInvalidCacheRef = false;
+
   if (!isa<ArraySectionExpr, ArraySubscriptExpr>(CurVarExpr)) {
     Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_cache);
     return ExprError();
@@ -540,19 +563,19 @@ ExprResult SemaOpenACC::ActOnCacheVar(Expr *VarExpr) {
   if (const auto *DRE = dyn_cast<DeclRefExpr>(CurVarExpr)) {
     if (isa<VarDecl, NonTypeTemplateParmDecl>(
             DRE->getFoundDecl()->getCanonicalDecl()))
-      return VarExpr;
+      return WasParsingInvalidCacheRef ? ExprEmpty() : VarExpr;
   }
 
   if (const auto *ME = dyn_cast<MemberExpr>(CurVarExpr)) {
     if (isa<FieldDecl>(ME->getMemberDecl()->getCanonicalDecl())) {
-      return VarExpr;
+      return WasParsingInvalidCacheRef ? ExprEmpty() : VarExpr;
     }
   }
 
   // Nothing really we can do here, as these are dependent.  So just return 
they
   // are valid.
   if (isa<DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr>(CurVarExpr))
-    return VarExpr;
+    return WasParsingInvalidCacheRef ? ExprEmpty() : VarExpr;
 
   // There isn't really anything we can do in the case of a recovery expr, so
   // skip the diagnostic rather than produce a confusing diagnostic.
@@ -562,6 +585,45 @@ ExprResult SemaOpenACC::ActOnCacheVar(Expr *VarExpr) {
   Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_cache);
   return ExprError();
 }
+
+void SemaOpenACC::CheckDeclReference(SourceLocation Loc, Expr *E, Decl *D) {
+  if (!getLangOpts().OpenACC || !CacheInfo.ParsingCacheVarList || !D ||
+      D->isInvalidDecl())
+    return;
+  // A 'cache' variable reference MUST be declared before the 'acc.loop' we
+  // generate in codegen, so we have to mark it invalid here in some way.  We 
do
+  // so in a bit of a convoluted way as there is no good way to put this into
+  // the AST, so we store it in SemaOpenACC State.  We can check the Scope
+  // during parsing to make sure there is a 'loop' before the decl is
+  // declared(and skip during instantiation).
+  // We only diagnose this as a warning, as this isn't required by the standard
+  // (unless you take a VERY awkward reading of some awkward prose).
+
+  Scope *CurScope = SemaRef.getCurScope();
+
+  // if we are at TU level, we are either doing some EXTRA wacky, or are in a
+  // template instantiation, so just give up.
+  if (CurScope->getDepth() == 0)
+    return;
+
+  while (CurScope) {
+    // If we run into a loop construct scope, than this is 'correct' in that 
the
+    // declaration is outside of the loop.
+    if (CurScope->isOpenACCLoopConstructScope())
+      return;
+
+    if (CurScope->isDeclScope(D)) {
+      Diag(Loc, diag::warn_acc_cache_var_not_outside_loop);
+
+      CacheInfo.IsInvalidCacheRef = true;
+    }
+
+    CurScope = CurScope->getParent();
+  }
+  // If we don't find the decl at all, we assume that it must be outside of the
+  // loop (or we aren't in a loop!) so skip the diagnostic.
+}
+
 ExprResult SemaOpenACC::ActOnVar(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
                                  Expr *VarExpr) {
   // This has unique enough restrictions that we should split it to a separate

diff  --git a/clang/test/AST/ast-print-openacc-cache-construct.cpp 
b/clang/test/AST/ast-print-openacc-cache-construct.cpp
index e98ff3ca67600..26dd1333ee9ed 100644
--- a/clang/test/AST/ast-print-openacc-cache-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-cache-construct.cpp
@@ -2,8 +2,11 @@
 
 void foo() {
   int Array[5];
+#pragma acc loop
+  for(int i = 0; i < 5; ++i) {
   // CHECK: #pragma acc cache(readonly: Array[1], Array[1:2])
   #pragma acc cache(readonly:Array[1], Array[1:2])
   // CHECK: #pragma acc cache(Array[1], Array[1:2])
   #pragma acc cache(Array[1], Array[1:2])
+  }
 }

diff  --git a/clang/test/ParserOpenACC/parse-cache-construct.c 
b/clang/test/ParserOpenACC/parse-cache-construct.c
index d52a5e4ea3441..76d6b32a574e8 100644
--- a/clang/test/ParserOpenACC/parse-cache-construct.c
+++ b/clang/test/ParserOpenACC/parse-cache-construct.c
@@ -46,10 +46,11 @@ void func() {
     #pragma acc cache(invalid
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+3{{expected ')'}}
     // expected-note@+2{{to match this '('}}
-    // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+    // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
     #pragma acc cache(ArrayPtr
   }
 
@@ -58,11 +59,13 @@ void func() {
     #pragma acc cache(invalid)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
-    // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+    // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
     #pragma acc cache(ArrayPtr)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+5{{expected expression}}
     // expected-error@+4{{expected ']'}}
@@ -72,6 +75,7 @@ void func() {
     #pragma acc cache(ArrayPtr[
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+3{{expected expression}}
     // expected-error@+2{{expected ']'}}
@@ -79,6 +83,7 @@ void func() {
     #pragma acc cache(ArrayPtr[, 5)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+3{{expected expression}}
     // expected-error@+2{{expected ']'}}
@@ -86,10 +91,12 @@ void func() {
     #pragma acc cache(Array[)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     #pragma acc cache(Array[*readonly])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+5{{expected expression}}
     // expected-error@+4{{expected ']'}}
@@ -99,49 +106,59 @@ void func() {
     #pragma acc cache(Array[*readonly:
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
-    // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+    // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
     #pragma acc cache(readonly)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+2{{invalid tag 'devnum' on 'cache' directive}}
-    // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+    // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
     #pragma acc cache(devnum:ArrayPtr)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+2{{invalid tag 'invalid' on 'cache' directive}}
-    // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+    // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
     #pragma acc cache(invalid:ArrayPtr)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
-    // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+    // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
     #pragma acc cache(readonly:ArrayPtr)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     #pragma acc cache(readonly:ArrayPtr[5:1])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     #pragma acc cache(readonly:ArrayPtr[5:*readonly])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
-    // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+    // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
     #pragma acc cache(readonly:ArrayPtr[5:*readonly], Array)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     #pragma acc cache(readonly:ArrayPtr[5:*readonly], Array[*readonly:3])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
-    #pragma acc cache(readonly:ArrayPtr[5 + i:*readonly], Array[*readonly + 
i:3])
+    #pragma acc cache(readonly:ArrayPtr[5 + 0:*readonly], Array[*readonly + 
0:3])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+3{{expected expression}}
     // expected-error@+2{{expected ')'}}
@@ -149,26 +166,31 @@ void func() {
     #pragma acc cache(readonly:ArrayPtr[5:*readonly],
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+1{{expected expression}}
     #pragma acc cache(readonly:ArrayPtr[5:*readonly],)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-warning@+1{{left operand of comma operator has no effect}}
     #pragma acc cache(readonly:ArrayPtr[5,6:*readonly])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-warning@+1{{left operand of comma operator has no effect}}
     #pragma acc cache(readonly:ArrayPtr[5:3, *readonly], ArrayPtr[0])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
-  // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+  // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
     #pragma acc cache(readonly:s.foo)
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-warning@+1{{left operand of comma operator has no effect}}
     #pragma acc cache(readonly:s.Array[1,2])

diff  --git a/clang/test/ParserOpenACC/parse-cache-construct.cpp 
b/clang/test/ParserOpenACC/parse-cache-construct.cpp
index 948f2e30f149c..fe794707bb7da 100644
--- a/clang/test/ParserOpenACC/parse-cache-construct.cpp
+++ b/clang/test/ParserOpenACC/parse-cache-construct.cpp
@@ -8,28 +8,34 @@ char *getArrayPtr();
 template<typename T, int I>
 void func() {
   char *ArrayPtr = getArrayPtr();
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-warning@+1{{left operand of comma operator has no effect}}
-    #pragma acc cache(ArrayPtr[T::value + I:I + 3], T::array[(i + T::value, 
2): 4])
+    #pragma acc cache(ArrayPtr[T::value + I:I + 3], T::array[(T::value, 2): 2])
   }
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     #pragma acc cache(NS::NSArray[NS::NSInt])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     #pragma acc cache(NS::NSArray[NS::NSInt : NS::NSInt])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+1{{use of undeclared identifier 'NSArray'}}
     #pragma acc cache(NSArray[NS::NSInt : NS::NSInt])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+1{{use of undeclared identifier 'NSInt'}}
     #pragma acc cache(NS::NSArray[NSInt : NS::NSInt])
   }
 
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+1{{use of undeclared identifier 'NSInt'}}
     #pragma acc cache(NS::NSArray[NS::NSInt : NSInt])
@@ -53,43 +59,53 @@ struct HasMembersArray {
 void use() {
 
   Members s;
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     #pragma acc cache(s.array[s.value])
   }
   HasMembersArray Arrs;
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     #pragma acc cache(Arrs.MemArr[3].array[4])
   }
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     #pragma acc cache(Arrs.MemArr[3].array[1:4])
   }
+#pragma acc loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+1{{OpenACC sub-array is not allowed here}}
     #pragma acc cache(Arrs.MemArr[2:1].array[1:4])
   }
+#pragma acc parallel loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+1{{OpenACC sub-array is not allowed here}}
     #pragma acc cache(Arrs.MemArr[2:1].array[4])
   }
+#pragma acc parallel loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+2{{expected ']'}}
     // expected-note@+1{{to match this '['}}
     #pragma acc cache(Arrs.MemArr[3:4:].array[4])
   }
+#pragma acc parallel loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+1{{OpenACC sub-array is not allowed here}}
     #pragma acc cache(Arrs.MemArr[:].array[4])
   }
+#pragma acc parallel loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+1{{expected unqualified-id}}
     #pragma acc cache(Arrs.MemArr[::].array[4])
   }
+#pragma acc parallel loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+3{{expected expression}}
     // expected-error@+2{{expected ']'}}
     // expected-note@+1{{to match this '['}}
     #pragma acc cache(Arrs.MemArr[: :].array[4])
   }
+#pragma acc parallel loop
   for (int i = 0; i < 10; ++i) {
     // expected-error@+1{{OpenACC sub-array is not allowed here}}
     #pragma acc cache(Arrs.MemArr[3:].array[4])

diff  --git a/clang/test/SemaOpenACC/cache-construct-ast.cpp 
b/clang/test/SemaOpenACC/cache-construct-ast.cpp
index 44c80beed77a3..fd8cd049840c2 100644
--- a/clang/test/SemaOpenACC/cache-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/cache-construct-ast.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+// RUN: %clang_cc1 %s -fopenacc -Wno-openacc-cache-var-inside-loop -ast-dump | 
FileCheck %s
 
 // Test this with PCH.
 // RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
@@ -14,6 +14,21 @@ void use() {
   // CHECK-NEXT: DeclStmt
   // CHECK-NEXT: VarDecl{{.*}}Array 'int[5]'
 
+#pragma acc loop
+  // CHECK-NEXT: OpenACCLoopConstruct
+  for (int i = 0; i < 5; ++i) {
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}}i 'int'
+  // CHECK-NEXT: IntegerLiteral
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr
+  // CHECK-NEXT: IntegerLiteral
+  // CHECK-NEXT: UnaryOperator
+  // CHECK-NEXT: DeclRefExpr
+  // CHECK-NEXT: CompoundStmt
 #pragma acc cache(Array[1])
   // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
   // CHECK-NEXT: ArraySubscriptExpr{{.*}}'int' lvalue
@@ -33,6 +48,18 @@ _Pragma("acc cache(Array[1])")
   // CHECK-NEXT: DeclRefExpr{{.*}}'Array' 'int[5]'
   // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
   // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2
+
+  // No vars here, since they are bad references.
+#pragma acc cache (Array[i])
+  // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
+  }
+
+  // No vars here, since they are bad references.
+#pragma acc cache(Array[1:2])
+  // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
+
+  return;
+  // CHECK-NEXT: ReturnStmt
 }
 
 struct S {
@@ -42,6 +69,21 @@ struct S {
   void StructUse() {
     // CHECK: CXXMethodDecl{{.*}}StructUse 'void ()'
     // CHECK-NEXT: CompoundStmt
+#pragma acc loop
+  // CHECK-NEXT: OpenACCLoopConstruct
+  for (int i = 0; i < 5; ++i) {
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}}i 'int'
+  // CHECK-NEXT: IntegerLiteral
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr
+  // CHECK-NEXT: IntegerLiteral
+  // CHECK-NEXT: UnaryOperator
+  // CHECK-NEXT: DeclRefExpr
+  // CHECK-NEXT: CompoundStmt
 #pragma acc cache(Array[1])
     // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
     // CHECK-NEXT: ArraySubscriptExpr{{.*}}'int' lvalue
@@ -79,6 +121,7 @@ struct S {
     // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
     // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2
   }
+  }
 };
 
 template<typename T>
@@ -89,6 +132,21 @@ void templ_use() {
   // CHECK-NEXT: DeclStmt
   // CHECK-NEXT: VarDecl{{.*}}Array 'T[5]'
 
+#pragma acc loop
+  // CHECK-NEXT: OpenACCLoopConstruct
+  for (int i = 0; i < 5; ++i) {
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}}i 'int'
+  // CHECK-NEXT: IntegerLiteral
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr
+  // CHECK-NEXT: IntegerLiteral
+  // CHECK-NEXT: UnaryOperator
+  // CHECK-NEXT: DeclRefExpr
+  // CHECK-NEXT: CompoundStmt
 #pragma acc cache(Array[1])
   // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
   // CHECK-NEXT: ArraySubscriptExpr{{.*}}'T' lvalue
@@ -101,6 +159,17 @@ void templ_use() {
   // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
   // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2
 
+    // No vars here, since they are bad references.
+#pragma acc cache(Array[i])
+  // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
+  }
+  // No vars here, since they are bad references.
+#pragma acc cache(Array[1])
+  // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
+
+  return;
+  // CHECK-NEXT: ReturnStmt
+
   // Instantiation:
   // CHECK: FunctionDecl{{.*}} templ_use 'void ()' implicit_instantiation
   // CHECK-NEXT: TemplateArgument type 'int'
@@ -108,6 +177,21 @@ void templ_use() {
   // CHECK-NEXT: CompoundStmt
   // CHECK-NEXT: DeclStmt
   // CHECK-NEXT: VarDecl{{.*}}Array 'int[5]'
+
+  // CHECK-NEXT: OpenACCLoopConstruct
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}}i 'int'
+  // CHECK-NEXT: IntegerLiteral
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr
+  // CHECK-NEXT: IntegerLiteral
+  // CHECK-NEXT: UnaryOperator
+  // CHECK-NEXT: DeclRefExpr
+  // CHECK-NEXT: CompoundStmt
+  //
   // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
   // CHECK-NEXT: ArraySubscriptExpr{{.*}}'int' lvalue
   // CHECK-NEXT: ImplicitCastExpr{{.*}}'int *' <ArrayToPointerDecay>
@@ -119,6 +203,11 @@ void templ_use() {
   // CHECK-NEXT: DeclRefExpr{{.*}}'Array' 'int[5]'
   // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
   // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2
+
+  // No vars here, since they are bad references.
+  // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
+  // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache
+  // CHECK-NEXT: ReturnStmt
 }
 
 void foo() {

diff  --git a/clang/test/SemaOpenACC/cache-construct.cpp 
b/clang/test/SemaOpenACC/cache-construct.cpp
index 0a32754c82710..3196cd010ceae 100644
--- a/clang/test/SemaOpenACC/cache-construct.cpp
+++ b/clang/test/SemaOpenACC/cache-construct.cpp
@@ -5,13 +5,16 @@ void use() {
   int Array[5];
   int NotArray;
 
+#pragma acc loop
+  for (int i = 0; i < 5;++i) {
 #pragma acc cache(Array[1])
 #pragma acc cache(Array[1:2])
 
-  // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+  // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
 #pragma acc cache(Array)
-  // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+  // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
 #pragma acc cache(NotArray)
+  }
 }
 
 struct S {
@@ -20,16 +23,19 @@ struct S {
   int Array2D[5][5];
 
   void use() {
+#pragma acc loop
+  for (int i = 0; i < 5;++i) {
 #pragma acc cache(Array[1])
 #pragma acc cache(Array[1:2])
 #pragma acc cache(Array2D[1][1])
 #pragma acc cache(Array2D[1][1:2])
 
-  // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+  // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
 #pragma acc cache(Array)
-  // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+  // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
 #pragma acc cache(NotArray)
   }
+  }
 };
 
 template<typename T>
@@ -37,13 +43,16 @@ void templ_use() {
   T Array[5];
   T NotArray;
 
+#pragma acc loop
+  for (int i = 0; i < 5;++i) {
 #pragma acc cache(Array[1])
 #pragma acc cache(Array[1:2])
 
-  // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+  // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
 #pragma acc cache(Array)
-  // expected-error@+1{{OpenACC variable in cache directive is not a valid 
sub-array or array element}}
+  // expected-error@+1{{OpenACC variable in 'cache' directive is not a valid 
sub-array or array element}}
 #pragma acc cache(NotArray)
+  }
 }
 
 void foo() {

diff  --git a/clang/test/SemaOpenACC/cache-warn-invalid-varloc.cpp 
b/clang/test/SemaOpenACC/cache-warn-invalid-varloc.cpp
new file mode 100644
index 0000000000000..bc24263c669dc
--- /dev/null
+++ b/clang/test/SemaOpenACC/cache-warn-invalid-varloc.cpp
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+void foo() {
+  int Array[5];
+  // expected-warning@+1 2{{OpenACC variable in 'cache' directive was not 
declared outside of the associated 'loop' directive; directive has no effect}}
+  #pragma acc cache(readonly:Array[1], Array[1:2])
+}
+
+
+void foo2() {
+#pragma acc loop
+  for(int i = 0; i < 5; ++i) {
+    int Array[5];
+    // expected-warning@+1 2{{OpenACC variable in 'cache' directive was not 
declared outside of the associated 'loop' directive; directive has no effect}}
+    #pragma acc cache(readonly:Array[1], Array[1:2])
+  }
+}
+
+
+template<typename T>
+void foo3() {
+  T Array[5];
+  // expected-warning@+1 2{{OpenACC variable in 'cache' directive was not 
declared outside of the associated 'loop' directive; directive has no effect}}
+  #pragma acc cache(readonly:Array[1], Array[1:2])
+}
+
+void Inst() {
+  foo3<int>();
+}


        
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to