[flang-commits] [flang] f513bd8 - [flang] CUDA Fortran - part 4/5: definability and characteristics
Peter Klausler via flang-commits
flang-commits at lists.llvm.org
Wed May 31 14:39:39 PDT 2023
Author: Peter Klausler
Date: 2023-05-31T14:25:38-07:00
New Revision: f513bd808867757a92791fb61d893534cafbe0b3
URL: https://github.com/llvm/llvm-project/commit/f513bd808867757a92791fb61d893534cafbe0b3
DIFF: https://github.com/llvm/llvm-project/commit/f513bd808867757a92791fb61d893534cafbe0b3.diff
LOG: [flang] CUDA Fortran - part 4/5: definability and characteristics
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.
Differential Revision: https://reviews.llvm.org/D150163
Added:
flang/test/Semantics/cuf07.cuf
flang/test/Semantics/cuf10.cuf
flang/test/Semantics/definable05.cuf
Modified:
flang/include/flang/Evaluate/call.h
flang/include/flang/Evaluate/characteristics.h
flang/include/flang/Semantics/expression.h
flang/lib/Evaluate/characteristics.cpp
flang/lib/Evaluate/formatting.cpp
flang/lib/Semantics/check-call.cpp
flang/lib/Semantics/definable.cpp
flang/lib/Semantics/expression.cpp
flang/test/Parser/cuf-sanity-tree.CUF
flang/test/Parser/cuf-sanity-unparse.CUF
Removed:
################################################################################
diff --git a/flang/include/flang/Evaluate/call.h b/flang/include/flang/Evaluate/call.h
index 76983853c169d..f2c231647390b 100644
--- a/flang/include/flang/Evaluate/call.h
+++ b/flang/include/flang/Evaluate/call.h
@@ -209,6 +209,8 @@ struct ProcedureDesignator {
u;
};
+using Chevrons = std::vector<Expr<SomeType>>;
+
class ProcedureRef {
public:
CLASS_BOILERPLATE(ProcedureRef)
@@ -223,6 +225,10 @@ class ProcedureRef {
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;
@@ -250,6 +256,7 @@ class ProcedureRef {
protected:
ProcedureDesignator proc_;
ActualArguments arguments_;
+ Chevrons chevrons_;
bool hasAlternateReturns_;
};
diff --git a/flang/include/flang/Evaluate/characteristics.h b/flang/include/flang/Evaluate/characteristics.h
index 46cc6f23bddc0..824060f725d2c 100644
--- a/flang/include/flang/Evaluate/characteristics.h
+++ b/flang/include/flang/Evaluate/characteristics.h
@@ -220,6 +220,7 @@ struct DummyDataObject {
common::Intent intent{common::Intent::Default};
Attrs attrs;
common::IgnoreTKRSet ignoreTKR;
+ std::optional<common::CUDADataAttr> cudaDataAttr;
};
// 15.3.2.3
@@ -317,6 +318,7 @@ struct FunctionResult {
Attrs attrs;
std::variant<TypeAndShape, CopyableIndirection<Procedure>> u;
+ std::optional<common::CUDADataAttr> cudaDataAttr;
};
// 15.3.1
@@ -368,6 +370,8 @@ struct Procedure {
std::optional<FunctionResult> functionResult;
DummyArguments dummyArguments;
Attrs attrs;
+ std::optional<common::CUDASubprogramAttrs> cudaSubprogramAttrs;
};
+
} // namespace Fortran::evaluate::characteristics
#endif // FORTRAN_EVALUATE_CHARACTERISTICS_H_
diff --git a/flang/include/flang/Semantics/expression.h b/flang/include/flang/Semantics/expression.h
index 7cf7089715b52..a75314b5188dd 100644
--- a/flang/include/flang/Semantics/expression.h
+++ b/flang/include/flang/Semantics/expression.h
@@ -381,6 +381,7 @@ class ExpressionAnalyzer {
}
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 &);
diff --git a/flang/lib/Evaluate/characteristics.cpp b/flang/lib/Evaluate/characteristics.cpp
index 6b961ac9fae56..b22025c8844bc 100644
--- a/flang/lib/Evaluate/characteristics.cpp
+++ b/flang/lib/Evaluate/characteristics.cpp
@@ -265,7 +265,8 @@ llvm::raw_ostream &TypeAndShape::Dump(llvm::raw_ostream &o) const {
bool DummyDataObject::operator==(const DummyDataObject &that) const {
return type == that.type && attrs == that.attrs && intent == that.intent &&
- coshape == that.coshape;
+ coshape == that.coshape && cudaDataAttr == that.cudaDataAttr;
+ ;
}
static bool AreCompatibleDummyDataObjectShapes(const Shape &x, const Shape &y) {
@@ -325,6 +326,13 @@ bool DummyDataObject::IsCompatibleWith(
*whyNot = "incompatible !DIR$ IGNORE_TKR directives";
}
}
+ if (!attrs.test(Attr::Value) &&
+ !common::AreCompatibleCUDADataAttrs(
+ cudaDataAttr, actual.cudaDataAttr, ignoreTKR)) {
+ if (whyNot) {
+ *whyNot = "incompatible CUDA data attributes";
+ }
+ }
return true;
}
@@ -360,6 +368,14 @@ std::optional<DummyDataObject> DummyDataObject::Characterize(
});
result->intent = GetIntent(symbol.attrs());
result->ignoreTKR = GetIgnoreTKR(symbol);
+ if (object) {
+ result->cudaDataAttr = object->cudaDataAttr();
+ if (!result->cudaDataAttr &&
+ !result->attrs.test(DummyDataObject::Attr::Value) &&
+ semantics::IsCUDADeviceContext(&symbol.owner())) {
+ result->cudaDataAttr = common::CUDADataAttr::Device;
+ }
+ }
return result;
}
}
@@ -380,6 +396,8 @@ bool DummyDataObject::CanBePassedViaImplicitInterface() const {
return false; // 15.4.2.2(3)(b-d)
} else if (type.type().IsPolymorphic()) {
return false; // 15.4.2.2(3)(f)
+ } else if (cudaDataAttr) {
+ return false;
} else if (const auto *derived{GetDerivedTypeSpec(type.type())}) {
return derived->parameters().empty(); // 15.4.2.2(3)(e)
} else {
@@ -400,6 +418,9 @@ llvm::raw_ostream &DummyDataObject::Dump(llvm::raw_ostream &o) const {
sep = ',';
}
}
+ if (cudaDataAttr) {
+ o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr);
+ }
if (!ignoreTKR.empty()) {
ignoreTKR.Dump(o << ' ', common::EnumToString);
}
@@ -522,6 +543,7 @@ static std::optional<Procedure> CharacterizeProcedure(
return std::nullopt;
}
}
+ result.cudaSubprogramAttrs = subp.cudaSubprogramAttrs();
return result;
},
[&](const semantics::ProcEntityDetails &proc)
@@ -554,6 +576,10 @@ static std::optional<Procedure> CharacterizeProcedure(
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};
@@ -844,13 +870,14 @@ 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,
@@ -859,6 +886,7 @@ static std::optional<FunctionResult> CharacterizeFunctionResult(
{semantics::Attr::CONTIGUOUS, FunctionResult::Attr::Contiguous},
{semantics::Attr::POINTER, FunctionResult::Attr::Pointer},
});
+ result.cudaDataAttr = object->cudaDataAttr();
return result;
}
} else if (auto maybeProc{
@@ -887,6 +915,8 @@ bool FunctionResult::IsAssumedLengthCharacter() const {
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)
@@ -953,6 +983,10 @@ bool FunctionResult::IsCompatibleWith(
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()) {
@@ -1033,6 +1067,9 @@ llvm::raw_ostream &FunctionResult::Dump(llvm::raw_ostream &o) const {
},
},
u);
+ if (cudaDataAttr) {
+ o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr);
+ }
return o;
}
@@ -1045,7 +1082,8 @@ 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,
@@ -1078,6 +1116,10 @@ bool Procedure::IsCompatibleWith(const Procedure &actual, std::string *whyNot,
}
} 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";
@@ -1200,6 +1242,10 @@ 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;
@@ -1227,7 +1273,11 @@ llvm::raw_ostream &Procedure::Dump(llvm::raw_ostream &o) const {
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
@@ -1329,6 +1379,9 @@ bool DistinguishUtils::Distinguishable(
if (pos2 >= 0 && pos2 <= name2) {
return true; // distinguishable based on C1514 rule 4
}
+ if (proc1.cudaSubprogramAttrs != proc2.cudaSubprogramAttrs) {
+ return true;
+ }
return false;
}
@@ -1456,6 +1509,9 @@ bool DistinguishUtils::Distinguishable(
} 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)) &&
@@ -1494,6 +1550,9 @@ bool DistinguishUtils::Distinguishable(
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) {
diff --git a/flang/lib/Evaluate/formatting.cpp b/flang/lib/Evaluate/formatting.cpp
index f9548e119f1a5..84dd4be76cd9c 100644
--- a/flang/lib/Evaluate/formatting.cpp
+++ b/flang/lib/Evaluate/formatting.cpp
@@ -135,6 +135,18 @@ llvm::raw_ostream &ProcedureRef::AsFortran(llvm::raw_ostream &o) const {
}
}
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()) {
diff --git a/flang/lib/Semantics/check-call.cpp b/flang/lib/Semantics/check-call.cpp
index 7b4e6e245c945..e5a338b64bf68 100644
--- a/flang/lib/Semantics/check-call.cpp
+++ b/flang/lib/Semantics/check-call.cpp
@@ -196,7 +196,8 @@ static void CheckExplicitDataArg(const characteristics::DummyDataObject &dummy,
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()};
@@ -628,6 +629,46 @@ static void CheckExplicitDataArg(const characteristics::DummyDataObject &dummy,
}
}
}
+
+ // 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,
@@ -819,7 +860,7 @@ static void CheckExplicitInterfaceArg(evaluate::ActualArgument &arg,
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
diff --git a/flang/lib/Semantics/definable.cpp b/flang/lib/Semantics/definable.cpp
index 675becd32c266..abb5f35c28eae 100644
--- a/flang/lib/Semantics/definable.cpp
+++ b/flang/lib/Semantics/definable.cpp
@@ -134,6 +134,33 @@ static std::optional<parser::Message> WhyNotDefinableBase(parser::CharBlock at,
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;
}
diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp
index 28c4ba16ae926..a6b54dd11d21b 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -2653,6 +2653,9 @@ void ExpressionAnalyzer::CheckForBadRecursion(
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);
}
@@ -2719,6 +2722,55 @@ bool ExpressionAnalyzer::CheckIsValidForwardReference(
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};
@@ -2730,17 +2782,17 @@ MaybeExpr ExpressionAnalyzer::Analyze(const parser::FunctionReference &funcRef,
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)}) {
@@ -2785,17 +2837,40 @@ void ExpressionAnalyzer::Analyze(const parser::CallStmt &callStmt) {
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;
}
}
@@ -3697,14 +3772,13 @@ MaybeExpr ExpressionAnalyzer::MakeFunctionRef(parser::CharBlock callSite,
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);
diff --git a/flang/test/Parser/cuf-sanity-tree.CUF b/flang/test/Parser/cuf-sanity-tree.CUF
index 2ce042bcdbc1b..f6cf9bbdd6b0c 100644
--- a/flang/test/Parser/cuf-sanity-tree.CUF
+++ b/flang/test/Parser/cuf-sanity-tree.CUF
@@ -106,6 +106,9 @@ include "cuf-sanity-common"
!CHECK: | | | | Name = 'attrs'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
+!CHECK: | | | | DeclarationConstruct -> SpecificationConstruct -> OtherSpecificationStmt -> CUDAAttributesStmt
+!CHECK: | | | | | CUDADataAttr = Device
+!CHECK: | | | | | Name = 'devx1'
!CHECK: | | | | DeclarationConstruct -> SpecificationConstruct -> TypeDeclarationStmt
!CHECK: | | | | | DeclarationTypeSpec -> IntrinsicTypeSpec -> Real
!CHECK: | | | | | AttrSpec -> CUDADataAttr = Device
@@ -159,27 +162,36 @@ include "cuf-sanity-common"
!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'
diff --git a/flang/test/Parser/cuf-sanity-unparse.CUF b/flang/test/Parser/cuf-sanity-unparse.CUF
index 3bd838a75e3ea..d4495c4fddccf 100644
--- a/flang/test/Parser/cuf-sanity-unparse.CUF
+++ b/flang/test/Parser/cuf-sanity-unparse.CUF
@@ -23,6 +23,7 @@ include "cuf-sanity-common"
!CHECK: ATTRIBUTES(GLOBAL) CLUSTER_DIMS(1_4, 2_4, 3_4) SUBROUTINE cdsub
!CHECK: END SUBROUTINE
!CHECK: ATTRIBUTES(DEVICE) SUBROUTINE attrs
+!CHECK: ATTRIBUTES(DEVICE) devx1
!CHECK: REAL, DEVICE :: devx2
!CHECK: END SUBROUTINE
!CHECK: SUBROUTINE test
@@ -33,9 +34,9 @@ include "cuf-sanity-common"
!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
diff --git a/flang/test/Semantics/cuf07.cuf b/flang/test/Semantics/cuf07.cuf
new file mode 100644
index 0000000000000..b520b5da51264
--- /dev/null
+++ b/flang/test/Semantics/cuf07.cuf
@@ -0,0 +1,26 @@
+! RUN: %python %S/test_errors.py %s %flang_fc1
+module m
+ real, allocatable :: xa
+ real, allocatable, managed :: ma
+ contains
+ attributes(device) subroutine devsubr
+ real, device, allocatable :: da
+ real, allocatable, managed :: dma
+ allocate(da) ! ok
+ deallocate(da) ! ok
+ allocate(dma) ! ok
+ deallocate(dma) ! ok
+ !ERROR: Name in ALLOCATE statement is not definable
+ !BECAUSE: 'xa' is a host variable and is not definable in a device subprogram
+ allocate(xa)
+ !ERROR: Name in DEALLOCATE statement is not definable
+ !BECAUSE: 'xa' is a host variable and is not definable in a device subprogram
+ deallocate(xa)
+ !ERROR: Name in ALLOCATE statement is not definable
+ !BECAUSE: 'ma' is a host-associated allocatable and is not definable in a device subprogram
+ allocate(ma)
+ !ERROR: Name in DEALLOCATE statement is not definable
+ !BECAUSE: 'ma' is a host-associated allocatable and is not definable in a device subprogram
+ deallocate(ma)
+ end subroutine
+end module
diff --git a/flang/test/Semantics/cuf10.cuf b/flang/test/Semantics/cuf10.cuf
new file mode 100644
index 0000000000000..0d05222d446df
--- /dev/null
+++ b/flang/test/Semantics/cuf10.cuf
@@ -0,0 +1,17 @@
+! RUN: %python %S/test_errors.py %s %flang_fc1
+module m
+ real, device :: a(4,8)
+ real, managed, allocatable :: b(:,:)
+ contains
+ attributes(global) subroutine kernel(a,b,c,n,m)
+ integer, value :: n
+ integer, intent(in) :: m
+ real a(n,m), c(n,m)
+ real, managed :: b(n,m)
+ end
+ subroutine test
+ allocate(b(4,8))
+ !ERROR: dummy argument 'm=' has ATTRIBUTES(DEVICE) but its associated actual argument has no CUDA data attribute
+ call kernel<<<1,32>>>(a,b,b,4,8)
+ end
+end
diff --git a/flang/test/Semantics/definable05.cuf b/flang/test/Semantics/definable05.cuf
new file mode 100644
index 0000000000000..5af3ca9244345
--- /dev/null
+++ b/flang/test/Semantics/definable05.cuf
@@ -0,0 +1,31 @@
+! RUN: %python %S/test_errors.py %s %flang_fc1
+module m
+ real, constant :: rc
+ !ERROR: Object 'rcp' with ATTRIBUTES(CONSTANT) may not be allocatable, pointer, or target
+ real, constant, pointer :: rcp
+ !ERROR: Object 'rct' with ATTRIBUTES(CONSTANT) may not be allocatable, pointer, or target
+ real, constant, target :: rct
+ real, device, pointer :: dp(:)
+ real, device, target :: dt(100)
+ contains
+ attributes(device) subroutine devsub
+ !ERROR: Left-hand side of assignment is not definable
+ !BECAUSE: 'rc' has ATTRIBUTES(CONSTANT) and is not definable in a device subprogram
+ rc = 1.
+ !ERROR: The left-hand side of a pointer assignment is not definable
+ !BECAUSE: 'dp' is a pointer and may not be associated in a device subprogram
+ dp => dt
+ end
+ attributes(global) subroutine globsub
+ !ERROR: Left-hand side of assignment is not definable
+ !BECAUSE: 'rc' has ATTRIBUTES(CONSTANT) and is not definable in a device subprogram
+ rc = 1.
+ !ERROR: The left-hand side of a pointer assignment is not definable
+ !BECAUSE: 'dp' is a pointer and may not be associated in a device subprogram
+ dp => dt
+ end
+ subroutine hostsub
+ rc = 1.
+ dp => dt
+ end
+end
More information about the flang-commits
mailing list