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

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

1 files changed, 50 insertions, 23 deletions

diff --git a/flang/lib/Parser/executable-parsers.cpp b/flang/lib/Parser/executable-parsers.cpp
index 92e7d25..56ca3ed 100644
--- a/flang/lib/Parser/executable-parsers.cpp
+++ b/flang/lib/Parser/executable-parsers.cpp
@@ -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