Skip to content

Commit

Permalink
[flang] CUDA Fortran - part 4/5: definability and characteristics
Browse files Browse the repository at this point in the history 
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
  • Loading branch information
@klausler
klausler committed May 31, 2023
1 parent f924104 commit f513bd8
Show file tree
Hide file tree
Showing 13 changed files with 344 additions and 32 deletions.
7 changes: 7 additions & 0 deletions flang/include/flang/Evaluate/call.h
View file
Expand Up @@ -209,6 +209,8 @@ struct ProcedureDesignator {
u;
};

using Chevrons = std::vector<Expr<SomeType>>;

class ProcedureRef {
public:
CLASS_BOILERPLATE(ProcedureRef)
Expand All @@ -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;
Expand Down Expand Up @@ -250,6 +256,7 @@ class ProcedureRef {
protected:
ProcedureDesignator proc_;
ActualArguments arguments_;
Chevrons chevrons_;
bool hasAlternateReturns_;
};

Expand Down
4 changes: 4 additions & 0 deletions flang/include/flang/Evaluate/characteristics.h
View file
Expand Up @@ -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
Expand Down Expand Up @@ -317,6 +318,7 @@ struct FunctionResult {

Attrs attrs;
std::variant<TypeAndShape, CopyableIndirection<Procedure>> u;
std::optional<common::CUDADataAttr> cudaDataAttr;
};

// 15.3.1
Expand Down Expand Up @@ -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_
1 change: 1 addition & 0 deletions flang/include/flang/Semantics/expression.h
View file
Expand Up @@ -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 &);

Expand Down
69 changes: 64 additions & 5 deletions flang/lib/Evaluate/characteristics.cpp
View file
Expand Up @@ -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) {
Expand Down Expand Up @@ -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;
}

Expand Down Expand Up @@ -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;
}
}
Expand All @@ -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 {
Expand All @@ -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);
}
Expand Down Expand Up @@ -522,6 +543,7 @@ static std::optional<Procedure> CharacterizeProcedure(
return std::nullopt;
}
}
result.cudaSubprogramAttrs = subp.cudaSubprogramAttrs();
return result;
},
[&](const semantics::ProcEntityDetails &proc)
Expand Down Expand Up @@ -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};
Expand Down Expand Up @@ -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,
Expand All @@ -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{
Expand Down Expand Up @@ -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)
Expand Down Expand Up @@ -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()) {
Expand Down Expand Up @@ -1033,6 +1067,9 @@ llvm::raw_ostream &FunctionResult::Dump(llvm::raw_ostream &o) const {
},
},
u);
if (cudaDataAttr) {
o << " cudaDataAttr: " << common::EnumToString(*cudaDataAttr);
}
return o;
}

Expand All @@ -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,
Expand Down Expand Up @@ -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";
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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;
}

Expand Down Expand Up @@ -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)) &&
Expand Down Expand Up @@ -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) {
Expand Down
12 changes: 12 additions & 0 deletions flang/lib/Evaluate/formatting.cpp
View file
Expand Up @@ -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()) {
Expand Down
45 changes: 43 additions & 2 deletions flang/lib/Semantics/check-call.cpp
View file
Expand Up @@ -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()};
Expand Down Expand Up @@ -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,
Expand Down Expand Up @@ -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
Expand Down
27 changes: 27 additions & 0 deletions flang/lib/Semantics/definable.cpp
View file
Expand Up @@ -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;
}

Expand Down

0 comments on commit f513bd8

Please sign in to comment.