Index: flang/include/flang/Evaluate/call.h =================================================================== --- flang/include/flang/Evaluate/call.h +++ flang/include/flang/Evaluate/call.h @@ -209,6 +209,8 @@ u; }; +using Chevrons = std::vector>; + class ProcedureRef { public: CLASS_BOILERPLATE(ProcedureRef) @@ -223,6 +225,10 @@ 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 @@ protected: ProcedureDesignator proc_; ActualArguments arguments_; + Chevrons chevrons_; bool hasAlternateReturns_; }; Index: flang/include/flang/Evaluate/characteristics.h =================================================================== --- flang/include/flang/Evaluate/characteristics.h +++ flang/include/flang/Evaluate/characteristics.h @@ -220,6 +220,7 @@ common::Intent intent{common::Intent::Default}; Attrs attrs; common::IgnoreTKRSet ignoreTKR; + std::optional cudaDataAttr; }; // 15.3.2.3 @@ -317,6 +318,7 @@ Attrs attrs; std::variant> u; + std::optional cudaDataAttr; }; // 15.3.1 @@ -368,6 +370,8 @@ std::optional functionResult; DummyArguments dummyArguments; Attrs attrs; + std::optional cudaSubprogramAttrs; }; + } // namespace Fortran::evaluate::characteristics #endif // FORTRAN_EVALUATE_CHARACTERISTICS_H_ Index: flang/include/flang/Evaluate/traverse.h =================================================================== --- flang/include/flang/Evaluate/traverse.h +++ flang/include/flang/Evaluate/traverse.h @@ -38,6 +38,7 @@ // expression of an ASSOCIATE (or related) construct entity. #include "expression.h" +#include "flang/Common/indirection.h" #include "flang/Semantics/symbol.h" #include "flang/Semantics/type.h" #include @@ -53,6 +54,10 @@ Result operator()(const common::Indirection &x) const { return visitor_(x.value()); } + template + Result operator()(const common::ForwardOwningPointer &p) const { + return visitor_(p.get()); + } template Result operator()(const SymbolRef x) const { return visitor_(*x); } @@ -76,13 +81,17 @@ return visitor_.Default(); } } - template - Result operator()(const std::variant &u) const { - return common::visit(visitor_, u); + template + Result operator()(const std::variant &u) const { + return common::visit([=](const auto &y) { return visitor_(y); }, u); } template Result operator()(const std::vector &x) const { return CombineContents(x); } + template + Result operator()(const std::pair &x) const { + return Combine(x.first, x.second); + } // Leaves Result operator()(const BOZLiteralConstant &) const { @@ -233,14 +242,25 @@ template Result operator()(const Expr &x) const { return visitor_(x.u); } + Result operator()(const Assignment &x) const { + return Combine(x.lhs, x.rhs, x.u); + } + Result operator()(const Assignment::Intrinsic &) const { + return visitor_.Default(); + ; + } + Result operator()(const GenericExprWrapper &x) const { return visitor_(x.v); } + Result operator()(const GenericAssignmentWrapper &x) const { + return visitor_(x.v); + } private: template Result CombineRange(ITER iter, ITER end) const { if (iter == end) { return visitor_.Default(); } else { - Result result{visitor_(*iter++)}; - for (; iter != end; ++iter) { + Result result{visitor_(*iter)}; + for (++iter; iter != end; ++iter) { result = visitor_.Combine(std::move(result), visitor_(*iter)); } return result; Index: flang/include/flang/Evaluate/type.h =================================================================== --- flang/include/flang/Evaluate/type.h +++ flang/include/flang/Evaluate/type.h @@ -22,6 +22,7 @@ #include "integer.h" #include "logical.h" #include "real.h" +#include "flang/Common/Fortran-features.h" #include "flang/Common/Fortran.h" #include "flang/Common/idioms.h" #include "flang/Common/real.h" @@ -472,8 +473,10 @@ std::optional ComparisonType( const DynamicType &, const DynamicType &); -bool IsInteroperableIntrinsicType( - const DynamicType &, bool checkCharLength = true); +bool IsInteroperableIntrinsicType(const DynamicType &, + const common::LanguageFeatureControl * = nullptr, + bool checkCharLength = true); +bool IsCUDAIntrinsicType(const DynamicType &); // Determine whether two derived type specs are sufficiently identical // to be considered the "same" type even if declared separately. Index: flang/include/flang/Semantics/expression.h =================================================================== --- flang/include/flang/Semantics/expression.h +++ flang/include/flang/Semantics/expression.h @@ -381,6 +381,7 @@ } 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 &); Index: flang/include/flang/Semantics/type.h =================================================================== --- flang/include/flang/Semantics/type.h +++ flang/include/flang/Semantics/type.h @@ -446,7 +446,8 @@ return const_cast(this)->AsDerived(); } -bool IsInteroperableIntrinsicType(const DeclTypeSpec &); +bool IsInteroperableIntrinsicType( + const DeclTypeSpec &, const common::LanguageFeatureControl &); } // namespace Fortran::semantics #endif // FORTRAN_SEMANTICS_TYPE_H_ Index: flang/lib/Evaluate/characteristics.cpp =================================================================== --- flang/lib/Evaluate/characteristics.cpp +++ flang/lib/Evaluate/characteristics.cpp @@ -265,7 +265,8 @@ 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 @@ *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 @@ }); 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 @@ 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 @@ sep = ','; } } + if (cudaDataAttr) { + o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr); + } return o; } @@ -519,6 +540,7 @@ return std::nullopt; } } + result.cudaSubprogramAttrs = subp.cudaSubprogramAttrs(); return result; }, [&](const semantics::ProcEntityDetails &proc) @@ -551,6 +573,10 @@ 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}; @@ -841,13 +867,14 @@ 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, @@ -856,6 +883,7 @@ {semantics::Attr::CONTIGUOUS, FunctionResult::Attr::Contiguous}, {semantics::Attr::POINTER, FunctionResult::Attr::Pointer}, }); + result.cudaDataAttr = object->cudaDataAttr(); return result; } } else if (auto maybeProc{ @@ -884,6 +912,8 @@ 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) @@ -950,6 +980,10 @@ 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()) { @@ -1030,6 +1064,9 @@ }, }, u); + if (cudaDataAttr) { + o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr); + } return o; } @@ -1042,7 +1079,8 @@ 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, @@ -1075,6 +1113,10 @@ } } 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"; @@ -1197,6 +1239,10 @@ // 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; @@ -1224,7 +1270,11 @@ 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 @@ -1326,6 +1376,9 @@ if (pos2 >= 0 && pos2 <= name2) { return true; // distinguishable based on C1514 rule 4 } + if (proc1.cudaSubprogramAttrs != proc2.cudaSubprogramAttrs) { + return true; + } return false; } @@ -1453,6 +1506,9 @@ } 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)) && @@ -1491,6 +1547,9 @@ 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) { Index: flang/lib/Evaluate/formatting.cpp =================================================================== --- flang/lib/Evaluate/formatting.cpp +++ flang/lib/Evaluate/formatting.cpp @@ -135,6 +135,18 @@ } } 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()) { Index: flang/lib/Evaluate/type.cpp =================================================================== --- flang/lib/Evaluate/type.cpp +++ flang/lib/Evaluate/type.cpp @@ -734,14 +734,16 @@ } } -bool IsInteroperableIntrinsicType( - const DynamicType &type, bool checkCharLength) { +bool IsInteroperableIntrinsicType(const DynamicType &type, + const common::LanguageFeatureControl *features, + bool checkCharLength) { switch (type.category()) { case TypeCategory::Integer: return true; case TypeCategory::Real: case TypeCategory::Complex: - return type.kind() >= 4; // no short or half floats + return (features && features->IsEnabled(common::LanguageFeature::CUDA)) || + type.kind() >= 4; // no short or half floats case TypeCategory::Logical: return type.kind() == 1; // C_BOOL case TypeCategory::Character: @@ -755,4 +757,21 @@ } } +bool IsCUDAIntrinsicType(const DynamicType &type) { + switch (type.category()) { + case TypeCategory::Integer: + case TypeCategory::Logical: + return type.kind() <= 8; + case TypeCategory::Real: + return type.kind() >= 2 && type.kind() <= 8; + case TypeCategory::Complex: + return type.kind() == 2 || type.kind() == 4 || type.kind() == 8; + case TypeCategory::Character: + return type.kind() == 1; + default: + // Derived types are tested in Semantics/check-declarations.cpp + return false; + } +} + } // namespace Fortran::evaluate Index: flang/lib/Parser/unparse.cpp =================================================================== --- flang/lib/Parser/unparse.cpp +++ flang/lib/Parser/unparse.cpp @@ -1689,7 +1689,7 @@ Put('('), Walk(std::get>(x.v.t), ", "), Put(')'); } void Unparse(const CallStmt &x) { // R1521 - if (asFortran_ && x.typedCall.get() && !x.chevrons /*CUDA todo*/) { + if (asFortran_ && x.typedCall.get()) { Put(' '); asFortran_->call(out_, *x.typedCall); Put('\n'); Index: flang/lib/Semantics/CMakeLists.txt =================================================================== --- flang/lib/Semantics/CMakeLists.txt +++ flang/lib/Semantics/CMakeLists.txt @@ -10,6 +10,7 @@ check-call.cpp check-case.cpp check-coarray.cpp + check-cuda.cpp check-data.cpp check-deallocate.cpp check-declarations.cpp Index: flang/lib/Semantics/canonicalize-acc.cpp =================================================================== --- flang/lib/Semantics/canonicalize-acc.cpp +++ flang/lib/Semantics/canonicalize-acc.cpp @@ -65,7 +65,7 @@ const auto &outer{std::get>(x.t)}; if (outer->IsDoConcurrent()) { - return; // Tile is not allowed on DO CONURRENT + return; // Tile is not allowed on DO CONCURRENT } for (const parser::DoConstruct *loop{&*outer}; loop && tileArgNb > 0; --tileArgNb) { Index: flang/lib/Semantics/check-allocate.cpp =================================================================== --- flang/lib/Semantics/check-allocate.cpp +++ flang/lib/Semantics/check-allocate.cpp @@ -31,6 +31,8 @@ bool gotTypeSpec{false}; bool gotSource{false}; bool gotMold{false}; + bool gotStream{false}; + bool gotPinned{false}; }; class AllocationCheckerHelper { @@ -179,8 +181,22 @@ parserSourceExpr = &mold.v.value(); info.gotMold = true; }, - [](const parser::AllocOpt::Stream &) { /* CUDA coming */ }, - [](const parser::AllocOpt::Pinned &) { /* CUDA coming */ }, + [&](const parser::AllocOpt::Stream &stream) { // CUDA + if (info.gotStream) { + context.Say( + "STREAM may not be duplicated in a ALLOCATE statement"_err_en_US); + stopCheckingAllocate = true; + } + info.gotStream = true; + }, + [&](const parser::AllocOpt::Pinned &pinned) { // CUDA + if (info.gotPinned) { + context.Say( + "PINNED may not be duplicated in a ALLOCATE statement"_err_en_US); + stopCheckingAllocate = true; + } + info.gotPinned = true; + }, }, allocOpt.u); } @@ -569,12 +585,13 @@ return false; } context.CheckIndexVarRedefine(name_); + const Scope &subpScope{ + GetProgramUnitContaining(context.FindScope(name_.source))}; if (allocateObject_.typedExpr && allocateObject_.typedExpr->v) { - if (auto whyNot{ - WhyNotDefinable(name_.source, context.FindScope(name_.source), - {DefinabilityFlag::PointerDefinition, - DefinabilityFlag::AcceptAllocatable}, - *allocateObject_.typedExpr->v)}) { + if (auto whyNot{WhyNotDefinable(name_.source, subpScope, + {DefinabilityFlag::PointerDefinition, + DefinabilityFlag::AcceptAllocatable}, + *allocateObject_.typedExpr->v)}) { context .Say(name_.source, "Name in ALLOCATE statement is not definable"_err_en_US) Index: flang/lib/Semantics/check-call.cpp =================================================================== --- flang/lib/Semantics/check-call.cpp +++ flang/lib/Semantics/check-call.cpp @@ -190,7 +190,8 @@ characteristics::TypeAndShape &actualType, bool isElemental, evaluate::FoldingContext &context, const Scope *scope, const evaluate::SpecificIntrinsic *intrinsic, - bool allowActualArgumentConversions) { + bool allowActualArgumentConversions, + const characteristics::Procedure &procedure) { // Basic type & rank checking parser::ContextualMessages &messages{context.messages()}; @@ -618,6 +619,46 @@ } } } + + // 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, @@ -808,7 +849,7 @@ object.type.Rank() == 0 && proc.IsElemental()}; CheckExplicitDataArg(object, dummyName, *expr, *type, isElemental, context, scope, intrinsic, - allowActualArgumentConversions); + allowActualArgumentConversions, proc); } else if (object.type.type().IsTypelessIntrinsicArgument() && IsBOZLiteral(*expr)) { // ok Index: flang/lib/Semantics/check-cuda.h =================================================================== --- /dev/null +++ flang/lib/Semantics/check-cuda.h @@ -0,0 +1,50 @@ +//===-- lib/Semantics/check-cuda.h ------------------------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef FORTRAN_SEMANTICS_CHECK_CUDA_H_ +#define FORTRAN_SEMANTICS_CHECK_CUDA_H_ + +#include "flang/Semantics/semantics.h" +#include + +namespace Fortran::parser { +struct Program; +class Messages; +class Name; +class CharBlock; +class ExecutionPartConstruct; +class ExecutableConstruct; +class ActionStmt; +class IfConstruct; +class CUFKernelDoConstruct; +class SubroutineSubprogram; +class FunctionSubprogram; +class SeparateModuleSubprogram; +} // namespace Fortran::parser + +namespace Fortran::semantics { + +class SemanticsContext; + +class CUDAChecker : public virtual BaseChecker { +public: + explicit CUDAChecker(SemanticsContext &c) : context_{c} {} + void Enter(const parser::SubroutineSubprogram &); + void Enter(const parser::FunctionSubprogram &); + void Enter(const parser::SeparateModuleSubprogram &); + void Enter(const parser::CUFKernelDoConstruct &); + +private: + SemanticsContext &context_; +}; + +bool CanonicalizeCUDA(parser::Program &); + +} // namespace Fortran::semantics + +#endif // FORTRAN_SEMANTICS_CHECK_CUDA_H_ Index: flang/lib/Semantics/check-cuda.cpp =================================================================== --- /dev/null +++ flang/lib/Semantics/check-cuda.cpp @@ -0,0 +1,416 @@ +//===-- lib/Semantics/check-cuda.cpp ----------------------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "check-cuda.h" +#include "flang/Common/template.h" +#include "flang/Evaluate/fold.h" +#include "flang/Evaluate/traverse.h" +#include "flang/Parser/parse-tree-visitor.h" +#include "flang/Parser/parse-tree.h" +#include "flang/Parser/tools.h" +#include "flang/Semantics/expression.h" +#include "flang/Semantics/symbol.h" + +// Once labeled DO constructs have been canonicalized and their parse subtrees +// transformed into parser::DoConstructs, scan the parser::Blocks of the program +// and merge adjacent CUFKernelDoConstructs and DoConstructs whenever the +// CUFKernelDoConstruct doesn't already have an embedded DoConstruct. Also +// emit errors about improper or missing DoConstructs. + +namespace Fortran::parser { +struct Mutator { + template bool Pre(A &) { return true; } + template void Post(A &) {} + bool Pre(Block &); +}; + +bool Mutator::Pre(Block &block) { + for (auto iter{block.begin()}; iter != block.end(); ++iter) { + if (auto *kernel{Unwrap(*iter)}) { + auto &nested{std::get>(kernel->t)}; + if (!nested) { + if (auto next{iter}; ++next != block.end()) { + if (auto *doConstruct{Unwrap(*next)}) { + nested = std::move(*doConstruct); + block.erase(next); + } + } + } + } else { + Walk(*iter, *this); + } + } + return false; +} +} // namespace Fortran::parser + +namespace Fortran::semantics { + +bool CanonicalizeCUDA(parser::Program &program) { + parser::Mutator mutator; + parser::Walk(program, mutator); + return true; +} + +using MaybeMsg = std::optional; + +// Traverses an evaluate::Expr<> in search of unsupported operations +// on the device. + +struct DeviceExprChecker + : public evaluate::AnyTraverse { + using Result = MaybeMsg; + using Base = evaluate::AnyTraverse; + DeviceExprChecker() : Base(*this) {} + using Base::operator(); + Result operator()(const evaluate::ProcedureDesignator &x) const { + if (const Symbol * sym{x.GetInterfaceSymbol()}) { + const auto *subp{ + sym->GetUltimate().detailsIf()}; + if (subp) { + if (auto attrs{subp->cudaSubprogramAttrs()}) { + if (*attrs == common::CUDASubprogramAttrs::HostDevice || + *attrs == common::CUDASubprogramAttrs::Device) { + return {}; + } + } + } + } else if (x.GetSpecificIntrinsic()) { + // TODO(CUDA): Check for unsupported intrinsics here + return {}; + } + return parser::MessageFormattedText( + "'%s' may not be called in device code"_err_en_US, x.GetName()); + } +}; + +template static MaybeMsg CheckUnwrappedExpr(const A &x) { + if (const auto *expr{parser::Unwrap(x)}) { + return DeviceExprChecker{}(expr->typedExpr); + } + return {}; +} + +template +static void CheckUnwrappedExpr( + SemanticsContext &context, SourceName at, const A &x) { + if (const auto *expr{parser::Unwrap(x)}) { + if (auto msg{DeviceExprChecker{}(expr->typedExpr)}) { + context.Say(at, std::move(*msg)); + } + } +} + +template struct ActionStmtChecker { + template static MaybeMsg WhyNotOk(const A &x) { + if constexpr (ConstraintTrait) { + return WhyNotOk(x.thing); + } else if constexpr (WrapperTrait) { + return WhyNotOk(x.v); + } else if constexpr (UnionTrait) { + return WhyNotOk(x.u); + } else if constexpr (TupleTrait) { + return WhyNotOk(x.t); + } else { + return parser::MessageFormattedText{ + "Statement may not appear in device code"_err_en_US}; + } + } + template + static MaybeMsg WhyNotOk(const common::Indirection &x) { + return WhyNotOk(x.value()); + } + template + static MaybeMsg WhyNotOk(const std::variant &x) { + return common::visit([](const auto &x) { return WhyNotOk(x); }, x); + } + template + static MaybeMsg WhyNotOk(const std::tuple &x) { + if constexpr (J == sizeof...(As)) { + return {}; + } else if (auto msg{WhyNotOk(std::get(x))}) { + return msg; + } else { + return WhyNotOk<(J + 1)>(x); + } + } + template static MaybeMsg WhyNotOk(const std::list &x) { + for (const auto &y : x) { + if (MaybeMsg result{WhyNotOk(y)}) { + return result; + } + } + return {}; + } + template static MaybeMsg WhyNotOk(const std::optional &x) { + if (x) { + return WhyNotOk(*x); + } else { + return {}; + } + } + template + static MaybeMsg WhyNotOk(const parser::UnlabeledStatement &x) { + return WhyNotOk(x.statement); + } + template + static MaybeMsg WhyNotOk(const parser::Statement &x) { + return WhyNotOk(x.statement); + } + static MaybeMsg WhyNotOk(const parser::AllocateStmt &) { + return {}; // AllocateObjects are checked elsewhere + } + static MaybeMsg WhyNotOk(const parser::AllocateCoarraySpec &) { + return parser::MessageFormattedText( + "A coarray may not be allocated on the device"_err_en_US); + } + static MaybeMsg WhyNotOk(const parser::DeallocateStmt &) { + return {}; // AllocateObjects are checked elsewhere + } + static MaybeMsg WhyNotOk(const parser::AssignmentStmt &x) { + return DeviceExprChecker{}(x.typedAssignment); + } + static MaybeMsg WhyNotOk(const parser::CallStmt &x) { + return DeviceExprChecker{}(x.typedCall); + } + static MaybeMsg WhyNotOk(const parser::ContinueStmt &) { return {}; } + static MaybeMsg WhyNotOk(const parser::IfStmt &x) { + if (auto result{ + CheckUnwrappedExpr(std::get(x.t))}) { + return result; + } + return WhyNotOk( + std::get>(x.t) + .statement); + } + static MaybeMsg WhyNotOk(const parser::NullifyStmt &x) { + for (const auto &y : x.v) { + if (MaybeMsg result{DeviceExprChecker{}(y.typedExpr)}) { + return result; + } + } + return {}; + } + static MaybeMsg WhyNotOk(const parser::PointerAssignmentStmt &x) { + return DeviceExprChecker{}(x.typedAssignment); + } +}; + +template class DeviceContextChecker { +public: + explicit DeviceContextChecker(SemanticsContext &c) : context_{c} {} + void CheckSubprogram(const parser::Name &name, const parser::Block &body) { + if (name.symbol) { + const auto *subp{ + name.symbol->GetUltimate().detailsIf()}; + if (subp && subp->moduleInterface()) { + subp = subp->moduleInterface() + ->GetUltimate() + .detailsIf(); + } + if (subp && + subp->cudaSubprogramAttrs().value_or( + common::CUDASubprogramAttrs::Host) != + common::CUDASubprogramAttrs::Host) { + Check(body); + } + } + } + void Check(const parser::Block &block) { + for (const auto &epc : block) { + Check(epc); + } + } + +private: + void Check(const parser::ExecutionPartConstruct &epc) { + common::visit( + common::visitors{ + [&](const parser::ExecutableConstruct &x) { Check(x); }, + [&](const parser::Statement> + &x) { + context_.Say(x.source, + "Device code may not contain an ENTRY statement"_err_en_US); + }, + [](const parser::Statement> + &) {}, + [](const parser::Statement> + &) {}, + [](const parser::Statement< + common::Indirection> &) {}, + [](const parser::ErrorRecovery &) {}, + }, + epc.u); + } + void Check(const parser::ExecutableConstruct &ec) { + common::visit( + common::visitors{ + [&](const parser::Statement &stmt) { + Check(stmt.statement, stmt.source); + }, + [&](const common::Indirection &x) { + if (const std::optional &control{ + x.value().GetLoopControl()}) { + common::visit([&](const auto &y) { Check(y); }, control->u); + } + Check(std::get(x.value().t)); + }, + [&](const common::Indirection &x) { + Check(std::get(x.value().t)); + }, + [&](const common::Indirection &x) { + Check(x.value()); + }, + [&](const auto &x) { + if (auto source{parser::GetSource(x)}) { + context_.Say(*source, + "Statement may not appear in device code"_err_en_US); + } + }, + }, + ec.u); + } + void Check(const parser::ActionStmt &stmt, const parser::CharBlock &source) { + common::visit( + common::visitors{ + [&](const auto &x) { + if (auto msg{ActionStmtChecker::WhyNotOk(x)}) { + context_.Say(source, std::move(*msg)); + } + }, + }, + stmt.u); + } + void Check(const parser::IfConstruct &ic) { + const auto &ifS{std::get>(ic.t)}; + CheckUnwrappedExpr(context_, ifS.source, + std::get(ifS.statement.t)); + Check(std::get(ic.t)); + for (const auto &eib : + std::get>(ic.t)) { + const auto &eIfS{std::get>(eib.t)}; + CheckUnwrappedExpr(context_, eIfS.source, + std::get(eIfS.statement.t)); + Check(std::get(eib.t)); + } + if (const auto &eb{ + std::get>(ic.t)}) { + Check(std::get(eb->t)); + } + } + void Check(const parser::LoopControl::Bounds &bounds) { + Check(bounds.lower); + Check(bounds.upper); + if (bounds.step) { + Check(*bounds.step); + } + } + void Check(const parser::LoopControl::Concurrent &x) { + const auto &header{std::get(x.t)}; + for (const auto &cc : + std::get>(header.t)) { + Check(std::get<1>(cc.t)); + Check(std::get<2>(cc.t)); + if (const auto &step{ + std::get>(cc.t)}) { + Check(*step); + } + } + if (const auto &mask{ + std::get>(header.t)}) { + Check(*mask); + } + } + void Check(const parser::ScalarLogicalExpr &x) { + Check(DEREF(parser::Unwrap(x))); + } + void Check(const parser::ScalarIntExpr &x) { + Check(DEREF(parser::Unwrap(x))); + } + void Check(const parser::ScalarExpr &x) { + Check(DEREF(parser::Unwrap(x))); + } + void Check(const parser::Expr &expr) { + if (MaybeMsg msg{DeviceExprChecker{}(expr.typedExpr)}) { + context_.Say(expr.source, std::move(*msg)); + } + } + + SemanticsContext &context_; +}; + +void CUDAChecker::Enter(const parser::SubroutineSubprogram &x) { + DeviceContextChecker{context_}.CheckSubprogram( + std::get( + std::get>(x.t).statement.t), + std::get(x.t).v); +} + +void CUDAChecker::Enter(const parser::FunctionSubprogram &x) { + DeviceContextChecker{context_}.CheckSubprogram( + std::get( + std::get>(x.t).statement.t), + std::get(x.t).v); +} + +void CUDAChecker::Enter(const parser::SeparateModuleSubprogram &x) { + DeviceContextChecker{context_}.CheckSubprogram( + std::get>(x.t).statement.v, + std::get(x.t).v); +} + +// !$CUF KERNEL DO semantic checks + +static int DoConstructTightNesting( + const parser::DoConstruct *doConstruct, const parser::Block *&innerBlock) { + if (!doConstruct || !doConstruct->IsDoNormal()) { + return 0; + } + innerBlock = &std::get(doConstruct->t); + if (innerBlock->size() == 1) { + if (const auto *execConstruct{ + std::get_if(&innerBlock->front().u)}) { + if (const auto *next{ + std::get_if>( + &execConstruct->u)}) { + return 1 + DoConstructTightNesting(&next->value(), innerBlock); + } + } + } + return 1; +} + +void CUDAChecker::Enter(const parser::CUFKernelDoConstruct &x) { + auto source{std::get(x.t).source}; + const auto &directive{std::get(x.t)}; + std::int64_t depth{1}; + if (auto expr{AnalyzeExpr(context_, + std::get>( + directive.t))}) { + depth = evaluate::ToInt64(expr).value_or(0); + if (depth <= 0) { + context_.Say(source, + "!$CUF KERNEL DO (%jd): loop nesting depth must be positive"_err_en_US, + std::intmax_t{depth}); + depth = 1; + } + } + const parser::DoConstruct *doConstruct{common::GetPtrFromOptional( + std::get>(x.t))}; + const parser::Block *innerBlock{nullptr}; + if (DoConstructTightNesting(doConstruct, innerBlock) < depth) { + context_.Say(source, + "!$CUF KERNEL DO (%jd) must be followed by a DO construct with tightly nested outer levels of counted DO loops"_err_en_US, + std::intmax_t{depth}); + } + if (innerBlock) { + DeviceContextChecker{context_}.Check(*innerBlock); + } +} + +} // namespace Fortran::semantics Index: flang/lib/Semantics/check-deallocate.cpp =================================================================== --- flang/lib/Semantics/check-deallocate.cpp +++ flang/lib/Semantics/check-deallocate.cpp @@ -19,10 +19,13 @@ void DeallocateChecker::Leave(const parser::DeallocateStmt &deallocateStmt) { for (const parser::AllocateObject &allocateObject : std::get>(deallocateStmt.t)) { + parser::CharBlock source; + const Symbol *symbol{nullptr}; common::visit( common::visitors{ [&](const parser::Name &name) { - auto const *symbol{name.symbol}; + source = name.source; + symbol = name.symbol; if (context_.HasError(symbol)) { // already reported an error } else if (!IsVariableName(*symbol)) { @@ -58,9 +61,10 @@ [&](const parser::StructureComponent &structureComponent) { // Only perform structureComponent checks if it was successfully // analyzed by expression analysis. + source = structureComponent.component.source; + symbol = structureComponent.component.symbol; if (const auto *expr{GetExpr(context_, allocateObject)}) { - if (const Symbol *symbol{structureComponent.component.symbol}) { - auto source{structureComponent.component.source}; + if (symbol) { if (!IsAllocatableOrPointer(*symbol)) { // C932 context_.Say(source, "Component in DEALLOCATE statement must have the ALLOCATABLE or POINTER attribute"_err_en_US); Index: flang/lib/Semantics/check-declarations.cpp =================================================================== --- flang/lib/Semantics/check-declarations.cpp +++ flang/lib/Semantics/check-declarations.cpp @@ -114,6 +114,19 @@ } return msg; } + template parser::Message *WarnIfNotInModuleFile(A &&...x) { + if (FindModuleFileContaining(context_.FindScope(messages_.at()))) { + return nullptr; + } + return messages_.Say(std::forward(x)...); + } + template + parser::Message *WarnIfNotInModuleFile(parser::CharBlock source, A &&...x) { + if (FindModuleFileContaining(context_.FindScope(source))) { + return nullptr; + } + return messages_.Say(source, std::forward(x)...); + } bool IsResultOkToDiffer(const FunctionResult &); void CheckGlobalName(const Symbol &); void CheckExplicitSave(const Symbol &); @@ -216,9 +229,8 @@ void CheckHelper::Check(const Symbol &symbol) { if (symbol.name().size() > common::maxNameLen && - &symbol == &symbol.GetUltimate() && - !FindModuleFileContaining(symbol.owner())) { - messages_.Say(symbol.name(), + &symbol == &symbol.GetUltimate()) { + WarnIfNotInModuleFile(symbol.name(), "%s has length %d, which is greater than the maximum name length " "%d"_port_en_US, symbol.name(), symbol.name().size(), common::maxNameLen); @@ -615,6 +627,7 @@ WarnMissingFinal(symbol); const DeclTypeSpec *type{details.type()}; const DerivedTypeSpec *derived{type ? type->AsDerived() : nullptr}; + bool isComponent{symbol.owner().IsDerivedType()}; if (!details.coshape().empty()) { bool isDeferredCoshape{details.coshape().CanBeDeferredShape()}; if (IsAllocatable(symbol)) { @@ -623,7 +636,7 @@ " coshape"_err_en_US, symbol.name()); } - } else if (symbol.owner().IsDerivedType()) { // C746 + } else if (isComponent) { // C746 std::string deferredMsg{ isDeferredCoshape ? "" : " and have a deferred coshape"}; messages_.Say("Component '%s' is a coarray and must have the ALLOCATABLE" @@ -727,7 +740,7 @@ if (IsPassedViaDescriptor(symbol)) { if (IsAllocatableOrPointer(symbol)) { if (inExplicitInterface) { - messages_.Say( + WarnIfNotInModuleFile( "!DIR$ IGNORE_TKR should not apply to an allocatable or pointer"_warn_en_US); } else { messages_.Say( @@ -735,10 +748,10 @@ } } else if (ignoreTKR.test(common::IgnoreTKR::Rank)) { if (ignoreTKR.count() == 1 && evaluate::IsAssumedRank(symbol)) { - messages_.Say( + WarnIfNotInModuleFile( "!DIR$ IGNORE_TKR(R) is not meaningful for an assumed-rank array"_warn_en_US); } else if (inExplicitInterface) { - messages_.Say( + WarnIfNotInModuleFile( "!DIR$ IGNORE_TKR(R) should not apply to a dummy argument passed via descriptor"_warn_en_US); } else { messages_.Say( @@ -806,9 +819,8 @@ messages_.Say("A dummy argument must not be initialized"_err_en_US); } else if (IsFunctionResult(symbol)) { messages_.Say("A function result must not be initialized"_err_en_US); - } else if (IsInBlankCommon(symbol) && - !FindModuleFileContaining(symbol.owner())) { - messages_.Say( + } else if (IsInBlankCommon(symbol)) { + WarnIfNotInModuleFile( "A variable in blank COMMON should not be initialized"_port_en_US); } } @@ -848,6 +860,156 @@ "'%s' is a data object and may not be EXTERNAL"_err_en_US, symbol.name()); } + + // Check CUDA attributes and special circumstances of being in device + // subprograms + const Scope &progUnit{GetProgramUnitContaining(symbol)}; + const auto *subpDetails{!isComponent && progUnit.symbol() + ? progUnit.symbol()->detailsIf() + : nullptr}; + bool inDeviceSubprogram{IsCUDADeviceContext(&symbol.owner())}; + if (inDeviceSubprogram) { + if (IsSaved(symbol)) { + WarnIfNotInModuleFile( + "'%s' should not have the SAVE attribute or initialization in a device subprogram"_warn_en_US, + symbol.name()); + } + if (IsPointer(symbol)) { + WarnIfNotInModuleFile( + "Pointer '%s' may not be associated in a device subprogram"_warn_en_US, + symbol.name()); + } + if (details.isDummy() && + details.cudaDataAttr().value_or(common::CUDADataAttr::Device) != + common::CUDADataAttr::Device && + details.cudaDataAttr().value_or(common::CUDADataAttr::Device) != + common::CUDADataAttr::Managed) { + WarnIfNotInModuleFile( + "Dummy argument '%s' may not have ATTRIBUTES(%s) in a device subprogram"_warn_en_US, + symbol.name(), + parser::ToUpperCaseLetters( + common::EnumToString(*details.cudaDataAttr()))); + } + } + if (details.cudaDataAttr()) { + if (auto dyType{evaluate::DynamicType::From(symbol)}) { + if (dyType->category() != TypeCategory::Derived) { + if (!IsCUDAIntrinsicType(*dyType)) { + messages_.Say( + "'%s' has intrinsic type '%s' that is not available on the device"_err_en_US, + symbol.name(), dyType->AsFortran()); + } + } + } + auto attr{*details.cudaDataAttr()}; + switch (attr) { + case common::CUDADataAttr::Constant: + if (IsAllocatableOrPointer(symbol) || symbol.attrs().test(Attr::TARGET)) { + messages_.Say( + "Object '%s' with ATTRIBUTES(CONSTANT) may not be allocatable, pointer, or target"_err_en_US, + symbol.name()); + } else if (auto shape{evaluate::GetShape(foldingContext_, symbol)}; + !shape || + !evaluate::AsConstantExtents(foldingContext_, *shape)) { + messages_.Say( + "Object '%s' with ATTRIBUTES(CONSTANT) must have constant array bounds"_err_en_US, + symbol.name()); + } + break; + case common::CUDADataAttr::Device: + if (isComponent && !IsAllocatable(symbol)) { + messages_.Say( + "Component '%s' with ATTRIBUTES(DEVICE) must also be allocatable"_err_en_US, + symbol.name()); + } + break; + case common::CUDADataAttr::Managed: + if (!IsAutomatic(symbol) && !IsAllocatable(symbol) && + !details.isDummy()) { + messages_.Say( + "Object '%s' with ATTRIBUTES(MANAGED) must also be allocatable, automatic, or a dummy argument"_err_en_US, + symbol.name()); + } + break; + case common::CUDADataAttr::Pinned: + if (inDeviceSubprogram) { + WarnIfNotInModuleFile( + "Object '%s' with ATTRIBUTES(PINNED) may not be declared in a device subprogram"_warn_en_US, + symbol.name()); + } else if (IsPointer(symbol)) { + WarnIfNotInModuleFile( + "Object '%s' with ATTRIBUTES(PINNED) may not be a pointer"_warn_en_US, + symbol.name()); + } else if (!IsAllocatable(symbol)) { + WarnIfNotInModuleFile( + "Object '%s' with ATTRIBUTES(PINNED) should also be allocatable"_warn_en_US, + symbol.name()); + } + break; + case common::CUDADataAttr::Shared: + if (IsAllocatableOrPointer(symbol) || symbol.attrs().test(Attr::TARGET)) { + messages_.Say( + "Object '%s' with ATTRIBUTES(SHARED) may not be allocatable, pointer, or target"_err_en_US, + symbol.name()); + } else if (!inDeviceSubprogram) { + messages_.Say( + "Object '%s' with ATTRIBUTES(SHARED) must be declared in a device subprogram"_err_en_US, + symbol.name()); + } + break; + case common::CUDADataAttr::Texture: + messages_.Say( + "ATTRIBUTES(TEXTURE) is obsolete and no longer supported"_err_en_US); + break; + } + if (attr != common::CUDADataAttr::Pinned) { + if (details.commonBlock()) { + messages_.Say( + "Object '%s' with ATTRIBUTES(%s) may not be in COMMON"_err_en_US, + symbol.name(), + parser::ToUpperCaseLetters(common::EnumToString(attr))); + } else if (FindEquivalenceSet(symbol)) { + messages_.Say( + "Object '%s' with ATTRIBUTES(%s) may not be in an equivalence group"_err_en_US, + symbol.name(), + parser::ToUpperCaseLetters(common::EnumToString(attr))); + } + } + if (subpDetails /* not a module variable */ && IsSaved(symbol) && + !inDeviceSubprogram && !IsAllocatable(symbol) && + attr == common::CUDADataAttr::Device) { + messages_.Say( + "Saved object '%s' in host code may not have ATTRIBUTES(DEVICE) unless allocatable"_err_en_US, + symbol.name(), + parser::ToUpperCaseLetters(common::EnumToString(attr))); + } + if (isComponent) { + if (attr == common::CUDADataAttr::Device) { + const DeclTypeSpec *type{symbol.GetType()}; + if (const DerivedTypeSpec * + derived{type ? type->AsDerived() : nullptr}) { + DirectComponentIterator directs{*derived}; + if (auto iter{std::find_if(directs.begin(), directs.end(), + [](const Symbol &) { return false; })}) { + messages_.Say( + "Derived type component '%s' may not have ATTRIBUTES(DEVICE) as it has a direct device component '%s'"_err_en_US, + symbol.name(), iter.BuildResultDesignatorName()); + } + } + } else if (attr == common::CUDADataAttr::Constant || + attr == common::CUDADataAttr::Shared) { + messages_.Say( + "Derived type component '%s' may not have ATTRIBUTES(%s)"_err_en_US, + symbol.name(), + parser::ToUpperCaseLetters(common::EnumToString(attr))); + } + } else if (!subpDetails && symbol.owner().kind() != Scope::Kind::Module && + symbol.owner().kind() != Scope::Kind::MainProgram) { + messages_.Say( + "ATTRIBUTES(%s) may apply only to module, host subprogram, or device subprogram data"_err_en_US, + parser::ToUpperCaseLetters(common::EnumToString(attr))); + } + } } void CheckHelper::CheckPointerInitialization(const Symbol &symbol) { @@ -919,6 +1081,9 @@ bool canBeAssumedShape{arraySpec.CanBeAssumedShape()}; bool canBeAssumedSize{arraySpec.CanBeAssumedSize()}; bool isAssumedRank{arraySpec.IsAssumedRank()}; + bool isCUDAShared{ + GetCUDADataAttr(&symbol).value_or(common::CUDADataAttr::Device) == + common::CUDADataAttr::Shared}; std::optional msg; if (symbol.test(Symbol::Flag::CrayPointee) && !isExplicit && !canBeAssumedSize) { @@ -948,12 +1113,12 @@ } } else if (canBeAssumedShape && !canBeDeferred) { msg = "Assumed-shape array '%s' must be a dummy argument"_err_en_US; - } else if (canBeAssumedSize && !canBeImplied) { // C833 + } else if (canBeAssumedSize && !canBeImplied && !isCUDAShared) { // C833 msg = "Assumed-size array '%s' must be a dummy argument"_err_en_US; } else if (isAssumedRank) { // C837 msg = "Assumed-rank array '%s' must be a dummy argument"_err_en_US; } else if (canBeImplied) { - if (!IsNamedConstant(symbol)) { // C835, C836 + if (!IsNamedConstant(symbol) && !isCUDAShared) { // C835, C836 msg = "Implied-shape array '%s' must be a named constant or a " "dummy argument"_err_en_US; } @@ -1187,6 +1352,50 @@ } CheckExternal(symbol); CheckModuleProcedureDef(symbol); + auto cudaAttrs{details.cudaSubprogramAttrs()}; + if (cudaAttrs && + (*cudaAttrs == common::CUDASubprogramAttrs::Global || + *cudaAttrs == common::CUDASubprogramAttrs::Grid_Global) && + details.isFunction()) { + messages_.Say(symbol.name(), + "A function may not have ATTRIBUTES(GLOBAL) or ATTRIBUTES(GRID_GLOBAL)"_err_en_US); + } + if (cudaAttrs && *cudaAttrs != common::CUDASubprogramAttrs::Host) { + // CUDA device subprogram checks + if (symbol.attrs().HasAny({Attr::RECURSIVE, Attr::PURE, Attr::ELEMENTAL})) { + messages_.Say(symbol.name(), + "A device subprogram may not be RECURSIVE, PURE, or ELEMENTAL"_err_en_US); + } + if (ClassifyProcedure(symbol) == ProcedureDefinitionClass::Internal) { + messages_.Say(symbol.name(), + "A device subprogram may not be an internal subprogram"_err_en_US); + } else if ((*cudaAttrs == common::CUDASubprogramAttrs::Device || + *cudaAttrs == common::CUDASubprogramAttrs::HostDevice) && + (symbol.owner().kind() != Scope::Kind::Module || + details.isInterface())) { + messages_.Say(symbol.name(), + "An ATTRIBUTES(DEVICE) subprogram must be a top-level module procedure"_err_en_US); + } + } + if ((!details.cudaLaunchBounds().empty() || + !details.cudaClusterDims().empty()) && + !(cudaAttrs && + (*cudaAttrs == common::CUDASubprogramAttrs::Global || + *cudaAttrs == common::CUDASubprogramAttrs::Grid_Global))) { + messages_.Say(symbol.name(), + "A subroutine may not have LAUNCH_BOUNDS() or CLUSTER_DIMS() unless it has ATTRIBUTES(GLOBAL) or ATTRIBUTES(GRID_GLOBAL)"_err_en_US); + } + if (!IsStmtFunction(symbol)) { + if (const Scope * outerDevice{FindCUDADeviceContext(&symbol.owner())}; + outerDevice && outerDevice->symbol()) { + if (auto *msg{messages_.Say(symbol.name(), + "'%s' may not be an internal procedure of CUDA device subprogram '%s'"_err_en_US, + symbol.name(), outerDevice->symbol()->name())}) { + msg->Attach(outerDevice->symbol()->name(), + "Containing CUDA device subprogram"_en_US); + } + } + } } void CheckHelper::CheckExternal(const Symbol &symbol) { @@ -1215,7 +1424,7 @@ if (chars->HasExplicitInterface()) { std::string whyNot; if (!chars->IsCompatibleWith(*globalChars, &whyNot)) { - msg = messages_.Say( + msg = WarnIfNotInModuleFile( "The global subprogram '%s' is not compatible with its local procedure declaration (%s)"_warn_en_US, global->name(), whyNot); } @@ -1241,7 +1450,7 @@ if (auto previousChars{Characterize(previous)}) { std::string whyNot; if (!chars->IsCompatibleWith(*previousChars, &whyNot)) { - if (auto *msg{messages_.Say( + if (auto *msg{WarnIfNotInModuleFile( "The external interface '%s' is not compatible with an earlier definition (%s)"_warn_en_US, symbol.name(), whyNot)}) { evaluate::AttachDeclaration(msg, previous); @@ -1628,12 +1837,14 @@ return true; // OK } bool isFatal{msg->IsFatal()}; - SayWithDeclaration( - specific, std::move(*msg), MakeOpName(opName), specific.name()); + if (isFatal || !FindModuleFileContaining(specific.owner())) { + SayWithDeclaration( + specific, std::move(*msg), MakeOpName(opName), specific.name()); + } if (isFatal) { context_.SetError(specific); } - return false; + return !isFatal; } // If the number of arguments is wrong for this intrinsic operator, return @@ -1694,15 +1905,24 @@ dataObject == nullptr) { msg = "In %s function '%s', dummy argument '%s' must be a" " data object"_err_en_US; + } else if (dataObject->intent == common::Intent::Out) { + msg = + "In %s function '%s', dummy argument '%s' may not be INTENT(OUT)"_err_en_US; } else if (dataObject->intent != common::Intent::In && !dataObject->attrs.test(DummyDataObject::Attr::Value)) { - msg = "In %s function '%s', dummy argument '%s' must have INTENT(IN)" - " or VALUE attribute"_err_en_US; + msg = + "In %s function '%s', dummy argument '%s' should have INTENT(IN) or VALUE attribute"_warn_en_US; } if (msg) { - SayWithDeclaration(symbol, std::move(*msg), - parser::ToUpperCaseLetters(opName.ToString()), symbol.name(), arg.name); - return false; + bool isFatal{msg->IsFatal()}; + if (isFatal || !FindModuleFileContaining(symbol.owner())) { + SayWithDeclaration(symbol, std::move(*msg), + parser::ToUpperCaseLetters(opName.ToString()), symbol.name(), + arg.name); + } + if (isFatal) { + return false; + } } return true; } @@ -1748,17 +1968,23 @@ " may not be OPTIONAL"_err_en_US; } else if (const auto *dataObject{std::get_if(&arg.u)}) { if (pos == 0) { - if (dataObject->intent != common::Intent::Out && + if (dataObject->intent == common::Intent::In) { + msg = "In defined assignment subroutine '%s', first dummy argument '%s'" + " may not have INTENT(IN)"_err_en_US; + } else if (dataObject->intent != common::Intent::Out && dataObject->intent != common::Intent::InOut) { msg = "In defined assignment subroutine '%s', first dummy argument '%s'" - " must have INTENT(OUT) or INTENT(INOUT)"_err_en_US; + " should have INTENT(OUT) or INTENT(INOUT)"_warn_en_US; } } else if (pos == 1) { - if (dataObject->intent != common::Intent::In && + if (dataObject->intent == common::Intent::Out) { + msg = "In defined assignment subroutine '%s', second dummy" + " argument '%s' may not have INTENT(OUT)"_err_en_US; + } else if (dataObject->intent != common::Intent::In && !dataObject->attrs.test(DummyDataObject::Attr::Value)) { msg = "In defined assignment subroutine '%s', second dummy" - " argument '%s' must have INTENT(IN) or VALUE attribute"_err_en_US; + " argument '%s' should have INTENT(IN) or VALUE attribute"_warn_en_US; } else if (dataObject->attrs.test(DummyDataObject::Attr::Pointer)) { msg = "In defined assignment subroutine '%s', second dummy argument '%s' must not be a pointer"_err_en_US; @@ -1774,9 +2000,14 @@ " must be a data object"_err_en_US; } if (msg) { - SayWithDeclaration(symbol, std::move(*msg), symbol.name(), arg.name); - context_.SetError(symbol); - return false; + bool isFatal{msg->IsFatal()}; + if (isFatal || !FindModuleFileContaining(symbol.owner())) { + SayWithDeclaration(symbol, std::move(*msg), symbol.name(), arg.name); + } + if (isFatal) { + context_.SetError(symbol); + return false; + } } return true; } @@ -1809,10 +2040,10 @@ if (!derivedDetails->finals().empty() && !derivedDetails->GetFinalForRank(rank)) { if (auto *msg{derivedSym == initialDerivedSym - ? messages_.Say(symbol.name(), + ? WarnIfNotInModuleFile(symbol.name(), "'%s' of derived type '%s' does not have a FINAL subroutine for its rank (%d)"_warn_en_US, symbol.name(), derivedSym->name(), rank) - : messages_.Say(symbol.name(), + : WarnIfNotInModuleFile(symbol.name(), "'%s' of derived type '%s' extended from '%s' does not have a FINAL subroutine for its rank (%d)"_warn_en_US, symbol.name(), initialDerivedSym->name(), derivedSym->name(), rank)}) { @@ -2431,15 +2662,17 @@ type->category() == DeclTypeSpec::Character && type->characterTypeSpec().length().isDeferred()) { // ok; F'2018 18.3.6 p2(6) - } else if (derived || IsInteroperableIntrinsicType(*type)) { + } else if (derived || + IsInteroperableIntrinsicType(*type, context_.languageFeatures())) { // F'2018 18.3.6 p2(4,5) - } else if (type->category() == DeclTypeSpec::Logical && IsDummy(symbol) && - evaluate::GetRank(*shape) == 0) { - // Special exception: LOGICAL scalar dummy arguments can be converted - // before a call -- & after if not INTENT(IN) -- without loss of - // information, and are accepted by some older compilers. - messages_.Say(symbol.name(), - "A BIND(C) LOGICAL dummy argument should have the interoperable KIND=C_BOOL"_port_en_US); + } else if (type->category() == DeclTypeSpec::Logical) { + if (IsDummy(symbol)) { + WarnIfNotInModuleFile(symbol.name(), + "A BIND(C) LOGICAL dummy argument should have the interoperable KIND=C_BOOL"_port_en_US); + } else { + WarnIfNotInModuleFile(symbol.name(), + "A BIND(C) LOGICAL object should have the interoperable KIND=C_BOOL"_port_en_US); + } } else if (symbol.attrs().test(Attr::VALUE)) { messages_.Say(symbol.name(), "A BIND(C) VALUE dummy argument must have an interoperable type"_err_en_US); @@ -2451,12 +2684,13 @@ } } if (IsOptional(symbol) && !symbol.attrs().test(Attr::VALUE)) { - messages_.Say(symbol.name(), + WarnIfNotInModuleFile(symbol.name(), "An interoperable procedure with an OPTIONAL dummy argument might not be portable"_port_en_US); } } else if (const auto *proc{symbol.detailsIf()}) { - if (!proc->procInterface() || - !proc->procInterface()->attrs().test(Attr::BIND_C)) { + if (!proc->isDummy() && + (!proc->procInterface() || + !proc->procInterface()->attrs().test(Attr::BIND_C))) { messages_.Say(symbol.name(), "An interface name with BIND attribute must be specified if the BIND attribute is specified in a procedure declaration statement"_err_en_US); context_.SetError(symbol); @@ -2508,10 +2742,21 @@ } context_.SetError(symbol); } - } else if (!IsInteroperableIntrinsicType(*type)) { - messages_.Say(component->name(), - "Each component of an interoperable derived type must have an interoperable type"_err_en_US); - context_.SetError(symbol); + } else if (!IsInteroperableIntrinsicType( + *type, context_.languageFeatures())) { + auto maybeDyType{evaluate::DynamicType::From(*type)}; + if (type->category() == DeclTypeSpec::Logical) { + WarnIfNotInModuleFile(component->name(), + "A LOGICAL component of a BIND(C) type should have the interoperable KIND=C_BOOL"_port_en_US); + } else if (type->category() == DeclTypeSpec::Character && + maybeDyType && maybeDyType->kind() == 1) { + WarnIfNotInModuleFile(component->name(), + "A CHARACTER component of a BIND(C) type should have length 1"_port_en_US); + } else { + messages_.Say(component->name(), + "Each component of an interoperable derived type must have an interoperable type"_err_en_US); + context_.SetError(symbol); + } } } if (auto extents{ @@ -2523,9 +2768,8 @@ } } } - if (derived->componentNames().empty() && - !FindModuleFileContaining(symbol.owner())) { // C1805 - messages_.Say(symbol.name(), + if (derived->componentNames().empty()) { // C1805 + WarnIfNotInModuleFile(symbol.name(), "A derived type with the BIND attribute is empty"_port_en_US); } } Index: flang/lib/Semantics/definable.cpp =================================================================== --- flang/lib/Semantics/definable.cpp +++ flang/lib/Semantics/definable.cpp @@ -134,6 +134,33 @@ 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; } Index: flang/lib/Semantics/expression.cpp =================================================================== --- flang/lib/Semantics/expression.cpp +++ flang/lib/Semantics/expression.cpp @@ -2616,6 +2616,9 @@ 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); } @@ -2682,6 +2685,55 @@ 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}; @@ -2693,17 +2745,17 @@ 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}; + 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)}) { @@ -2748,17 +2800,40 @@ 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; } } @@ -3660,14 +3735,13 @@ 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); Index: flang/lib/Semantics/semantics.cpp =================================================================== --- flang/lib/Semantics/semantics.cpp +++ flang/lib/Semantics/semantics.cpp @@ -16,6 +16,7 @@ #include "check-arithmeticif.h" #include "check-case.h" #include "check-coarray.h" +#include "check-cuda.h" #include "check-data.h" #include "check-deallocate.h" #include "check-declarations.h" @@ -158,12 +159,14 @@ }; using StatementSemanticsPass1 = ExprChecker; -using StatementSemanticsPass2 = SemanticsVisitor; +using StatementSemanticsPass2 = SemanticsVisitor; +using StatementSemanticsPass3 = + SemanticsVisitor; static bool PerformStatementSemantics( SemanticsContext &context, parser::Program &program) { @@ -174,6 +177,11 @@ StatementSemanticsPass1{context}.Walk(program); StatementSemanticsPass2 pass2{context}; pass2.Walk(program); + if (context.languageFeatures().IsEnabled(common::LanguageFeature::OpenACC) || + context.languageFeatures().IsEnabled(common::LanguageFeature::OpenMP) || + context.languageFeatures().IsEnabled(common::LanguageFeature::CUDA)) { + StatementSemanticsPass3{context}.Walk(program); + } if (!context.AnyFatalError()) { pass2.CompileDataInitializationsIntoInitializers(); } @@ -515,6 +523,7 @@ parser::CanonicalizeDo(program_) && // force line break CanonicalizeAcc(context_.messages(), program_) && CanonicalizeOmp(context_.messages(), program_) && + CanonicalizeCUDA(program_) && PerformStatementSemantics(context_, program_) && ModFileWriter{context_}.WriteAll(); } @@ -556,7 +565,7 @@ if (scope.derivedTypeSpec()) { os << " instantiation of " << *scope.derivedTypeSpec(); } - os << '\n'; + os << " sourceRange=" << scope.sourceRange().size() << " bytes\n"; ++indent; for (const auto &pair : scope) { const auto &symbol{*pair.second}; Index: flang/lib/Semantics/type.cpp =================================================================== --- flang/lib/Semantics/type.cpp +++ flang/lib/Semantics/type.cpp @@ -797,9 +797,10 @@ return o << x.AsFortran(); } -bool IsInteroperableIntrinsicType(const DeclTypeSpec &type) { +bool IsInteroperableIntrinsicType( + const DeclTypeSpec &type, const common::LanguageFeatureControl &features) { auto dyType{evaluate::DynamicType::From(type)}; - return dyType && IsInteroperableIntrinsicType(*dyType); + return dyType && IsInteroperableIntrinsicType(*dyType, &features); } } // namespace Fortran::semantics Index: flang/module/__cuda_builtins.f90 =================================================================== --- /dev/null +++ flang/module/__cuda_builtins.f90 @@ -0,0 +1,19 @@ +!===-- module/__cuda_builtins.f90 ------------------------------------------===! +! +! Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +! See https://llvm.org/LICENSE.txt for license information. +! SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +! +!===------------------------------------------------------------------------===! + +! These CUDA predefined variables are automatically available in device +! subprograms. + +module __CUDA_builtins + use __Fortran_builtins, only: & + threadIdx => __builtin_threadIdx, & + blockDim => __builtin_blockDim, & + blockIdx => __builtin_blockIdx, & + gridDim => __builtin_gridDim, & + warpsize => __builtin_warpsize +end module Index: flang/module/__fortran_builtins.f90 =================================================================== --- flang/module/__fortran_builtins.f90 +++ flang/module/__fortran_builtins.f90 @@ -73,4 +73,23 @@ type(__builtin_team_type) :: team_type end type + interface operator(==) + module procedure __builtin_c_ptr_eq + end interface + interface operator(/=) + module procedure __builtin_c_ptr_eq + end interface + +contains + + elemental logical function __builtin_c_ptr_eq(x, y) + type(__builtin_c_ptr), intent(in) :: x, y + __builtin_c_ptr_eq = x%__address == y%__address + end function + + elemental logical function __builtin_c_ptr_ne(x, y) + type(__builtin_c_ptr), intent(in) :: x, y + __builtin_c_ptr_ne = x%__address /= y%__address + end function + end module Index: flang/module/iso_c_binding.f90 =================================================================== --- flang/module/iso_c_binding.f90 +++ flang/module/iso_c_binding.f90 @@ -15,7 +15,8 @@ c_ptr => __builtin_c_ptr, & c_funptr => __builtin_c_funptr, & c_sizeof => sizeof, & - c_loc => __builtin_c_loc + c_loc => __builtin_c_loc, & + operator(==), operator(/=) type(c_ptr), parameter :: c_null_ptr = c_ptr(0) type(c_funptr), parameter :: c_null_funptr = c_funptr(0) Index: flang/test/Parser/cuf-sanity-common =================================================================== --- flang/test/Parser/cuf-sanity-common +++ flang/test/Parser/cuf-sanity-common @@ -18,7 +18,7 @@ attributes(global) launch_bounds(1, 2) subroutine lbsub; end attributes(global) cluster_dims(1, 2, 3) subroutine cdsub; end attributes(device) subroutine attrs - attributes(device) :: devx1 + attributes(device)::devx1 real, device :: devx2 end subroutine subroutine test Index: flang/test/Parser/cuf-sanity-tree.CUF =================================================================== --- flang/test/Parser/cuf-sanity-tree.CUF +++ flang/test/Parser/cuf-sanity-tree.CUF @@ -106,6 +106,9 @@ !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 @@ !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' Index: flang/test/Parser/cuf-sanity-unparse.CUF =================================================================== --- flang/test/Parser/cuf-sanity-unparse.CUF +++ flang/test/Parser/cuf-sanity-unparse.CUF @@ -23,6 +23,7 @@ !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,10 @@ !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 + Index: flang/test/Semantics/bind-c06.f90 =================================================================== --- flang/test/Semantics/bind-c06.f90 +++ flang/test/Semantics/bind-c06.f90 @@ -65,7 +65,7 @@ end type type, bind(c) :: t10 - !ERROR: Each component of an interoperable derived type must have an interoperable type + !WARNING: A CHARACTER component of a BIND(C) type should have length 1 character(len=2) x end type type, bind(c) :: t11 @@ -73,7 +73,7 @@ character(kind=2) x end type type, bind(c) :: t12 - !ERROR: Each component of an interoperable derived type must have an interoperable type + !PORTABILITY: A LOGICAL component of a BIND(C) type should have the interoperable KIND=C_BOOL logical(kind=8) x end type type, bind(c) :: t13 Index: flang/test/Semantics/cuf02.cuf =================================================================== --- /dev/null +++ flang/test/Semantics/cuf02.cuf @@ -0,0 +1,49 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +module m + interface + !ERROR: An ATTRIBUTES(DEVICE) subprogram must be a top-level module procedure + attributes(device) subroutine exts1 + end + end interface + contains + !ERROR: A device subprogram may not be RECURSIVE, PURE, or ELEMENTAL + recursive attributes(device) subroutine s1 + end + !ERROR: A device subprogram may not be RECURSIVE, PURE, or ELEMENTAL + pure attributes(device) subroutine s2 + end + !ERROR: A device subprogram may not be RECURSIVE, PURE, or ELEMENTAL + elemental attributes(device) subroutine s3 + end + subroutine s4 + contains + !ERROR: A device subprogram may not be an internal subprogram + attributes(device) subroutine inner + end + end + attributes(device) subroutine s5 ! nvfortran crashes on this one + contains + !ERROR: 'inner' may not be an internal procedure of CUDA device subprogram 's5' + subroutine inner + end + end + attributes(device) subroutine s6 + stmtfunc(x) = x + 1. ! ok + end + !ERROR: A function may not have ATTRIBUTES(GLOBAL) or ATTRIBUTES(GRID_GLOBAL) + attributes(global) real function f1 + end + !ERROR: A device subprogram may not be RECURSIVE, PURE, or ELEMENTAL + recursive attributes(global) subroutine s7 + end + !ERROR: A device subprogram may not be RECURSIVE, PURE, or ELEMENTAL + pure attributes(global) subroutine s8 + end + !ERROR: A device subprogram may not be RECURSIVE, PURE, or ELEMENTAL + elemental attributes(global) subroutine s9 + end +end + +!ERROR: An ATTRIBUTES(DEVICE) subprogram must be a top-level module procedure +attributes(device) subroutine exts1 +end Index: flang/test/Semantics/cuf03.cuf =================================================================== --- /dev/null +++ flang/test/Semantics/cuf03.cuf @@ -0,0 +1,59 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +! Exercise CUDA data attribute checks +module m + real, constant :: mc ! ok + real, constant :: mci = 1. ! ok + !ERROR: Object 'mcl' with ATTRIBUTES(CONSTANT) may not be allocatable, pointer, or target + real, constant, allocatable :: mcl + !ERROR: Object 'mcp' with ATTRIBUTES(CONSTANT) may not be allocatable, pointer, or target + real, constant, pointer :: mcp + !ERROR: Object 'mct' with ATTRIBUTES(CONSTANT) may not be allocatable, pointer, or target + real, constant, target :: mct + real, device :: md ! ok + real, device :: mdi = 1. + real, device, allocatable :: mdl ! ok + real, device, pointer :: mdp ! ok at module level + real, device, target :: mdt ! ok + !ERROR: Object 'ms' with ATTRIBUTES(SHARED) must be declared in a device subprogram + real, shared :: ms + !ERROR: Object 'msi' with ATTRIBUTES(SHARED) must be declared in a device subprogram + real, shared :: msi = 1. + !ERROR: Object 'msl' with ATTRIBUTES(SHARED) may not be allocatable, pointer, or target + real, shared, allocatable :: msl + !ERROR: Object 'msp' with ATTRIBUTES(SHARED) may not be allocatable, pointer, or target + real, shared, pointer :: msp + !ERROR: Object 'mst' with ATTRIBUTES(SHARED) may not be allocatable, pointer, or target + real, shared, target :: mst + !ERROR: Object 'msa' with ATTRIBUTES(SHARED) must be declared in a device subprogram + real, shared :: msa(*) + !ERROR: Object 'mm' with ATTRIBUTES(MANAGED) must also be allocatable, automatic, or a dummy argument + real, managed :: mm + !ERROR: Object 'mmi' with ATTRIBUTES(MANAGED) must also be allocatable, automatic, or a dummy argument + real, managed :: mmi = 1. + real, managed, allocatable :: mml ! ok + !ERROR: Object 'mmp' with ATTRIBUTES(MANAGED) must also be allocatable, automatic, or a dummy argument + real, managed, pointer :: mmp ! ok + !ERROR: Object 'mmt' with ATTRIBUTES(MANAGED) must also be allocatable, automatic, or a dummy argument + real, managed, target :: mmt + !WARNING: Object 'mp' with ATTRIBUTES(PINNED) should also be allocatable + real, pinned :: mp + !WARNING: Object 'mpi' with ATTRIBUTES(PINNED) should also be allocatable + real, pinned :: mpi = 1. + real, pinned, allocatable :: mpl ! ok + !ERROR: Object 'mpp' with ATTRIBUTES(PINNED) may not be a pointer + real, pinned, pointer :: mpp + !WARNING: Object 'mpt' with ATTRIBUTES(PINNED) should also be allocatable + real, pinned, target :: mpt ! ok + !ERROR: ATTRIBUTES(TEXTURE) is obsolete and no longer supported + real, texture, pointer :: mt + !ERROR: 'bigint' has intrinsic type 'INTEGER(16)' that is not available on the device + integer(16), device :: bigint + contains + attributes(device) subroutine devsubr(n,da) + integer, intent(in) :: n + real, device :: da(*) ! ok + real, managed :: ma(n) ! ok + !WARNING: Pointer 'dp' may not be associated in a device subprogram + real, device, pointer :: dp + end subroutine +end module Index: flang/test/Semantics/cuf04.cuf =================================================================== --- /dev/null +++ flang/test/Semantics/cuf04.cuf @@ -0,0 +1,24 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +! CUDA Fortran section 2.5.6 restrictions +module m + contains + attributes(device) subroutine devsubr(n) + integer, intent(in) :: n + !WARNING: 'x1' should not have the SAVE attribute or initialization in a device subprogram + real, save :: x1 + !WARNING: 'x2' should not have the SAVE attribute or initialization in a device subprogram + real :: x2 = 1. + !ERROR: Device subprogram 'devsubr' cannot call itself + if (n > 0) call devsubr(n-1) + end subroutine + attributes(global) subroutine globsubr + end subroutine + subroutine boring + end subroutine + subroutine test + !ERROR: 'globsubr' is a kernel subroutine and must be called with kernel launch parameters in chevrons + call globsubr + !ERROR: Kernel launch parameters in chevrons may not be used unless calling a kernel subroutine + call boring<<<1,2>>> + end subroutine +end module Index: flang/test/Semantics/cuf05.cuf =================================================================== --- /dev/null +++ flang/test/Semantics/cuf05.cuf @@ -0,0 +1,12 @@ +! RUN: %flang_fc1 -fdebug-dump-symbols %s 2>&1 | FileCheck %s +module m + implicit none + contains + attributes(device) subroutine devsubr() + !CHECK: blockdim, PUBLIC: Use from __builtin_blockdim in __fortran_builtins + !CHECK: blockidx, PUBLIC: Use from __builtin_blockidx in __fortran_builtins + !CHECK: griddim, PUBLIC: Use from __builtin_griddim in __fortran_builtins + !CHECK: threadidx, PUBLIC: Use from __builtin_threadidx in __fortran_builtins + !CHECK: warpsize, PARAMETER, PUBLIC: Use from __builtin_warpsize in __fortran_builtins + end subroutine +end module Index: flang/test/Semantics/cuf06.cuf =================================================================== --- /dev/null +++ flang/test/Semantics/cuf06.cuf @@ -0,0 +1,15 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +module m + use, intrinsic :: __fortran_builtins, only: __builtin_dim3 + contains + attributes(global) subroutine kernel + end subroutine + subroutine test + call kernel<<< 1, 32 >>> ! ok + call kernel<<< __builtin_dim3(1,1), __builtin_dim3(32,1,1) >>> ! ok + !ERROR: Kernel launch grid parameter must be either integer or TYPE(dim3) + call kernel<<< 1.d0, 32 >>> + !ERROR: Kernel launch block parameter must be either integer or TYPE(dim3) + call kernel<<< 1, "abc" >>> + end +end module Index: flang/test/Semantics/cuf07.cuf =================================================================== --- /dev/null +++ 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 Index: flang/test/Semantics/cuf08.cuf =================================================================== --- /dev/null +++ flang/test/Semantics/cuf08.cuf @@ -0,0 +1,22 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +module m + contains + !ERROR: A subroutine may not have LAUNCH_BOUNDS() or CLUSTER_DIMS() unless it has ATTRIBUTES(GLOBAL) or ATTRIBUTES(GRID_GLOBAL) + launch_bounds(1,2) subroutine bad1; end + !ERROR: A subroutine may not have LAUNCH_BOUNDS() or CLUSTER_DIMS() unless it has ATTRIBUTES(GLOBAL) or ATTRIBUTES(GRID_GLOBAL) + cluster_dims(1,2,3) subroutine bad2; end + attributes(global) launch_bounds(1,2) subroutine good1; end + !ERROR: LAUNCH_BOUNDS() may only appear once + attributes(global) launch_bounds(1,2) launch_bounds(3,4) subroutine bad3; end + !ERROR: Operands of LAUNCH_BOUNDS() must be two integer constants + attributes(global) launch_bounds(1) subroutine bad4; end + !ERROR: Operands of LAUNCH_BOUNDS() must be two integer constants + attributes(global) launch_bounds(1,2,3) subroutine bad5; end + attributes(global) cluster_dims(1,2,3) subroutine good2; end + !ERROR: CLUSTER_DIMS() may only appear once + attributes(global) cluster_dims(1,2,3) cluster_dims(4,5,6) subroutine bad6; end + !ERROR: Operands of CLUSTER_DIMS() must be three integer constants + attributes(global) cluster_dims(1) subroutine bad7; end + !ERROR: Operands of CLUSTER_DIMS() must be three integer constants + attributes(global) cluster_dims(1,2,3,4) subroutine bad8; end +end module Index: flang/test/Semantics/cuf09.cuf =================================================================== --- /dev/null +++ flang/test/Semantics/cuf09.cuf @@ -0,0 +1,76 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +module m + contains + attributes(device) subroutine devsub + !ERROR: Statement may not appear in device code + !$cuf kernel do <<< 1, 2 >>> + do k=1,10 + end do + end +end + +program main + !ERROR: !$CUF KERNEL DO (1) must be followed by a DO construct with tightly nested outer levels of counted DO loops + !$cuf kernel do <<< 1, 2 >>> + do while (.false.) + end do + !ERROR: !$CUF KERNEL DO (1) must be followed by a DO construct with tightly nested outer levels of counted DO loops + !$cuf kernel do <<< 1, 2 >>> + do + exit + end do + !ERROR: !$CUF KERNEL DO (1) must be followed by a DO construct with tightly nested outer levels of counted DO loops + !$cuf kernel do <<< 1, 2 >>> + do concurrent (j=1:10) + end do + !$cuf kernel do <<< 1, 2 >>> + do 1 j=1,10 +1 continue ! ok + !$cuf kernel do <<< 1, 2 >>> + do j=1,10 + end do ! ok + !$cuf kernel do <<< 1, 2 >>> + do j=1,10 + !ERROR: Statement may not appear in device code + !$cuf kernel do <<< 1, 2 >>> + do k=1,10 + end do + end do + !ERROR: !$CUF KERNEL DO (-1): loop nesting depth must be positive + !$cuf kernel do (-1) <<< 1, 2 >>> + do j=1,10 + end do + !ERROR: !$CUF KERNEL DO (1) must be followed by a DO construct with tightly nested outer levels of counted DO loops + !$cuf kernel do <<< 1, 2 >>> + continue + !ERROR: !$CUF KERNEL DO (2) must be followed by a DO construct with tightly nested outer levels of counted DO loops + !$cuf kernel do (2) <<< 1, 2 >>> + do j=1,10 + end do + !ERROR: !$CUF KERNEL DO (2) must be followed by a DO construct with tightly nested outer levels of counted DO loops + !$cuf kernel do (2) <<< 1, 2 >>> + do j=1,10 + continue + end do + !ERROR: !$CUF KERNEL DO (2) must be followed by a DO construct with tightly nested outer levels of counted DO loops + !$cuf kernel do (2) <<< 1, 2 >>> + do j=1,10 + do k=1,10 + end do + continue + end do + !$cuf kernel do <<< 1, 2 >>> + do j = 1, 10 + !ERROR: 'foo' may not be called in device code + call foo + !ERROR: 'bar' may not be called in device code + x = bar() + !ERROR: 'ifunc' may not be called in device code + if (ifunc() /= 0) continue + !ERROR: 'ifunc' may not be called in device code + if (ifunc() /= 0) then + !ERROR: 'ifunc' may not be called in device code + else if (ifunc() /= 1) then + end if + end do +end Index: flang/test/Semantics/cuf10.cuf =================================================================== --- /dev/null +++ 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 Index: flang/test/Semantics/definable05.cuf =================================================================== --- /dev/null +++ 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 Index: flang/test/Semantics/resolve65.f90 =================================================================== --- flang/test/Semantics/resolve65.f90 +++ flang/test/Semantics/resolve65.f90 @@ -5,6 +5,9 @@ implicit none type :: t contains + !ERROR: Generic 'assignment(=)' may not have specific procedures 't%assign_t4' and 't%assign_t5' as their interfaces are not distinguishable + !ERROR: Generic 'assignment(=)' may not have specific procedures 't%assign_t4' and 't%assign_t6' as their interfaces are not distinguishable + !ERROR: Generic 'assignment(=)' may not have specific procedures 't%assign_t5' and 't%assign_t6' as their interfaces are not distinguishable !ERROR: Defined assignment procedure 'binding' must be a subroutine generic :: assignment(=) => binding procedure :: binding => assign_t1 @@ -12,10 +15,14 @@ procedure :: assign_t2 procedure :: assign_t3 !ERROR: Defined assignment subroutine 'assign_t2' must have two dummy arguments - !ERROR: In defined assignment subroutine 'assign_t3', second dummy argument 'y' must have INTENT(IN) or VALUE attribute - !ERROR: In defined assignment subroutine 'assign_t4', first dummy argument 'x' must have INTENT(OUT) or INTENT(INOUT) - generic :: assignment(=) => assign_t, assign_t2, assign_t3, assign_t4 + !WARNING: In defined assignment subroutine 'assign_t3', second dummy argument 'y' should have INTENT(IN) or VALUE attribute + !WARNING: In defined assignment subroutine 'assign_t4', first dummy argument 'x' should have INTENT(OUT) or INTENT(INOUT) + !ERROR: In defined assignment subroutine 'assign_t5', first dummy argument 'x' may not have INTENT(IN) + !ERROR: In defined assignment subroutine 'assign_t6', second dummy argument 'y' may not have INTENT(OUT) + generic :: assignment(=) => assign_t, assign_t2, assign_t3, assign_t4, assign_t5, assign_t6 procedure :: assign_t4 + procedure :: assign_t5 + procedure :: assign_t6 end type type :: t2 contains @@ -41,7 +48,15 @@ end subroutine assign_t4(x, y) class(t) :: x - integer, intent(in) :: y + integer, intent(in) :: y + end + subroutine assign_t5(x, y) + class(t), intent(in) :: x + integer, intent(in) :: y + end + subroutine assign_t6(x, y) + class(t), intent(out) :: x + integer, intent(out) :: y end end Index: flang/test/Semantics/resolve67.f90 =================================================================== --- flang/test/Semantics/resolve67.f90 +++ flang/test/Semantics/resolve67.f90 @@ -41,15 +41,16 @@ end end interface interface operator(<) - !ERROR: In OPERATOR(<) function 'lt1', dummy argument 'x' must have INTENT(IN) or VALUE attribute + !WARNING: In OPERATOR(<) function 'lt1', dummy argument 'x' should have INTENT(IN) or VALUE attribute !ERROR: In OPERATOR(<) function 'lt1', dummy argument 'y' may not be OPTIONAL logical function lt1(x, y) logical :: x real, value, optional :: y end + !ERROR: In OPERATOR(<) function 'lt2', dummy argument 'x' may not be INTENT(OUT) !ERROR: In OPERATOR(<) function 'lt2', dummy argument 'y' must be a data object logical function lt2(x, y) - logical, intent(in) :: x + logical, intent(out) :: x intent(in) :: y interface subroutine y()