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

[Peter Klausler via flang-commits]

https://github.com/klausler created https://github.com/llvm/llvm-project/pull/72310

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

>From a015b1d24c1b98cada17cdba2aafdac3729e2fd0 Mon Sep 17 00:00:00 2001
From: Peter Klausler <pklausler at nvidia.com>
Date: Tue, 14 Nov 2023 12:50:49 -0800
Subject: [PATCH] [flang][openacc] Disable CUDA argument checks in OpenACC
 regions

Checks for CUDA Fortran data attribute compatibility don't need
to be applied in OpenACC regions.
---
 flang/include/flang/Semantics/scope.h |  2 +-
 flang/include/flang/Semantics/tools.h |  3 ++-
 flang/lib/Semantics/check-call.cpp    |  7 +++++--
 flang/lib/Semantics/resolve-names.cpp | 11 ++++++++++-
 flang/lib/Semantics/tools.cpp         |  8 ++++++++
 flang/test/Semantics/cuf10.cuf        | 11 ++++++++++-
 6 files changed, 36 insertions(+), 6 deletions(-)

diff --git a/flang/include/flang/Semantics/scope.h b/flang/include/flang/Semantics/scope.h
index d2d42d0a79af8c5..21072772d184b66 100644
--- a/flang/include/flang/Semantics/scope.h
+++ b/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
diff --git a/flang/include/flang/Semantics/tools.h b/flang/include/flang/Semantics/tools.h
index 633787f45e85255..b245081847016b6 100644
--- a/flang/include/flang/Semantics/tools.h
+++ b/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 &);
diff --git a/flang/lib/Semantics/check-call.cpp b/flang/lib/Semantics/check-call.cpp
index efc2cb0a291ddce..b3f3b74b04ee111 100644
--- a/flang/lib/Semantics/check-call.cpp
+++ b/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>()
diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp
index 8f15f2f51da7c89..5fa31d6b4f3f77d 100644
--- a/flang/lib/Semantics/resolve-names.cpp
+++ b/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:
diff --git a/flang/lib/Semantics/tools.cpp b/flang/lib/Semantics/tools.cpp
index 7d6ab2c83cc5952..39d6fdc97512aaa 100644
--- a/flang/lib/Semantics/tools.cpp
+++ b/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
diff --git a/flang/test/Semantics/cuf10.cuf b/flang/test/Semantics/cuf10.cuf
index 0d05222d446df1f..047503b3cca4eac 100644
--- a/flang/test/Semantics/cuf10.cuf
+++ b/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