[flang] Parse REDUCE clauses in !$CUF KERNEL DO

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

@@ -236,6 +236,7 @@ class ParseTreeDumper {
   NODE(parser, CUFKernelDoConstruct)
   NODE(CUFKernelDoConstruct, StarOrExpr)
   NODE(CUFKernelDoConstruct, Directive)
+  NODE(parser, CUFReduction)
   NODE(parser, CycleStmt)
   NODE(parser, DataComponentDefStmt)
   NODE(parser, DataIDoObject)

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

@@ -4303,20 +4303,32 @@ struct OpenACCConstruct {
 };
 
 // CUF-kernel-do-construct ->
-//     !$CUF KERNEL DO [ (scalar-int-constant-expr) ] <<< grid, block [, stream]
-//     >>> do-construct
+//   !$CUF KERNEL DO [ (scalar-int-constant-expr) ]
+//      <<< grid, block [, stream] >>>
+//      [ cuf-reduction... ]
+//      do-construct
 // star-or-expr -> * | scalar-int-expr
 // grid -> * | scalar-int-expr | ( star-or-expr-list )
 // block -> * | scalar-int-expr | ( star-or-expr-list )
 // stream -> 0, scalar-int-expr | STREAM = scalar-int-expr
+// cuf-reduction -> [ REDUCE | REDUCTION ] (
+//                  acc-reduction-op : scalar-variable-list )
+
+struct CUFReduction {
+  TUPLE_CLASS_BOILERPLATE(CUFReduction);
+  using Operator = AccReductionOperator;
+  std::tuple<Operator, std::list<Scalar<Variable>>> t;
+};
+
 struct CUFKernelDoConstruct {
   TUPLE_CLASS_BOILERPLATE(CUFKernelDoConstruct);
   WRAPPER_CLASS(StarOrExpr, std::optional<ScalarIntExpr>);
   struct Directive {
     TUPLE_CLASS_BOILERPLATE(Directive);
     CharBlock source;
     std::tuple<std::optional<ScalarIntConstantExpr>, std::list<StarOrExpr>,
-        std::list<StarOrExpr>, std::optional<ScalarIntExpr>>
+        std::list<StarOrExpr>, std::optional<ScalarIntExpr>,
+        std::list<CUFReduction>>
         t;
   };
   std::tuple<Directive, std::optional<DoConstruct>> t;

flang/lib/Parser/executable-parsers.cpp

@@ -538,25 +538,34 @@ 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
+// CUF-kernel-do-construct ->
+//   !$CUF KERNEL DO [ (scalar-int-constant-expr) ]
+//      <<< grid, block [, stream] >>>
+//      [ cuf-reduction... ]
+//      do-construct
 // star-or-expr -> * | scalar-int-expr
 // grid -> * | scalar-int-expr | ( star-or-expr-list )
 // block -> * | scalar-int-expr | ( star-or-expr-list )
-// stream -> ( 0, | STREAM = ) scalar-int-expr
+// stream -> 0, scalar-int-expr | STREAM = scalar-int-expr
+// cuf-reduction -> [ REDUCTION | REDUCE ] (
+//                  acc-reduction-op : scalar-variable-list )
+
 constexpr auto starOrExpr{construct<CUFKernelDoConstruct::StarOrExpr>(
     "*" >> pure<std::optional<ScalarIntExpr>>() ||
     applyFunction(presentOptional<ScalarIntExpr>, scalarIntExpr))};
 constexpr auto gridOrBlock{parenthesized(nonemptyList(starOrExpr)) ||
     applyFunction(singletonList<CUFKernelDoConstruct::StarOrExpr>, starOrExpr)};
+
+TYPE_PARSER(("REDUCTION"_tok || "REDUCE"_tok) >>
+    parenthesized(construct<CUFReduction>(Parser<CUFReduction::Operator>{},
+        ":" >> nonemptyList(scalar(variable)))))
+
 TYPE_PARSER(sourced(beginDirective >> "$CUF KERNEL DO"_tok >>
     construct<CUFKernelDoConstruct::Directive>(
         maybe(parenthesized(scalarIntConstantExpr)), "<<<" >> gridOrBlock,
         "," >> gridOrBlock,
-        maybe((", 0 ,"_tok || ", STREAM ="_tok) >> scalarIntExpr) / ">>>" /
-            endDirective)))
+        maybe((", 0 ,"_tok || ", STREAM ="_tok) >> scalarIntExpr) / ">>>",
+        many(Parser<CUFReduction>{}) / endDirective)))
 TYPE_CONTEXT_PARSER("!$CUF KERNEL DO construct"_en_US,
     extension<LanguageFeature::CUDA>(construct<CUFKernelDoConstruct>(
         Parser<CUFKernelDoConstruct::Directive>{},

flang/lib/Parser/openacc-parsers.cpp

@@ -19,9 +19,9 @@
 // OpenACC Directives and Clauses
 namespace Fortran::parser {
 
-constexpr auto startAccLine = skipStuffBeforeStatement >>
-    ("!$ACC "_sptok || "C$ACC "_sptok || "*$ACC "_sptok);
-constexpr auto endAccLine = space >> endOfLine;
+constexpr auto startAccLine{skipStuffBeforeStatement >>
+    ("!$ACC "_sptok || "C$ACC "_sptok || "*$ACC "_sptok)};
+constexpr auto endAccLine{space >> endOfLine};
 
 // Autogenerated clauses parser. Information is taken from ACC.td and the
 // parser is generated by tablegen.

flang/lib/Parser/unparse.cpp

@@ -2705,7 +2705,6 @@ class UnparseVisitor {
   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
@@ -2736,6 +2735,31 @@ class UnparseVisitor {
   WALK_NESTED_ENUM(OmpOrderClause, Type) // OMP order-type
   WALK_NESTED_ENUM(OmpOrderModifier, Kind) // OMP order-modifier
 #undef WALK_NESTED_ENUM
+  void Unparse(const AccReductionOperator::Operator x) {
+    switch (x) {
+    case AccReductionOperator::Operator::Plus:
+      Word("+");
+      break;
+    case AccReductionOperator::Operator::Multiply:
+      Word("*");
+      break;
+    case AccReductionOperator::Operator::And:
+      Word(".AND.");
+      break;
+    case AccReductionOperator::Operator::Or:
+      Word(".OR.");
+      break;
+    case AccReductionOperator::Operator::Eqv:
+      Word(".EQV.");
+      break;
+    case AccReductionOperator::Operator::Neqv:
+      Word(".NEQV.");
+      break;
+    default:
+      Word(AccReductionOperator::EnumToString(x));
+      break;
+    }
+  }
 
   void Unparse(const CUFKernelDoConstruct::StarOrExpr &x) {
     if (x.v) {
@@ -2768,13 +2792,19 @@ class UnparseVisitor {
     if (const auto &stream{std::get<3>(x.t)}) {
       Word(",STREAM="), Walk(*stream);
     }
-    Word(">>>\n");
+    Word(">>>");
+    Walk(" ", std::get<std::list<CUFReduction>>(x.t), " ");
+    Word("\n");
   }
-
   void Unparse(const CUFKernelDoConstruct &x) {
     Walk(std::get<CUFKernelDoConstruct::Directive>(x.t));
     Walk(std::get<std::optional<DoConstruct>>(x.t));
   }
+  void Unparse(const CUFReduction &x) {
+    Word("REDUCE(");
+    Walk(std::get<CUFReduction::Operator>(x.t));
+    Walk(":", std::get<std::list<Scalar<Variable>>>(x.t), ",", ")");
+  }
 
   void Done() const { CHECK(indent_ == 0); }
 

flang/lib/Semantics/check-cuda.cpp

@@ -463,6 +463,46 @@ static int DoConstructTightNesting(
   return 1;
 }
 
+static void CheckReduce(
+    SemanticsContext &context, const parser::CUFReduction &reduce) {
+  auto op{std::get<parser::CUFReduction::Operator>(reduce.t).v};
+  for (const auto &var :
+      std::get<std::list<parser::Scalar<parser::Variable>>>(reduce.t)) {
+    if (const auto &typedExprPtr{var.thing.typedExpr};
+        typedExprPtr && typedExprPtr->v) {
+      const auto &expr{*typedExprPtr->v};
+      if (auto type{expr.GetType()}) {
+        auto cat{type->category()};
+        bool isOk{false};
+        switch (op) {
+        case parser::AccReductionOperator::Operator::Plus:
+        case parser::AccReductionOperator::Operator::Multiply:
+        case parser::AccReductionOperator::Operator::Max:
+        case parser::AccReductionOperator::Operator::Min:
+          isOk = cat == TypeCategory::Integer || cat == TypeCategory::Real;
+          break;
+        case parser::AccReductionOperator::Operator::Iand:
+        case parser::AccReductionOperator::Operator::Ior:
+        case parser::AccReductionOperator::Operator::Ieor:
+          isOk = cat == TypeCategory::Integer;
+          break;
+        case parser::AccReductionOperator::Operator::And:
+        case parser::AccReductionOperator::Operator::Or:
+        case parser::AccReductionOperator::Operator::Eqv:
+        case parser::AccReductionOperator::Operator::Neqv:
+          isOk = cat == TypeCategory::Logical;
+          break;
+        }
+        if (!isOk) {
+          context.Say(var.thing.GetSource(),
+              "!$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type %s"_err_en_US,
+              type->AsFortran());
+        }
+      }
+    }
+  }
+}
+
 void CUDAChecker::Enter(const parser::CUFKernelDoConstruct &x) {
   auto source{std::get<parser::CUFKernelDoConstruct::Directive>(x.t).source};
   const auto &directive{std::get<parser::CUFKernelDoConstruct::Directive>(x.t)};
@@ -489,6 +529,10 @@ void CUDAChecker::Enter(const parser::CUFKernelDoConstruct &x) {
   if (innerBlock) {
     DeviceContextChecker<true>{context_}.Check(*innerBlock);
   }
+  for (const auto &reduce :
+      std::get<std::list<parser::CUFReduction>>(directive.t)) {
+    CheckReduce(context_, reduce);
+  }
 }
 
 void CUDAChecker::Enter(const parser::AssignmentStmt &x) {

flang/lib/Semantics/resolve-directives.h

@@ -21,7 +21,7 @@ class SemanticsContext;
 
 // Name resolution for OpenACC and OpenMP directives
 void ResolveAccParts(
-    SemanticsContext &, const parser::ProgramUnit &, Scope *topScope = {});
+    SemanticsContext &, const parser::ProgramUnit &, Scope *topScope);
 void ResolveOmpParts(SemanticsContext &, const parser::ProgramUnit &);
 void ResolveOmpTopLevelParts(SemanticsContext &, const parser::Program &);
 

flang/lib/Semantics/resolve-names.cpp

@@ -8941,7 +8941,7 @@ bool ResolveNamesVisitor::Pre(const parser::ProgramUnit &x) {
   FinishSpecificationParts(root);
   ResolveExecutionParts(root);
   FinishExecutionParts(root);
-  ResolveAccParts(context(), x);
+  ResolveAccParts(context(), x, /*topScope=*/nullptr);
   ResolveOmpParts(context(), x);
   return false;
 }

flang/test/Parser/cuf-sanity-common

@@ -23,12 +23,19 @@ module m
   end subroutine
   subroutine test
     logical isPinned
+    real a(10), x, y, z
     !$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
+    !$cuf kernel do <<<*, *>>> reduce(+:x,y) reduce(*:z)
+    do j = 1, 10
+      x = x + a(j)
+      y = y + a(j)
+      z = z * a(j)
+    end do
     call globalsub<<<1, 2>>>
     call globalsub<<<1, 2, 3>>>
     call globalsub<<<1, 2, 3, 4>>>

flang/test/Parser/cuf-sanity-unparse.CUF

@@ -34,6 +34,12 @@ include "cuf-sanity-common"
 !CHECK:   !$CUF KERNEL DO <<<1_4,(2_4,3_4),STREAM=1_4>>>
 !CHECK:   DO j=1_4,10_4
 !CHECK:   END DO
+!CHECK:   !$CUF KERNEL DO <<<*,*>>> REDUCE(+:x,y) REDUCE(*:z)
+!CHECK:   DO j=1_4,10_4
+!CHECK:    x=x+a(int(j,kind=8))
+!CHECK:    y=y+a(int(j,kind=8))
+!CHECK:    z=z*a(int(j,kind=8))
+!CHECK:   END DO
 !CHECK:    CALL globalsub<<<1_4,2_4>>>()
 !CHECK:    CALL globalsub<<<1_4,2_4,3_4>>>()
 !CHECK:    CALL globalsub<<<1_4,2_4,3_4,4_4>>>()

flang/test/Semantics/reduce.cuf

@@ -0,0 +1,72 @@
+! RUN: %python %S/test_errors.py %s %flang_fc1
+subroutine s(n,m,a,l)
+  integer, intent(in) :: n
+  integer, intent(in) :: m(n)
+  real, intent(in) :: a(n)
+  logical, intent(in) :: l(n)
+  integer j, mr
+  real ar
+  logical lr
+!$cuf kernel do <<<*,*>>> reduce (+:mr,ar)
+  do j=1,n; mr = mr + m(j); ar = ar + a(j); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type LOGICAL(4)
+!$cuf kernel do <<<*,*>>> reduce (+:lr)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (*:mr,ar)
+  do j=1,n; mr = mr * m(j); ar = ar * a(j); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type LOGICAL(4)
+!$cuf kernel do <<<*,*>>> reduce (*:lr)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (max:mr,ar)
+  do j=1,n; mr = max(mr,m(j)); ar = max(ar,a(j)); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type LOGICAL(4)
+!$cuf kernel do <<<*,*>>> reduce (max:lr)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (min:mr,ar)
+  do j=1,n; mr = min(mr,m(j)); ar = min(ar,a(j)); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type LOGICAL(4)
+!$cuf kernel do <<<*,*>>> reduce (min:lr)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (iand:mr)
+  do j=1,n; mr = iand(mr,m(j)); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type REAL(4)
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type LOGICAL(4)
+!$cuf kernel do <<<*,*>>> reduce (iand:ar,lr)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (ieor:mr)
+  do j=1,n; mr = ieor(mr,m(j)); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type REAL(4)
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type LOGICAL(4)
+!$cuf kernel do <<<*,*>>> reduce (ieor:ar,lr)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (ior:mr)
+  do j=1,n; mr = ior(mr,m(j)); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type REAL(4)
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type LOGICAL(4)
+!$cuf kernel do <<<*,*>>> reduce (ior:ar,lr)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (.and.:lr)
+  do j=1,n; lr = lr .and. l(j); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type INTEGER(4)
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type REAL(4)
+!$cuf kernel do <<<*,*>>> reduce (.and.:mr,ar)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (.eqv.:lr)
+  do j=1,n; lr = lr .eqv. l(j); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type INTEGER(4)
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type REAL(4)
+!$cuf kernel do <<<*,*>>> reduce (.eqv.:mr,ar)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (.neqv.:lr)
+  do j=1,n; lr = lr .neqv. l(j); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type INTEGER(4)
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type REAL(4)
+!$cuf kernel do <<<*,*>>> reduce (.neqv.:mr,ar)
+  do j=1,n; end do
+!$cuf kernel do <<<*,*>>> reduce (.or.:lr)
+  do j=1,n; lr = lr .or. l(j); end do
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type INTEGER(4)
+!ERROR: !$CUF KERNEL DO REDUCE operation is not acceptable for a variable with type REAL(4)
+!$cuf kernel do <<<*,*>>> reduce (.or.:mr,ar)
+  do j=1,n; end do
+end