From 504bff6995bdf21ea1903995938c8ebe7f61497f Mon Sep 17 00:00:00 2001 From: Peter Klausler 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>> t; +}; + struct CUFKernelDoConstruct { TUPLE_CLASS_BOILERPLATE(CUFKernelDoConstruct); WRAPPER_CLASS(StarOrExpr, std::optional); @@ -4316,7 +4327,8 @@ struct CUFKernelDoConstruct { TUPLE_CLASS_BOILERPLATE(Directive); CharBlock source; std::tuple, std::list, - std::list, std::optional> + std::list, std::optional, + std::list> t; }; std::tuple> 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("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( "*" >> pure>() || applyFunction(presentOptional, scalarIntExpr))}; constexpr auto gridOrBlock{parenthesized(nonemptyList(starOrExpr)) || applyFunction(singletonList, starOrExpr)}; + +TYPE_PARSER(("REDUCTION"_tok || "REDUCE"_tok) >> + parenthesized(construct(Parser{}, + ":" >> nonemptyList(scalar(variable))))) + TYPE_PARSER(sourced(beginDirective >> "$CUF KERNEL DO"_tok >> construct( maybe(parenthesized(scalarIntConstantExpr)), "<<<" >> gridOrBlock, "," >> gridOrBlock, - maybe((", 0 ,"_tok || ", STREAM ="_tok) >> scalarIntExpr) / ">>>" / - endDirective))) + maybe((", 0 ,"_tok || ", STREAM ="_tok) >> scalarIntExpr) / ">>>", + many(Parser{}) / endDirective))) TYPE_CONTEXT_PARSER("!$CUF KERNEL DO construct"_en_US, extension(construct( Parser{}, 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>(x.t), " "); + Word("\n"); } - void Unparse(const CUFKernelDoConstruct &x) { Walk(std::get(x.t)); Walk(std::get>(x.t)); } + void Unparse(const CUFReduction &x) { + Word("REDUCE("); + Walk(std::get(x.t)); + Walk(":", std::get>>(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(reduce.t).v}; + for (const auto &var : + std::get>>(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(x.t).source}; const auto &directive{std::get(x.t)}; @@ -489,6 +529,10 @@ void CUDAChecker::Enter(const parser::CUFKernelDoConstruct &x) { if (innerBlock) { DeviceContextChecker{context_}.Check(*innerBlock); } + for (const auto &reduce : + std::get>(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