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/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/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-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/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/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/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/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()