diff --git a/flang/include/flang/Semantics/scope.h b/flang/include/flang/Semantics/scope.h index d2d42d0a79af8..21072772d184b 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 633787f45e852..b245081847016 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 efc2cb0a291dd..b3f3b74b04ee1 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 actualDataAttr, dummyDataAttr; if (const auto *actualObject{actualLastSymbol ? actualLastSymbol->detailsIf() diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp index 0e59c3dd3b1af..0f69c1ccd2850 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 7d6ab2c83cc59..39d6fdc97512a 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 0d05222d446df..047503b3cca4e 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