[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.

Author: erichkeane

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

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.

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 *

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};
 }
 

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) {

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)

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

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])
+  }
 }