[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, 19 insertions, 0 deletions

diff --git a/flang/lib/Common/Fortran.cpp b/flang/lib/Common/Fortran.cpp
index e8d8fef..27ff31e 100644
--- a/flang/lib/Common/Fortran.cpp
+++ b/flang/lib/Common/Fortran.cpp
@@ -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