[flang] CUDA Fortran - part 2/5: symbols & scopes

[klausler]

Summary

Add representations of CUDA Fortran data and subprogram attributes
to the symbol table and scopes of semantics. Set them in name
resolution, and emit them to module files.

Depends on https://reviews.llvm.org/D150159.

[klausler]

Remove redundantly added tests.

[clementval]

LGTM

[wangzpgi]

flang/lib/Semantics/resolve-names.cpp

3526

Can launch_bounds take third argument of max cluster rank for Hopper?

3551

Can cluster_dim take variable argument such as "cluster_dim(x,y,z)", where x, y, and z are integer variables?

[wangzpgi]

flang/include/flang/Semantics/symbol.h

152 ↗

(On Diff #520536)

misspelling: LANCH_BOUNDS -> LAUNCH_BOUNDS

[klausler]

flang/lib/Semantics/resolve-names.cpp

3526

Not according to https://docs.nvidia.com/hpc-sdk/compilers/cuda-fortran-prog-guide/index.html#cfref-new-subs-func-attr which has "The form used in CUDA Fortran is:

attributes(global) launch_bounds(maxTPB, minBPM) subroutine sub(...)

where maxTPB is the maxThreadsPerBlock, the maximum number of threads per block with which the application will ever launch, and minBPM is the desired minimum number of resident blocks per multiprocessor. Both values must be numeric constants." in 3.1.2.

3551

No, not according to https://docs.nvidia.com/hpc-sdk/compilers/cuda-fortran-prog-guide/index.html#cfref-new-subs-func-attr , which has "Beginning with the 23.3 release, support for thread block clusters is enabled for Hopper (cc90) and later targets. To specify the dimensions of the cluster, use the cluster_dims syntax and specify each x, y, and z dimension. Values must be numeric constants." in 3.1.3.

[klausler]

Fix misspelling in comment.

[klausler]

Generate diff against part 1 rather than main (oops).

[wangzpgi]

flang/lib/Semantics/resolve-names.cpp

3551

It mentioned at the end of 3.1.2: "Also, as part of the Hopper support, the launch_bounds syntax has been extended to accept a third argument, an upper bound on the cluster size."

[klausler]

flang/lib/Semantics/resolve-names.cpp

3551

That's for LAUNCH_BOUNDS(), not CLUSTER_DIMS().

[klausler]

Allow third parameter to LAUNCH_BOUNDS().

[klausler]

Upload correct diff (sorry).

[wangzpgi]

flang/lib/Semantics/resolve-names.cpp

3551

Yes, allowing third argument was for launch_bounds, not cluster_dims, I commented at the wrong place. Not quite sure how this Phabricator review works.

For cluster_dims as well as launch_bounds, Brent has requested to take maybe const expressions. He would like to use module parameters, for global routines that are in that module. Neither Michael nor I implemented that in nvfortran.

[wangzpgi]

Is there checks for CUDA Version that some features are allowed for certain CUDA Version?

[klausler]

In D150161#4331054, @wangzpgi wrote:

Is there checks for CUDA Version that some features are allowed for certain CUDA Version?

Not yet.

flang/lib/Semantics/resolve-names.cpp

3551

That's what I implemented -- any expression that can be folded to a constant value at compilation time will work. They don't have to integer literals.

[wangzpgi]

flang/lib/Semantics/resolve-names.cpp

3551

Great. Thanks.

[klausler]

Rebase to current head.

[klausler]

Rebasing to current head in preparation for merging.

flang/include/flang/Evaluate/call.h

   std::variant<SpecificIntrinsic, SymbolRef,
       common::CopyableIndirection<Component>>
       u;
 };
 
-using Chevrons = std::vector<Expr<SomeType>>;
-
 class ProcedureRef {
 public:
   CLASS_BOILERPLATE(ProcedureRef)
   ProcedureRef(ProcedureDesignator &&p, ActualArguments &&a,
       bool hasAlternateReturns = false)

 
   ProcedureDesignator &proc() { return proc_; }
   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<Expr<SubscriptInteger>> LEN() const;
   int Rank() const;
   bool IsElemental() const { return proc_.IsElemental(); }
   bool hasAlternateReturns() const { return hasAlternateReturns_; }

   llvm::raw_ostream &AsFortran(llvm::raw_ostream &) const;
 
 protected:
   ProcedureDesignator proc_;
   ActualArguments arguments_;
-  Chevrons chevrons_;
   bool hasAlternateReturns_;
 };
 
 template <typename A> class FunctionRef : public ProcedureRef {
 public:

flang/include/flang/Evaluate/characteristics.h

   TypeAndShape type;
   std::vector<Expr<SubscriptInteger>> coshape;
   common::Intent intent{common::Intent::Default};
   Attrs attrs;
   common::IgnoreTKRSet ignoreTKR;
-  std::optional<common::CUDADataAttr> cudaDataAttr;
 };
 
 // 15.3.2.3
 struct DummyProcedure {
   ENUM_CLASS(Attr, Pointer, Optional)

 
   llvm::raw_ostream &Dump(llvm::raw_ostream &) const;
 
   Attrs attrs;
   std::variant<TypeAndShape, CopyableIndirection<Procedure>> u;
-  std::optional<common::CUDADataAttr> cudaDataAttr;
 };
 
 // 15.3.1
 struct Procedure {
   ENUM_CLASS(

   llvm::raw_ostream &Dump(llvm::raw_ostream &) const;
 
   std::optional<FunctionResult> functionResult;
   DummyArguments dummyArguments;
   Attrs attrs;
-  std::optional<common::CUDASubprogramAttrs> cudaSubprogramAttrs;
 };
-
 } // namespace Fortran::evaluate::characteristics
 #endif // FORTRAN_EVALUATE_CHARACTERISTICS_H_

flang/include/flang/Evaluate/type.h

 #include "complex.h"
 #include "formatting.h"
 #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"
 #include "flang/Common/template.h"
 #include <cinttypes>

 // type to which they must be converted in order to be compared with
 // intrinsic OPERATOR(==) or .EQV.
 std::optional<DynamicType> 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.
 bool AreSameDerivedType(
     const semantics::DerivedTypeSpec &x, const semantics::DerivedTypeSpec &y);

flang/include/flang/Semantics/expression.h

   template <typename T> T Fold(T &&expr) {
     return evaluate::Fold(foldingContext_, std::move(expr));
   }
   bool CheckIsValidForwardReference(const semantics::DerivedTypeSpec &);
   MaybeExpr AnalyzeComplex(MaybeExpr &&re, MaybeExpr &&im, const char *what);
-  std::optional<Chevrons> AnalyzeChevrons(const parser::CallStmt &);
 
   MaybeExpr IterativelyAnalyzeSubexpressions(const parser::Expr &);
 
   semantics::SemanticsContext &context_;
   FoldingContext &foldingContext_{context_.foldingContext()};

flang/include/flang/Semantics/type.h

 }
 inline const DerivedTypeSpec *DeclTypeSpec::AsDerived() const {
   return const_cast<DeclTypeSpec *>(this)->AsDerived();
 }
 
-bool IsInteroperableIntrinsicType(
-    const DeclTypeSpec &, const common::LanguageFeatureControl &);
+bool IsInteroperableIntrinsicType(const DeclTypeSpec &);
 
 } // namespace Fortran::semantics
 #endif // FORTRAN_SEMANTICS_TYPE_H_

flang/lib/Evaluate/characteristics.cpp

   return o;
 }
 
 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) {
   int n{GetRank(x)};
   if (n != GetRank(y)) {

   if (ignoreTKR != actual.ignoreTKR) {
     if (whyNot) {
       *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;
 }
 
 static common::Intent GetIntent(const semantics::Attrs &attrs) {
   if (attrs.test(semantics::Attr::INTENT_IN)) {

               {Attr::POINTER, DummyDataObject::Attr::Pointer},
               {Attr::TARGET, DummyDataObject::Attr::Target},
           });
       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;
     }
   }
   return std::nullopt;
 }

                      TypeAndShape::Attr::Coarray})
                  .any()) {
     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 {
     return true;
   }

     for (const auto &expr : coshape) {
       expr.AsFortran(o << sep);
       sep = ',';
     }
   }
-  if (cudaDataAttr) {
-    o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr);
-  }
   return o;
 }
 
 DummyProcedure::DummyProcedure(Procedure &&p)
     : procedure{new Procedure{std::move(p)}} {}

                     std::move(argCharacteristics.value()));
               } else {
                 return std::nullopt;
               }
             }
-            result.cudaSubprogramAttrs = subp.cudaSubprogramAttrs();
             return result;
           },
           [&](const semantics::ProcEntityDetails &proc)
               -> std::optional<Procedure> {
             if (symbol.attrs().test(semantics::Attr::INTRINSIC)) {

               result.attrs.set(Procedure::Attr::ImplicitInterface);
               const semantics::DeclTypeSpec *type{proc.type()};
               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};
                 } else {
                   return std::nullopt;

 FunctionResult::FunctionResult(TypeAndShape &&t) : u{std::move(t)} {}
 FunctionResult::FunctionResult(Procedure &&p) : u{std::move(p)} {}
 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<FunctionResult> CharacterizeFunctionResult(
     const semantics::Symbol &symbol, FoldingContext &context,
     semantics::UnorderedSymbolSet seenProcs) {
-  if (const auto *object{symbol.detailsIf<semantics::ObjectEntityDetails>()}) {
+  if (symbol.has<semantics::ObjectEntityDetails>()) {
     if (auto type{TypeAndShape::Characterize(symbol, context)}) {
       FunctionResult result{std::move(*type)};
       CopyAttrs<FunctionResult, FunctionResult::Attr>(symbol, result,
           {
               {semantics::Attr::ALLOCATABLE, FunctionResult::Attr::Allocatable},
               {semantics::Attr::CONTIGUOUS, FunctionResult::Attr::Contiguous},
               {semantics::Attr::POINTER, FunctionResult::Attr::Pointer},
           });
-      result.cudaDataAttr = object->cudaDataAttr();
       return result;
     }
   } else if (auto maybeProc{
                  CharacterizeProcedure(symbol, context, seenProcs)}) {
     FunctionResult result{std::move(*maybeProc)};

 }
 
 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)
     } else {
       const DynamicType &type{typeAndShape->type()};

   }
   if (attrs != actualAttrs) {
     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<TypeAndShape>(&u)}) {
     if (const auto *actualTypeShape{std::get_if<TypeAndShape>(&actual.u)}) {
       if (ifaceTypeShape->Rank() != actualTypeShape->Rank()) {
         if (whyNot) {
           *whyNot = "function results have distinct ranks";

                     [&](const CopyableIndirection<Procedure> &p) {
                       p.value().Dump(o << " procedure(") << ')';
                     },
                 },
       u);
-  if (cudaDataAttr) {
-    o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr);
-  }
   return o;
 }
 
 Procedure::Procedure(FunctionResult &&fr, DummyArguments &&args, Attrs a)
     : functionResult{std::move(fr)}, dummyArguments{std::move(args)}, attrs{a} {

     : dummyArguments{std::move(args)}, attrs{a} {}
 Procedure::~Procedure() {}
 
 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,
     const SpecificIntrinsic *specificIntrinsic) const {
   // 15.5.2.9(1): if dummy is not pure, actual need not be.

       *whyNot =
           "incompatible procedures: one is a function, the other a subroutine";
     }
   } 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";
     }
   } else {

 
 bool Procedure::CanBeCalledViaImplicitInterface() const {
   // TODO: Pass back information on why we return false
   if (attrs.test(Attr::Elemental) || attrs.test(Attr::BindC)) {
     return false; // 15.4.2.2(5,6)
-  } else if (cudaSubprogramAttrs &&
-      *cudaSubprogramAttrs != common::CUDASubprogramAttrs::Host &&
-      *cudaSubprogramAttrs != common::CUDASubprogramAttrs::Global) {
-    return false;
   } else if (IsFunction() &&
       !functionResult->CanBeReturnedViaImplicitInterface()) {
     return false;
   } else {
     for (const DummyArgument &arg : dummyArguments) {

   char sep{'('};
   for (const auto &dummy : dummyArguments) {
     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
 class DistinguishUtils {
 public:

   int pos2{FindFirstToDistinguishByPosition(args2, args1)};
   int name2{FindLastToDistinguishByName(args2, args1)};
   if (pos2 >= 0 && pos2 <= name2) {
     return true; // distinguishable based on C1514 rule 4
   }
-  if (proc1.cudaSubprogramAttrs != proc2.cudaSubprogramAttrs) {
-    return true;
-  }
   return false;
 }
 
 // C1514 rule 3: Procedures are distinguishable if both have a passed-object
 // dummy argument and those are distinguishable.

       y.intent != common::Intent::In) {
     return true;
   } 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)) &&
       (y.attrs.test(Attr::Allocatable) || y.attrs.test(Attr::Pointer)) &&
       (x.type.type().IsUnlimitedPolymorphic() !=

 bool DistinguishUtils::Distinguishable(
     const FunctionResult &x, const FunctionResult &y) const {
   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) {
             return Distinguishable(
                 z, std::get<TypeAndShape>(y.u), common::IgnoreTKRSet{});

flang/lib/Evaluate/formatting.cpp

       arg->AsFortran(o) << '%';
       break;
     }
   }
   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()) {
       arg->AsFortran(o << separator);
       separator = ',';

flang/lib/Evaluate/type.cpp

   default:
     return std::nullopt;
   }
 }
 
-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:
     if (checkCharLength && type.knownLength().value_or(0) != 1) {
       return false;

     // Derived types are tested in Semantics/check-declarations.cpp
     return false;
   }
 }
 
-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

flang/lib/Parser/unparse.cpp

   void Unparse(const FunctionReference &x) { // R1520
     Walk(std::get<ProcedureDesignator>(x.v.t));
     Put('('), Walk(std::get<std::list<ActualArgSpec>>(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');
     } else {
       const auto &pd{std::get<ProcedureDesignator>(x.call.t)};

flang/lib/Semantics/check-call.cpp

 static void CheckExplicitDataArg(const characteristics::DummyDataObject &dummy,
     const std::string &dummyName, evaluate::Expr<evaluate::SomeType> &actual,
     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()};
   CheckCharacterActual(actual, dummy, actualType, context, messages);
   bool dummyIsAllocatable{

             "Any pointer associated with TARGET %s during this call must not be used afterwards, as '%s' is not a target"_warn_en_US,
             dummyName, actual.AsFortran());
       }
     }
   }
-
-  // CUDA
-  if (!intrinsic &&
-      !dummy.attrs.test(characteristics::DummyDataObject::Attr::Value)) {
-    std::optional<common::CUDADataAttr> actualDataAttr, dummyDataAttr;
-    if (const auto *actualObject{actualLastSymbol
-                ? actualLastSymbol->detailsIf<ObjectEntityDetails>()
-                : 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<common::CUDADataAttr> 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,
     const characteristics::Procedure &proc,
     const characteristics::DummyProcedure &dummy, const std::string &dummyName,

                   arg.set_dummyIntent(object.intent);
                   bool isElemental{
                       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
                 } else if (object.type.type().IsTypelessIntrinsicArgument() &&
                     evaluate::IsNullObjectPointer(*expr)) {

flang/lib/Semantics/check-declarations.cpp

     if (msg && messages_.at().begin() != symbol.name().begin()) {
       evaluate::AttachDeclaration(*msg, symbol);
     }
     return msg;
   }
-  template <typename... A> parser::Message *WarnIfNotInModuleFile(A &&...x) {
-    if (FindModuleFileContaining(context_.FindScope(messages_.at()))) {
-      return nullptr;
-    }
-    return messages_.Say(std::forward<A>(x)...);
-  }
-  template <typename... A>
-  parser::Message *WarnIfNotInModuleFile(parser::CharBlock source, A &&...x) {
-    if (FindModuleFileContaining(context_.FindScope(source))) {
-      return nullptr;
-    }
-    return messages_.Say(source, std::forward<A>(x)...);
-  }
   bool IsResultOkToDiffer(const FunctionResult &);
   void CheckGlobalName(const Symbol &);
   void CheckExplicitSave(const Symbol &);
   void CheckBindC(const Symbol &);
   void CheckBindCFunctionResult(const Symbol &);

   }
 }
 
 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);
   }
   if (context_.HasError(symbol)) {

   }
   CheckAssumedTypeEntity(symbol, details);
   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)) {
       if (!isDeferredCoshape) { // C827
         messages_.Say("'%s' is an ALLOCATABLE coarray and must have a deferred"
                       " 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"
                     " attribute%s"_err_en_US,
           symbol.name(), deferredMsg);

             "!DIR$ IGNORE_TKR(R) may not apply in an ELEMENTAL procedure"_err_en_US);
       }
       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(
                 "!DIR$ IGNORE_TKR may not apply to an allocatable or pointer"_err_en_US);
           }
         } 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(
                 "!DIR$ IGNORE_TKR(R) may not apply to a dummy argument passed via descriptor"_err_en_US);
           }

           "An automatic variable or component must not be initialized"_err_en_US);
     } else if (IsDummy(symbol)) {
       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);
     }
   }
   if (symbol.owner().kind() == Scope::Kind::BlockData) {
     if (IsAllocatable(symbol)) {

   if (symbol.attrs().test(Attr::EXTERNAL)) {
     SayWithDeclaration(symbol,
         "'%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<SubprogramDetails>()
-          : 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) {
   if (IsPointer(symbol) && !context_.HasError(symbol) &&
       !scopeIsUninstantiatedPDT_) {

   bool canBeDeferred{arraySpec.CanBeDeferredShape()};
   bool canBeImplied{arraySpec.CanBeImpliedShape()};
   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<parser::MessageFixedText> msg;
   if (symbol.test(Symbol::Flag::CrayPointee) && !isExplicit &&
       !canBeAssumedSize) {
     msg = "Cray pointee '%s' must have explicit shape or"
           " assumed size"_err_en_US;

     if (canBeImplied && !canBeAssumedSize) { // C836
       msg = "Dummy array argument '%s' may not have implied shape"_err_en_US;
     }
   } 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;
     }
   } else if (IsNamedConstant(symbol)) {
     if (!isExplicit && !canBeImplied) {

           "A function interface may not declare an assumed-length CHARACTER(*) result"_err_en_US);
     }
   }
   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) {
   if (IsExternal(symbol)) {
     std::string interfaceName{symbol.name().ToString()};

         } else if (auto chars{Characterize(symbol)}) {
           if (auto globalChars{Characterize(*global)}) {
             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);
               }
             } else if (!globalChars->CanBeCalledViaImplicitInterface()) {
               msg = messages_.Say(

       const Symbol &previous{*iter->second};
       if (auto chars{Characterize(symbol)}) {
         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);
               evaluate::AttachDeclaration(msg, symbol);
             }

     msg = "%s function '%s' conflicts with intrinsic operator"_err_en_US;
   } else {
     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
 // false and return the error message in msg.
 std::optional<parser::MessageFixedText> CheckHelper::CheckNumberOfArgs(

           " OPTIONAL"_err_en_US;
   } else if (const auto *dataObject{std::get_if<DummyDataObject>(&arg.u)};
              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;
 }
 
 // Check if this procedure can be used for defined assignment (see 15.4.3.4.3).

   if (arg.IsOptional()) {
     msg = "In defined assignment subroutine '%s', dummy argument '%s'"
           " may not be OPTIONAL"_err_en_US;
   } else if (const auto *dataObject{std::get_if<DummyDataObject>(&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;
       } else if (dataObject->attrs.test(DummyDataObject::Attr::Allocatable)) {
         msg =

   } else {
     msg = "In defined assignment subroutine '%s', dummy argument '%s'"
           " 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;
 }
 
 // Report a conflicting attribute error if symbol has both of these attributes

   while (const auto *derivedDetails{
       derivedSym ? derivedSym->detailsIf<DerivedTypeDetails>() : nullptr}) {
     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)}) {
         msg->Attach(derivedSym->name(),
             "Declaration of derived type '%s'"_en_US, derivedSym->name());

         // ok
       } else if (IsAllocatableOrPointer(symbol) &&
           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);
         context_.SetError(symbol);
       } else {

             "A BIND(C) object must have an interoperable type"_err_en_US);
         context_.SetError(symbol);
       }
     }
     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<ProcEntityDetails>()}) {
-    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);
     }
   } else if (const auto *subp{symbol.detailsIf<SubprogramDetails>()}) {

                 msg->Attach(derived->typeSymbol().name(),
                     "Non-interoperable component type"_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);
-            }
+          } 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{
                 evaluate::GetConstantExtents(foldingContext_, component)};
             extents && evaluate::GetSize(*extents) == 0) {

               "An array component of an interoperable type must have at least one element"_err_en_US);
           context_.SetError(symbol);
         }
       }
     }
-    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);
     }
   }
 }
 

flang/lib/Semantics/definable.cpp

       return BlameSymbol(at,
           "'%s' is externally visible via '%s' and not definable in a pure subprogram"_en_US,
           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;
 }
 
 static std::optional<parser::Message> WhyNotDefinableLast(parser::CharBlock at,
     const Scope &scope, DefinabilityFlags flags, const Symbol &original) {

flang/lib/Semantics/expression.cpp

       } else if (IsAssumedLengthCharacter(proc) && IsExternal(proc)) {
         // TODO: Also catch assumed PDT type parameters
         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);
     }
   }
 }

     return false;
   }
   return true;
 }
 
-std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
-    const parser::CallStmt &call) {
-  Chevrons result;
-  auto checkLaunchArg{[&](const Expr<SomeType> &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<parser::StructureConstructor> *structureConstructor) {
   const parser::Call &call{funcRef.v};
   auto restorer{GetContextualMessages().SetLocation(funcRef.source)};
   ArgumentAnalyzer analyzer{*this, funcRef.source, true /* isProcedureCall */};

     analyzer.Analyze(arg, false /* not subroutine call */);
   }
   if (analyzer.fatalErrors()) {
     return std::nullopt;
   }
-  bool mightBeStructureConstructor{structureConstructor != nullptr};
-  if (std::optional<CalleeAndArguments> callee{GetCalleeAndArguments(
-          std::get<parser::ProcedureDesignator>(call.t), analyzer.GetActuals(),
-          false /* not subroutine */, mightBeStructureConstructor)}) {
+  if (std::optional<CalleeAndArguments> callee{
+          GetCalleeAndArguments(std::get<parser::ProcedureDesignator>(call.t),
+              analyzer.GetActuals(), false /* not subroutine */,
+              true /* might be structure constructor */)}) {
     if (auto *proc{std::get_if<ProcedureDesignator>(&callee->u)}) {
       return MakeFunctionRef(
           funcRef.source, std::move(*proc), std::move(callee->arguments));
     }
     CHECK(std::holds_alternative<semantics::SymbolRef>(callee->u));
     const Symbol &symbol{*std::get<semantics::SymbolRef>(callee->u)};
-    if (mightBeStructureConstructor) {
+    if (structureConstructor) {
       // Structure constructor misparsed as function reference?
       const auto &designator{std::get<parser::ProcedureDesignator>(call.t)};
       if (const auto *name{std::get_if<parser::Name>(&designator.u)}) {
         semantics::Scope &scope{context_.FindScope(name->source)};
         semantics::DerivedTypeSpec dtSpec{name->source, symbol.GetUltimate()};

   ArgumentAnalyzer analyzer{*this, callStmt.source, true /* isProcedureCall */};
   const auto &actualArgList{std::get<std::list<parser::ActualArgSpec>>(call.t)};
   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<CalleeAndArguments> callee{
             GetCalleeAndArguments(std::get<parser::ProcedureDesignator>(call.t),
                 analyzer.GetActuals(), true /* subroutine */)}) {
       ProcedureDesignator *proc{std::get_if<ProcedureDesignator>(&callee->u)};
       CHECK(proc);
-      bool isKernel{false};
-      if (const Symbol * procSym{proc->GetSymbol()}) {
-        const Symbol &ultimate{procSym->GetUltimate()};
-        if (const auto *subpDetails{
-                ultimate.detailsIf<semantics::SubprogramDetails>()}) {
-          if (auto attrs{subpDetails->cudaSubprogramAttrs()}) {
-            isKernel = *attrs == common::CUDASubprogramAttrs::Global ||
-                *attrs == common::CUDASubprogramAttrs::Grid_Global;
-          }
-        } else if (const auto *procDetails{
-                       ultimate.detailsIf<semantics::ProcEntityDetails>()}) {
-          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;
       }
     }
     if (!context_.AnyFatalError()) {
       std::string buf;

     }
   }
   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<SomeType>{std::move(procRef)};
+        return Expr<SomeType>{
+            ProcedureRef{std::move(proc), std::move(arguments)}};
       } else {
         // Not a procedure pointer, so type and shape are known.
         return TypedWrapper<FunctionRef, ProcedureRef>(
-            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);
     }
   }

flang/lib/Semantics/resolve-names.cpp

       bounds.push_back(*value);
     } else {
       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<SubprogramDetails>()
                      : nullptr}) {
     if (subp->cudaLaunchBounds().empty()) {
       subp->set_cudaLaunchBounds(std::move(bounds));

flang/lib/Semantics/type.cpp

 
 llvm::raw_ostream &operator<<(llvm::raw_ostream &o, const DeclTypeSpec &x) {
   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

flang/module/__cuda_builtins.f90

-!===-- 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

flang/test/Parser/cuf-sanity-tree.CUF

 !CHECK: | | | SubroutineStmt
 !CHECK: | | | | PrefixSpec -> Attributes -> CUDASubprogramAttrs = Device
 !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
 !CHECK: | | | | | EntityDecl
 !CHECK: | | | | | | Name = 'devx2'

 !CHECK: | | | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
 !CHECK: | | | | | | | | Scalar -> Expr = '10_4'
 !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'
 !CHECK: | | | | | | | Name = 'pa'
 !CHECK: | | | | | | AllocateShapeSpec

flang/test/Parser/cuf-sanity-unparse.CUF

 !CHECK:  ATTRIBUTES(GLOBAL) LAUNCH_BOUNDS(1_4, 2_4) SUBROUTINE lbsub
 !CHECK:  END SUBROUTINE
 !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
 !CHECK:   LOGICAL ispinned
 !CHECK:   !$CUF KERNEL DO (1_4) <<<*,*,STREAM=1_4>>>
 !CHECK:   DO j=1_4,10_4
 !CHECK:   END DO
 !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
-

flang/test/Semantics/bind-c06.f90

     type(t8) :: y
     integer :: z
   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
     !ERROR: Each component of an interoperable derived type must have an interoperable type
     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
     !ERROR: Each component of an interoperable derived type must have an interoperable type
     real(kind=2) x

flang/test/Semantics/cuf02.cuf

-! 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

flang/test/Semantics/cuf03.cuf

-! 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

flang/test/Semantics/cuf07.cuf

-! 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

flang/test/Semantics/cuf08.cuf

-! 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

flang/test/Semantics/cuf10.cuf

-! 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

flang/test/Semantics/definable05.cuf

-! 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

flang/test/Semantics/resolve65.f90

 
 module m1
   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
     procedure :: assign_t
     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
     procedure, nopass :: assign_t
     !ERROR: Defined assignment procedure 'assign_t' may not have NOPASS attribute

     class(t), intent(out) :: x
     real :: y
   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
   end
 end
 
 module m2
   type :: t

flang/test/Semantics/resolve67.f90

     character(*) function divide(x, y)
       character(*), intent(in) :: x, y
     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()
         end
       end interface