shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

[flang] Parse REDUCE clauses in !$CUF KERNEL DO (#92154)

Browse Source 
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.
This commit is contained in:
Peter Klausler 2024-05-15 16:28:58 -07:00 committed by GitHub
parent 463f58a564
commit 5bbb63bd6d
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
11 changed files with 199 additions and 18 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
flang/include/flang/Parser/dump-parse-tree.h
View File

@ -236,6 +236,7 @@ public:
NODE(parser, CUFKernelDoConstruct)
NODE(CUFKernelDoConstruct, StarOrExpr)
NODE(CUFKernelDoConstruct, Directive)
NODE(parser, CUFReduction)
NODE(parser, CycleStmt)
NODE(parser, DataComponentDefStmt)
NODE(parser, DataIDoObject)

18
flang/include/flang/Parser/parse-tree.h
View File

@ -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;

23
flang/lib/Parser/executable-parsers.cpp
View File

@ -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>{},

6
flang/lib/Parser/openacc-parsers.cpp
View File

@ -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.

36
flang/lib/Parser/unparse.cpp
View File

@ -2705,7 +2705,6 @@ public:
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 @@ public:
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 @@ public:
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); }

44
flang/lib/Semantics/check-cuda.cpp
View File

@ -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) {

2
flang/lib/Semantics/resolve-directives.h
View File

@ -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 &);

2
flang/lib/Semantics/resolve-names.cpp
View File

@ -8940,7 +8940,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;
}

7
flang/test/Parser/cuf-sanity-common
View File

@ -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>>>

6
flang/test/Parser/cuf-sanity-unparse.CUF
View File

@ -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>>>()

72
flang/test/Semantics/reduce.cuf Normal file
View File

@ -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