From f513bd808867757a92791fb61d893534cafbe0b3 Mon Sep 17 00:00:00 2001 From: Peter Klausler Date: Sat, 6 May 2023 15:03:39 -0700 Subject: [PATCH] [flang] CUDA Fortran - part 4/5: definability and characteristics Extend the definability and procedure characteristics checking infrastructure in semantics to check for context-dependent CUDA object definability violations and problems with CUDA attribute incompatibility in procedure interfaces. Depends on https://reviews.llvm.org/D150159, https://reviews.llvm.org/D150161, & https://reviews.llvm.org/D150162. Differential Revision: https://reviews.llvm.org/D150163 --- flang/include/flang/Evaluate/call.h | 7 ++ .../include/flang/Evaluate/characteristics.h | 4 + flang/include/flang/Semantics/expression.h | 1 + flang/lib/Evaluate/characteristics.cpp | 69 +++++++++++++- flang/lib/Evaluate/formatting.cpp | 12 +++ flang/lib/Semantics/check-call.cpp | 45 ++++++++- flang/lib/Semantics/definable.cpp | 27 ++++++ flang/lib/Semantics/expression.cpp | 94 +++++++++++++++++-- flang/test/Parser/cuf-sanity-tree.CUF | 36 ++++--- flang/test/Parser/cuf-sanity-unparse.CUF | 7 +- flang/test/Semantics/cuf07.cuf | 26 +++++ flang/test/Semantics/cuf10.cuf | 17 ++++ flang/test/Semantics/definable05.cuf | 31 ++++++ 13 files changed, 344 insertions(+), 32 deletions(-) create mode 100644 flang/test/Semantics/cuf07.cuf create mode 100644 flang/test/Semantics/cuf10.cuf create mode 100644 flang/test/Semantics/definable05.cuf diff --git a/flang/include/flang/Evaluate/call.h b/flang/include/flang/Evaluate/call.h index 76983853c169d..f2c231647390b 100644 --- a/flang/include/flang/Evaluate/call.h +++ b/flang/include/flang/Evaluate/call.h @@ -209,6 +209,8 @@ struct ProcedureDesignator { u; }; +using Chevrons = std::vector>; + class ProcedureRef { public: CLASS_BOILERPLATE(ProcedureRef) @@ -223,6 +225,10 @@ class ProcedureRef { const ProcedureDesignator &proc() const { return proc_; } ActualArguments &arguments() { return arguments_; } const ActualArguments &arguments() const { return arguments_; } + // CALL subr <<< kernel launch >>> (...); not function + Chevrons &chevrons() { return chevrons_; } + const Chevrons &chevrons() const { return chevrons_; } + void set_chevrons(Chevrons &&chevrons) { chevrons_ = std::move(chevrons); } std::optional> LEN() const; int Rank() const; @@ -250,6 +256,7 @@ class ProcedureRef { protected: ProcedureDesignator proc_; ActualArguments arguments_; + Chevrons chevrons_; bool hasAlternateReturns_; }; diff --git a/flang/include/flang/Evaluate/characteristics.h b/flang/include/flang/Evaluate/characteristics.h index 46cc6f23bddc0..824060f725d2c 100644 --- a/flang/include/flang/Evaluate/characteristics.h +++ b/flang/include/flang/Evaluate/characteristics.h @@ -220,6 +220,7 @@ struct DummyDataObject { common::Intent intent{common::Intent::Default}; Attrs attrs; common::IgnoreTKRSet ignoreTKR; + std::optional cudaDataAttr; }; // 15.3.2.3 @@ -317,6 +318,7 @@ struct FunctionResult { Attrs attrs; std::variant> u; + std::optional cudaDataAttr; }; // 15.3.1 @@ -368,6 +370,8 @@ struct Procedure { std::optional functionResult; DummyArguments dummyArguments; Attrs attrs; + std::optional cudaSubprogramAttrs; }; + } // namespace Fortran::evaluate::characteristics #endif // FORTRAN_EVALUATE_CHARACTERISTICS_H_ diff --git a/flang/include/flang/Semantics/expression.h b/flang/include/flang/Semantics/expression.h index 7cf7089715b52..a75314b5188dd 100644 --- a/flang/include/flang/Semantics/expression.h +++ b/flang/include/flang/Semantics/expression.h @@ -381,6 +381,7 @@ class ExpressionAnalyzer { } bool CheckIsValidForwardReference(const semantics::DerivedTypeSpec &); MaybeExpr AnalyzeComplex(MaybeExpr &&re, MaybeExpr &&im, const char *what); + std::optional AnalyzeChevrons(const parser::CallStmt &); MaybeExpr IterativelyAnalyzeSubexpressions(const parser::Expr &); diff --git a/flang/lib/Evaluate/characteristics.cpp b/flang/lib/Evaluate/characteristics.cpp index 6b961ac9fae56..b22025c8844bc 100644 --- a/flang/lib/Evaluate/characteristics.cpp +++ b/flang/lib/Evaluate/characteristics.cpp @@ -265,7 +265,8 @@ llvm::raw_ostream &TypeAndShape::Dump(llvm::raw_ostream &o) const { bool DummyDataObject::operator==(const DummyDataObject &that) const { return type == that.type && attrs == that.attrs && intent == that.intent && - coshape == that.coshape; + coshape == that.coshape && cudaDataAttr == that.cudaDataAttr; + ; } static bool AreCompatibleDummyDataObjectShapes(const Shape &x, const Shape &y) { @@ -325,6 +326,13 @@ bool DummyDataObject::IsCompatibleWith( *whyNot = "incompatible !DIR$ IGNORE_TKR directives"; } } + if (!attrs.test(Attr::Value) && + !common::AreCompatibleCUDADataAttrs( + cudaDataAttr, actual.cudaDataAttr, ignoreTKR)) { + if (whyNot) { + *whyNot = "incompatible CUDA data attributes"; + } + } return true; } @@ -360,6 +368,14 @@ std::optional DummyDataObject::Characterize( }); result->intent = GetIntent(symbol.attrs()); result->ignoreTKR = GetIgnoreTKR(symbol); + if (object) { + result->cudaDataAttr = object->cudaDataAttr(); + if (!result->cudaDataAttr && + !result->attrs.test(DummyDataObject::Attr::Value) && + semantics::IsCUDADeviceContext(&symbol.owner())) { + result->cudaDataAttr = common::CUDADataAttr::Device; + } + } return result; } } @@ -380,6 +396,8 @@ bool DummyDataObject::CanBePassedViaImplicitInterface() const { return false; // 15.4.2.2(3)(b-d) } else if (type.type().IsPolymorphic()) { return false; // 15.4.2.2(3)(f) + } else if (cudaDataAttr) { + return false; } else if (const auto *derived{GetDerivedTypeSpec(type.type())}) { return derived->parameters().empty(); // 15.4.2.2(3)(e) } else { @@ -400,6 +418,9 @@ llvm::raw_ostream &DummyDataObject::Dump(llvm::raw_ostream &o) const { sep = ','; } } + if (cudaDataAttr) { + o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr); + } if (!ignoreTKR.empty()) { ignoreTKR.Dump(o << ' ', common::EnumToString); } @@ -522,6 +543,7 @@ static std::optional CharacterizeProcedure( return std::nullopt; } } + result.cudaSubprogramAttrs = subp.cudaSubprogramAttrs(); return result; }, [&](const semantics::ProcEntityDetails &proc) @@ -554,6 +576,10 @@ static std::optional CharacterizeProcedure( if (symbol.test(semantics::Symbol::Flag::Subroutine)) { // ignore any implicit typing result.attrs.set(Procedure::Attr::Subroutine); + if (proc.isCUDAKernel()) { + result.cudaSubprogramAttrs = + common::CUDASubprogramAttrs::Global; + } } else if (type) { if (auto resultType{DynamicType::From(*type)}) { result.functionResult = FunctionResult{*resultType}; @@ -844,13 +870,14 @@ FunctionResult::FunctionResult(Procedure &&p) : u{std::move(p)} {} FunctionResult::~FunctionResult() {} bool FunctionResult::operator==(const FunctionResult &that) const { - return attrs == that.attrs && u == that.u; + return attrs == that.attrs && cudaDataAttr == that.cudaDataAttr && + u == that.u; } static std::optional CharacterizeFunctionResult( const semantics::Symbol &symbol, FoldingContext &context, semantics::UnorderedSymbolSet seenProcs) { - if (symbol.has()) { + if (const auto *object{symbol.detailsIf()}) { if (auto type{TypeAndShape::Characterize(symbol, context)}) { FunctionResult result{std::move(*type)}; CopyAttrs(symbol, result, @@ -859,6 +886,7 @@ static std::optional CharacterizeFunctionResult( {semantics::Attr::CONTIGUOUS, FunctionResult::Attr::Contiguous}, {semantics::Attr::POINTER, FunctionResult::Attr::Pointer}, }); + result.cudaDataAttr = object->cudaDataAttr(); return result; } } else if (auto maybeProc{ @@ -887,6 +915,8 @@ bool FunctionResult::IsAssumedLengthCharacter() const { bool FunctionResult::CanBeReturnedViaImplicitInterface() const { if (attrs.test(Attr::Pointer) || attrs.test(Attr::Allocatable)) { return false; // 15.4.2.2(4)(b) + } else if (cudaDataAttr) { + return false; } else if (const auto *typeAndShape{GetTypeAndShape()}) { if (typeAndShape->Rank() > 0) { return false; // 15.4.2.2(4)(a) @@ -953,6 +983,10 @@ bool FunctionResult::IsCompatibleWith( if (whyNot) { *whyNot = "function results have incompatible attributes"; } + } else if (cudaDataAttr != actual.cudaDataAttr) { + if (whyNot) { + *whyNot = "function results have incompatible CUDA data attributes"; + } } else if (const auto *ifaceTypeShape{std::get_if(&u)}) { if (const auto *actualTypeShape{std::get_if(&actual.u)}) { if (ifaceTypeShape->Rank() != actualTypeShape->Rank()) { @@ -1033,6 +1067,9 @@ llvm::raw_ostream &FunctionResult::Dump(llvm::raw_ostream &o) const { }, }, u); + if (cudaDataAttr) { + o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr); + } return o; } @@ -1045,7 +1082,8 @@ Procedure::~Procedure() {} bool Procedure::operator==(const Procedure &that) const { return attrs == that.attrs && functionResult == that.functionResult && - dummyArguments == that.dummyArguments; + dummyArguments == that.dummyArguments && + cudaSubprogramAttrs == that.cudaSubprogramAttrs; } bool Procedure::IsCompatibleWith(const Procedure &actual, std::string *whyNot, @@ -1078,6 +1116,10 @@ bool Procedure::IsCompatibleWith(const Procedure &actual, std::string *whyNot, } } else if (functionResult && actual.functionResult && !functionResult->IsCompatibleWith(*actual.functionResult, whyNot)) { + } else if (cudaSubprogramAttrs != actual.cudaSubprogramAttrs) { + if (whyNot) { + *whyNot = "incompatible CUDA subprogram attributes"; + } } else if (dummyArguments.size() != actual.dummyArguments.size()) { if (whyNot) { *whyNot = "distinct numbers of dummy arguments"; @@ -1200,6 +1242,10 @@ bool Procedure::CanBeCalledViaImplicitInterface() const { // TODO: Pass back information on why we return false if (attrs.test(Attr::Elemental) || attrs.test(Attr::BindC)) { return false; // 15.4.2.2(5,6) + } else if (cudaSubprogramAttrs && + *cudaSubprogramAttrs != common::CUDASubprogramAttrs::Host && + *cudaSubprogramAttrs != common::CUDASubprogramAttrs::Global) { + return false; } else if (IsFunction() && !functionResult->CanBeReturnedViaImplicitInterface()) { return false; @@ -1227,7 +1273,11 @@ llvm::raw_ostream &Procedure::Dump(llvm::raw_ostream &o) const { dummy.Dump(o << sep); sep = ','; } - return o << (sep == '(' ? "()" : ")"); + o << (sep == '(' ? "()" : ")"); + if (cudaSubprogramAttrs) { + o << " cudaSubprogramAttrs: " << common::EnumToString(*cudaSubprogramAttrs); + } + return o; } // Utility class to determine if Procedures, etc. are distinguishable @@ -1329,6 +1379,9 @@ bool DistinguishUtils::Distinguishable( if (pos2 >= 0 && pos2 <= name2) { return true; // distinguishable based on C1514 rule 4 } + if (proc1.cudaSubprogramAttrs != proc2.cudaSubprogramAttrs) { + return true; + } return false; } @@ -1456,6 +1509,9 @@ bool DistinguishUtils::Distinguishable( } else if (y.attrs.test(Attr::Allocatable) && x.attrs.test(Attr::Pointer) && x.intent != common::Intent::In) { return true; + } else if (!common::AreCompatibleCUDADataAttrs( + x.cudaDataAttr, y.cudaDataAttr, x.ignoreTKR | y.ignoreTKR)) { + return true; } else if (features_.IsEnabled( common::LanguageFeature::DistinguishableSpecifics) && (x.attrs.test(Attr::Allocatable) || x.attrs.test(Attr::Pointer)) && @@ -1494,6 +1550,9 @@ bool DistinguishUtils::Distinguishable( if (x.u.index() != y.u.index()) { return true; // one is data object, one is procedure } + if (x.cudaDataAttr != y.cudaDataAttr) { + return true; + } return common::visit( common::visitors{ [&](const TypeAndShape &z) { diff --git a/flang/lib/Evaluate/formatting.cpp b/flang/lib/Evaluate/formatting.cpp index f9548e119f1a5..84dd4be76cd9c 100644 --- a/flang/lib/Evaluate/formatting.cpp +++ b/flang/lib/Evaluate/formatting.cpp @@ -135,6 +135,18 @@ llvm::raw_ostream &ProcedureRef::AsFortran(llvm::raw_ostream &o) const { } } proc_.AsFortran(o); + if (!chevrons_.empty()) { + bool first{true}; + for (const auto &expr : chevrons_) { + if (first) { + expr.AsFortran(o << "<<<"); + first = false; + } else { + expr.AsFortran(o << ","); + } + } + o << ">>>"; + } char separator{'('}; for (const auto &arg : arguments_) { if (arg && !arg->isPassedObject()) { diff --git a/flang/lib/Semantics/check-call.cpp b/flang/lib/Semantics/check-call.cpp index 7b4e6e245c945..e5a338b64bf68 100644 --- a/flang/lib/Semantics/check-call.cpp +++ b/flang/lib/Semantics/check-call.cpp @@ -196,7 +196,8 @@ static void CheckExplicitDataArg(const characteristics::DummyDataObject &dummy, characteristics::TypeAndShape &actualType, bool isElemental, SemanticsContext &context, evaluate::FoldingContext &foldingContext, const Scope *scope, const evaluate::SpecificIntrinsic *intrinsic, - bool allowActualArgumentConversions) { + bool allowActualArgumentConversions, + const characteristics::Procedure &procedure) { // Basic type & rank checking parser::ContextualMessages &messages{foldingContext.messages()}; @@ -628,6 +629,46 @@ static void CheckExplicitDataArg(const characteristics::DummyDataObject &dummy, } } } + + // CUDA + if (!intrinsic && + !dummy.attrs.test(characteristics::DummyDataObject::Attr::Value)) { + std::optional actualDataAttr, dummyDataAttr; + if (const auto *actualObject{actualLastSymbol + ? actualLastSymbol->detailsIf() + : nullptr}) { + actualDataAttr = actualObject->cudaDataAttr(); + } + dummyDataAttr = dummy.cudaDataAttr; + // Treat MANAGED like DEVICE for nonallocatable nonpointer arguments to + // device subprograms + if (procedure.cudaSubprogramAttrs.value_or( + common::CUDASubprogramAttrs::Host) != + common::CUDASubprogramAttrs::Host && + !dummy.attrs.test( + characteristics::DummyDataObject::Attr::Allocatable) && + !dummy.attrs.test(characteristics::DummyDataObject::Attr::Pointer)) { + if (!dummyDataAttr || *dummyDataAttr == common::CUDADataAttr::Managed) { + dummyDataAttr = common::CUDADataAttr::Device; + } + if ((!actualDataAttr && FindCUDADeviceContext(scope)) || + (actualDataAttr && + *actualDataAttr == common::CUDADataAttr::Managed)) { + actualDataAttr = common::CUDADataAttr::Device; + } + } + if (!common::AreCompatibleCUDADataAttrs( + dummyDataAttr, actualDataAttr, dummy.ignoreTKR)) { + auto toStr{[](std::optional x) { + return x ? "ATTRIBUTES("s + + parser::ToUpperCaseLetters(common::EnumToString(*x)) + ")"s + : "no CUDA data attribute"s; + }}; + messages.Say( + "%s has %s but its associated actual argument has %s"_err_en_US, + dummyName, toStr(dummyDataAttr), toStr(actualDataAttr)); + } + } } static void CheckProcedureArg(evaluate::ActualArgument &arg, @@ -819,7 +860,7 @@ static void CheckExplicitInterfaceArg(evaluate::ActualArgument &arg, object.type.Rank() == 0 && proc.IsElemental()}; CheckExplicitDataArg(object, dummyName, *expr, *type, isElemental, context, foldingContext, scope, intrinsic, - allowActualArgumentConversions); + allowActualArgumentConversions, proc); } else if (object.type.type().IsTypelessIntrinsicArgument() && IsBOZLiteral(*expr)) { // ok diff --git a/flang/lib/Semantics/definable.cpp b/flang/lib/Semantics/definable.cpp index 675becd32c266..abb5f35c28eae 100644 --- a/flang/lib/Semantics/definable.cpp +++ b/flang/lib/Semantics/definable.cpp @@ -134,6 +134,33 @@ static std::optional WhyNotDefinableBase(parser::CharBlock at, original, visible->name()); } } + if (const Scope * deviceContext{FindCUDADeviceContext(&scope)}) { + bool isOwnedByDeviceCode{deviceContext->Contains(ultimate.owner())}; + if (isPointerDefinition && !acceptAllocatable) { + return BlameSymbol(at, + "'%s' is a pointer and may not be associated in a device subprogram"_err_en_US, + original); + } else if (auto cudaDataAttr{GetCUDADataAttr(&ultimate)}) { + if (*cudaDataAttr == common::CUDADataAttr::Constant) { + return BlameSymbol(at, + "'%s' has ATTRIBUTES(CONSTANT) and is not definable in a device subprogram"_err_en_US, + original); + } else if (acceptAllocatable && !isOwnedByDeviceCode) { + return BlameSymbol(at, + "'%s' is a host-associated allocatable and is not definable in a device subprogram"_err_en_US, + original); + } else if (*cudaDataAttr != common::CUDADataAttr::Device && + *cudaDataAttr != common::CUDADataAttr::Managed) { + return BlameSymbol(at, + "'%s' is not device or managed data and is not definable in a device subprogram"_err_en_US, + original); + } + } else if (!isOwnedByDeviceCode) { + return BlameSymbol(at, + "'%s' is a host variable and is not definable in a device subprogram"_err_en_US, + original); + } + } return std::nullopt; } diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp index 28c4ba16ae926..a6b54dd11d21b 100644 --- a/flang/lib/Semantics/expression.cpp +++ b/flang/lib/Semantics/expression.cpp @@ -2653,6 +2653,9 @@ void ExpressionAnalyzer::CheckForBadRecursion( msg = Say( // 15.6.2.1(3) "Assumed-length CHARACTER(*) function '%s' cannot call itself"_err_en_US, callSite); + } else if (FindCUDADeviceContext(scope)) { + msg = Say( + "Device subprogram '%s' cannot call itself"_err_en_US, callSite); } AttachDeclaration(msg, proc); } @@ -2719,6 +2722,55 @@ bool ExpressionAnalyzer::CheckIsValidForwardReference( return true; } +std::optional ExpressionAnalyzer::AnalyzeChevrons( + const parser::CallStmt &call) { + Chevrons result; + auto checkLaunchArg{[&](const Expr &expr, const char *which) { + if (auto dyType{expr.GetType()}) { + if (dyType->category() == TypeCategory::Integer) { + return true; + } + if (dyType->category() == TypeCategory::Derived && + !dyType->IsPolymorphic() && + IsBuiltinDerivedType(&dyType->GetDerivedTypeSpec(), "dim3")) { + return true; + } + } + Say("Kernel launch %s parameter must be either integer or TYPE(dim3)"_err_en_US, + which); + return false; + }}; + if (const auto &chevrons{call.chevrons}) { + if (auto expr{Analyze(std::get<0>(chevrons->t))}; + expr && checkLaunchArg(*expr, "grid")) { + result.emplace_back(*expr); + } else { + return std::nullopt; + } + if (auto expr{Analyze(std::get<1>(chevrons->t))}; + expr && checkLaunchArg(*expr, "block")) { + result.emplace_back(*expr); + } else { + return std::nullopt; + } + if (const auto &maybeExpr{std::get<2>(chevrons->t)}) { + if (auto expr{Analyze(*maybeExpr)}) { + result.emplace_back(*expr); + } else { + return std::nullopt; + } + } + if (const auto &maybeExpr{std::get<3>(chevrons->t)}) { + if (auto expr{Analyze(*maybeExpr)}) { + result.emplace_back(*expr); + } else { + return std::nullopt; + } + } + } + return std::move(result); +} + MaybeExpr ExpressionAnalyzer::Analyze(const parser::FunctionReference &funcRef, std::optional *structureConstructor) { const parser::Call &call{funcRef.v}; @@ -2730,17 +2782,17 @@ MaybeExpr ExpressionAnalyzer::Analyze(const parser::FunctionReference &funcRef, if (analyzer.fatalErrors()) { return std::nullopt; } - if (std::optional callee{ - GetCalleeAndArguments(std::get(call.t), - analyzer.GetActuals(), false /* not subroutine */, - true /* might be structure constructor */)}) { + bool mightBeStructureConstructor{structureConstructor != nullptr}; + if (std::optional callee{GetCalleeAndArguments( + std::get(call.t), analyzer.GetActuals(), + false /* not subroutine */, mightBeStructureConstructor)}) { if (auto *proc{std::get_if(&callee->u)}) { return MakeFunctionRef( funcRef.source, std::move(*proc), std::move(callee->arguments)); } CHECK(std::holds_alternative(callee->u)); const Symbol &symbol{*std::get(callee->u)}; - if (structureConstructor) { + if (mightBeStructureConstructor) { // Structure constructor misparsed as function reference? const auto &designator{std::get(call.t)}; if (const auto *name{std::get_if(&designator.u)}) { @@ -2785,17 +2837,40 @@ void ExpressionAnalyzer::Analyze(const parser::CallStmt &callStmt) { for (const auto &arg : actualArgList) { analyzer.Analyze(arg, true /* is subroutine call */); } - if (!analyzer.fatalErrors()) { + auto chevrons{AnalyzeChevrons(callStmt)}; + if (!analyzer.fatalErrors() && chevrons) { if (std::optional callee{ GetCalleeAndArguments(std::get(call.t), analyzer.GetActuals(), true /* subroutine */)}) { ProcedureDesignator *proc{std::get_if(&callee->u)}; CHECK(proc); + bool isKernel{false}; + if (const Symbol * procSym{proc->GetSymbol()}) { + const Symbol &ultimate{procSym->GetUltimate()}; + if (const auto *subpDetails{ + ultimate.detailsIf()}) { + if (auto attrs{subpDetails->cudaSubprogramAttrs()}) { + isKernel = *attrs == common::CUDASubprogramAttrs::Global || + *attrs == common::CUDASubprogramAttrs::Grid_Global; + } + } else if (const auto *procDetails{ + ultimate.detailsIf()}) { + isKernel = procDetails->isCUDAKernel(); + } + if (isKernel && chevrons->empty()) { + Say("'%s' is a kernel subroutine and must be called with kernel launch parameters in chevrons"_err_en_US, + procSym->name()); + } + } + if (!isKernel && !chevrons->empty()) { + Say("Kernel launch parameters in chevrons may not be used unless calling a kernel subroutine"_err_en_US); + } if (CheckCall(callStmt.source, *proc, callee->arguments)) { callStmt.typedCall.Reset( new ProcedureRef{std::move(*proc), std::move(callee->arguments), HasAlternateReturns(callee->arguments)}, ProcedureRef::Deleter); + DEREF(callStmt.typedCall.get()).set_chevrons(std::move(*chevrons)); return; } } @@ -3697,14 +3772,13 @@ MaybeExpr ExpressionAnalyzer::MakeFunctionRef(parser::CharBlock callSite, if (auto chars{CheckCall(callSite, proc, arguments)}) { if (chars->functionResult) { const auto &result{*chars->functionResult}; + ProcedureRef procRef{std::move(proc), std::move(arguments)}; if (result.IsProcedurePointer()) { - return Expr{ - ProcedureRef{std::move(proc), std::move(arguments)}}; + return Expr{std::move(procRef)}; } else { // Not a procedure pointer, so type and shape are known. return TypedWrapper( - DEREF(result.GetTypeAndShape()).type(), - ProcedureRef{std::move(proc), std::move(arguments)}); + DEREF(result.GetTypeAndShape()).type(), std::move(procRef)); } } else { Say("Function result characteristics are not known"_err_en_US); diff --git a/flang/test/Parser/cuf-sanity-tree.CUF b/flang/test/Parser/cuf-sanity-tree.CUF index 2ce042bcdbc1b..f6cf9bbdd6b0c 100644 --- a/flang/test/Parser/cuf-sanity-tree.CUF +++ b/flang/test/Parser/cuf-sanity-tree.CUF @@ -106,6 +106,9 @@ include "cuf-sanity-common" !CHECK: | | | | Name = 'attrs' !CHECK: | | | SpecificationPart !CHECK: | | | | ImplicitPart -> +!CHECK: | | | | DeclarationConstruct -> SpecificationConstruct -> OtherSpecificationStmt -> CUDAAttributesStmt +!CHECK: | | | | | CUDADataAttr = Device +!CHECK: | | | | | Name = 'devx1' !CHECK: | | | | DeclarationConstruct -> SpecificationConstruct -> TypeDeclarationStmt !CHECK: | | | | | DeclarationTypeSpec -> IntrinsicTypeSpec -> Real !CHECK: | | | | | AttrSpec -> CUDADataAttr = Device @@ -159,27 +162,36 @@ include "cuf-sanity-common" !CHECK: | | | | | | | | | LiteralConstant -> IntLiteralConstant = '10' !CHECK: | | | | | | Block !CHECK: | | | | | | EndDoStmt -> -!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub()' +!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4>>>()' !CHECK: | | | | | Call !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' !CHECK: | | | | | Chevrons -!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '1' -!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '2' -!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub()' +!CHECK: | | | | | | Scalar -> Expr = '1_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1' +!CHECK: | | | | | | Scalar -> Expr = '2_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2' +!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4,3_4>>>()' !CHECK: | | | | | Call !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' !CHECK: | | | | | Chevrons -!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '1' -!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '2' -!CHECK: | | | | | | Scalar -> Integer -> Expr -> LiteralConstant -> IntLiteralConstant = '3' -!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub()' +!CHECK: | | | | | | Scalar -> Expr = '1_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1' +!CHECK: | | | | | | Scalar -> Expr = '2_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2' +!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3' +!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4,3_4,4_4>>>()' !CHECK: | | | | | Call !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' !CHECK: | | | | | Chevrons -!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '1' -!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '2' -!CHECK: | | | | | | Scalar -> Integer -> Expr -> LiteralConstant -> IntLiteralConstant = '3' -!CHECK: | | | | | | Scalar -> Integer -> Expr -> LiteralConstant -> IntLiteralConstant = '4' +!CHECK: | | | | | | Scalar -> Expr = '1_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1' +!CHECK: | | | | | | Scalar -> Expr = '2_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2' +!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3' +!CHECK: | | | | | | Scalar -> Integer -> Expr = '4_4' +!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '4' !CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> AllocateStmt !CHECK: | | | | | Allocation !CHECK: | | | | | | AllocateObject = 'pa' diff --git a/flang/test/Parser/cuf-sanity-unparse.CUF b/flang/test/Parser/cuf-sanity-unparse.CUF index 3bd838a75e3ea..d4495c4fddccf 100644 --- a/flang/test/Parser/cuf-sanity-unparse.CUF +++ b/flang/test/Parser/cuf-sanity-unparse.CUF @@ -23,6 +23,7 @@ include "cuf-sanity-common" !CHECK: ATTRIBUTES(GLOBAL) CLUSTER_DIMS(1_4, 2_4, 3_4) SUBROUTINE cdsub !CHECK: END SUBROUTINE !CHECK: ATTRIBUTES(DEVICE) SUBROUTINE attrs +!CHECK: ATTRIBUTES(DEVICE) devx1 !CHECK: REAL, DEVICE :: devx2 !CHECK: END SUBROUTINE !CHECK: SUBROUTINE test @@ -33,9 +34,9 @@ include "cuf-sanity-common" !CHECK: !$CUF KERNEL DO <<<1_4,(2_4,3_4),STREAM=1_4>>> !CHECK: DO j=1_4,10_4 !CHECK: END DO -!CHECK: CALL globalsub<<<1,2>>> -!CHECK: CALL globalsub<<<1,2,3>>> -!CHECK: CALL globalsub<<<1,2,3,4>>> +!CHECK: CALL globalsub<<<1_4,2_4>>> +!CHECK: CALL globalsub<<<1_4,2_4,3_4>>> +!CHECK: CALL globalsub<<<1_4,2_4,3_4,4_4>>> !CHECK: ALLOCATE(pa(32_4), STREAM=1_4, PINNED=ispinned) !CHECK: END SUBROUTINE !CHECK: END MODULE diff --git a/flang/test/Semantics/cuf07.cuf b/flang/test/Semantics/cuf07.cuf new file mode 100644 index 0000000000000..b520b5da51264 --- /dev/null +++ b/flang/test/Semantics/cuf07.cuf @@ -0,0 +1,26 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +module m + real, allocatable :: xa + real, allocatable, managed :: ma + contains + attributes(device) subroutine devsubr + real, device, allocatable :: da + real, allocatable, managed :: dma + allocate(da) ! ok + deallocate(da) ! ok + allocate(dma) ! ok + deallocate(dma) ! ok + !ERROR: Name in ALLOCATE statement is not definable + !BECAUSE: 'xa' is a host variable and is not definable in a device subprogram + allocate(xa) + !ERROR: Name in DEALLOCATE statement is not definable + !BECAUSE: 'xa' is a host variable and is not definable in a device subprogram + deallocate(xa) + !ERROR: Name in ALLOCATE statement is not definable + !BECAUSE: 'ma' is a host-associated allocatable and is not definable in a device subprogram + allocate(ma) + !ERROR: Name in DEALLOCATE statement is not definable + !BECAUSE: 'ma' is a host-associated allocatable and is not definable in a device subprogram + deallocate(ma) + end subroutine +end module diff --git a/flang/test/Semantics/cuf10.cuf b/flang/test/Semantics/cuf10.cuf new file mode 100644 index 0000000000000..0d05222d446df --- /dev/null +++ b/flang/test/Semantics/cuf10.cuf @@ -0,0 +1,17 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +module m + real, device :: a(4,8) + real, managed, allocatable :: b(:,:) + contains + attributes(global) subroutine kernel(a,b,c,n,m) + integer, value :: n + integer, intent(in) :: m + real a(n,m), c(n,m) + real, managed :: b(n,m) + end + subroutine test + 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) + end +end diff --git a/flang/test/Semantics/definable05.cuf b/flang/test/Semantics/definable05.cuf new file mode 100644 index 0000000000000..5af3ca9244345 --- /dev/null +++ b/flang/test/Semantics/definable05.cuf @@ -0,0 +1,31 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +module m + real, constant :: rc + !ERROR: Object 'rcp' with ATTRIBUTES(CONSTANT) may not be allocatable, pointer, or target + real, constant, pointer :: rcp + !ERROR: Object 'rct' with ATTRIBUTES(CONSTANT) may not be allocatable, pointer, or target + real, constant, target :: rct + real, device, pointer :: dp(:) + real, device, target :: dt(100) + contains + attributes(device) subroutine devsub + !ERROR: Left-hand side of assignment is not definable + !BECAUSE: 'rc' has ATTRIBUTES(CONSTANT) and is not definable in a device subprogram + rc = 1. + !ERROR: The left-hand side of a pointer assignment is not definable + !BECAUSE: 'dp' is a pointer and may not be associated in a device subprogram + dp => dt + end + attributes(global) subroutine globsub + !ERROR: Left-hand side of assignment is not definable + !BECAUSE: 'rc' has ATTRIBUTES(CONSTANT) and is not definable in a device subprogram + rc = 1. + !ERROR: The left-hand side of a pointer assignment is not definable + !BECAUSE: 'dp' is a pointer and may not be associated in a device subprogram + dp => dt + end + subroutine hostsub + rc = 1. + dp => dt + end +end