[flang] Parse REDUCE clauses in !$CUF KERNEL DO #92154
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
flang
|
@llvm/pr-subscribers-flang-parser @llvm/pr-subscribers-flang-semantics Author: Peter Klausler (klausler) ChangesA !$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. Full diff: https://github.com/llvm/llvm-project/pull/92154.diff 9 Files Affected:
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>>>()
|
razvanlupusoru
approved these changes
May 14, 2024
ImanHosseini
approved these changes
May 14, 2024
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.
mub-at-arm
pushed a commit
to mub-at-arm/llvm-project
that referenced
this pull request
May 16, 2024
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.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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.