[flang-commits] [flang] [flang] Parse REDUCE clauses in !$CUF KERNEL DO (PR #92154)
Peter Klausler via flang-commits
flang-commits at lists.llvm.org
Wed May 15 14:27:22 PDT 2024
https://github.com/klausler updated https://github.com/llvm/llvm-project/pull/92154
>From 504bff6995bdf21ea1903995938c8ebe7f61497f 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/check-cuda.cpp | 44 ++++++++++++
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 ++
flang/test/Semantics/reduce.cuf | 72 ++++++++++++++++++++
11 files changed, 199 insertions(+), 18 deletions(-)
create mode 100644 flang/test/Semantics/reduce.cuf
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..0a40aa8b8f616 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-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>);
@@ -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..382a593416872 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-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>{},
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..1639e900903fe 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<Variable>>>(x.t), ",", ")");
+ }
void Done() const { CHECK(indent_ == 0); }
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index 96ab902392633..45217ed2e3ccd 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/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) {
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>>>()
diff --git a/flang/test/Semantics/reduce.cuf b/flang/test/Semantics/reduce.cuf
new file mode 100644
index 0000000000000..95ff2e87c09b4
--- /dev/null
+++ b/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
More information about the flang-commits
mailing list