[flang-commits] [flang] [flang] Parse REDUCE clauses in !$CUF KERNEL DO (PR #92154)
Peter Klausler via flang-commits
flang-commits at lists.llvm.org
Tue May 14 11:02:06 PDT 2024
https://github.com/klausler created https://github.com/llvm/llvm-project/pull/92154
A !$CUF KERNEL DO directive is allowed to have advisory REDUCE clauses similar to those in OpenACC and DO CONCURRENT. Parse and represent them. Semantic validation will follow.
>From 2627f1e71d9b244c0efe10531e233559d2e7efe0 Mon Sep 17 00:00:00 2001
From: Peter Klausler <pklausler at nvidia.com>
Date: Tue, 14 May 2024 10:59:12 -0700
Subject: [PATCH] [flang] Parse REDUCE clauses in !$CUF KERNEL DO
A !$CUF KERNEL DO directive is allowed to have advisory REDUCE
clauses similar to those in OpenACC and DO CONCURRENT. Parse
and represent them. Semantic validation will follow.
---
flang/include/flang/Parser/dump-parse-tree.h | 1 +
flang/include/flang/Parser/parse-tree.h | 18 ++++++++--
flang/lib/Parser/executable-parsers.cpp | 23 +++++++++----
flang/lib/Parser/openacc-parsers.cpp | 6 ++--
flang/lib/Parser/unparse.cpp | 36 ++++++++++++++++++--
flang/lib/Semantics/resolve-directives.h | 2 +-
flang/lib/Semantics/resolve-names.cpp | 2 +-
flang/test/Parser/cuf-sanity-common | 7 ++++
flang/test/Parser/cuf-sanity-unparse.CUF | 6 ++++
9 files changed, 83 insertions(+), 18 deletions(-)
diff --git a/flang/include/flang/Parser/dump-parse-tree.h b/flang/include/flang/Parser/dump-parse-tree.h
index 477d391277ee2..68ae50c312cde 100644
--- a/flang/include/flang/Parser/dump-parse-tree.h
+++ b/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)
diff --git a/flang/include/flang/Parser/parse-tree.h b/flang/include/flang/Parser/parse-tree.h
index c063544583790..39cc514b5dfd6 100644
--- a/flang/include/flang/Parser/parse-tree.h
+++ b/flang/include/flang/Parser/parse-tree.h
@@ -4303,12 +4303,23 @@ 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-designator-list )
+
+struct CUFReduction {
+ TUPLE_CLASS_BOILERPLATE(CUFReduction);
+ using Operator = AccReductionOperator;
+ std::tuple<Operator, std::list<Scalar<Designator>>> t;
+};
+
struct CUFKernelDoConstruct {
TUPLE_CLASS_BOILERPLATE(CUFKernelDoConstruct);
WRAPPER_CLASS(StarOrExpr, std::optional<ScalarIntExpr>);
@@ -4316,7 +4327,8 @@ struct CUFKernelDoConstruct {
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;
diff --git a/flang/lib/Parser/executable-parsers.cpp b/flang/lib/Parser/executable-parsers.cpp
index 07a570bd61e99..c1f34cc85ea28 100644
--- a/flang/lib/Parser/executable-parsers.cpp
+++ b/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-designator-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(designator)))))
+
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>{},
diff --git a/flang/lib/Parser/openacc-parsers.cpp b/flang/lib/Parser/openacc-parsers.cpp
index 946b33d0084a9..3d919e29a2482 100644
--- a/flang/lib/Parser/openacc-parsers.cpp
+++ b/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.
diff --git a/flang/lib/Parser/unparse.cpp b/flang/lib/Parser/unparse.cpp
index 3398b395f198f..a9498659138a1 100644
--- a/flang/lib/Parser/unparse.cpp
+++ b/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<Designator>>>(x.t), ",", ")");
+ }
void Done() const { CHECK(indent_ == 0); }
diff --git a/flang/lib/Semantics/resolve-directives.h b/flang/lib/Semantics/resolve-directives.h
index 4aef8ad6c4008..5a890c26aa334 100644
--- a/flang/lib/Semantics/resolve-directives.h
+++ b/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 &);
diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp
index e2875081b732c..121745f9b13d6 100644
--- a/flang/lib/Semantics/resolve-names.cpp
+++ b/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;
}
diff --git a/flang/test/Parser/cuf-sanity-common b/flang/test/Parser/cuf-sanity-common
index b097a6aa30045..9d73204e3f5f6 100644
--- a/flang/test/Parser/cuf-sanity-common
+++ b/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>>>
diff --git a/flang/test/Parser/cuf-sanity-unparse.CUF b/flang/test/Parser/cuf-sanity-unparse.CUF
index b6921e74fc05a..d4be347dd044e 100644
--- a/flang/test/Parser/cuf-sanity-unparse.CUF
+++ b/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>>>()
More information about the flang-commits
mailing list