[flang][openacc] Disable CUDA argument checks in OpenACC regions (#72310)

Checks for CUDA Fortran data attribute compatibility don't need to be
applied in OpenACC regions.

Author: klausler

flang/include/flang/Semantics/scope.h

@@ -61,7 +61,7 @@ class Scope {
 public:
   ENUM_CLASS(Kind, Global, IntrinsicModules, Module, MainProgram, Subprogram,
       BlockData, DerivedType, BlockConstruct, Forall, OtherConstruct,
-      ImpliedDos)
+      OpenACCConstruct, ImpliedDos)
   using ImportKind = common::ImportKind;
 
   // Create the Global scope -- the root of the scope tree

flang/include/flang/Semantics/tools.h

@@ -44,7 +44,8 @@ const Scope &GetProgramUnitOrBlockConstructContaining(const Symbol &);
 const Scope *FindModuleContaining(const Scope &);
 const Scope *FindModuleFileContaining(const Scope &);
 const Scope *FindPureProcedureContaining(const Scope &);
-const Scope *FindPureProcedureContaining(const Symbol &);
+const Scope *FindOpenACCConstructContaining(const Scope *);
+
 const Symbol *FindPointerComponent(const Scope &);
 const Symbol *FindPointerComponent(const DerivedTypeSpec &);
 const Symbol *FindPointerComponent(const DeclTypeSpec &);

flang/lib/Semantics/check-call.cpp

@@ -856,9 +856,12 @@ static void CheckExplicitDataArg(const characteristics::DummyDataObject &dummy,
     }
   }
 
-  // CUDA
+  // CUDA specific checks
+  // TODO: These are disabled in OpenACC constructs, which may not be
+  // correct when the target is not a GPU.
   if (!intrinsic &&
-      !dummy.attrs.test(characteristics::DummyDataObject::Attr::Value)) {
+      !dummy.attrs.test(characteristics::DummyDataObject::Attr::Value) &&
+      !FindOpenACCConstructContaining(scope)) {
     std::optional<common::CUDADataAttr> actualDataAttr, dummyDataAttr;
     if (const auto *actualObject{actualLastSymbol
                 ? actualLastSymbol->detailsIf<ObjectEntityDetails>()

flang/lib/Semantics/resolve-names.cpp

@@ -1332,6 +1332,8 @@ class AccVisitor : public virtual DeclarationVisitor {
 
   bool Pre(const parser::OpenACCBlockConstruct &);
   void Post(const parser::OpenACCBlockConstruct &);
+  bool Pre(const parser::OpenACCCombinedConstruct &);
+  void Post(const parser::OpenACCCombinedConstruct &);
   bool Pre(const parser::AccBeginBlockDirective &x) {
     AddAccSourceRange(x.source);
     return true;
@@ -1377,7 +1379,7 @@ void AccVisitor::AddAccSourceRange(const parser::CharBlock &source) {
 
 bool AccVisitor::Pre(const parser::OpenACCBlockConstruct &x) {
   if (NeedsScope(x)) {
-    PushScope(Scope::Kind::OtherConstruct, nullptr);
+    PushScope(Scope::Kind::OpenACCConstruct, nullptr);
   }
   return true;
 }
@@ -1388,6 +1390,13 @@ void AccVisitor::Post(const parser::OpenACCBlockConstruct &x) {
   }
 }
 
+bool AccVisitor::Pre(const parser::OpenACCCombinedConstruct &x) {
+  PushScope(Scope::Kind::OpenACCConstruct, nullptr);
+  return true;
+}
+
+void AccVisitor::Post(const parser::OpenACCCombinedConstruct &x) { PopScope(); }
+
 // Create scopes for OpenMP constructs
 class OmpVisitor : public virtual DeclarationVisitor {
 public:

flang/lib/Semantics/tools.cpp

@@ -108,6 +108,14 @@ const Scope *FindPureProcedureContaining(const Scope &start) {
   }
 }
 
+const Scope *FindOpenACCConstructContaining(const Scope *scope) {
+  return scope ? FindScopeContaining(*scope,
+                     [](const Scope &s) {
+                       return s.kind() == Scope::Kind::OpenACCConstruct;
+                     })
+               : nullptr;
+}
+
 // 7.5.2.4 "same derived type" test -- rely on IsTkCompatibleWith() and its
 // infrastructure to detect and handle comparisons on distinct (but "same")
 // sequence/bind(C) derived types

flang/test/Semantics/cuf10.cuf

@@ -1,4 +1,4 @@
-! RUN: %python %S/test_errors.py %s %flang_fc1
+! RUN: %python %S/test_errors.py %s %flang_fc1 -fopenacc
 module m
   real, device :: a(4,8)
   real, managed, allocatable :: b(:,:)
@@ -9,9 +9,18 @@ module m
     real a(n,m), c(n,m)
     real, managed :: b(n,m)
   end
+  attributes(device) subroutine devsub(a,n)
+    integer, value :: n
+    real, device :: a(n)
+  end
   subroutine test
+    real c(4)
     allocate(b(4,8))
     !ERROR: dummy argument 'm=' has ATTRIBUTES(DEVICE) but its associated actual argument has no CUDA data attribute
     call kernel<<<1,32>>>(a,b,b,4,8)
+    !$acc parallel loop copy(c)
+    do j = 1, 1
+      call devsub(c,4) ! not checked in OpenACC construct
+    end do
   end
 end