[flang] CUDA Fortran - part 4/5: definability and characteristics

[klausler]

Summary

Extend the definability and procedure characteristics checking
infrastructure in semantics to check for context-dependent CUDA object
definability violations and problems with CUDA attribute incompatibility
in procedure interfaces.

Depends on https://reviews.llvm.org/D150159,
https://reviews.llvm.org/D150161, & https://reviews.llvm.org/D150162.

[klausler]

Remove redundantly added tests.

[klausler]

Remove redundantly added tests.

[klausler]

Fix compilation issue with more recent GCC.

[clementval]

LGTM

[klausler]

Rebase to current head

[klausler]

Rebase to current head.

[klausler]

Rebase to current head.

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/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/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;
+      coshape == that.coshape && cudaDataAttr == that.cudaDataAttr;
+  ;
 }
 
 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);
+  }
   if (!ignoreTKR.empty()) {
     ignoreTKR.Dump(o << ' ', common::EnumToString);
   }
   return o;
 }

                     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 && u == that.u;
+  return attrs == that.attrs && cudaDataAttr == that.cudaDataAttr &&
+      u == that.u;
 }
 
 static std::optional<FunctionResult> CharacterizeFunctionResult(
     const semantics::Symbol &symbol, FoldingContext &context,
     semantics::UnorderedSymbolSet seenProcs) {
-  if (symbol.has<semantics::ObjectEntityDetails>()) {
+  if (const auto *object{symbol.detailsIf<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;
+      dummyArguments == that.dummyArguments &&
+      cudaSubprogramAttrs == that.cudaSubprogramAttrs;
 }
 
 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 = ',';
   }
-  return o << (sep == '(' ? "()" : ")");
+  o << (sep == '(' ? "()" : ")");
+  if (cudaSubprogramAttrs) {
+    o << " cudaSubprogramAttrs: " << common::EnumToString(*cudaSubprogramAttrs);
+  }
+  return o;
 }
 
 // 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/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() && !x.chevrons /*CUDA todo*/) {
+    if (asFortran_ && x.typedCall.get()) {
       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,
     SemanticsContext &context, evaluate::FoldingContext &foldingContext,
     const Scope *scope, const evaluate::SpecificIntrinsic *intrinsic,
-    bool allowActualArgumentConversions) {
+    bool allowActualArgumentConversions,
+    const characteristics::Procedure &procedure) {
 
   // Basic type & rank checking
   parser::ContextualMessages &messages{foldingContext.messages()};
   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, foldingContext, scope, intrinsic,
-                      allowActualArgumentConversions);
+                      allowActualArgumentConversions, proc);
                 } else if (object.type.type().IsTypelessIntrinsicArgument() &&
                     IsBOZLiteral(*expr)) {
                   // ok
                 } else if (object.type.type().IsTypelessIntrinsicArgument() &&
                     evaluate::IsNullObjectPointer(*expr)) {

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;
   }
-  if (std::optional<CalleeAndArguments> callee{
-          GetCalleeAndArguments(std::get<parser::ProcedureDesignator>(call.t),
-              analyzer.GetActuals(), false /* not subroutine */,
-              true /* might be structure constructor */)}) {
+  bool mightBeStructureConstructor{structureConstructor != nullptr};
+  if (std::optional<CalleeAndArguments> callee{GetCalleeAndArguments(
+          std::get<parser::ProcedureDesignator>(call.t), analyzer.GetActuals(),
+          false /* not subroutine */, mightBeStructureConstructor)}) {
     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 (structureConstructor) {
+    if (mightBeStructureConstructor) {
       // 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 */);
   }
-  if (!analyzer.fatalErrors()) {
+  auto chevrons{AnalyzeChevrons(callStmt)};
+  if (!analyzer.fatalErrors() && chevrons) {
     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>{
-            ProcedureRef{std::move(proc), std::move(arguments)}};
+        return Expr<SomeType>{std::move(procRef)};
       } else {
         // Not a procedure pointer, so type and shape are known.
         return TypedWrapper<FunctionRef, ProcedureRef>(
-            DEREF(result.GetTypeAndShape()).type(),
-            ProcedureRef{std::move(proc), std::move(arguments)});
+            DEREF(result.GetTypeAndShape()).type(), std::move(procRef));
       }
     } else {
       Say("Function result characteristics are not known"_err_en_US);
     }
   }

flang/lib/Semantics/symbol.cpp

     os << " cudaDataAttr: " << common::EnumToString(*x.cudaDataAttr());
   }
   if (x.cudaDataAttr()) {
     os << " cudaDataAttr: " << common::EnumToString(*x.cudaDataAttr());
   }
+  if (x.cudaDataAttr()) {
+    os << " cudaDataAttr: " << common::EnumToString(*x.cudaDataAttr());
+  }
   return os;
 }
 
 llvm::raw_ostream &operator<<(
     llvm::raw_ostream &os, const AssocEntityDetails &x) {

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()'
+!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4>>>()'
 !CHECK: | | | | | Call
 !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub'
 !CHECK: | | | | | Chevrons
-!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '1'
-!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '2'
-!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub()'
+!CHECK: | | | | | | Scalar -> Expr = '1_4'
+!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
+!CHECK: | | | | | | Scalar -> Expr = '2_4'
+!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
+!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4,3_4>>>()'
 !CHECK: | | | | | Call
 !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub'
 !CHECK: | | | | | Chevrons
-!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '1'
-!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '2'
-!CHECK: | | | | | | Scalar -> Integer -> Expr -> LiteralConstant -> IntLiteralConstant = '3'
-!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub()'
+!CHECK: | | | | | | Scalar -> Expr = '1_4'
+!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
+!CHECK: | | | | | | Scalar -> Expr = '2_4'
+!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
+!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4'
+!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3'
+!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4,3_4,4_4>>>()'
 !CHECK: | | | | | Call
 !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub'
 !CHECK: | | | | | Chevrons
-!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '1'
-!CHECK: | | | | | | Scalar -> Expr -> LiteralConstant -> IntLiteralConstant = '2'
-!CHECK: | | | | | | Scalar -> Integer -> Expr -> LiteralConstant -> IntLiteralConstant = '3'
-!CHECK: | | | | | | Scalar -> Integer -> Expr -> LiteralConstant -> IntLiteralConstant = '4'
+!CHECK: | | | | | | Scalar -> Expr = '1_4'
+!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
+!CHECK: | | | | | | Scalar -> Expr = '2_4'
+!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
+!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4'
+!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3'
+!CHECK: | | | | | | Scalar -> Integer -> Expr = '4_4'
+!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '4'
 !CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> AllocateStmt
 !CHECK: | | | | | Allocation
 !CHECK: | | | | | | AllocateObject = 'pa'
 !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,2>>>
-!CHECK:    CALL globalsub<<<1,2,3>>>
-!CHECK:    CALL globalsub<<<1,2,3,4>>>
+!CHECK:    CALL globalsub<<<1_4,2_4>>>()
+!CHECK:    CALL globalsub<<<1_4,2_4,3_4>>>()
+!CHECK:    CALL globalsub<<<1_4,2_4,3_4,4_4>>>()
 !CHECK:   ALLOCATE(pa(32_4), STREAM=1_4, PINNED=ispinned)
 !CHECK:  END SUBROUTINE
 !CHECK: END MODULE
+

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