Index: flang/include/flang/Evaluate/call.h =================================================================== --- flang/include/flang/Evaluate/call.h +++ flang/include/flang/Evaluate/call.h @@ -209,8 +209,6 @@ u; }; -using Chevrons = std::vector>; - class ProcedureRef { public: CLASS_BOILERPLATE(ProcedureRef) @@ -225,10 +223,6 @@ 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; @@ -256,7 +250,6 @@ 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,7 +220,6 @@ common::Intent intent{common::Intent::Default}; Attrs attrs; common::IgnoreTKRSet ignoreTKR; - std::optional cudaDataAttr; }; // 15.3.2.3 @@ -318,7 +317,6 @@ Attrs attrs; std::variant> u; - std::optional cudaDataAttr; }; // 15.3.1 @@ -370,8 +368,6 @@ 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,7 +22,6 @@ #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" @@ -473,10 +472,8 @@ std::optional ComparisonType( const DynamicType &, const DynamicType &); -bool IsInteroperableIntrinsicType(const DynamicType &, - const common::LanguageFeatureControl * = nullptr, - bool checkCharLength = true); -bool IsCUDAIntrinsicType(const DynamicType &); +bool IsInteroperableIntrinsicType( + const DynamicType &, bool checkCharLength = true); // 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,7 +381,6 @@ } 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,8 +446,7 @@ return const_cast(this)->AsDerived(); } -bool IsInteroperableIntrinsicType( - const DeclTypeSpec &, const common::LanguageFeatureControl &); +bool IsInteroperableIntrinsicType(const DeclTypeSpec &); } // 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,8 +265,7 @@ bool DummyDataObject::operator==(const DummyDataObject &that) const { return type == that.type && attrs == that.attrs && intent == that.intent && - coshape == that.coshape && cudaDataAttr == that.cudaDataAttr; - ; + coshape == that.coshape; } static bool AreCompatibleDummyDataObjectShapes(const Shape &x, const Shape &y) { @@ -326,13 +325,6 @@ *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; } @@ -368,14 +360,6 @@ }); 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; } } @@ -396,8 +380,6 @@ 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 { @@ -418,9 +400,6 @@ sep = ','; } } - if (cudaDataAttr) { - o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr); - } return o; } @@ -540,7 +519,6 @@ return std::nullopt; } } - result.cudaSubprogramAttrs = subp.cudaSubprogramAttrs(); return result; }, [&](const semantics::ProcEntityDetails &proc) @@ -573,10 +551,6 @@ 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}; @@ -867,14 +841,13 @@ FunctionResult::~FunctionResult() {} bool FunctionResult::operator==(const FunctionResult &that) const { - return attrs == that.attrs && cudaDataAttr == that.cudaDataAttr && - u == that.u; + return attrs == that.attrs && u == that.u; } static std::optional CharacterizeFunctionResult( const semantics::Symbol &symbol, FoldingContext &context, semantics::UnorderedSymbolSet seenProcs) { - if (const auto *object{symbol.detailsIf()}) { + if (symbol.has()) { if (auto type{TypeAndShape::Characterize(symbol, context)}) { FunctionResult result{std::move(*type)}; CopyAttrs(symbol, result, @@ -883,7 +856,6 @@ {semantics::Attr::CONTIGUOUS, FunctionResult::Attr::Contiguous}, {semantics::Attr::POINTER, FunctionResult::Attr::Pointer}, }); - result.cudaDataAttr = object->cudaDataAttr(); return result; } } else if (auto maybeProc{ @@ -912,8 +884,6 @@ 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) @@ -980,10 +950,6 @@ 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()) { @@ -1064,9 +1030,6 @@ }, }, u); - if (cudaDataAttr) { - o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr); - } return o; } @@ -1079,8 +1042,7 @@ bool Procedure::operator==(const Procedure &that) const { return attrs == that.attrs && functionResult == that.functionResult && - dummyArguments == that.dummyArguments && - cudaSubprogramAttrs == that.cudaSubprogramAttrs; + dummyArguments == that.dummyArguments; } bool Procedure::IsCompatibleWith(const Procedure &actual, std::string *whyNot, @@ -1113,10 +1075,6 @@ } } 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"; @@ -1239,10 +1197,6 @@ // 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; @@ -1270,11 +1224,7 @@ dummy.Dump(o << sep); sep = ','; } - o << (sep == '(' ? "()" : ")"); - if (cudaSubprogramAttrs) { - o << " cudaSubprogramAttrs: " << common::EnumToString(*cudaSubprogramAttrs); - } - return o; + return o << (sep == '(' ? "()" : ")"); } // Utility class to determine if Procedures, etc. are distinguishable @@ -1376,9 +1326,6 @@ if (pos2 >= 0 && pos2 <= name2) { return true; // distinguishable based on C1514 rule 4 } - if (proc1.cudaSubprogramAttrs != proc2.cudaSubprogramAttrs) { - return true; - } return false; } @@ -1506,9 +1453,6 @@ } 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)) && @@ -1547,9 +1491,6 @@ 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,18 +135,6 @@ } } 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,15 +734,14 @@ } } -bool IsInteroperableIntrinsicType(const DynamicType &type, - const common::LanguageFeatureControl *features, bool checkCharLength) { +bool IsInteroperableIntrinsicType( + const DynamicType &type, bool checkCharLength) { switch (type.category()) { case TypeCategory::Integer: return true; case TypeCategory::Real: case TypeCategory::Complex: - return (features && features->IsEnabled(common::LanguageFeature::CUDA)) || - type.kind() >= 4; // no short or half floats + return type.kind() >= 4; // no short or half floats case TypeCategory::Logical: return type.kind() == 1; // C_BOOL case TypeCategory::Character: @@ -756,21 +755,4 @@ } } -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()) { + if (asFortran_ && x.typedCall.get() && !x.chevrons /*CUDA todo*/) { 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,8 +190,7 @@ characteristics::TypeAndShape &actualType, bool isElemental, evaluate::FoldingContext &context, const Scope *scope, const evaluate::SpecificIntrinsic *intrinsic, - bool allowActualArgumentConversions, - const characteristics::Procedure &procedure) { + bool allowActualArgumentConversions) { // Basic type & rank checking parser::ContextualMessages &messages{context.messages()}; @@ -619,46 +618,6 @@ } } } - - // 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, @@ -849,7 +808,7 @@ object.type.Rank() == 0 && proc.IsElemental()}; CheckExplicitDataArg(object, dummyName, *expr, *type, isElemental, context, scope, intrinsic, - allowActualArgumentConversions, proc); + allowActualArgumentConversions); } 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,19 +114,6 @@ } 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 &); @@ -229,8 +216,9 @@ void CheckHelper::Check(const Symbol &symbol) { if (symbol.name().size() > common::maxNameLen && - &symbol == &symbol.GetUltimate()) { - WarnIfNotInModuleFile(symbol.name(), + &symbol == &symbol.GetUltimate() && + !FindModuleFileContaining(symbol.owner())) { + messages_.Say(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); @@ -627,7 +615,6 @@ 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)) { @@ -636,7 +623,7 @@ " coshape"_err_en_US, symbol.name()); } - } else if (isComponent) { // C746 + } else if (symbol.owner().IsDerivedType()) { // C746 std::string deferredMsg{ isDeferredCoshape ? "" : " and have a deferred coshape"}; messages_.Say("Component '%s' is a coarray and must have the ALLOCATABLE" @@ -740,7 +727,7 @@ if (IsPassedViaDescriptor(symbol)) { if (IsAllocatableOrPointer(symbol)) { if (inExplicitInterface) { - WarnIfNotInModuleFile( + messages_.Say( "!DIR$ IGNORE_TKR should not apply to an allocatable or pointer"_warn_en_US); } else { messages_.Say( @@ -748,10 +735,10 @@ } } else if (ignoreTKR.test(common::IgnoreTKR::Rank)) { if (ignoreTKR.count() == 1 && evaluate::IsAssumedRank(symbol)) { - WarnIfNotInModuleFile( + messages_.Say( "!DIR$ IGNORE_TKR(R) is not meaningful for an assumed-rank array"_warn_en_US); } else if (inExplicitInterface) { - WarnIfNotInModuleFile( + messages_.Say( "!DIR$ IGNORE_TKR(R) should not apply to a dummy argument passed via descriptor"_warn_en_US); } else { messages_.Say( @@ -819,8 +806,9 @@ 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)) { - WarnIfNotInModuleFile( + } else if (IsInBlankCommon(symbol) && + !FindModuleFileContaining(symbol.owner())) { + messages_.Say( "A variable in blank COMMON should not be initialized"_port_en_US); } } @@ -860,156 +848,6 @@ "'%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) { @@ -1081,9 +919,6 @@ 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) { @@ -1113,12 +948,12 @@ } } else if (canBeAssumedShape && !canBeDeferred) { msg = "Assumed-shape array '%s' must be a dummy argument"_err_en_US; - } else if (canBeAssumedSize && !canBeImplied && !isCUDAShared) { // C833 + } else if (canBeAssumedSize && !canBeImplied) { // 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) && !isCUDAShared) { // C835, C836 + if (!IsNamedConstant(symbol)) { // C835, C836 msg = "Implied-shape array '%s' must be a named constant or a " "dummy argument"_err_en_US; } @@ -1352,50 +1187,6 @@ } 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) { @@ -1424,7 +1215,7 @@ if (chars->HasExplicitInterface()) { std::string whyNot; if (!chars->IsCompatibleWith(*globalChars, &whyNot)) { - msg = WarnIfNotInModuleFile( + msg = messages_.Say( "The global subprogram '%s' is not compatible with its local procedure declaration (%s)"_warn_en_US, global->name(), whyNot); } @@ -1450,7 +1241,7 @@ if (auto previousChars{Characterize(previous)}) { std::string whyNot; if (!chars->IsCompatibleWith(*previousChars, &whyNot)) { - if (auto *msg{WarnIfNotInModuleFile( + if (auto *msg{messages_.Say( "The external interface '%s' is not compatible with an earlier definition (%s)"_warn_en_US, symbol.name(), whyNot)}) { evaluate::AttachDeclaration(msg, previous); @@ -1837,14 +1628,12 @@ return true; // OK } bool isFatal{msg->IsFatal()}; - if (isFatal || !FindModuleFileContaining(specific.owner())) { - SayWithDeclaration( - specific, std::move(*msg), MakeOpName(opName), specific.name()); - } + SayWithDeclaration( + specific, std::move(*msg), MakeOpName(opName), specific.name()); if (isFatal) { context_.SetError(specific); } - return !isFatal; + return false; } // If the number of arguments is wrong for this intrinsic operator, return @@ -1905,24 +1694,15 @@ 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' should have INTENT(IN) or VALUE attribute"_warn_en_US; + msg = "In %s function '%s', dummy argument '%s' must have INTENT(IN)" + " or VALUE attribute"_err_en_US; } if (msg) { - 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; - } + SayWithDeclaration(symbol, std::move(*msg), + parser::ToUpperCaseLetters(opName.ToString()), symbol.name(), arg.name); + return false; } return true; } @@ -1968,23 +1748,17 @@ " 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::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 && + if (dataObject->intent != common::Intent::Out && dataObject->intent != common::Intent::InOut) { msg = "In defined assignment subroutine '%s', first dummy argument '%s'" - " should have INTENT(OUT) or INTENT(INOUT)"_warn_en_US; + " must have INTENT(OUT) or INTENT(INOUT)"_err_en_US; } } else if (pos == 1) { - 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 && + if (dataObject->intent != common::Intent::In && !dataObject->attrs.test(DummyDataObject::Attr::Value)) { msg = "In defined assignment subroutine '%s', second dummy" - " argument '%s' should have INTENT(IN) or VALUE attribute"_warn_en_US; + " argument '%s' must have INTENT(IN) or VALUE attribute"_err_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; @@ -2000,14 +1774,9 @@ " must be a data object"_err_en_US; } if (msg) { - 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; - } + SayWithDeclaration(symbol, std::move(*msg), symbol.name(), arg.name); + context_.SetError(symbol); + return false; } return true; } @@ -2040,10 +1809,10 @@ if (!derivedDetails->finals().empty() && !derivedDetails->GetFinalForRank(rank)) { if (auto *msg{derivedSym == initialDerivedSym - ? WarnIfNotInModuleFile(symbol.name(), + ? messages_.Say(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) - : WarnIfNotInModuleFile(symbol.name(), + : messages_.Say(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)}) { @@ -2662,17 +2431,15 @@ type->category() == DeclTypeSpec::Character && type->characterTypeSpec().length().isDeferred()) { // ok; F'2018 18.3.6 p2(6) - } else if (derived || - IsInteroperableIntrinsicType(*type, context_.languageFeatures())) { + } else if (derived || IsInteroperableIntrinsicType(*type)) { // F'2018 18.3.6 p2(4,5) - } 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 (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 (symbol.attrs().test(Attr::VALUE)) { messages_.Say(symbol.name(), "A BIND(C) VALUE dummy argument must have an interoperable type"_err_en_US); @@ -2684,13 +2451,12 @@ } } if (IsOptional(symbol) && !symbol.attrs().test(Attr::VALUE)) { - WarnIfNotInModuleFile(symbol.name(), + messages_.Say(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->isDummy() && - (!proc->procInterface() || - !proc->procInterface()->attrs().test(Attr::BIND_C))) { + if (!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); @@ -2742,21 +2508,10 @@ } 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); - } + } 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); } } if (auto extents{ @@ -2768,8 +2523,9 @@ } } } - if (derived->componentNames().empty()) { // C1805 - WarnIfNotInModuleFile(symbol.name(), + if (derived->componentNames().empty() && + !FindModuleFileContaining(symbol.owner())) { // C1805 + messages_.Say(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,33 +134,6 @@ 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,9 +2616,6 @@ 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); } @@ -2685,55 +2682,6 @@ 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}; @@ -2745,17 +2693,17 @@ if (analyzer.fatalErrors()) { return std::nullopt; } - bool mightBeStructureConstructor{structureConstructor != nullptr}; - if (std::optional callee{GetCalleeAndArguments( - std::get(call.t), analyzer.GetActuals(), - false /* not subroutine */, mightBeStructureConstructor)}) { + if (std::optional callee{ + GetCalleeAndArguments(std::get(call.t), + analyzer.GetActuals(), false /* not subroutine */, + true /* might be structure constructor */)}) { 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 (mightBeStructureConstructor) { + if (structureConstructor) { // Structure constructor misparsed as function reference? const auto &designator{std::get(call.t)}; if (const auto *name{std::get_if(&designator.u)}) { @@ -2800,40 +2748,17 @@ for (const auto &arg : actualArgList) { analyzer.Analyze(arg, true /* is subroutine call */); } - auto chevrons{AnalyzeChevrons(callStmt)}; - if (!analyzer.fatalErrors() && chevrons) { + if (!analyzer.fatalErrors()) { 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; } } @@ -3735,13 +3660,14 @@ 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{std::move(procRef)}; + return Expr{ + ProcedureRef{std::move(proc), std::move(arguments)}}; } else { // Not a procedure pointer, so type and shape are known. return TypedWrapper( - DEREF(result.GetTypeAndShape()).type(), std::move(procRef)); + DEREF(result.GetTypeAndShape()).type(), + ProcedureRef{std::move(proc), std::move(arguments)}); } } else { Say("Function result characteristics are not known"_err_en_US); Index: flang/lib/Semantics/resolve-names.cpp =================================================================== --- flang/lib/Semantics/resolve-names.cpp +++ flang/lib/Semantics/resolve-names.cpp @@ -3521,9 +3521,9 @@ ok = false; } } - if (!ok || bounds.size() != 2) { + if (!ok || bounds.size() < 2 || bounds.size() > 3) { Say(currStmtSource().value(), - "Operands of LAUNCH_BOUNDS() must be two integer constants"_err_en_US); + "Operands of LAUNCH_BOUNDS() must be 2 or 3 integer constants"_err_en_US); } else if (auto *subp{currScope().symbol() ? currScope().symbol()->detailsIf() : nullptr}) { Index: flang/lib/Semantics/type.cpp =================================================================== --- flang/lib/Semantics/type.cpp +++ flang/lib/Semantics/type.cpp @@ -797,10 +797,9 @@ return o << x.AsFortran(); } -bool IsInteroperableIntrinsicType( - const DeclTypeSpec &type, const common::LanguageFeatureControl &features) { +bool IsInteroperableIntrinsicType(const DeclTypeSpec &type) { auto dyType{evaluate::DynamicType::From(type)}; - return dyType && IsInteroperableIntrinsicType(*dyType, &features); + return dyType && IsInteroperableIntrinsicType(*dyType); } } // namespace Fortran::semantics Index: flang/module/__cuda_builtins.f90 =================================================================== --- flang/module/__cuda_builtins.f90 +++ /dev/null @@ -1,19 +0,0 @@ -!===-- 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,9 +106,6 @@ !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 @@ -162,36 +159,27 @@ !CHECK: | | | | | | | | | LiteralConstant -> IntLiteralConstant = '10' !CHECK: | | | | | | Block !CHECK: | | | | | | EndDoStmt -> -!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4>>>()' +!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub()' !CHECK: | | | | | Call !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' !CHECK: | | | | | Chevrons -!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: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '1' +!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '2' +!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub()' !CHECK: | | | | | Call !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' !CHECK: | | | | | Chevrons -!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: | | | | | | 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: | | | | | Call !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' !CHECK: | | | | | Chevrons -!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: | | | | | | 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: | | | | 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,7 +23,6 @@ !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 @@ -34,10 +33,9 @@ !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_4,2_4>>>() -!CHECK: CALL globalsub<<<1_4,2_4,3_4>>>() -!CHECK: CALL globalsub<<<1_4,2_4,3_4,4_4>>>() +!CHECK: CALL globalsub<<<1,2>>> +!CHECK: CALL globalsub<<<1,2,3>>> +!CHECK: CALL globalsub<<<1,2,3,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 - !WARNING: A CHARACTER component of a BIND(C) type should have length 1 + !ERROR: Each component of an interoperable derived type must have an interoperable type character(len=2) x end type type, bind(c) :: t11 @@ -73,7 +73,7 @@ character(kind=2) x end type type, bind(c) :: t12 - !PORTABILITY: A LOGICAL component of a BIND(C) type should have the interoperable KIND=C_BOOL + !ERROR: Each component of an interoperable derived type must have an interoperable type logical(kind=8) x end type type, bind(c) :: t13 Index: flang/test/Semantics/cuf02.cuf =================================================================== --- flang/test/Semantics/cuf02.cuf +++ /dev/null @@ -1,49 +0,0 @@ -! 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 =================================================================== --- flang/test/Semantics/cuf03.cuf +++ /dev/null @@ -1,59 +0,0 @@ -! 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 =================================================================== --- flang/test/Semantics/cuf07.cuf +++ /dev/null @@ -1,26 +0,0 @@ -! 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 =================================================================== --- flang/test/Semantics/cuf08.cuf +++ /dev/null @@ -1,22 +0,0 @@ -! 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 =================================================================== --- flang/test/Semantics/cuf10.cuf +++ /dev/null @@ -1,17 +0,0 @@ -! 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 =================================================================== --- flang/test/Semantics/definable05.cuf +++ /dev/null @@ -1,31 +0,0 @@ -! 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,9 +5,6 @@ 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 @@ -15,14 +12,10 @@ procedure :: assign_t2 procedure :: assign_t3 !ERROR: Defined assignment subroutine 'assign_t2' must have two dummy arguments - !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 + !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 procedure :: assign_t4 - procedure :: assign_t5 - procedure :: assign_t6 end type type :: t2 contains @@ -48,15 +41,7 @@ end subroutine assign_t4(x, y) class(t) :: x - 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 + integer, intent(in) :: y end end Index: flang/test/Semantics/resolve67.f90 =================================================================== --- flang/test/Semantics/resolve67.f90 +++ flang/test/Semantics/resolve67.f90 @@ -41,16 +41,15 @@ end end interface interface operator(<) - !WARNING: In OPERATOR(<) function 'lt1', dummy argument 'x' should have INTENT(IN) or VALUE attribute + !ERROR: In OPERATOR(<) function 'lt1', dummy argument 'x' must 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(out) :: x + logical, intent(in) :: x intent(in) :: y interface subroutine y()