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

[klausler]

Summary

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.

[clementval]

I'm not a CUDA Fortran expert but this looks good to me.

flang/lib/Lower/Allocatable.cpp

371

TODO(loc, "cuda stream"); so it triggers the not yet implemented error.

[tschuett]

driver change to enable that feature for *.cuf and *.CUF source files or did you added it to lit?

[awarzynski]

Driver logic LGTM, thanks!

[tschuett]

Sorry. I missed the frontend changes.

[klausler]

Address review comment about TODO() usage in lowering.

[klausler]

If anybody still wishes to review this patch, please let me know, and I will happily wait for you to do so before I merge it.

[MehdiChinoune]

Shouldn't an RFC be opened before start upstreaming CUDA stuff.
Is the community OK with supporting this Nvidia/PGI language extension/dialect?
If for example Intel invents a new extension based OneAPI/SYCL, Should Flang supports it?
There should be a policy for this kind of changes.

[clementval]

In D150159#4338550, @MehdiChinoune wrote:

Shouldn't an RFC be opened before start upstreaming CUDA stuff.
Is the community OK with supporting this Nvidia/PGI language extension/dialect?
If for example Intel invents a new extension based OneAPI/SYCL, Should Flang supports it?
There should be a policy for this kind of changes.

clang has cuda support.

[klausler]

In D150159#4338556, @clementval wrote:

In D150159#4338550, @MehdiChinoune wrote:

Shouldn't an RFC be opened before start upstreaming CUDA stuff.
Is the community OK with supporting this Nvidia/PGI language extension/dialect?
If for example Intel invents a new extension based OneAPI/SYCL, Should Flang supports it?
There should be a policy for this kind of changes.

clang has cuda support.

And so does Classic Flang.

Note that this particular extension is enabled only by means of a distinct source file suffix.

And yes, I would absolutely welcome support for more Fortran language extensions. We already support as many CDC, Cray, DEC, IBM, Intel, NVIDIA/PGI, and probably other extensions that I know about and that have survived to the present day, and others are busily working on adding more, like IBM PPC vector extensions. The more Fortran applications that LLVM Flang can correctly compile unmodified, the better.

[awarzynski]

In D150159#4338550, @MehdiChinoune wrote:

Shouldn't an RFC be opened before start upstreaming CUDA stuff.

CUDA is very widely adopted and I doubt that anyone would object to this. The changes in the driver are very non-intrusive and I see no need to discuss this more widely. Definitely not objecting to an RFC (in fact, a short PSA would be nice), but I'm happy for this to land without further advertising.

If for example Intel invents a new extension based OneAPI/SYCL, Should Flang supports it?

If there is a critical mass of users that would like to use a particular extension then it should be supported. It doesn't matter whether it comes from "company A" or "company B". As for SYCL specifically, isn't that a purely C++ extension? What's the relevance here?

[MehdiChinoune]

There is a big difference between CUDA and CUDA-Fortran.
1 - CUDA could be used by any compiler MSVC, Intel, IBM, GNU... as It doesn't require any support from the compiler. It's just libraries+headers.
While CUDA-Fortran requires a support from the compiler itself.
GCC, Intel, Cray, Fujitsu ... compilers do not support it because It's just an Nvidia/PGI dialect, Flang-Classic is just an open-sourced Frontend of PGI compiler.
2 - CUDA-Fortran requires a proprietary CUDA SDK to be installed which make it different from other dialects/extensions you mentioned above. those are language extensions they don't need any SDK to be installed.

We are talking here about a company pushing for a dialect not widely used by the community.
If you fellow many discussions online in Fortran community, you know for sure that PGI compiler is one of the least/worst used compiler.
I think Nvidia should keep CUDA-Fortran support in the commercial/proprietary compiler that will be based on flang.

[MehdiChinoune]

Could you at least make CUDA-Fortran support behind a CMake option.
As a package maintainer on a platform where CUDA is not supported (I am sure I am not alone), I don't want to fix code that will never benefit my platform.
You know this a community project, CUDA is not supported everywhere.

[everythingfunctional]

I'm fairly sympathetic to @MehdiChinoune 's argument here. You're essentially asking an open-source community to support your proprietary technology. Maybe it's widely enough used that it deserves to be (I'm at least open to the argument), but you should definitely be open and upfront about wanting to do that, and ask the community through the proper procedures first. I think it would have a been a violation of the community's trust if it had been approved and made it into the official repository without anybody outside Nvidia having a chance to comment. Maybe I missed it, but I don't remember hearing about this plan before the patches got submitted. And based on the number of reviewers listed here, it doesn't look like you were trying to publicise it very broadly.

[klausler]

In D150159#4344586, @MehdiChinoune wrote:

There is a big difference between CUDA and CUDA-Fortran.
1 - CUDA could be used by any compiler MSVC, Intel, IBM, GNU... as It doesn't require any support from the compiler. It's just libraries+headers.
While CUDA-Fortran requires a support from the compiler itself.

This is actually not the case. A CUDA C(++) program needs to be compiled by a C(++) compiler capable of handling the same features (function attributes, data attributes, kernel launch, &c.) that are present in CUDA Fortran.

GCC, Intel, Cray, Fujitsu ... compilers do not support it because It's just an Nvidia/PGI dialect, Flang-Classic is just an open-sourced Frontend of PGI compiler.

IBM's XLF Fortran compiler also supports CUDA Fortran.

2 - CUDA-Fortran requires a proprietary CUDA SDK to be installed which make it different from other dialects/extensions you mentioned above. those are language extensions they don't need any SDK to be installed.

It's no different from the CUDA support in Clang or the PTX support in the LLVM code generator in that respect.

We are talking here about a company pushing for a dialect not widely used by the community.
If you fellow many discussions online in Fortran community, you know for sure that PGI compiler is one of the least/worst used compiler.
I think Nvidia should keep CUDA-Fortran support in the commercial/proprietary compiler that will be based on flang.

I would much rather have our source in the open and available to all, which is why LLVM Flang exists in the first place rather than being a closed-source NVIDIA project to improve its Fortran product.

[klausler]

In D150159#4346067, @everythingfunctional wrote:

I'm fairly sympathetic to @MehdiChinoune 's argument here. You're essentially asking an open-source community to support your proprietary technology. Maybe it's widely enough used that it deserves to be (I'm at least open to the argument), but you should definitely be open and upfront about wanting to do that, and ask the community through the proper procedures first. I think it would have a been a violation of the community's trust if it had been approved and made it into the official repository without anybody outside Nvidia having a chance to comment. Maybe I missed it, but I don't remember hearing about this plan before the patches got submitted. And based on the number of reviewers listed here, it doesn't look like you were trying to publicise it very broadly.

I've left these reviews open for comment for a long time, and will continue to do so to collect relevant comments from people who actually participate in the project.

[sscalpone]

In D150159#4344598, @MehdiChinoune wrote:

Could you at least make CUDA-Fortran support behind a CMake option.
As a package maintainer on a platform where CUDA is not supported (I am sure I am not alone), I don't want to fix code that will never benefit my platform.
You know this a community project, CUDA is not supported everywhere.

Hi Mehdi,

It's a good idea to guard CUDA Fortran dependencies on CUDA in a similar way to how clang's CUDA C++ dependencies on CUDA are guarded.

This patch is front end work; because there's no lowering yet, there's no need to add such CMake options yet. The feature is gated from Fortran users by way of a special required file extension (.cuf).

CUDA can work with Fortran using modules, interfaces, and libraries similarly to how CUDA can work with C++ using headers and libraries. But both CUDA C++ and CUDA Fortran are language extensions, e.g. chevron syntax.

Generally regarding extensions, the Flang policy is spelled out in https://github.com/llvm/llvm-project/blob/48c3ae5cc3d4612283018ea597db3f7214aabfd9/flang/docs/Extensions.md.

As a general principle, this compiler will accept by default and without complaint many legacy features, extensions to the standard language, and features that have been deleted from the standard, so long as the recognition of those features would not cause a standard-conforming program to be rejected or misinterpreted.

Adding support for CUDA Fortran follows both this spirit and intent as did the addition of OpenMP and OpenACC.

[MehdiChinoune]

In D150159#4346788, @sscalpone wrote:

Hi Mehdi,

It's a good idea to guard CUDA Fortran dependencies on CUDA in a similar way to how clang's CUDA C++ dependencies on CUDA are guarded.

This patch is front end work; because there's no lowering yet, there's no need to add such CMake options yet. The feature is gated from Fortran users by way of a special required file extension (.cuf).

I was talking about flang code. These changes could make flang fails to build. So we have to fix it because a feature we don't use.

Adding support for CUDA Fortran follows both this spirit and intent as did the addition of OpenMP and OpenACC.

OpenMP and OpenACC are open standards.

[everythingfunctional]

In D150159#4346752, @klausler wrote:

In D150159#4346067, @everythingfunctional wrote:

I'm fairly sympathetic to @MehdiChinoune 's argument here. You're essentially asking an open-source community to support your proprietary technology. Maybe it's widely enough used that it deserves to be (I'm at least open to the argument), but you should definitely be open and upfront about wanting to do that, and ask the community through the proper procedures first. I think it would have a been a violation of the community's trust if it had been approved and made it into the official repository without anybody outside Nvidia having a chance to comment. Maybe I missed it, but I don't remember hearing about this plan before the patches got submitted. And based on the number of reviewers listed here, it doesn't look like you were trying to publicise it very broadly.

I've left these reviews open for comment for a long time, and will continue to do so to collect relevant comments from people who actually participate in the project.

But my point was that you have not actually solicited feedback from everybody who "actually participates," only those actively monitoring every submission to Phabricator. Some people can't actively monitor every submission, since many people contribute on a voluntary or part time basis, and I think their opinions matter too.

[klausler]

In D150159#4346968, @MehdiChinoune wrote:

I was talking about flang code. These changes could make flang fails to build. So we have to fix it because a feature we don't use.

I couldn't find any CMakeLists.txt configuration enablement switches for CUDA in Clang.

[klausler]

Rebase to current head.

[tschuett]

Is CUDA Fortran an extrension of ISO Fortran? Do you need to mention it in Extensions.md?

[kiranchandramohan]

I believe it is fine to support CUDA Fortran in Flang. This is broadly in line with CUDA support in Clang. I believe there is sufficient documentation in the open that specifies its syntax and semantics (https://docs.nvidia.com/hpc-sdk/compilers/cuda-fortran-prog-guide/index.html). Also, it can be seen that we included cuf as part of the file extensions that we wanted to test initially.

I am not a CUDA Fortran expert but the patch looks OK to me.

[klausler]

Rebase to current llvm-project/main in preparation for merging.

flang/docs/ParserCombinators.md

 * `skipMany(p)` is the same as `many(p)`, but it discards the results.
 * `maybe(p)` tries to match p, returning an `std::optional<T>` value.
   It always succeeds.
 * `defaulted(p)` matches p, and when p fails it returns a
   default-constructed instance of p's resultType.  It always succeeds.
 * `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
   message that `p` generates will be tagged with `msg` as its
   context.  Contexts may nest.
 * `withMessage(msg, p)` succeeds if `p` does, and if it does not,
   it discards the messages from `p` and fails with the specified message.
 * `recovery(p, q)` is equivalent to `p || q`, except that error messages
* `"..."_tok` match the content of the string, skipping spaces before and
   free form source in pedantic mode if the next character could also
   be part of an identifier -- add a trailing blank to avoid this.
 * `"..."_sptok` is a string match in which the spaces are required in
    free form source.
 * `"..."_id` is a string match for a complete identifier (not a prefix of
    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.
 
 ### Debugging Parser
 Last, a string literal `"..."_debug` denotes a parser that emits the string to
 `llvm::errs` and succeeds.  It is useful for tracing while debugging a parser but should
 obviously not be committed for production code.

flang/include/flang/Common/Fortran-features.h

ENUM_CLASS(LanguageFeature, BackslashEscapes, OldDebugLines,
     XOROperator, PunctuationInNames, OptionalFreeFormSpace, BOZExtensions,
     EmptyStatement, AlternativeNE, ExecutionPartNamelist, DECStructures,
     DoubleComplex, Byte, StarKind, ExponentMatchingKindParam, QuadPrecision,
     SlashInitialization, TripletInArrayConstructor, MissingColons,
     SignedComplexLiteral, OldStyleParameter, ComplexConstructor, PercentLOC,
     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,
     OldLabelDoEndStatements, LogicalIntegerAssignment, EmptySourceFile,
     ProgramReturn, ImplicitNoneTypeNever, ImplicitNoneTypeAlways,
     ForwardRefImplicitNone, OpenAccessAppend, BOZAsDefaultInteger,
     DistinguishableSpecifics, DefaultSave, PointerInSeqType, NonCharacterFormat,
     SaveMainProgram, SaveBigMainProgramVariables,

 
 class LanguageFeatureControl {
 public:
   LanguageFeatureControl() {
     // These features must be explicitly enabled by command line options.
     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);
     disable_.set(LanguageFeature::SaveMainProgram);
     // These features, if enabled, conflict with valid standard usage,
     // so there are disabled here by default.
     disable_.set(LanguageFeature::BackslashEscapes);
     disable_.set(LanguageFeature::LogicalAbbreviations);

flang/include/flang/Common/Fortran.h

 #define FORTRAN_COMMON_FORTRAN_H_
 
 // Fortran language concepts that are used in many phases are defined
 // once here to avoid redundancy and needless translation.
 
 #include "enum-set.h"
 #include "idioms.h"
 #include <cinttypes>
+#include <optional>
 #include <string>
 
 namespace Fortran::common {
 
 // Fortran has five kinds of intrinsic data types, plus the derived types.
 ENUM_CLASS(TypeCategory, Integer, Real, Complex, Character, Logical, Derived)
 ENUM_CLASS(VectorElementCategory, Integer, Unsigned, Real)
 

 };
 
 // Fortran label. Must be in [1..99999].
 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};
 
 // !DIR$ IGNORE_TKR [[(letters) name] ... letters
 // "A" expands to all of TKRDM
 ENUM_CLASS(IgnoreTKR,
     Type, // T - don't check type category
     Kind, // K - don't check kind
     Rank, // R - don't check ranks
     Device, // D - don't check host/device residence
     Managed, // M - don't check managed storage
     Contiguous) // C - legacy; disabled NVFORTRAN's convention that leading
                 // dimension of assumed-shape was contiguous
 using IgnoreTKRSet = EnumSet<IgnoreTKR, 8>;
 // IGNORE_TKR(A) = IGNORE_TKR(TKRDM)
 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_

flang/include/flang/Common/indirection.h

if (p_) {
       deleter_(p_);
     }
   }
 
   A &operator*() const { return *p_; }
   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;
     return result;
   }
 
   void Reset(A *p = nullptr) {
     if (p_) {

flang/include/flang/Common/template.h

 };
 
 template <template <typename> class PREDICATE, typename TUPLEorVARIANT>
 constexpr int SearchMembers{
     OverMembers<SearchMembersHelper<PREDICATE>::template Scanner,
         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>
 std::optional<A> JoinOptional(std::optional<std::optional<A>> &&x) {
   if (x) {
     return std::move(*x);
   }
   return std::nullopt;

flang/include/flang/Frontend/FrontendOptions.h

 /// \param suffix The file extension
 /// \return True if the file extension should be processed as free form
 bool isFreeFormSuffix(llvm::StringRef suffix);
 
 /// \param suffix The file extension
 /// \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,
 
   /// MLIR: we accept this so that we can run the optimizer on it, and compile
   /// it to LLVM IR, assembly or object code.
   MLIR,
 
   /// LLVM IR: we accept this so that we can run the optimizer on it,
class FrontendInputFile {
   bool isFixedForm = false;
 
   /// Must this file be preprocessed? Note that in Flang the preprocessor is
   /// always run. This flag is used to control whether predefined and command
   /// line preprocessor macros are enabled or not. In practice, this is
   /// 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)
       : file(file.str()), kind(inKind) {
 
     // Based on the extension, decide whether this is a fixed or free form
     // file.
     auto pathDotIndex{file.rfind(".")};
     std::string pathSuffix{file.substr(pathDotIndex + 1)};
     isFixedForm = isFixedFormSuffix(pathSuffix);
     mustBePreprocessed = isToBePreprocessed(pathSuffix);
+    isCUDAFortran = isCUDAFortranSuffix(pathSuffix);
   }
 
   FrontendInputFile(const llvm::MemoryBuffer *memBuf, InputKind inKind)
       : buffer(memBuf), kind(inKind) {}
 
   InputKind getKind() const { return kind; }
 
   bool isEmpty() const { return file.empty() && buffer == nullptr; }
   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());
     return file;
   }
 
   const llvm::MemoryBuffer *getBuffer() const {
     assert(buffer && "Requested buffer, but it is empty!");

flang/include/flang/Parser/dump-parse-tree.h

 //===-- include/flang/Parser/dump-parse-tree.h ------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef FORTRAN_PARSER_DUMP_PARSE_TREE_H_
 #define FORTRAN_PARSER_DUMP_PARSE_TREE_H_
 
 #include "format-specification.h"
 #include "parse-tree-visitor.h"
 #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"
 #include <string>
 #include <type_traits>
 
 namespace Fortran::parser {
 
static std::string GetNodeName(const T::E &x) { \
     return #E " = "s + std::string{T::EnumToString(x)}; \
   }
 #define NODE(T1, T2) NODE_NAME(T1::T2, #T2)
   NODE_NAME(bool, "bool")
   NODE_NAME(int, "int")
   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)
   NODE(format, FormatItem)
   NODE(format, FormatSpecification)
   NODE(format, IntrinsicTypeDataEditDesc)
   NODE(format::IntrinsicTypeDataEditDesc, Kind)
   NODE(parser, Abstract)
#include "llvm/Frontend/OpenACC/ACC.inc"
   NODE(parser, ActualArg)
   NODE(ActualArg, PercentRef)
   NODE(ActualArg, PercentVal)
   NODE(parser, ActualArgSpec)
   NODE(AcValue, Triplet)
   NODE(parser, AllocOpt)
   NODE(AllocOpt, Mold)
   NODE(AllocOpt, Source)
+  NODE(AllocOpt, Stream)
+  NODE(AllocOpt, Pinned)
   NODE(parser, Allocatable)
   NODE(parser, AllocatableStmt)
   NODE(parser, AllocateCoarraySpec)
   NODE(parser, AllocateObject)
   NODE(parser, AllocateShapeSpec)
   NODE(parser, AllocateStmt)
   NODE(parser, Allocation)
   NODE(parser, AltReturnSpec)
#include "llvm/Frontend/OpenACC/ACC.inc"
   NODE(parser, BlockData)
   NODE(parser, BlockDataStmt)
   NODE(parser, BlockSpecificationPart)
   NODE(parser, BlockStmt)
   NODE(parser, BoundsRemapping)
   NODE(parser, BoundsSpec)
   NODE(parser, Call)
   NODE(parser, CallStmt)
+  NODE(CallStmt, Chevrons)
   NODE(parser, CaseConstruct)
   NODE(CaseConstruct, Case)
   NODE(parser, CaseSelector)
   NODE(parser, CaseStmt)
   NODE(parser, CaseValueRange)
   NODE(CaseValueRange, Range)
   NODE(parser, ChangeTeamConstruct)
   NODE(parser, ChangeTeamStmt)
#include "llvm/Frontend/OpenACC/ACC.inc"
   NODE(ConnectSpec, Newunit)
   NODE(ConnectSpec, Recl)
   NODE(parser, ContainsStmt)
   NODE(parser, Contiguous)
   NODE(parser, ContiguousStmt)
   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)
   NODE(parser, DataImpliedDo)
   NODE(parser, DataRef)
   NODE(parser, DataStmt)
   NODE(parser, DataStmtConstant)
   NODE(parser, DataStmtObject)
#include "llvm/Frontend/OpenMP/OMP.inc"
   NODE(parser, PositionOrFlushSpec)
   NODE(parser, PrefixSpec)
   NODE(PrefixSpec, Elemental)
   NODE(PrefixSpec, Impure)
   NODE(PrefixSpec, Module)
   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)
   NODE(parser, ProcAttrSpec)
   NODE(parser, ProcComponentAttrSpec)
   NODE(parser, ProcComponentDefStmt)
   NODE(parser, ProcComponentRef)
   NODE(parser, ProcDecl)

flang/include/flang/Parser/message.h

constexpr MessageFixedText(
       const char str[], std::size_t n, Severity severity = Severity::None)
       : text_{str, n}, severity_{severity} {}
   constexpr MessageFixedText(const MessageFixedText &) = default;
   constexpr MessageFixedText(MessageFixedText &&) = default;
   constexpr MessageFixedText &operator=(const MessageFixedText &) = default;
   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;
     return *this;
   }
   bool IsFatal() const {
     return severity_ == Severity::Error || severity_ == Severity::Todo;
   }

flang/include/flang/Parser/parse-tree-visitor.h

 }
 template <typename M> void Walk(Designator &x, M &mutator) {
   if (mutator.Pre(x)) {
     Walk(x.source, mutator);
     Walk(x.u, 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);
   }
 }
 template <typename V> void Walk(const PartRef &x, V &visitor) {
   if (visitor.Pre(x)) {
     Walk(x.name, visitor);
     Walk(x.subscripts, visitor);
     Walk(x.imageSelector, visitor);

flang/include/flang/Parser/parse-tree.h

 class DerivedTypeSpec;
 } // namespace Fortran::semantics
 
 // Expressions in the parse tree have owning pointers that can be set to
 // type-checked generic expression representations by semantic analysis.
 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
 // for their move constructor and move assignment operator=, and disable
 // their copy constructor and copy assignment operator=.
 #define COPY_AND_ASSIGN_BOILERPLATE(classname) \
   classname(classname &&) = default; \
   classname &operator=(classname &&) = default; \

 struct SeparateModuleSubprogram; // R1538
 struct EntryStmt; // R1541
 struct ReturnStmt; // R1542
 struct StmtFunctionStmt; // R1544
 
 // Directives, extensions, and deprecated statements
 struct CompilerDirective;
 struct BasedPointerStmt;
+struct CUDAAttributesStmt;
 struct StructureDef;
 struct ArithmeticIfStmt;
 struct AssignStmt;
 struct AssignedGotoStmt;
 struct PauseStmt;
 struct OpenACCConstruct;
 struct AccEndCombinedDirective;
 struct OpenACCDeclarativeConstruct;
 struct OpenMPConstruct;
 struct OpenMPDeclarativeConstruct;
 struct OmpEndLoopDirective;
+struct CUFKernelDoConstruct;
 
 // Cooked character stream locations
 using Location = const char *;
 
 // A parse tree node with provenance only
 struct Verbatim {
   // Allow a no-arg constructor for Verbatim so parsers can return `RESULT{}`.
   constexpr Verbatim() {}

 
 // R513 other-specification-stmt ->
 //        access-stmt | allocatable-stmt | asynchronous-stmt | bind-stmt |
 //        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
 // Extension: (Cray) based POINTER statement
+// Extension: CUDA data attribute statement
 struct OtherSpecificationStmt {
   UNION_CLASS_BOILERPLATE(OtherSpecificationStmt);
   std::variant<common::Indirection<AccessStmt>,
       common::Indirection<AllocatableStmt>,
       common::Indirection<AsynchronousStmt>, common::Indirection<BindStmt>,
       common::Indirection<CodimensionStmt>, common::Indirection<ContiguousStmt>,
       common::Indirection<DimensionStmt>, common::Indirection<ExternalStmt>,
       common::Indirection<IntentStmt>, common::Indirection<IntrinsicStmt>,
       common::Indirection<NamelistStmt>, common::Indirection<OptionalStmt>,
       common::Indirection<PointerStmt>, common::Indirection<ProtectedStmt>,
       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;
 };
 
 // R508 specification-construct ->
 //        derived-type-def | enum-def | generic-stmt | interface-block |
 //        parameter-stmt | procedure-declaration-stmt |
 //        other-specification-stmt | type-declaration-stmt
 struct SpecificationConstruct {
std::variant<common::Indirection<AllocateStmt>,
       common::Indirection<PauseStmt>>
       u;
 };
 
 // R514 executable-construct ->
 //        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>,
       common::Indirection<BlockConstruct>, common::Indirection<CaseConstruct>,
       common::Indirection<ChangeTeamConstruct>,
       common::Indirection<CriticalConstruct>,
       Statement<common::Indirection<LabelDoStmt>>,
       Statement<common::Indirection<EndDoStmt>>,
       common::Indirection<DoConstruct>, common::Indirection<IfConstruct>,
       common::Indirection<SelectRankConstruct>,
       common::Indirection<SelectTypeConstruct>,
       common::Indirection<WhereConstruct>, common::Indirection<ForallConstruct>,
       common::Indirection<CompilerDirective>,
       common::Indirection<OpenACCConstruct>,
       common::Indirection<AccEndCombinedDirective>,
       common::Indirection<OpenMPConstruct>,
-      common::Indirection<OmpEndLoopDirective>>
+      common::Indirection<OmpEndLoopDirective>,
+      common::Indirection<CUFKernelDoConstruct>>
       u;
 };
 
 // R510 execution-part-construct ->
 //        executable-construct | format-stmt | entry-stmt | data-stmt
 // Extension (PGI/Intel): also accept NAMELIST in execution part
 struct ExecutionPartConstruct {
   UNION_CLASS_BOILERPLATE(ExecutionPartConstruct);

 struct ComponentArraySpec {
   UNION_CLASS_BOILERPLATE(ComponentArraySpec);
   std::variant<std::list<ExplicitShapeSpec>, DeferredShapeSpecList> u;
 };
 
 // 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;
 };
 
 // R806 null-init -> function-reference   ... which must be NULL()
 WRAPPER_CLASS(NullInit, common::Indirection<Expr>);
 
 // R744 initial-data-target -> designator
 using InitialDataTarget = common::Indirection<Designator>;
struct IntentSpec {
   WRAPPER_CLASS_BOILERPLATE(IntentSpec, Intent);
 };
 
 // R802 attr-spec ->
 //        access-spec | ALLOCATABLE | ASYNCHRONOUS |
 //        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);
 EMPTY_CLASS(Optional);
 EMPTY_CLASS(Parameter);
 EMPTY_CLASS(Protected);
 EMPTY_CLASS(Save);
 EMPTY_CLASS(Target);
 EMPTY_CLASS(Value);
 EMPTY_CLASS(Volatile);
 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;
 };
 
 // R803 entity-decl ->
 //        object-name [( array-spec )] [lbracket coarray-spec rbracket]
 //          [* char-length] [initialization] |
 //        function-name [* char-length]
 struct EntityDecl {

 // R1165 sync-stat -> STAT = stat-variable | ERRMSG = errmsg-variable
 struct StatOrErrmsg {
   UNION_CLASS_BOILERPLATE(StatOrErrmsg);
   std::variant<StatVariable, MsgVariable> u;
 };
 
 // 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 ->
 //        ALLOCATE ( [type-spec ::] allocation-list [, alloc-opt-list] )
 struct AllocateStmt {
   TUPLE_CLASS_BOILERPLATE(AllocateStmt);
   std::tuple<std::optional<TypeSpec>, std::list<Allocation>,
       std::list<AllocOpt>>
struct ProcedureDeclarationStmt {
   TUPLE_CLASS_BOILERPLATE(ProcedureDeclarationStmt);
   std::tuple<std::optional<ProcInterface>, std::list<ProcAttrSpec>,
       std::list<ProcDecl>>
       t;
 };
 
 // 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);
   EMPTY_CLASS(Impure);
   EMPTY_CLASS(Module);
   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;
 };
 
 // R1532 suffix ->
 //         proc-language-binding-spec [RESULT ( result-name )] |
 //         RESULT ( result-name ) [proc-language-binding-spec]
 struct Suffix {
   BOILERPLATE(Suffix);

 };
 
 // R1523 actual-arg-spec -> [keyword =] actual-arg
 struct ActualArgSpec {
   TUPLE_CLASS_BOILERPLATE(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
 };
 
 // R1529 function-subprogram ->
 //         function-stmt [specification-part] [execution-part]
 //         [internal-subprogram-part] end-function-stmt
 struct FunctionSubprogram {
struct CompilerDirective {
   struct NameValue {
     TUPLE_CLASS_BOILERPLATE(NameValue);
     std::tuple<Name, std::optional<std::uint64_t>> t;
   };
   CharBlock source;
   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);
   std::tuple<ObjectName, ObjectName, std::optional<ArraySpec>> t;
 };
 WRAPPER_CLASS(BasedPointerStmt, std::list<BasedPointer>);
 
 struct Union;

 struct OpenACCConstruct {
   UNION_CLASS_BOILERPLATE(OpenACCConstruct);
   std::variant<OpenACCBlockConstruct, OpenACCCombinedConstruct,
       OpenACCLoopConstruct, OpenACCStandaloneConstruct, OpenACCCacheConstruct,
       OpenACCWaitConstruct, OpenACCAtomicConstruct>
       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_

flang/lib/Common/Fortran.cpp

if (tkr.test(IgnoreTKR::Managed)) {
     result += 'M';
   }
   if (tkr.test(IgnoreTKR::Contiguous)) {
     result += 'C';
   }
   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

flang/lib/Frontend/FrontendAction.cpp

bool FrontendAction::beginSourceFile(CompilerInstance &ci,
   // to decide whether to include them or not.
   if ((invoc.getPreprocessorOpts().macrosFlag == PPMacrosFlag::Include) ||
       (invoc.getPreprocessorOpts().macrosFlag == PPMacrosFlag::Unknown &&
        getCurrentInput().getMustBePreprocessed())) {
     invoc.setDefaultPredefinitions();
     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) {
     invoc.getFortranOpts().isFixedForm = getCurrentInput().getIsFixedForm();
   }
 
   if (!beginSourceFileAction()) {
     beginSourceFileCleanUp(*this, ci);

flang/lib/Frontend/FrontendOptions.cpp

bool Fortran::frontend::isFixedFormSuffix(llvm::StringRef suffix) {
   // Note: Keep this list in-sync with flang/test/lit.cfg.py
   return suffix == "f77" || suffix == "f" || suffix == "F" || suffix == "ff" ||
          suffix == "for" || suffix == "FOR" || suffix == "fpp" ||
          suffix == "FPP";
 }
 
 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) {
   if (isFixedFormSuffix(extension) || isFreeFormSuffix(extension)) {
     return Language::Fortran;
   }
 
   if (extension == "bc" || extension == "ll")
     return Language::LLVM_IR;
   if (extension == "fir" || extension == "mlir")
     return Language::MLIR;
 
   return Language::Unknown;
 }

flang/lib/Lower/Allocatable.cpp

for (const auto &allocOption :
                     statOrErr.u);
               },
               [&](const Fortran::parser::AllocOpt::Source &source) {
                 sourceExpr = Fortran::semantics::GetExpr(source.v.value());
               },
               [&](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);
   }
 
   void lowerAllocation(const Allocation &alloc) {
     fir::MutableBoxValue boxAddr =
         genMutableBoxValue(converter, loc, alloc.getAllocObj());
 

flang/lib/Lower/Bridge.cpp

mlir::func::FuncOp getFunc(llvm::StringRef name, mlir::FunctionType ty) {
     }
     return builder->createFunction(toLocation(), name, ty);
   }
 
   /// Lowering of CALL statement
   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()) {
       std::optional<mlir::Type> resultType;
       if (stmt.typedCall->hasAlternateReturns())
         resultType = builder->getIndexType();
       auto hlfirRes = Fortran::lower::convertCallToHLFIR(
           toLocation(), *this, *stmt.typedCall, resultType, localSymbols,
void genFIR(const Fortran::parser::CallStmt &stmt) {
     if (!res)
       return; // "Normal" subroutine call.
     // Call with alternate return specifiers.
     // The call returns an index that selects an alternate return branch target.
     llvm::SmallVector<int64_t> indexList;
     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)) {
         indexList.push_back(++index);
         labelList.push_back(altReturn->v);
       }
     }
     genMultiwayBranch(res, indexList, labelList, eval.nonNopSuccessor());

flang/lib/Lower/PFTBuilder.cpp

void analyzeBranches(lower::pft::Evaluation *parentConstruct,
                        std::list<lower::pft::Evaluation> &evaluationList) {
     lower::pft::Evaluation *lastConstructStmtEvaluation{};
     for (auto &eval : evaluationList) {
       eval.visit(common::visitors{
           // Action statements (except IO statements)
           [&](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 =
                       std::get_if<parser::AltReturnSpec>(&actual.u))
                 markBranchTarget(eval, altReturn->v);
             }
           },
           [&](const parser::CycleStmt &s) {

flang/lib/Parser/Fortran-parsers.cpp

 TYPE_PARSER(construct<DataComponentDefStmt>(declarationTypeSpec,
     optionalListBeforeColons(Parser<ComponentAttrSpec>{}),
     nonemptyList("expected component declarations"_err_en_US,
         Parser<ComponentOrFill>{})))
 
 // 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),
         kindOrLen >> construct<ErrorRecovery>())))
 
 // R739 component-decl ->
 //        component-name [( component-array-spec )]
 //        [lbracket coarray-spec rbracket] [* char-length]
extension<LanguageFeature::MissingColons>(
             withMessage("expected entity declarations"_err_en_US,
                 "," >> nonemptyList(entityDecl)))))
 
 // R802 attr-spec ->
 //        access-spec | ALLOCATABLE | ASYNCHRONOUS |
 //        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)) ||
     construct<AttrSpec>("CODIMENSION" >> coarraySpec) ||
     construct<AttrSpec>(contiguous) ||
     construct<AttrSpec>("DIMENSION" >> arraySpec) ||
     construct<AttrSpec>(construct<External>("EXTERNAL"_tok)) ||
     construct<AttrSpec>("INTENT" >> parenthesized(intentSpec)) ||
     construct<AttrSpec>(construct<Intrinsic>("INTRINSIC"_tok)) ||
     construct<AttrSpec>(languageBindingSpec) || construct<AttrSpec>(optional) ||
     construct<AttrSpec>(construct<Parameter>("PARAMETER"_tok)) ||
     construct<AttrSpec>(pointer) || construct<AttrSpec>(protectedAttr) ||
     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};
 
 // R803 entity-decl ->
 //        object-name [( array-spec )] [lbracket coarray-spec rbracket]
 //          [* char-length] [initialization] |
 //        function-name [* char-length]

 //        ALLOCATE ( [type-spec ::] allocation-list [, alloc-opt-list] )
 TYPE_CONTEXT_PARSER("ALLOCATE statement"_en_US,
     construct<AllocateStmt>("ALLOCATE (" >> maybe(typeSpec / "::"),
         nonemptyList(Parser<Allocation>{}),
         defaulted("," >> nonemptyList(Parser<AllocOpt>{})) / ")"))
 
 // 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))))
 
 // R932 allocation ->
 //        allocate-object [( allocate-shape-spec-list )]
 //        [lbracket allocate-coarray-spec rbracket]
 TYPE_PARSER(construct<Allocation>(Parser<AllocateObject>{},

 // R1165 sync-stat -> STAT = stat-variable | ERRMSG = errmsg-variable
 TYPE_PARSER(construct<StatOrErrmsg>("STAT =" >> statVariable) ||
     construct<StatOrErrmsg>("ERRMSG =" >> msgVariable))
 
 // Directives, extensions, and deprecated statements
 // !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) ||
         construct<CompilerDirective>(
             "DIR$" >> many(construct<CompilerDirective::NameValue>(name,
                           maybe(("="_tok || ":"_tok) >> digitString64))))) /
         endOfStmt)
 
 TYPE_PARSER(extension<LanguageFeature::CrayPointer>(
     "nonstandard usage: based POINTER"_port_en_US,
     construct<BasedPointerStmt>(
         "POINTER" >> nonemptyList("expected POINTER associations"_err_en_US,
                          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>(
     "STRUCTURE" >> maybe(sourced("/" >> name / "/")), optionalList(entityDecl)))
 
 constexpr auto nestedStructureDef{
     CONTEXT_PARSER("nested STRUCTURE definition"_en_US,
         construct<StructureDef>(statement(NestedStructureStmt{}),

flang/lib/Parser/basic-parsers.h

 // is a parser that optionally enabled, sets a strict conformance violation
 // flag, and may emit a warning message, if those are enabled.
 template <LanguageFeature LF, typename PA> class NonstandardParser {
 public:
   using resultType = typename PA::resultType;
   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)) {
         return std::nullopt;
       }
     }
     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_);
     }
     return result;
   }
 
 private:
   const PA parser_;
   const MessageFixedText message_;
 };
 
 template <LanguageFeature LF, typename PA>
 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.
 template <LanguageFeature LF, typename PA> class DeprecatedParser {
 public:
   using resultType = typename PA::resultType;
   constexpr DeprecatedParser(const DeprecatedParser &) = default;
   constexpr DeprecatedParser(PA parser) : parser_{parser} {}

flang/lib/Parser/executable-parsers.cpp

 //===-- lib/Parser/executable-parsers.cpp ---------------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 // 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"
 #include "token-parsers.h"
 #include "type-parser-implementation.h"
 #include "flang/Parser/characters.h"
 #include "flang/Parser/parse-tree.h"
 
 namespace Fortran::parser {
 
 // Fortran allows the statement with the corresponding label at the end of
 // a do-construct that begins with an old-style label-do-stmt to be a
 // new-style END DO statement; e.g., DO 10 I=1,N; ...; 10 END DO.  Usually,
 // END DO statements appear only at the ends of do-constructs that begin
 // with a nonlabel-do-stmt, so care must be taken to recognize this case and
 // essentially treat them like CONTINUE statements.
 
 // R514 executable-construct ->
 //        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>(indirect(compilerDirective)),
+    construct<ExecutableConstruct>(indirect(Parser<CUFKernelDoConstruct>{})))};
 
 // R510 execution-part-construct ->
 //        executable-construct | format-stmt | entry-stmt | data-stmt
 // Extension (PGI/Intel): also accept NAMELIST in execution part
 constexpr auto obsoleteExecutionPartConstruct{recovery(ignoredStatementPrefix >>
         fail<ExecutionPartConstruct>(
             "obsolete legacy extension is not supported"_err_en_US),
     construct<ExecutionPartConstruct>(construct<ErrorRecovery>(ok /
TYPE_PARSER(
     construct<LockStmt::LockStat>("ACQUIRED_LOCK =" >> scalarLogicalVariable) ||
     construct<LockStmt::LockStat>(statOrErrmsg))
 
 // R1181 unlock-stmt -> UNLOCK ( lock-variable [, sync-stat-list] )
 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

flang/lib/Parser/io-parsers.cpp

 TYPE_PARSER(first(construct<WaitSpec>(maybe("UNIT ="_tok) >> fileUnitNumber),
     construct<WaitSpec>("END =" >> endLabel),
     construct<WaitSpec>("EOR =" >> eorLabel),
     construct<WaitSpec>("ERR =" >> errLabel),
     construct<WaitSpec>("ID =" >> idExpr),
     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))};
 constexpr auto positionOrFlushSpecList{
     parenthesized(nonemptyList(positionOrFlushSpec)) || bareUnitNumberAsList};
 
 // R1224 backspace-stmt ->
 //         BACKSPACE file-unit-number | BACKSPACE ( position-spec-list )

flang/lib/Parser/misc-parsers.h

 constexpr auto star{construct<Star>("*"_tok)};
 constexpr auto allocatable{construct<Allocatable>("ALLOCATABLE"_tok)};
 constexpr auto contiguous{construct<Contiguous>("CONTIGUOUS"_tok)};
 constexpr auto optional{construct<Optional>("OPTIONAL"_tok)};
 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

flang/lib/Parser/parse-tree.cpp

return common::visit(
             return common::visit(
                 common::visitors{
                     [&](common::Indirection<Designator> &z) {
                       return WithSource(
                           z.value().source, Expr{std::move(z.value())});
                     },
                     [&](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);
           },
           [&](auto &) -> Expr { common::die("unexpected type"); },
       },
       std::get<ActualArg>(arg.t).u);
 }
 
 Designator FunctionReference::ConvertToArrayElementRef() {
   std::list<Expr> args;
   for (auto &arg : std::get<std::list<ActualArgSpec>>(v.t)) {
     args.emplace_back(ActualArgToExpr(arg));
   }
   return common::visit(
       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)));
           },
       },
       std::get<ProcedureDesignator>(v.t).u);
 }
 
 StructureConstructor FunctionReference::ConvertToStructureConstructor(
     const semantics::DerivedTypeSpec &derived) {
Statement<ActionStmt> StmtFunctionStmt::ConvertToAssignment() {
   for (const Name &arg : funcArgs) {
     actuals.emplace_back(std::optional<Keyword>{},
         ActualArg{Expr{WithSource(
             arg.source, Designator{DataRef{Name{arg.source, arg.symbol}}})}});
     source.ExtendToCover(arg.source);
   }
   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{
           AssignmentStmt{std::move(variable), std::move(funcExpr)}}}};
 }
 
 CharBlock Variable::GetSource() const {
   return common::visit(
       common::visitors{
           [&](const common::Indirection<Designator> &des) {
             return des.value().source;
           },
           [&](const common::Indirection<parser::FunctionReference> &call) {
-            return call.value().v.source;
+            return call.value().source;
           },
       },
       u);
 }
 
 llvm::raw_ostream &operator<<(llvm::raw_ostream &os, const Name &x) {
   return os << x.ToString();
 }
 } // namespace Fortran::parser

flang/lib/Parser/parsing.cpp

prescanner.set_fixedForm(options.isFixedForm)
       .AddCompilerDirectiveSentinel("dir$");
   if (options.features.IsEnabled(LanguageFeature::OpenACC)) {
     prescanner.AddCompilerDirectiveSentinel("$acc");
   }
   if (options.features.IsEnabled(LanguageFeature::OpenMP)) {
     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);
   if (currentCooked_->BufferedBytes() == 0 && !options.isModuleFile) {
     // Input is empty.  Append a newline so that any warning
     // message about nonstandard usage will have provenance.
     currentCooked_->Put('\n', range.start());
   }

flang/lib/Parser/preprocessor.cpp

if (dir.TokenAt(j).ToString() == "<") { // #include <foo>
       }
       if (k >= tokens) {
         prescanner.Say(dir.GetIntervalProvenanceRange(j, tokens - j),
             "#include: expected '>' at end of included file"_port_en_US);
       }
       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());
       }
     } else {
       prescanner.Say(dir.GetTokenProvenanceRange(j < tokens ? j : tokens - 1),
           "#include: expected name of file to include"_err_en_US);
       return;

flang/lib/Parser/prescan.cpp

if (directiveSentinel_[0] == '$' && directiveSentinel_[1] == '\0') {
       // OpenMP conditional compilation line.  Remove the sentinel and then
       // treat the line as if it were normal source.
       at_ += 2, column_ += 2;
       if (inFixedForm_) {
         LabelField(tokens);
       } 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, '!');
       ++at_, ++column_;
       for (const char *sp{directiveSentinel_}; *sp != '\0';
            ++sp, ++at_, ++column_) {
         EmitChar(tokens, *sp);
       }

flang/lib/Parser/program-parsers.cpp

 //        entry-stmt | stmt-function-stmt
 // N.B. These parsers incorporate recognition of some other statements that
 // may have been misplaced in the sequence of statements that are acceptable
 // as a specification part in order to improve error recovery.
 // Also note that many instances of specification-part in the standard grammar
 // 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>{} >>
         fail<DeclarationConstruct>("misplaced USE statement"_err_en_US) ||
     Parser<ImportStmt>{} >>
         fail<DeclarationConstruct>(
             "IMPORT statements must follow any USE statements and precede all other declarations"_err_en_US) ||
     Parser<ImplicitStmt>{} >>
first(construct<SpecificationConstruct>(indirect(Parser<DerivedTypeDef>{})),
         construct<SpecificationConstruct>(indirect(openmpDeclarativeConstruct)),
         construct<SpecificationConstruct>(indirect(compilerDirective))))
 
 // R513 other-specification-stmt ->
 //        access-stmt | allocatable-stmt | asynchronous-stmt | bind-stmt |
 //        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>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<AsynchronousStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<BindStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<CodimensionStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<ContiguousStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<DimensionStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<ExternalStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<IntentStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<IntrinsicStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<NamelistStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<OptionalStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<PointerStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<ProtectedStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<SaveStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<TargetStmt>{})),
     construct<OtherSpecificationStmt>(indirect(Parser<ValueStmt>{})),
     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]
 //         [internal-subprogram-part] end-program-stmt
 TYPE_CONTEXT_PARSER("main program"_en_US,
     construct<MainProgram>(maybe(statement(Parser<ProgramStmt>{})),
         specificationPart, executionPart, maybe(internalSubprogramPart),
         unterminatedStatement(Parser<EndProgramStmt>{})))

 // R1518 initial-proc-target -> procedure-name
 TYPE_PARSER(
     construct<ProcPointerInit>(nullInit) || construct<ProcPointerInit>(name))
 
 // R1519 intrinsic-stmt -> INTRINSIC [::] intrinsic-procedure-name-list
 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 ->
 //         procedure-name | proc-component-ref | data-ref % binding-name
 TYPE_PARSER(construct<ProcedureDesignator>(Parser<ProcComponentRef>{}) ||
     construct<ProcedureDesignator>(name))
 
 // R1523 actual-arg-spec -> [keyword =] actual-arg
extension<LanguageFeature::PercentRefAndVal>(
         construct<ActualArg>(
             construct<ActualArg::PercentVal>("%VAL" >> parenthesized(expr)))))
 
 // R1525 alt-return-spec -> * label
</