shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[flang] CUDA Fortran - part 1/5: parsing

Browse Source 
Begin upstreaming of CUDA Fortran support in LLVM Flang.

This first patch implements parsing for CUDA Fortran syntax,
including:
 - a new LanguageFeature enum value for CUDA Fortran
 - driver change to enable that feature for *.cuf and *.CUF source files
 - parse tree representation of CUDA Fortran syntax
 - dumping and unparsing of the parse tree
 - the actual parsers for CUDA Fortran syntax
 - prescanning support for !@CUF and !$CUF
 - basic sanity testing via unparsing and parse tree dumps

... along with any minimized changes elsewhere to make these
work, mostly no-op cases in common::visitors instances in
semantics and lowering to allow them to compile in the face
of new types in variant<> instances in the parse tree.

Because CUDA Fortran allows the kernel launch chevron syntax
("call foo<<<blocks, threads>>>()") only on CALL statements and
not on function references, the parse tree nodes for CallStmt,
FunctionReference, and their shared Call were rearranged a bit;
this caused a fair amount of one-line changes in many files.

More patches will follow that implement CUDA Fortran in the symbol
table and name resolution, and then semantic checking.

Differential Revision: https://reviews.llvm.org/D150159
This commit is contained in:
Peter Klausler 2023-05-06 15:03:39 -07:00
parent 5a0108947c
commit 4ad7279392
No known key found for this signature in database
38 changed files with 701 additions and 117 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
flang/docs/ParserCombinators.md
View File

@ -97,8 +97,9 @@ They are `constexpr`, so they should be viewed as type-safe macros.
* `nonemptySeparated(p, q)` repeatedly matches "p q p q p q ... p",
returning a `std::list<>` of only the values of the p's. It fails if
p immediately fails.
* `extension(p)` parses p if strict standard compliance is disabled,
or with a warning if nonstandard usage warnings are enabled.
* `extension<feature>([msg,]p)` parses p if strict standard compliance is
disabled, or with an optional warning when nonstandard usage warnings
are enabled.
* `deprecated(p)` parses p if strict standard compliance is disabled,
with a warning if deprecated usage warnings are enabled.
* `inContext(msg, p)` runs p within an error message context; any
@ -165,9 +166,9 @@ is built. All of the following parsers consume characters acquired from
a longer identifier or keyword).
* `parenthesized(p)` is shorthand for `"(" >> p / ")"`.
* `bracketed(p)` is shorthand for `"[" >> p / "]"`.
* `nonEmptyList(p)` matches a comma-separated list of one or more
* `nonemptyList(p)` matches a comma-separated list of one or more
instances of p.
* `nonEmptyList(errorMessage, p)` is equivalent to
* `nonemptyList(errorMessage, p)` is equivalent to
`withMessage(errorMessage, nonemptyList(p))`, which allows one to supply
a meaningful error message in the event of an empty list.
* `optionalList(p)` is the same thing, but can be empty, and always succeeds.

3
flang/include/flang/Common/Fortran-features.h
View File

@ -27,7 +27,7 @@ ENUM_CLASS(LanguageFeature, BackslashEscapes, OldDebugLines,
SignedPrimary, FileName, Carriagecontrol, Convert, Dispose,
IOListLeadingComma, AbbreviatedEditDescriptor, ProgramParentheses,
PercentRefAndVal, OmitFunctionDummies, CrayPointer, Hollerith, ArithmeticIF,
Assign, AssignedGOTO, Pause, OpenACC, OpenMP, CruftAfterAmpersand,
Assign, AssignedGOTO, Pause, OpenACC, OpenMP, CUDA, CruftAfterAmpersand,
ClassicCComments, AdditionalFormats, BigIntLiterals, RealDoControls,
EquivalenceNumericWithCharacter, EquivalenceNonDefaultNumeric,
EquivalenceSameNonSequence, AdditionalIntrinsics, AnonymousParents,
@ -54,6 +54,7 @@ public:
disable_.set(LanguageFeature::OldDebugLines);
disable_.set(LanguageFeature::OpenACC);
disable_.set(LanguageFeature::OpenMP);
disable_.set(LanguageFeature::CUDA); // !@cuf
disable_.set(LanguageFeature::ImplicitNoneTypeNever);
disable_.set(LanguageFeature::ImplicitNoneTypeAlways);
disable_.set(LanguageFeature::DefaultSave);

10
flang/include/flang/Common/Fortran.h
View File

@ -15,6 +15,7 @@
#include "enum-set.h"
#include "idioms.h"
#include <cinttypes>
#include <optional>
#include <string>
namespace Fortran::common {
@ -80,6 +81,12 @@ using Label = std::uint64_t;
// Fortran arrays may have up to 15 dimensions (See Fortran 2018 section 5.4.6).
static constexpr int maxRank{15};
// CUDA subprogram attribute combinations
ENUM_CLASS(CUDASubprogramAttrs, Host, Device, HostDevice, Global, Grid_Global)
// CUDA data attributes; mutually exclusive
ENUM_CLASS(CUDADataAttr, Constant, Device, Managed, Pinned, Shared, Texture)
// Fortran names may have up to 63 characters (See Fortran 2018 C601).
static constexpr int maxNameLen{63};
@ -99,5 +106,8 @@ static constexpr IgnoreTKRSet ignoreTKRAll{IgnoreTKR::Type, IgnoreTKR::Kind,
IgnoreTKR::Rank, IgnoreTKR::Device, IgnoreTKR::Managed};
std::string AsFortran(IgnoreTKRSet);
bool AreCompatibleCUDADataAttrs(
std::optional<CUDADataAttr>, std::optional<CUDADataAttr>, IgnoreTKRSet);
} // namespace Fortran::common
#endif // FORTRAN_COMMON_FORTRAN_H_

1
flang/include/flang/Common/indirection.h
View File

@ -148,6 +148,7 @@ public:
A *operator->() const { return p_; }
operator bool() const { return p_ != nullptr; }
A *get() { return p_; }
auto get() const { return reinterpret_cast<std::add_const_t<A> *>(p_); }
A *release() {
A *result{p_};
p_ = nullptr;

6
flang/include/flang/Common/template.h
View File

@ -94,8 +94,10 @@ constexpr int SearchMembers{
TUPLEorVARIANT>::value()};
template <typename A, typename TUPLEorVARIANT>
constexpr bool HasMember{
SearchMembers<MatchType<A>::template Match, TUPLEorVARIANT> >= 0};
constexpr int FindMember{
SearchMembers<MatchType<A>::template Match, TUPLEorVARIANT>};
template <typename A, typename TUPLEorVARIANT>
constexpr bool HasMember{FindMember<A, TUPLEorVARIANT> >= 0};
// std::optional<std::optional<A>> -> std::optional<A>
template <typename A>

9
flang/include/flang/Frontend/FrontendOptions.h
View File

@ -113,6 +113,10 @@ bool isFreeFormSuffix(llvm::StringRef suffix);
/// \return True if the file should be preprocessed
bool isToBePreprocessed(llvm::StringRef suffix);
/// \param suffix The file extension
/// \return True if the file contains CUDA Fortran
bool isCUDAFortranSuffix(llvm::StringRef suffix);
enum class Language : uint8_t {
Unknown,
@ -182,6 +186,9 @@ class FrontendInputFile {
/// sufficient to implement gfortran`s logic controlled with `-cpp/-nocpp`.
unsigned mustBePreprocessed : 1;
/// Whether to enable CUDA Fortran language extensions
bool isCUDAFortran{false};
public:
FrontendInputFile() = default;
FrontendInputFile(llvm::StringRef file, InputKind inKind)
@ -193,6 +200,7 @@ public:
std::string pathSuffix{file.substr(pathDotIndex + 1)};
isFixedForm = isFixedFormSuffix(pathSuffix);
mustBePreprocessed = isToBePreprocessed(pathSuffix);
isCUDAFortran = isCUDAFortranSuffix(pathSuffix);
}
FrontendInputFile(const llvm::MemoryBuffer *memBuf, InputKind inKind)
@ -204,6 +212,7 @@ public:
bool isFile() const { return (buffer == nullptr); }
bool getIsFixedForm() const { return isFixedForm; }
bool getMustBePreprocessed() const { return mustBePreprocessed; }
bool getIsCUDAFortran() const { return isCUDAFortran; }
llvm::StringRef getFile() const {
assert(isFile());

12
flang/include/flang/Parser/dump-parse-tree.h
View File

@ -14,6 +14,7 @@
#include "parse-tree.h"
#include "tools.h"
#include "unparse.h"
#include "flang/Common/Fortran.h"
#include "flang/Common/idioms.h"
#include "flang/Common/indirection.h"
#include "llvm/Support/raw_ostream.h"
@ -45,6 +46,8 @@ public:
NODE(std, string)
NODE(std, int64_t)
NODE(std, uint64_t)
NODE_ENUM(common, CUDADataAttr)
NODE_ENUM(common, CUDASubprogramAttrs)
NODE(format, ControlEditDesc)
NODE(format::ControlEditDesc, Kind)
NODE(format, DerivedTypeDataEditDesc)
@ -120,6 +123,8 @@ public:
NODE(parser, AllocOpt)
NODE(AllocOpt, Mold)
NODE(AllocOpt, Source)
NODE(AllocOpt, Stream)
NODE(AllocOpt, Pinned)
NODE(parser, Allocatable)
NODE(parser, AllocatableStmt)
NODE(parser, AllocateCoarraySpec)
@ -165,6 +170,7 @@ public:
NODE(parser, BoundsSpec)
NODE(parser, Call)
NODE(parser, CallStmt)
NODE(CallStmt, Chevrons)
NODE(parser, CaseConstruct)
NODE(CaseConstruct, Case)
NODE(parser, CaseSelector)
@ -216,6 +222,9 @@ public:
NODE(parser, ContinueStmt)
NODE(parser, CriticalConstruct)
NODE(parser, CriticalStmt)
NODE(parser, CUDAAttributesStmt)
NODE(parser, CUFKernelDoConstruct)
NODE(CUFKernelDoConstruct, Directive)
NODE(parser, CycleStmt)
NODE(parser, DataComponentDefStmt)
NODE(parser, DataIDoObject)
@ -610,6 +619,9 @@ public:
NODE(PrefixSpec, Non_Recursive)
NODE(PrefixSpec, Pure)
NODE(PrefixSpec, Recursive)
NODE(PrefixSpec, Attributes)
NODE(PrefixSpec, Launch_Bounds)
NODE(PrefixSpec, Cluster_Dims)
NODE(parser, PrintStmt)
NODE(parser, PrivateStmt)
NODE(parser, PrivateOrSequence)

1
flang/include/flang/Parser/message.h
View File

@ -54,6 +54,7 @@ public:
constexpr MessageFixedText &operator=(MessageFixedText &&) = default;
CharBlock text() const { return text_; }
bool empty() const { return text_.empty(); }
Severity severity() const { return severity_; }
MessageFixedText &set_severity(Severity severity) {
severity_ = severity;

24
flang/include/flang/Parser/parse-tree-visitor.h
View File

@ -568,17 +568,33 @@ template <typename M> void Walk(Designator &x, M &mutator) {
mutator.Post(x);
}
}
template <typename V> void Walk(const Call &x, V &visitor) {
template <typename V> void Walk(const FunctionReference &x, V &visitor) {
if (visitor.Pre(x)) {
Walk(x.source, visitor);
Walk(x.t, visitor);
Walk(x.v, visitor);
visitor.Post(x);
}
}
template <typename M> void Walk(Call &x, M &mutator) {
template <typename M> void Walk(FunctionReference &x, M &mutator) {
if (mutator.Pre(x)) {
Walk(x.source, mutator);
Walk(x.t, mutator);
Walk(x.v, mutator);
mutator.Post(x);
}
}
template <typename V> void Walk(const CallStmt &x, V &visitor) {
if (visitor.Pre(x)) {
Walk(x.source, visitor);
Walk(x.call, visitor);
Walk(x.chevrons, visitor);
visitor.Post(x);
}
}
template <typename M> void Walk(CallStmt &x, M &mutator) {
if (mutator.Pre(x)) {
Walk(x.source, mutator);
Walk(x.call, mutator);
Walk(x.chevrons, mutator);
mutator.Post(x);
}
}

90
flang/include/flang/Parser/parse-tree.h
View File

@ -67,7 +67,7 @@ class DerivedTypeSpec;
namespace Fortran::evaluate {
struct GenericExprWrapper; // forward definition, wraps Expr<SomeType>
struct GenericAssignmentWrapper; // forward definition, represent assignment
class ProcedureRef; // forward definition, represents a CALL statement
class ProcedureRef; // forward definition, represents a CALL or function ref
} // namespace Fortran::evaluate
// Most non-template classes in this file use these default definitions
@ -253,6 +253,7 @@ struct StmtFunctionStmt; // R1544
// Directives, extensions, and deprecated statements
struct CompilerDirective;
struct BasedPointerStmt;
struct CUDAAttributesStmt;
struct StructureDef;
struct ArithmeticIfStmt;
struct AssignStmt;
@ -264,6 +265,7 @@ struct OpenACCDeclarativeConstruct;
struct OpenMPConstruct;
struct OpenMPDeclarativeConstruct;
struct OmpEndLoopDirective;
struct CUFKernelDoConstruct;
// Cooked character stream locations
using Location = const char *;
@ -361,6 +363,7 @@ EMPTY_CLASS(ErrorRecovery);
// pointer-stmt | protected-stmt | save-stmt | target-stmt |
// volatile-stmt | value-stmt | common-stmt | equivalence-stmt
// Extension: (Cray) based POINTER statement
// Extension: CUDA data attribute statement
struct OtherSpecificationStmt {
UNION_CLASS_BOILERPLATE(OtherSpecificationStmt);
std::variant<common::Indirection<AccessStmt>,
@ -374,7 +377,8 @@ struct OtherSpecificationStmt {
common::Indirection<SaveStmt>, common::Indirection<TargetStmt>,
common::Indirection<ValueStmt>, common::Indirection<VolatileStmt>,
common::Indirection<CommonStmt>, common::Indirection<EquivalenceStmt>,
common::Indirection<BasedPointerStmt>>
common::Indirection<BasedPointerStmt>,
common::Indirection<CUDAAttributesStmt>>
u;
};
@ -507,7 +511,8 @@ struct ActionStmt {
// action-stmt | associate-construct | block-construct |
// case-construct | change-team-construct | critical-construct |
// do-construct | if-construct | select-rank-construct |
// select-type-construct | where-construct | forall-construct
// select-type-construct | where-construct | forall-construct |
// (CUDA) CUF-kernel-do-construct
struct ExecutableConstruct {
UNION_CLASS_BOILERPLATE(ExecutableConstruct);
std::variant<Statement<ActionStmt>, common::Indirection<AssociateConstruct>,
@ -524,7 +529,8 @@ struct ExecutableConstruct {
common::Indirection<OpenACCConstruct>,
common::Indirection<AccEndCombinedDirective>,
common::Indirection<OpenMPConstruct>,
common::Indirection<OmpEndLoopDirective>>
common::Indirection<OmpEndLoopDirective>,
common::Indirection<CUFKernelDoConstruct>>
u;
};
@ -977,14 +983,15 @@ struct ComponentArraySpec {
// R738 component-attr-spec ->
// access-spec | ALLOCATABLE |
// CODIMENSION lbracket coarray-spec rbracket |
// CONTIGUOUS | DIMENSION ( component-array-spec ) | POINTER
// CONTIGUOUS | DIMENSION ( component-array-spec ) | POINTER |
// (CUDA) CONSTANT | DEVICE | MANAGED | PINNED | SHARED | TEXTURE
EMPTY_CLASS(Allocatable);
EMPTY_CLASS(Pointer);
EMPTY_CLASS(Contiguous);
struct ComponentAttrSpec {
UNION_CLASS_BOILERPLATE(ComponentAttrSpec);
std::variant<AccessSpec, Allocatable, CoarraySpec, Contiguous,
ComponentArraySpec, Pointer, ErrorRecovery>
ComponentArraySpec, Pointer, common::CUDADataAttr, ErrorRecovery>
u;
};
@ -1337,7 +1344,8 @@ struct IntentSpec {
// CODIMENSION lbracket coarray-spec rbracket | CONTIGUOUS |
// DIMENSION ( array-spec ) | EXTERNAL | INTENT ( intent-spec ) |
// INTRINSIC | language-binding-spec | OPTIONAL | PARAMETER | POINTER |
// PROTECTED | SAVE | TARGET | VALUE | VOLATILE
// PROTECTED | SAVE | TARGET | VALUE | VOLATILE |
// (CUDA) CONSTANT | DEVICE | MANAGED | PINNED | SHARED | TEXTURE
EMPTY_CLASS(Asynchronous);
EMPTY_CLASS(External);
EMPTY_CLASS(Intrinsic);
@ -1352,7 +1360,8 @@ struct AttrSpec {
UNION_CLASS_BOILERPLATE(AttrSpec);
std::variant<AccessSpec, Allocatable, Asynchronous, CoarraySpec, Contiguous,
ArraySpec, External, IntentSpec, Intrinsic, LanguageBindingSpec, Optional,
Parameter, Pointer, Protected, Save, Target, Value, Volatile>
Parameter, Pointer, Protected, Save, Target, Value, Volatile,
common::CUDADataAttr>
u;
};
@ -1926,13 +1935,17 @@ struct StatOrErrmsg {
// R928 alloc-opt ->
// ERRMSG = errmsg-variable | MOLD = source-expr |
// SOURCE = source-expr | STAT = stat-variable
// SOURCE = source-expr | STAT = stat-variable |
// (CUDA) STREAM = scalar-int-expr
// PINNED = scalar-logical-variable
// R931 source-expr -> expr
struct AllocOpt {
UNION_CLASS_BOILERPLATE(AllocOpt);
WRAPPER_CLASS(Mold, common::Indirection<Expr>);
WRAPPER_CLASS(Source, common::Indirection<Expr>);
std::variant<Mold, Source, StatOrErrmsg> u;
WRAPPER_CLASS(Stream, common::Indirection<ScalarIntExpr>);
WRAPPER_CLASS(Pinned, common::Indirection<ScalarLogicalVariable>);
std::variant<Mold, Source, StatOrErrmsg, Stream, Pinned> u;
};
// R927 allocate-stmt ->
@ -3033,7 +3046,9 @@ struct ProcedureDeclarationStmt {
// R1527 prefix-spec ->
// declaration-type-spec | ELEMENTAL | IMPURE | MODULE |
// NON_RECURSIVE | PURE | RECURSIVE
// NON_RECURSIVE | PURE | RECURSIVE |
// (CUDA) ATTRIBUTES ( (DEVICE | GLOBAL | GRID_GLOBAL | HOST)... )
// LAUNCH_BOUNDS(expr-list) | CLUSTER_DIMS(expr-list)
struct PrefixSpec {
UNION_CLASS_BOILERPLATE(PrefixSpec);
EMPTY_CLASS(Elemental);
@ -3042,8 +3057,11 @@ struct PrefixSpec {
EMPTY_CLASS(Non_Recursive);
EMPTY_CLASS(Pure);
EMPTY_CLASS(Recursive);
WRAPPER_CLASS(Attributes, std::list<common::CUDASubprogramAttrs>);
WRAPPER_CLASS(Launch_Bounds, std::list<ScalarIntConstantExpr>);
WRAPPER_CLASS(Cluster_Dims, std::list<ScalarIntConstantExpr>);
std::variant<DeclarationTypeSpec, Elemental, Impure, Module, Non_Recursive,
Pure, Recursive>
Pure, Recursive, Attributes, Launch_Bounds, Cluster_Dims>
u;
};
@ -3172,23 +3190,39 @@ struct ActualArgSpec {
std::tuple<std::optional<Keyword>, ActualArg> t;
};
// R1520 function-reference -> procedure-designator ( [actual-arg-spec-list] )
// R1520 function-reference -> procedure-designator
// ( [actual-arg-spec-list] )
struct Call {
TUPLE_CLASS_BOILERPLATE(Call);
CharBlock source;
std::tuple<ProcedureDesignator, std::list<ActualArgSpec>> t;
};
struct FunctionReference {
WRAPPER_CLASS_BOILERPLATE(FunctionReference, Call);
CharBlock source;
Designator ConvertToArrayElementRef();
StructureConstructor ConvertToStructureConstructor(
const semantics::DerivedTypeSpec &);
};
// R1521 call-stmt -> CALL procedure-designator [( [actual-arg-spec-list] )]
// R1521 call-stmt -> CALL procedure-designator [ chevrons ]
// [( [actual-arg-spec-list] )]
// (CUDA) chevrons -> <<< scalar-expr, scalar-expr [,
// scalar-int-expr [, scalar-int-expr ] ] >>>
struct CallStmt {
WRAPPER_CLASS_BOILERPLATE(CallStmt, Call);
BOILERPLATE(CallStmt);
struct Chevrons {
TUPLE_CLASS_BOILERPLATE(Chevrons);
std::tuple<ScalarExpr, ScalarExpr, std::optional<ScalarIntExpr>,
std::optional<ScalarIntExpr>>
t;
};
explicit CallStmt(ProcedureDesignator &&pd, std::optional<Chevrons> &&ch,
std::list<ActualArgSpec> &&args)
: call{std::move(pd), std::move(args)}, chevrons{std::move(ch)} {}
Call call;
std::optional<Chevrons> chevrons;
CharBlock source;
mutable common::ForwardOwningPointer<evaluate::ProcedureRef>
typedCall; // filled by semantics
};
@ -3267,6 +3301,12 @@ struct CompilerDirective {
std::variant<std::list<IgnoreTKR>, LoopCount, std::list<NameValue>> u;
};
// (CUDA) ATTRIBUTE(attribute) [::] name-list
struct CUDAAttributesStmt {
TUPLE_CLASS_BOILERPLATE(CUDAAttributesStmt);
std::tuple<common::CUDADataAttr, std::list<Name>> t;
};
// Legacy extensions
struct BasedPointer {
TUPLE_CLASS_BOILERPLATE(BasedPointer);
@ -4213,5 +4253,23 @@ struct OpenACCConstruct {
u;
};
// CUF-kernel-do-construct ->
// !$CUF KERNEL DO [ (scalar-int-constant-expr) ] <<< grid, block [, stream]
// >>> do-construct
// grid -> * | scalar-int-expr | ( scalar-int-expr-list )
// block -> * | scalar-int-expr | ( scalar-int-expr-list )
// stream -> 0, scalar-int-expr | STREAM = scalar-int-expr
struct CUFKernelDoConstruct {
TUPLE_CLASS_BOILERPLATE(CUFKernelDoConstruct);
struct Directive {
TUPLE_CLASS_BOILERPLATE(Directive);
CharBlock source;
std::tuple<std::optional<ScalarIntConstantExpr>, std::list<ScalarIntExpr>,
std::list<ScalarIntExpr>, std::optional<ScalarIntExpr>>
t;
};
std::tuple<Directive, std::optional<DoConstruct>> t;
};
} // namespace Fortran::parser
#endif // FORTRAN_PARSER_PARSE_TREE_H_

19
flang/lib/Common/Fortran.cpp
View File

@ -97,4 +97,23 @@ std::string AsFortran(IgnoreTKRSet tkr) {
return result;
}
bool AreCompatibleCUDADataAttrs(std::optional<CUDADataAttr> x,
std::optional<CUDADataAttr> y, IgnoreTKRSet ignoreTKR) {
if (!x && !y) {
return true;
} else if (x && y && *x == *y) {
return true;
} else if (ignoreTKR.test(IgnoreTKR::Device) &&
x.value_or(CUDADataAttr::Device) == CUDADataAttr::Device &&
y.value_or(CUDADataAttr::Device) == CUDADataAttr::Device) {
return true;
} else if (ignoreTKR.test(IgnoreTKR::Managed) &&
x.value_or(CUDADataAttr::Managed) == CUDADataAttr::Managed &&
y.value_or(CUDADataAttr::Managed) == CUDADataAttr::Managed) {
return true;
} else {
return false;
}
}
} // namespace Fortran::common

4
flang/lib/Frontend/FrontendAction.cpp
View File

@ -86,6 +86,10 @@ bool FrontendAction::beginSourceFile(CompilerInstance &ci,
invoc.collectMacroDefinitions();
}
// Enable CUDA Fortran if source file is *.cuf/*.CUF.
invoc.getFortranOpts().features.Enable(Fortran::common::LanguageFeature::CUDA,
getCurrentInput().getIsCUDAFortran());
// Decide between fixed and free form (if the user didn't express any
// preference, use the file extension to decide)
if (invoc.getFrontendOpts().fortranForm == FortranForm::Unknown) {

11
flang/lib/Frontend/FrontendOptions.cpp
View File

@ -23,17 +23,22 @@ bool Fortran::frontend::isFixedFormSuffix(llvm::StringRef suffix) {
bool Fortran::frontend::isFreeFormSuffix(llvm::StringRef suffix) {
// Note: Keep this list in-sync with flang/test/lit.cfg.py
// TODO: Add Cuda Fortan files (i.e. `*.cuf` and `*.CUF`).
return suffix == "f90" || suffix == "F90" || suffix == "ff90" ||
suffix == "f95" || suffix == "F95" || suffix == "ff95" ||
suffix == "f03" || suffix == "F03" || suffix == "f08" ||
suffix == "F08" || suffix == "f18" || suffix == "F18";
suffix == "F08" || suffix == "f18" || suffix == "F18" ||
suffix == "cuf" || suffix == "CUF";
}
bool Fortran::frontend::isToBePreprocessed(llvm::StringRef suffix) {
return suffix == "F" || suffix == "FOR" || suffix == "fpp" ||
suffix == "FPP" || suffix == "F90" || suffix == "F95" ||
suffix == "F03" || suffix == "F08" || suffix == "F18";
suffix == "F03" || suffix == "F08" || suffix == "F18" ||
suffix == "CUF";
}
bool Fortran::frontend::isCUDAFortranSuffix(llvm::StringRef suffix) {
return suffix == "cuf" || suffix == "CUF";
}
InputKind FrontendOptions::getInputKindForExtension(llvm::StringRef extension) {

6
flang/lib/Lower/Allocatable.cpp
View File

@ -367,6 +367,12 @@ private:
[&](const Fortran::parser::AllocOpt::Mold &mold) {
moldExpr = Fortran::semantics::GetExpr(mold.v.value());
},
[&](const Fortran::parser::AllocOpt::Stream &) {
TODO(loc, "CUDA ALLOCATE(STREAM=)");
},
[&](const Fortran::parser::AllocOpt::Pinned &) {
TODO(loc, "CUDA ALLOCATE(PINNED=)");
},
},
allocOption.u);
}

4
flang/lib/Lower/Bridge.cpp
View File

@ -1321,7 +1321,7 @@ private:
void genFIR(const Fortran::parser::CallStmt &stmt) {
Fortran::lower::StatementContext stmtCtx;
Fortran::lower::pft::Evaluation &eval = getEval();
setCurrentPosition(stmt.v.source);
setCurrentPosition(stmt.source);
assert(stmt.typedCall && "Call was not analyzed");
mlir::Value res{};
if (lowerToHighLevelFIR()) {
@ -1348,7 +1348,7 @@ private:
llvm::SmallVector<Fortran::parser::Label> labelList;
int64_t index = 0;
for (const Fortran::parser::ActualArgSpec &arg :
std::get<std::list<Fortran::parser::ActualArgSpec>>(stmt.v.t)) {
std::get<std::list<Fortran::parser::ActualArgSpec>>(stmt.call.t)) {
const auto &actual = std::get<Fortran::parser::ActualArg>(arg.t);
if (const auto *altReturn =
std::get_if<Fortran::parser::AltReturnSpec>(&actual.u)) {

2
flang/lib/Lower/PFTBuilder.cpp
View File

@ -726,7 +726,7 @@ private:
[&](const parser::CallStmt &s) {
// Look for alternate return specifiers.
const auto &args =
std::get<std::list<parser::ActualArgSpec>>(s.v.t);
std::get<std::list<parser::ActualArgSpec>>(s.call.t);
for (const auto &arg : args) {
const auto &actual = std::get<parser::ActualArg>(arg.t);
if (const auto *altReturn =

39
flang/lib/Parser/Fortran-parsers.cpp
View File

@ -451,13 +451,16 @@ TYPE_PARSER(construct<DataComponentDefStmt>(declarationTypeSpec,
// R738 component-attr-spec ->
// access-spec | ALLOCATABLE |
// CODIMENSION lbracket coarray-spec rbracket |
// CONTIGUOUS | DIMENSION ( component-array-spec ) | POINTER
// CONTIGUOUS | DIMENSION ( component-array-spec ) | POINTER |
// CUDA-data-attr
TYPE_PARSER(construct<ComponentAttrSpec>(accessSpec) ||
construct<ComponentAttrSpec>(allocatable) ||
construct<ComponentAttrSpec>("CODIMENSION" >> coarraySpec) ||
construct<ComponentAttrSpec>(contiguous) ||
construct<ComponentAttrSpec>("DIMENSION" >> Parser<ComponentArraySpec>{}) ||
construct<ComponentAttrSpec>(pointer) ||
extension<LanguageFeature::CUDA>(
construct<ComponentAttrSpec>(Parser<common::CUDADataAttr>{})) ||
construct<ComponentAttrSpec>(recovery(
fail<ErrorRecovery>(
"type parameter definitions must appear before component declarations"_err_en_US),
@ -677,7 +680,8 @@ TYPE_PARSER(
// CODIMENSION lbracket coarray-spec rbracket | CONTIGUOUS |
// DIMENSION ( array-spec ) | EXTERNAL | INTENT ( intent-spec ) |
// INTRINSIC | language-binding-spec | OPTIONAL | PARAMETER | POINTER |
// PROTECTED | SAVE | TARGET | VALUE | VOLATILE
// PROTECTED | SAVE | TARGET | VALUE | VOLATILE |
// CUDA-data-attr
TYPE_PARSER(construct<AttrSpec>(accessSpec) ||
construct<AttrSpec>(allocatable) ||
construct<AttrSpec>(construct<Asynchronous>("ASYNCHRONOUS"_tok)) ||
@ -693,7 +697,17 @@ TYPE_PARSER(construct<AttrSpec>(accessSpec) ||
construct<AttrSpec>(save) ||
construct<AttrSpec>(construct<Target>("TARGET"_tok)) ||
construct<AttrSpec>(construct<Value>("VALUE"_tok)) ||
construct<AttrSpec>(construct<Volatile>("VOLATILE"_tok)))
construct<AttrSpec>(construct<Volatile>("VOLATILE"_tok)) ||
extension<LanguageFeature::CUDA>(
construct<AttrSpec>(Parser<common::CUDADataAttr>{})))
// CUDA-data-attr -> CONSTANT | DEVICE | MANAGED | PINNED | SHARED | TEXTURE
TYPE_PARSER("CONSTANT" >> pure(common::CUDADataAttr::Constant) ||
"DEVICE" >> pure(common::CUDADataAttr::Device) ||
"MANAGED" >> pure(common::CUDADataAttr::Managed) ||
"PINNED" >> pure(common::CUDADataAttr::Pinned) ||
"SHARED" >> pure(common::CUDADataAttr::Shared) ||
"TEXTURE" >> pure(common::CUDADataAttr::Texture))
// R804 object-name -> name
constexpr auto objectName{name};
@ -1181,13 +1195,20 @@ TYPE_CONTEXT_PARSER("ALLOCATE statement"_en_US,
// R928 alloc-opt ->
// ERRMSG = errmsg-variable | MOLD = source-expr |
// SOURCE = source-expr | STAT = stat-variable
// SOURCE = source-expr | STAT = stat-variable |
// (CUDA) STREAM = scalar-int-expr
// PINNED = scalar-logical-variable
// R931 source-expr -> expr
TYPE_PARSER(construct<AllocOpt>(
construct<AllocOpt::Mold>("MOLD =" >> indirect(expr))) ||
construct<AllocOpt>(
construct<AllocOpt::Source>("SOURCE =" >> indirect(expr))) ||
construct<AllocOpt>(statOrErrmsg))
construct<AllocOpt>(statOrErrmsg) ||
extension<LanguageFeature::CUDA>(
construct<AllocOpt>(construct<AllocOpt::Stream>(
"STREAM =" >> indirect(scalarIntExpr))) ||
construct<AllocOpt>(construct<AllocOpt::Pinned>(
"PINNED =" >> indirect(scalarLogicalVariable)))))
// R929 stat-variable -> scalar-int-variable
TYPE_PARSER(construct<StatVariable>(scalar(integer(variable))))
@ -1239,14 +1260,12 @@ TYPE_PARSER(construct<StatOrErrmsg>("STAT =" >> statVariable) ||
// !DIR$ IGNORE_TKR [ [(tkrdmac...)] name ]...
// !DIR$ LOOP COUNT (n1[, n2]...)
// !DIR$ name...
constexpr auto beginDirective{skipStuffBeforeStatement >> "!"_ch};
constexpr auto ignore_tkr{
"DIR$ IGNORE_TKR" >> optionalList(construct<CompilerDirective::IgnoreTKR>(
maybe(parenthesized(many(letter))), name))};
constexpr auto loopCount{
"DIR$ LOOP COUNT" >> construct<CompilerDirective::LoopCount>(
parenthesized(nonemptyList(digitString64)))};
TYPE_PARSER(beginDirective >>
sourced(construct<CompilerDirective>(ignore_tkr) ||
construct<CompilerDirective>(loopCount) ||
@ -1262,6 +1281,12 @@ TYPE_PARSER(extension<LanguageFeature::CrayPointer>(
construct<BasedPointer>("(" >> objectName / ",",
objectName, maybe(Parser<ArraySpec>{}) / ")")))))
// CUDA-attributes-stmt -> ATTRIBUTES (CUDA-data-attr) [::] name-list
TYPE_PARSER(extension<LanguageFeature::CUDA>(construct<CUDAAttributesStmt>(
"ATTRIBUTES" >> parenthesized(Parser<common::CUDADataAttr>{}),
defaulted(
maybe("::"_tok) >> nonemptyList("expected names"_err_en_US, name)))))
// Subtle: the name includes the surrounding slashes, which avoids
// clashes with other uses of the name in the same scope.
TYPE_PARSER(construct<StructureStmt>(

8
flang/lib/Parser/basic-parsers.h
View File

@ -852,6 +852,7 @@ public:
constexpr NonstandardParser(const NonstandardParser &) = default;
constexpr NonstandardParser(PA parser, MessageFixedText msg)
: parser_{parser}, message_{msg} {}
constexpr NonstandardParser(PA parser) : parser_{parser} {}
std::optional<resultType> Parse(ParseState &state) const {
if (UserState * ustate{state.userState()}) {
if (!ustate->features().IsEnabled(LF)) {
@ -860,7 +861,7 @@ public:
}
auto at{state.GetLocation()};
auto result{parser_.Parse(state)};
if (result) {
if (result && !message_.empty()) {
state.Nonstandard(
CharBlock{at, std::max(state.GetLocation(), at + 1)}, LF, message_);
}
@ -877,6 +878,11 @@ inline constexpr auto extension(MessageFixedText feature, PA parser) {
return NonstandardParser<LF, PA>(parser, feature);
}
template <LanguageFeature LF, typename PA>
inline constexpr auto extension(PA parser) {
return NonstandardParser<LF, PA>(parser);
}
// If a is a parser for some deprecated or deleted language feature LF,
// deprecated<LF>(a) is a parser that is optionally enabled, sets a strict
// conformance violation flag, and may emit a warning message, if enabled.

73
flang/lib/Parser/executable-parsers.cpp
View File

@ -9,6 +9,7 @@
// Per-type parsers for executable statements
#include "basic-parsers.h"
#include "debug-parser.h"
#include "expr-parsers.h"
#include "misc-parsers.h"
#include "stmt-parser.h"
@ -30,29 +31,31 @@ namespace Fortran::parser {
// action-stmt | associate-construct | block-construct |
// case-construct | change-team-construct | critical-construct |
// do-construct | if-construct | select-rank-construct |
// select-type-construct | where-construct | forall-construct
constexpr auto executableConstruct{
first(construct<ExecutableConstruct>(CapturedLabelDoStmt{}),
construct<ExecutableConstruct>(EndDoStmtForCapturedLabelDoStmt{}),
construct<ExecutableConstruct>(indirect(Parser<DoConstruct>{})),
// Attempt DO statements before assignment statements for better
// error messages in cases like "DO10I=1,(error)".
construct<ExecutableConstruct>(statement(actionStmt)),
construct<ExecutableConstruct>(indirect(Parser<AssociateConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<BlockConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<CaseConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<ChangeTeamConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<CriticalConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<IfConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<SelectRankConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<SelectTypeConstruct>{})),
construct<ExecutableConstruct>(indirect(whereConstruct)),
construct<ExecutableConstruct>(indirect(forallConstruct)),
construct<ExecutableConstruct>(indirect(ompEndLoopDirective)),
construct<ExecutableConstruct>(indirect(openmpConstruct)),
construct<ExecutableConstruct>(indirect(accEndCombinedDirective)),
construct<ExecutableConstruct>(indirect(openaccConstruct)),
construct<ExecutableConstruct>(indirect(compilerDirective)))};
// select-type-construct | where-construct | forall-construct |
// (CUDA) CUF-kernel-do-construct
constexpr auto executableConstruct{first(
construct<ExecutableConstruct>(CapturedLabelDoStmt{}),
construct<ExecutableConstruct>(EndDoStmtForCapturedLabelDoStmt{}),
construct<ExecutableConstruct>(indirect(Parser<DoConstruct>{})),
// Attempt DO statements before assignment statements for better
// error messages in cases like "DO10I=1,(error)".
construct<ExecutableConstruct>(statement(actionStmt)),
construct<ExecutableConstruct>(indirect(Parser<AssociateConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<BlockConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<CaseConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<ChangeTeamConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<CriticalConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<IfConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<SelectRankConstruct>{})),
construct<ExecutableConstruct>(indirect(Parser<SelectTypeConstruct>{})),
construct<ExecutableConstruct>(indirect(whereConstruct)),
construct<ExecutableConstruct>(indirect(forallConstruct)),
construct<ExecutableConstruct>(indirect(ompEndLoopDirective)),
construct<ExecutableConstruct>(indirect(openmpConstruct)),
construct<ExecutableConstruct>(indirect(accEndCombinedDirective)),
construct<ExecutableConstruct>(indirect(openaccConstruct)),
construct<ExecutableConstruct>(indirect(compilerDirective)),
construct<ExecutableConstruct>(indirect(Parser<CUFKernelDoConstruct>{})))};
// R510 execution-part-construct ->
// executable-construct | format-stmt | entry-stmt | data-stmt
@ -525,4 +528,28 @@ TYPE_CONTEXT_PARSER("UNLOCK statement"_en_US,
construct<UnlockStmt>("UNLOCK (" >> lockVariable,
defaulted("," >> nonemptyList(statOrErrmsg)) / ")"))
// CUF-kernel-do-construct -> CUF-kernel-do-directive do-construct
// CUF-kernel-do-directive ->
// !$CUF KERNEL DO [ (scalar-int-constant-expr) ] <<< grid, block [, stream]
// >>> do-construct
// grid -> * | scalar-int-expr | ( scalar-int-expr-list )
// block -> * | scalar-int-expr | ( scalar-int-expr-list )
// stream -> ( 0, | STREAM = ) scalar-int-expr
TYPE_PARSER(sourced(beginDirective >> "$CUF KERNEL DO"_tok >>
construct<CUFKernelDoConstruct::Directive>(
maybe(parenthesized(scalarIntConstantExpr)),
"<<<" >>
("*" >> pure<std::list<ScalarIntExpr>>() ||
parenthesized(nonemptyList(scalarIntExpr)) ||
applyFunction(singletonList<ScalarIntExpr>, scalarIntExpr)),
"," >> ("*" >> pure<std::list<ScalarIntExpr>>() ||
parenthesized(nonemptyList(scalarIntExpr)) ||
applyFunction(singletonList<ScalarIntExpr>, scalarIntExpr)),
maybe((", 0 ,"_tok || ", STREAM ="_tok) >> scalarIntExpr) / ">>>" /
endDirective)))
TYPE_CONTEXT_PARSER("!$CUF KERNEL DO construct"_en_US,
extension<LanguageFeature::CUDA>(construct<CUFKernelDoConstruct>(
Parser<CUFKernelDoConstruct::Directive>{},
maybe(Parser<DoConstruct>{}))))
} // namespace Fortran::parser

5
flang/lib/Parser/io-parsers.cpp
View File

@ -301,11 +301,6 @@ TYPE_PARSER(first(construct<WaitSpec>(maybe("UNIT ="_tok) >> fileUnitNumber),
construct<WaitSpec>("IOMSG =" >> msgVariable),
construct<WaitSpec>("IOSTAT =" >> statVariable)))
template <typename A> common::IfNoLvalue<std::list<A>, A> singletonList(A &&x) {
std::list<A> result;
result.push_front(std::move(x));
return result;
}
constexpr auto bareUnitNumberAsList{
applyFunction(singletonList<PositionOrFlushSpec>,
construct<PositionOrFlushSpec>(fileUnitNumber))};

5
flang/lib/Parser/misc-parsers.h
View File

@ -52,5 +52,10 @@ constexpr auto pointer{construct<Pointer>("POINTER"_tok)};
constexpr auto protectedAttr{construct<Protected>("PROTECTED"_tok)};
constexpr auto save{construct<Save>("SAVE"_tok)};
template <typename A> common::IfNoLvalue<std::list<A>, A> singletonList(A &&x) {
std::list<A> result;
result.emplace_back(std::move(x));
return result;
}
} // namespace Fortran::parser
#endif

13
flang/lib/Parser/parse-tree.cpp
View File

@ -132,7 +132,7 @@ static Expr ActualArgToExpr(ActualArgSpec &arg) {
},
[&](common::Indirection<FunctionReference> &z) {
return WithSource(
z.value().v.source, Expr{std::move(z.value())});
z.value().source, Expr{std::move(z.value())});
},
},
y.value().u);
@ -151,10 +151,10 @@ Designator FunctionReference::ConvertToArrayElementRef() {
common::visitors{
[&](const Name &name) {
return WithSource(
v.source, MakeArrayElementRef(name, std::move(args)));
source, MakeArrayElementRef(name, std::move(args)));
},
[&](ProcComponentRef &pcr) {
return WithSource(v.source,
return WithSource(source,
MakeArrayElementRef(std::move(pcr.v.thing), std::move(args)));
},
},
@ -226,9 +226,10 @@ Statement<ActionStmt> StmtFunctionStmt::ConvertToAssignment() {
}
CHECK(*source.end() == ')');
source = CharBlock{source.begin(), source.end() + 1};
FunctionReference funcRef{WithSource(source,
FunctionReference funcRef{
Call{ProcedureDesignator{Name{funcName.source, funcName.symbol}},
std::move(actuals)})};
std::move(actuals)}};
funcRef.source = source;
auto variable{Variable{common::Indirection{std::move(funcRef)}}};
return Statement{std::nullopt,
ActionStmt{common::Indirection{
@ -242,7 +243,7 @@ CharBlock Variable::GetSource() const {
return des.value().source;
},
[&](const common::Indirection<parser::FunctionReference> &call) {
return call.value().v.source;
return call.value().source;
},
},
u);

5
flang/lib/Parser/parsing.cpp
View File

@ -84,6 +84,11 @@ const SourceFile *Parsing::Prescan(const std::string &path, Options options) {
prescanner.AddCompilerDirectiveSentinel("$omp");
prescanner.AddCompilerDirectiveSentinel("$"); // OMP conditional line
}
if (options.features.IsEnabled(LanguageFeature::CUDA)) {
prescanner.AddCompilerDirectiveSentinel("$cuf");
prescanner.AddCompilerDirectiveSentinel("@cuf");
preprocessor.Define("_CUDA", "1");
}
ProvenanceRange range{allSources.AddIncludedFile(
*sourceFile, ProvenanceRange{}, options.isModuleFile)};
prescanner.Prescan(range);

9
flang/lib/Parser/preprocessor.cpp
View File

@ -622,11 +622,12 @@ void Preprocessor::Directive(const TokenSequence &dir, Prescanner &prescanner) {
TokenSequence braced{dir, j + 1, k - j - 1};
include = braced.ToString();
j = k;
} else if ((include = dir.TokenAt(j).ToString()).substr(0, 1) == "\"" &&
include.substr(include.size() - 1, 1) == "\"") { // #include "foo"
} else if (((include = dir.TokenAt(j).ToString()).substr(0, 1) == "\"" ||
include.substr(0, 1) == "'") &&
include.substr(include.size() - 1, 1) == include.substr(0, 1)) {
// #include "foo" and #include 'foo'
include = include.substr(1, include.size() - 2);
// #include "foo" starts search in directory of file containing
// the directive
// Start search in directory of file containing the directive
auto prov{dir.GetTokenProvenanceRange(dirOffset).start()};
if (const auto *currentFile{allSources_.GetSourceFile(prov)}) {
prependPath = DirectoryName(currentFile->path());

11
flang/lib/Parser/prescan.cpp
View File

@ -127,6 +127,17 @@ void Prescanner::Statement() {
} else {
SkipSpaces();
}
} else if (directiveSentinel_[0] == '@' && directiveSentinel_[1] == 'c' &&
directiveSentinel_[2] == 'u' && directiveSentinel_[3] == 'f' &&
directiveSentinel_[4] == '\0') {
// CUDA conditional compilation line. Remove the sentinel and then
// treat the line as if it were normal source.
at_ += 5, column_ += 5;
if (inFixedForm_) {
LabelField(tokens);
} else {
SkipSpaces();
}
} else {
// Compiler directive. Emit normalized sentinel.
EmitChar(tokens, '!');

55
flang/lib/Parser/program-parsers.cpp
View File

@ -81,10 +81,10 @@ TYPE_CONTEXT_PARSER("specification part"_en_US,
// are in contexts that impose constraints on the kinds of statements that
// are allowed, and so we have a variant production for declaration-construct
// that implements those constraints.
constexpr auto execPartLookAhead{
first(actionStmt >> ok, openaccConstruct >> ok, openmpConstruct >> ok,
"ASSOCIATE ("_tok, "BLOCK"_tok, "SELECT"_tok, "CHANGE TEAM"_sptok,
"CRITICAL"_tok, "DO"_tok, "IF ("_tok, "WHERE ("_tok, "FORALL ("_tok)};
constexpr auto execPartLookAhead{first(actionStmt >> ok, openaccConstruct >> ok,
openmpConstruct >> ok, "ASSOCIATE ("_tok, "BLOCK"_tok, "SELECT"_tok,
"CHANGE TEAM"_sptok, "CRITICAL"_tok, "DO"_tok, "IF ("_tok, "WHERE ("_tok,
"FORALL ("_tok, "!$CUF"_tok)};
constexpr auto declErrorRecovery{
stmtErrorRecoveryStart >> !execPartLookAhead >> skipStmtErrorRecovery};
constexpr auto misplacedSpecificationStmt{Parser<UseStmt>{} >>
@ -168,7 +168,8 @@ TYPE_CONTEXT_PARSER("specification construct"_en_US,
// codimension-stmt | contiguous-stmt | dimension-stmt | external-stmt |
// intent-stmt | intrinsic-stmt | namelist-stmt | optional-stmt |
// pointer-stmt | protected-stmt | save-stmt | target-stmt |
// volatile-stmt | value-stmt | common-stmt | equivalence-stmt
// volatile-stmt | value-stmt | common-stmt | equivalence-stmt |
// (CUDA) CUDA-attributes-stmt
TYPE_PARSER(first(
construct<OtherSpecificationStmt>(indirect(Parser<AccessStmt>{})),
construct<OtherSpecificationStmt>(indirect(Parser<AllocatableStmt>{})),
@ -190,7 +191,8 @@ TYPE_PARSER(first(
construct<OtherSpecificationStmt>(indirect(Parser<VolatileStmt>{})),
construct<OtherSpecificationStmt>(indirect(Parser<CommonStmt>{})),
construct<OtherSpecificationStmt>(indirect(Parser<EquivalenceStmt>{})),
construct<OtherSpecificationStmt>(indirect(Parser<BasedPointerStmt>{}))))
construct<OtherSpecificationStmt>(indirect(Parser<BasedPointerStmt>{})),
construct<OtherSpecificationStmt>(indirect(Parser<CUDAAttributesStmt>{}))))
// R1401 main-program ->
// [program-stmt] [specification-part] [execution-part]
@ -422,16 +424,25 @@ TYPE_PARSER(
TYPE_PARSER(
"INTRINSIC" >> maybe("::"_tok) >> construct<IntrinsicStmt>(listOfNames))
// R1520 function-reference -> procedure-designator ( [actual-arg-spec-list] )
// R1520 function-reference -> procedure-designator
// ( [actual-arg-spec-list] )
TYPE_CONTEXT_PARSER("function reference"_en_US,
construct<FunctionReference>(
sourced(construct<Call>(Parser<ProcedureDesignator>{},
sourced(construct<FunctionReference>(
construct<Call>(Parser<ProcedureDesignator>{},
parenthesized(optionalList(actualArgSpec))))) /
!"["_tok)
// R1521 call-stmt -> CALL procedure-designator [( [actual-arg-spec-list] )]
// R1521 call-stmt -> CALL procedure-designator [chevrons]
/// [( [actual-arg-spec-list] )]
// (CUDA) chevrons -> <<< scalar-expr, scalar-expr [, scalar-int-expr
// [, scalar-int-expr ] ] >>>
TYPE_PARSER(extension<LanguageFeature::CUDA>(
"<<<" >> construct<CallStmt::Chevrons>(scalarExpr, "," >> scalarExpr,
maybe("," >> scalarIntExpr), maybe("," >> scalarIntExpr)) /
">>>"))
TYPE_PARSER(construct<CallStmt>(
sourced(construct<Call>("CALL" >> Parser<ProcedureDesignator>{},
sourced(construct<CallStmt>("CALL" >> Parser<ProcedureDesignator>{},
maybe(Parser<CallStmt::Chevrons>{}),
defaulted(parenthesized(optionalList(actualArgSpec)))))))
// R1522 procedure-designator ->
@ -467,7 +478,13 @@ TYPE_PARSER(construct<AltReturnSpec>(star >> label))
// R1527 prefix-spec ->
// declaration-type-spec | ELEMENTAL | IMPURE | MODULE |
// NON_RECURSIVE | PURE | RECURSIVE
// NON_RECURSIVE | PURE | RECURSIVE |
// (CUDA) ATTRIBUTES ( (DEVICE | GLOBAL | GRID_GLOBAL | HOST)... ) |
// LAUNCH_BOUNDS(expr-list) | CLUSTER_DIMS(expr-list)
TYPE_PARSER(first("DEVICE" >> pure(common::CUDASubprogramAttrs::Device),
"GLOBAL" >> pure(common::CUDASubprogramAttrs::Global),
"GRID_GLOBAL" >> pure(common::CUDASubprogramAttrs::Grid_Global),
"HOST" >> pure(common::CUDASubprogramAttrs::Host)))
TYPE_PARSER(first(construct<PrefixSpec>(declarationTypeSpec),
construct<PrefixSpec>(construct<PrefixSpec::Elemental>("ELEMENTAL"_tok)),
construct<PrefixSpec>(construct<PrefixSpec::Impure>("IMPURE"_tok)),
@ -475,7 +492,19 @@ TYPE_PARSER(first(construct<PrefixSpec>(declarationTypeSpec),
construct<PrefixSpec>(
construct<PrefixSpec::Non_Recursive>("NON_RECURSIVE"_tok)),
construct<PrefixSpec>(construct<PrefixSpec::Pure>("PURE"_tok)),
construct<PrefixSpec>(construct<PrefixSpec::Recursive>("RECURSIVE"_tok))))
construct<PrefixSpec>(construct<PrefixSpec::Recursive>("RECURSIVE"_tok)),
extension<LanguageFeature::CUDA>(
construct<PrefixSpec>(construct<PrefixSpec::Attributes>("ATTRIBUTES" >>
parenthesized(
optionalList(Parser<common::CUDASubprogramAttrs>{}))))),
extension<LanguageFeature::CUDA>(construct<PrefixSpec>(
construct<PrefixSpec::Launch_Bounds>("LAUNCH_BOUNDS" >>
parenthesized(nonemptyList(
"expected launch bounds"_err_en_US, scalarIntConstantExpr))))),
extension<LanguageFeature::CUDA>(construct<PrefixSpec>(
construct<PrefixSpec::Cluster_Dims>("CLUSTER_DIMS" >>
parenthesized(nonemptyList("expected cluster dimensions"_err_en_US,
scalarIntConstantExpr)))))))
// R1529 function-subprogram ->
// function-stmt [specification-part] [execution-part]

4
flang/lib/Parser/stmt-parser.h
View File

@ -105,5 +105,9 @@ constexpr auto progUnitEndStmtErrorRecovery{
(many(!"END"_tok >> SkipPast<'\n'>{}) >>
("END"_tok >> SkipTo<'\n'>{} || consumedAllInput)) >>
missingOptionalName};
constexpr auto beginDirective{skipStuffBeforeStatement >> "!"_ch};
constexpr auto endDirective{space >> endOfLine};
} // namespace Fortran::parser
#endif // FORTRAN_PARSER_STMT_PARSER_H_

70
flang/lib/Parser/unparse.cpp
View File

@ -820,6 +820,8 @@ public:
common::visit(common::visitors{
[&](const AllocOpt::Mold &) { Word("MOLD="); },
[&](const AllocOpt::Source &) { Word("SOURCE="); },
[&](const AllocOpt::Stream &) { Word("STREAM="); },
[&](const AllocOpt::Pinned &) { Word("PINNED="); },
[](const StatOrErrmsg &) {},
},
x.u);
@ -1685,19 +1687,26 @@ public:
void Unparse(const IntrinsicStmt &x) { // R1519
Word("INTRINSIC :: "), Walk(x.v, ", ");
}
void Unparse(const CallStmt::Chevrons &x) { // CUDA
Walk(std::get<0>(x.t)); // grid
Word(","), Walk(std::get<1>(x.t)); // block
Walk(",", std::get<2>(x.t)); // bytes
Walk(",", std::get<3>(x.t)); // stream
}
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.v.t)};
const auto &args{std::get<std::list<ActualArgSpec>>(x.v.t)};
const auto &pd{std::get<ProcedureDesignator>(x.call.t)};
Word("CALL "), Walk(pd);
Walk("<<<", x.chevrons, ">>>");
const auto &args{std::get<std::list<ActualArgSpec>>(x.call.t)};
if (args.empty()) {
if (std::holds_alternative<ProcComponentRef>(pd.u)) {
Put("()"); // pgf90 crashes on CALL to tbp without parentheses
@ -1726,6 +1735,15 @@ public:
void Post(const PrefixSpec::Non_Recursive) { Word("NON_RECURSIVE"); }
void Post(const PrefixSpec::Pure) { Word("PURE"); }
void Post(const PrefixSpec::Recursive) { Word("RECURSIVE"); }
void Unparse(const PrefixSpec::Attributes &x) {
Word("ATTRIBUTES("), Walk(x.v), Word(")");
}
void Unparse(const PrefixSpec::Launch_Bounds &x) {
Word("LAUNCH_BOUNDS("), Walk(x.v), Word(")");
}
void Unparse(const PrefixSpec::Cluster_Dims &x) {
Word("CLUSTER_DIMS("), Walk(x.v), Word(")");
}
void Unparse(const FunctionStmt &x) { // R1530
Walk("", std::get<std::list<PrefixSpec>>(x.t), " ", " ");
Word("FUNCTION "), Walk(std::get<Name>(x.t)), Put("(");
@ -1870,9 +1888,6 @@ public:
Walk(std::get<std::optional<AccDataModifier>>(x.t), ":");
Walk(std::get<AccObjectList>(x.t));
}
void Unparse(const AccDataModifier::Modifier &x) {
Word(AccDataModifier::EnumToString(x));
}
void Unparse(const AccBindClause &x) {
common::visit(common::visitors{
[&](const Name &y) { Put('('), Walk(y), Put(')'); },
@ -1966,9 +1981,6 @@ public:
x.u);
}
void Unparse(const AccObjectList &x) { Walk(x.v, ","); }
void Unparse(const AccReductionOperator::Operator &x) {
Word(AccReductionOperator::EnumToString(x));
}
void Unparse(const AccObjectListWithReduction &x) {
Walk(std::get<AccReductionOperator>(x.t));
Put(":");
@ -2613,6 +2625,10 @@ public:
Walk("(", std::get<std::optional<ArraySpec>>(x.t), ")"), Put(')');
}
void Unparse(const BasedPointerStmt &x) { Walk("POINTER ", x.v, ","); }
void Unparse(const CUDAAttributesStmt &x) {
Word("ATTRIBUTES("), Walk(std::get<common::CUDADataAttr>(x.t));
Word(") "), Walk(std::get<std::list<Name>>(x.t), ", ");
}
void Post(const StructureField &x) {
if (const auto *def{std::get_if<Statement<DataComponentDefStmt>>(&x.u)}) {
for (const auto &item :
@ -2658,8 +2674,12 @@ public:
#define WALK_NESTED_ENUM(CLASS, ENUM) \
void Unparse(const CLASS::ENUM &x) { Word(CLASS::EnumToString(x)); }
WALK_NESTED_ENUM(AccDataModifier, Modifier)
WALK_NESTED_ENUM(AccessSpec, Kind) // R807
WALK_NESTED_ENUM(AccReductionOperator, Operator)
WALK_NESTED_ENUM(common, TypeParamAttr) // R734
WALK_NESTED_ENUM(common, CUDADataAttr) // CUDA
WALK_NESTED_ENUM(common, CUDASubprogramAttrs) // CUDA
WALK_NESTED_ENUM(IntentSpec, Intent) // R826
WALK_NESTED_ENUM(ImplicitStmt, ImplicitNoneNameSpec) // R866
WALK_NESTED_ENUM(ConnectSpec::CharExpr, Kind) // R1205
@ -2686,6 +2706,38 @@ public:
WALK_NESTED_ENUM(OmpOrderModifier, Kind) // OMP order-modifier
#undef WALK_NESTED_ENUM
void Unparse(const CUFKernelDoConstruct::Directive &x) {
Word("!$CUF KERNEL DO");
Walk(" (", std::get<std::optional<ScalarIntConstantExpr>>(x.t), ")");
Word(" <<<");
const auto &grid{std::get<1>(x.t)};
if (grid.empty()) {
Word("*");
} else if (grid.size() == 1) {
Walk(grid.front());
} else {
Walk("(", grid, ",", ")");
}
Word(",");
const auto &block{std::get<2>(x.t)};
if (block.empty()) {
Word("*");
} else if (block.size() == 1) {
Walk(block.front());
} else {
Walk("(", block, ",", ")");
}
if (const auto &stream{std::get<3>(x.t)}) {
Word(",STREAM="), Walk(*stream);
}
Word(">>>\n");
}
void Unparse(const CUFKernelDoConstruct &x) {
Walk(std::get<CUFKernelDoConstruct::Directive>(x.t));
Walk(std::get<std::optional<DoConstruct>>(x.t));
}
void Done() const { CHECK(indent_ == 0); }
private:

2
flang/lib/Semantics/check-allocate.cpp
View File

@ -179,6 +179,8 @@ static std::optional<AllocateCheckerInfo> CheckAllocateOptions(
parserSourceExpr = &mold.v.value();
info.gotMold = true;
},
[](const parser::AllocOpt::Stream &) { /* CUDA coming */ },
[](const parser::AllocOpt::Pinned &) { /* CUDA coming */ },
},
allocOpt.u);
}

2
flang/lib/Semantics/check-do-forall.cpp
View File

@ -975,7 +975,7 @@ static void CheckIfArgIsDoVar(const evaluate::ActualArgument &arg,
void DoForallChecker::Leave(const parser::CallStmt &callStmt) {
if (const auto &typedCall{callStmt.typedCall}) {
const auto &parsedArgs{
std::get<std::list<parser::ActualArgSpec>>(callStmt.v.t)};
std::get<std::list<parser::ActualArgSpec>>(callStmt.call.t)};
auto parsedArgIter{parsedArgs.begin()};
const evaluate::ActualArguments &checkedArgs{typedCall->arguments()};
for (const auto &checkedOptionalArg : checkedArgs) {

18
flang/lib/Semantics/expression.cpp
View File

@ -2722,8 +2722,8 @@ bool ExpressionAnalyzer::CheckIsValidForwardReference(
MaybeExpr ExpressionAnalyzer::Analyze(const parser::FunctionReference &funcRef,
std::optional<parser::StructureConstructor> *structureConstructor) {
const parser::Call &call{funcRef.v};
auto restorer{GetContextualMessages().SetLocation(call.source)};
ArgumentAnalyzer analyzer{*this, call.source, true /* isProcedureCall */};
auto restorer{GetContextualMessages().SetLocation(funcRef.source)};
ArgumentAnalyzer analyzer{*this, funcRef.source, true /* isProcedureCall */};
for (const auto &arg : std::get<std::list<parser::ActualArgSpec>>(call.t)) {
analyzer.Analyze(arg, false /* not subroutine call */);
}
@ -2736,7 +2736,7 @@ MaybeExpr ExpressionAnalyzer::Analyze(const parser::FunctionReference &funcRef,
true /* might be structure constructor */)}) {
if (auto *proc{std::get_if<ProcedureDesignator>(&callee->u)}) {
return MakeFunctionRef(
call.source, std::move(*proc), std::move(callee->arguments));
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)};
@ -2778,9 +2778,9 @@ static bool HasAlternateReturns(const evaluate::ActualArguments &args) {
}
void ExpressionAnalyzer::Analyze(const parser::CallStmt &callStmt) {
const parser::Call &call{callStmt.v};
auto restorer{GetContextualMessages().SetLocation(call.source)};
ArgumentAnalyzer analyzer{*this, call.source, true /* isProcedureCall */};
const parser::Call &call{callStmt.call};
auto restorer{GetContextualMessages().SetLocation(callStmt.source)};
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 */);
@ -2791,7 +2791,7 @@ void ExpressionAnalyzer::Analyze(const parser::CallStmt &callStmt) {
analyzer.GetActuals(), true /* subroutine */)}) {
ProcedureDesignator *proc{std::get_if<ProcedureDesignator>(&callee->u)};
CHECK(proc);
if (CheckCall(call.source, *proc, callee->arguments)) {
if (CheckCall(callStmt.source, *proc, callee->arguments)) {
callStmt.typedCall.Reset(
new ProcedureRef{std::move(*proc), std::move(callee->arguments),
HasAlternateReturns(callee->arguments)},
@ -3284,7 +3284,7 @@ static bool CheckFuncRefToArrayElement(semantics::SemanticsContext &context,
} else if (name->symbol->Rank() == 0) {
if (const Symbol *function{
semantics::IsFunctionResultWithSameNameAsFunction(*name->symbol)}) {
auto &msg{context.Say(funcRef.v.source,
auto &msg{context.Say(funcRef.source,
function->flags().test(Symbol::Flag::StmtFunction)
? "Recursive call to statement function '%s' is not allowed"_err_en_US
: "Recursive call to '%s' requires a distinct RESULT in its declaration"_err_en_US,
@ -3295,7 +3295,7 @@ static bool CheckFuncRefToArrayElement(semantics::SemanticsContext &context,
return false;
} else {
if (std::get<std::list<parser::ActualArgSpec>>(funcRef.v.t).empty()) {
auto &msg{context.Say(funcRef.v.source,
auto &msg{context.Say(funcRef.source,
"Reference to array '%s' with empty subscript list"_err_en_US,
name->source)};
if (name->symbol) {

4
flang/lib/Semantics/resolve-names.cpp
View File

@ -6869,7 +6869,7 @@ bool ResolveNamesVisitor::Pre(const parser::FunctionReference &x) {
return false;
}
bool ResolveNamesVisitor::Pre(const parser::CallStmt &x) {
HandleCall(Symbol::Flag::Subroutine, x.v);
HandleCall(Symbol::Flag::Subroutine, x.call);
return false;
}
@ -8085,7 +8085,7 @@ public:
resolver_.NoteExecutablePartCall(Symbol::Flag::Function, fr.v);
}
void Post(const parser::CallStmt &cs) {
resolver_.NoteExecutablePartCall(Symbol::Flag::Subroutine, cs.v);
resolver_.NoteExecutablePartCall(Symbol::Flag::Subroutine, cs.call);
}
private:

5
flang/lib/Semantics/tools.cpp
View File

@ -933,11 +933,12 @@ public:
}
bool operator()(const parser::CallStmt &stmt) {
const auto &procedureDesignator{
std::get<parser::ProcedureDesignator>(stmt.v.t)};
std::get<parser::ProcedureDesignator>(stmt.call.t)};
if (auto *name{std::get_if<parser::Name>(&procedureDesignator.u)}) {
// TODO: also ensure that the procedure is, in fact, an intrinsic
if (name->source == "move_alloc") {
const auto &args{std::get<std::list<parser::ActualArgSpec>>(stmt.v.t)};
const auto &args{
std::get<std::list<parser::ActualArgSpec>>(stmt.call.t)};
if (!args.empty()) {
const parser::ActualArg &actualArg{
std::get<parser::ActualArg>(args.front().t)};

37
flang/test/Parser/cuf-sanity-common Normal file
View File

@ -0,0 +1,37 @@
! Common source for CUF parse tree and unparsing tests.
!@cuf subroutine atcuf;
end
#ifdef _CUDA
subroutine cudadefd;
end
#endif
module m
real, allocatable, pinned ::pa(:)
contains
attributes(device) subroutine devicesub; end
attributes(device) real function devicefunc(); devicefunc = 1.; end
attributes(global) subroutine globalsub; end
attributes(grid_global) subroutine gridglobalsub; end
attributes(host) subroutine hostsub; end
attributes(global) launch_bounds(1, 2) subroutine lbsub; end
attributes(global) cluster_dims(1, 2, 3) subroutine cdsub; end
attributes(device) subroutine attrs
! enable with name resolution: attributes(device) :: devx1
real, device :: devx2
end subroutine
subroutine test
logical isPinned
!$cuf kernel do(1) <<<*, *, stream = 1>>>
do j = 1, 10
end do
!$cuf kernel do <<<1, (2, 3), stream = 1>>>
do j = 1, 10
end do
call globalsub<<<1, 2>>>
call globalsub<<<1, 2, 3>>>
call globalsub<<<1, 2, 3, 4>>>
allocate(pa(32), stream = 1, pinned = isPinned)
end subroutine
end module

195
flang/test/Parser/cuf-sanity-tree.CUF Normal file
View File

@ -0,0 +1,195 @@
! RUN: %flang_fc1 -fdebug-dump-parse-tree %s 2>&1 | FileCheck %s
include "cuf-sanity-common"
!CHECK: Program -> ProgramUnit -> SubroutineSubprogram
!CHECK: | SubroutineStmt
!CHECK: | | Name = 'atcuf'
!CHECK: | SpecificationPart
!CHECK: | | ImplicitPart ->
!CHECK: | ExecutionPart -> Block
!CHECK: | EndSubroutineStmt ->
!CHECK: ProgramUnit -> SubroutineSubprogram
!CHECK: | SubroutineStmt
!CHECK: | | Name = 'cudadefd'
!CHECK: | SpecificationPart
!CHECK: | | ImplicitPart ->
!CHECK: | ExecutionPart -> Block
!CHECK: | EndSubroutineStmt ->
!CHECK: ProgramUnit -> Module
!CHECK: | ModuleStmt -> Name = 'm'
!CHECK: | SpecificationPart
!CHECK: | | ImplicitPart ->
!CHECK: | | DeclarationConstruct -> SpecificationConstruct -> TypeDeclarationStmt
!CHECK: | | | DeclarationTypeSpec -> IntrinsicTypeSpec -> Real
!CHECK: | | | AttrSpec -> Allocatable
!CHECK: | | | AttrSpec -> CUDADataAttr = Pinned
!CHECK: | | | EntityDecl
!CHECK: | | | | Name = 'pa'
!CHECK: | | | | ArraySpec -> DeferredShapeSpecList -> int
!CHECK: | ModuleSubprogramPart
!CHECK: | | ContainsStmt
!CHECK: | | ModuleSubprogram -> SubroutineSubprogram
!CHECK: | | | SubroutineStmt
!CHECK: | | | | PrefixSpec -> Attributes -> CUDASubprogramAttrs = Device
!CHECK: | | | | Name = 'devicesub'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
!CHECK: | | | ExecutionPart -> Block
!CHECK: | | | EndSubroutineStmt ->
!CHECK: | | ModuleSubprogram -> FunctionSubprogram
!CHECK: | | | FunctionStmt
!CHECK: | | | | PrefixSpec -> Attributes -> CUDASubprogramAttrs = Device
!CHECK: | | | | PrefixSpec -> DeclarationTypeSpec -> IntrinsicTypeSpec -> Real
!CHECK: | | | | Name = 'devicefunc'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
!CHECK: | | | ExecutionPart -> Block
!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> AssignmentStmt = 'devicefunc=1._4'
!CHECK: | | | | | Variable = 'devicefunc'
!CHECK: | | | | | | Designator -> DataRef -> Name = 'devicefunc'
!CHECK: | | | | | Expr = '1._4'
!CHECK: | | | | | | LiteralConstant -> RealLiteralConstant
!CHECK: | | | | | | | Real = '1.'
!CHECK: | | | EndFunctionStmt ->
!CHECK: | | ModuleSubprogram -> SubroutineSubprogram
!CHECK: | | | SubroutineStmt
!CHECK: | | | | PrefixSpec -> Attributes -> CUDASubprogramAttrs = Global
!CHECK: | | | | Name = 'globalsub'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
!CHECK: | | | ExecutionPart -> Block
!CHECK: | | | EndSubroutineStmt ->
!CHECK: | | ModuleSubprogram -> SubroutineSubprogram
!CHECK: | | | SubroutineStmt
!CHECK: | | | | PrefixSpec -> Attributes -> CUDASubprogramAttrs = Grid_Global
!CHECK: | | | | Name = 'gridglobalsub'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
!CHECK: | | | ExecutionPart -> Block
!CHECK: | | | EndSubroutineStmt ->
!CHECK: | | ModuleSubprogram -> SubroutineSubprogram
!CHECK: | | | SubroutineStmt
!CHECK: | | | | PrefixSpec -> Attributes -> CUDASubprogramAttrs = Host
!CHECK: | | | | Name = 'hostsub'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
!CHECK: | | | ExecutionPart -> Block
!CHECK: | | | EndSubroutineStmt ->
!CHECK: | | ModuleSubprogram -> SubroutineSubprogram
!CHECK: | | | SubroutineStmt
!CHECK: | | | | PrefixSpec -> Attributes -> CUDASubprogramAttrs = Global
!CHECK: | | | | PrefixSpec -> Launch_Bounds -> Scalar -> Integer -> Constant -> Expr = '1_4'
!CHECK: | | | | | LiteralConstant -> IntLiteralConstant = '1'
!CHECK: | | | | Scalar -> Integer -> Constant -> Expr = '2_4'
!CHECK: | | | | | LiteralConstant -> IntLiteralConstant = '2'
!CHECK: | | | | Name = 'lbsub'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
!CHECK: | | | ExecutionPart -> Block
!CHECK: | | | EndSubroutineStmt ->
!CHECK: | | ModuleSubprogram -> SubroutineSubprogram
!CHECK: | | | SubroutineStmt
!CHECK: | | | | PrefixSpec -> Attributes -> CUDASubprogramAttrs = Global
!CHECK: | | | | PrefixSpec -> Cluster_Dims -> Scalar -> Integer -> Constant -> Expr = '1_4'
!CHECK: | | | | | LiteralConstant -> IntLiteralConstant = '1'
!CHECK: | | | | Scalar -> Integer -> Constant -> Expr = '2_4'
!CHECK: | | | | | LiteralConstant -> IntLiteralConstant = '2'
!CHECK: | | | | Scalar -> Integer -> Constant -> Expr = '3_4'
!CHECK: | | | | | LiteralConstant -> IntLiteralConstant = '3'
!CHECK: | | | | Name = 'cdsub'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
!CHECK: | | | ExecutionPart -> Block
!CHECK: | | | EndSubroutineStmt ->
!CHECK: | | ModuleSubprogram -> SubroutineSubprogram
!CHECK: | | | SubroutineStmt
!CHECK: | | | | PrefixSpec -> Attributes -> CUDASubprogramAttrs = Device
!CHECK: | | | | Name = 'attrs'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
!CHECK: | | | | DeclarationConstruct -> SpecificationConstruct -> TypeDeclarationStmt
!CHECK: | | | | | DeclarationTypeSpec -> IntrinsicTypeSpec -> Real
!CHECK: | | | | | AttrSpec -> CUDADataAttr = Device
!CHECK: | | | | | EntityDecl
!CHECK: | | | | | | Name = 'devx2'
!CHECK: | | | ExecutionPart -> Block
!CHECK: | | | EndSubroutineStmt ->
!CHECK: | | ModuleSubprogram -> SubroutineSubprogram
!CHECK: | | | SubroutineStmt
!CHECK: | | | | Name = 'test'
!CHECK: | | | SpecificationPart
!CHECK: | | | | ImplicitPart ->
!CHECK: | | | | DeclarationConstruct -> SpecificationConstruct -> TypeDeclarationStmt
!CHECK: | | | | | DeclarationTypeSpec -> IntrinsicTypeSpec -> Logical
!CHECK: | | | | | EntityDecl
!CHECK: | | | | | | Name = 'ispinned'
!CHECK: | | | ExecutionPart -> Block
!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> CUFKernelDoConstruct
!CHECK: | | | | | Directive
!CHECK: | | | | | | Scalar -> Integer -> Constant -> Expr = '1_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
!CHECK: | | | | | | Scalar -> Integer -> Expr = '1_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
!CHECK: | | | | | DoConstruct
!CHECK: | | | | | | NonLabelDoStmt
!CHECK: | | | | | | | LoopControl -> LoopBounds
!CHECK: | | | | | | | | Scalar -> Name = 'j'
!CHECK: | | | | | | | | Scalar -> Expr = '1_4'
!CHECK: | | | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
!CHECK: | | | | | | | | Scalar -> Expr = '10_4'
!CHECK: | | | | | | | | | LiteralConstant -> IntLiteralConstant = '10'
!CHECK: | | | | | | Block
!CHECK: | | | | | | EndDoStmt ->
!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> CUFKernelDoConstruct
!CHECK: | | | | | Directive
!CHECK: | | | | | | Scalar -> Integer -> Expr = '1_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
!CHECK: | | | | | | Scalar -> Integer -> Expr = '2_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3'
!CHECK: | | | | | | Scalar -> Integer -> Expr = '1_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
!CHECK: | | | | | DoConstruct
!CHECK: | | | | | | NonLabelDoStmt
!CHECK: | | | | | | | LoopControl -> LoopBounds
!CHECK: | | | | | | | | Scalar -> Name = 'j'
!CHECK: | | | | | | | | Scalar -> Expr = '1_4'
!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: | | | | | 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: | | | | | 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: | | | | | 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: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> AllocateStmt
!CHECK: | | | | | Allocation
!CHECK: | | | | | | AllocateObject = 'pa'
!CHECK: | | | | | | | Name = 'pa'
!CHECK: | | | | | | AllocateShapeSpec
!CHECK: | | | | | | | Scalar -> Integer -> Expr = '32_4'
!CHECK: | | | | | | | | LiteralConstant -> IntLiteralConstant = '32'
!CHECK: | | | | | AllocOpt -> Stream -> Scalar -> Integer -> Expr = '1_4'
!CHECK: | | | | | | LiteralConstant -> IntLiteralConstant = '1'
!CHECK: | | | | | AllocOpt -> Pinned -> Scalar -> Logical -> Variable = 'ispinned'
!CHECK: | | | | | | Designator -> DataRef -> Name = 'ispinned'
!CHECK: | | | EndSubroutineStmt ->
!CHECK: | EndModuleStmt ->

41
flang/test/Parser/cuf-sanity-unparse.CUF Normal file
View File

@ -0,0 +1,41 @@
! RUN: %flang_fc1 -fdebug-unparse %s 2>&1 | FileCheck %s
include "cuf-sanity-common"
!CHECK: SUBROUTINE atcuf
!CHECK: END SUBROUTINE
!CHECK: SUBROUTINE cudadefd
!CHECK: END SUBROUTINE
!CHECK: MODULE m
!CHECK: REAL, ALLOCATABLE, PINNED :: pa(:)
!CHECK: CONTAINS
!CHECK: ATTRIBUTES(DEVICE) SUBROUTINE devicesub
!CHECK: END SUBROUTINE
!CHECK: ATTRIBUTES(DEVICE) REAL FUNCTION devicefunc()
!CHECK: devicefunc=1._4
!CHECK: END FUNCTION
!CHECK: ATTRIBUTES(GLOBAL) SUBROUTINE globalsub
!CHECK: END SUBROUTINE
!CHECK: ATTRIBUTES(GRID_GLOBAL) SUBROUTINE gridglobalsub
!CHECK: END SUBROUTINE
!CHECK: ATTRIBUTES(HOST) SUBROUTINE hostsub
!CHECK: END SUBROUTINE
!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: 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: ALLOCATE(pa(32_4), STREAM=1_4, PINNED=ispinned)
!CHECK: END SUBROUTINE
!CHECK: END MODULE

3
flang/test/lib/lit.local.cfg
View File

@ -15,7 +15,8 @@ config.suffixes = [
".ff95",
".fpp",
".FPP",
".cuf" ".CUF",
".cuf",
".CUF",
".f18",
".F18",
".f03",

3
flang/test/lit.cfg.py
View File

@ -42,7 +42,8 @@ config.suffixes = [
".ff95",
".fpp",
".FPP",
".cuf" ".CUF",
".cuf",
".CUF",
".f18",
".F18",
".f03",