Skip to content

Commit 60fa2b0

Browse files
klauslerklausler
authored
[flang] Parse !$CUF KERNEL DO <<< (*) (#85338)
Accept and represent asterisks within the parenthesized grid and block specification lists.
1 parent 5e21fa2 commit 60fa2b0

File tree

9 files changed

+63
-29
lines changed

9 files changed

+63
-29
lines changed

‎flang/include/flang/Parser/dump-parse-tree.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -233,6 +233,7 @@ class ParseTreeDumper {
233233
NODE(parser, CriticalStmt)
234234
NODE(parser, CUDAAttributesStmt)
235235
NODE(parser, CUFKernelDoConstruct)
236+
NODE(CUFKernelDoConstruct, StarOrExpr)
236237
NODE(CUFKernelDoConstruct, Directive)
237238
NODE(parser, CycleStmt)
238239
NODE(parser, DataComponentDefStmt)

‎flang/include/flang/Parser/parse-tree.h

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -4297,16 +4297,18 @@ struct OpenACCConstruct {
42974297
// CUF-kernel-do-construct ->
42984298
// !$CUF KERNEL DO [ (scalar-int-constant-expr) ] <<< grid, block [, stream]
42994299
// >>> do-construct
4300-
// grid -> * | scalar-int-expr | ( scalar-int-expr-list )
4301-
// block -> * | scalar-int-expr | ( scalar-int-expr-list )
4300+
// star-or-expr -> * | scalar-int-expr
4301+
// grid -> * | scalar-int-expr | ( star-or-expr-list )
4302+
// block -> * | scalar-int-expr | ( star-or-expr-list )
43024303
// stream -> 0, scalar-int-expr | STREAM = scalar-int-expr
43034304
struct CUFKernelDoConstruct {
43044305
TUPLE_CLASS_BOILERPLATE(CUFKernelDoConstruct);
4306+
WRAPPER_CLASS(StarOrExpr, std::optional<ScalarIntExpr>);
43054307
struct Directive {
43064308
TUPLE_CLASS_BOILERPLATE(Directive);
43074309
CharBlock source;
4308-
std::tuple<std::optional<ScalarIntConstantExpr>, std::list<ScalarIntExpr>,
4309-
std::list<ScalarIntExpr>, std::optional<ScalarIntExpr>>
4310+
std::tuple<std::optional<ScalarIntConstantExpr>, std::list<StarOrExpr>,
4311+
std::list<StarOrExpr>, std::optional<ScalarIntExpr>>
43104312
t;
43114313
};
43124314
std::tuple<Directive, std::optional<DoConstruct>> t;

‎flang/lib/Lower/Bridge.cpp

Lines changed: 21 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -2508,19 +2508,32 @@ class FirConverter : public Fortran::lower::AbstractConverter {
25082508
if (nestedLoops > 1)
25092509
n = builder->getIntegerAttr(builder->getI64Type(), nestedLoops);
25102510

2511-
const std::list<Fortran::parser::ScalarIntExpr> &grid = std::get<1>(dir.t);
2512-
const std::list<Fortran::parser::ScalarIntExpr> &block = std::get<2>(dir.t);
2511+
const std::list<Fortran::parser::CUFKernelDoConstruct::StarOrExpr> &grid =
2512+
std::get<1>(dir.t);
2513+
const std::list<Fortran::parser::CUFKernelDoConstruct::StarOrExpr> &block =
2514+
std::get<2>(dir.t);
25132515
const std::optional<Fortran::parser::ScalarIntExpr> &stream =
25142516
std::get<3>(dir.t);
25152517

25162518
llvm::SmallVector<mlir::Value> gridValues;
2517-
for (const Fortran::parser::ScalarIntExpr &expr : grid)
2518-
gridValues.push_back(fir::getBase(
2519-
genExprValue(*Fortran::semantics::GetExpr(expr), stmtCtx)));
2519+
for (const Fortran::parser::CUFKernelDoConstruct::StarOrExpr &expr : grid) {
2520+
if (expr.v) {
2521+
gridValues.push_back(fir::getBase(
2522+
genExprValue(*Fortran::semantics::GetExpr(*expr.v), stmtCtx)));
2523+
} else {
2524+
// TODO: '*'
2525+
}
2526+
}
25202527
llvm::SmallVector<mlir::Value> blockValues;
2521-
for (const Fortran::parser::ScalarIntExpr &expr : block)
2522-
blockValues.push_back(fir::getBase(
2523-
genExprValue(*Fortran::semantics::GetExpr(expr), stmtCtx)));
2528+
for (const Fortran::parser::CUFKernelDoConstruct::StarOrExpr &expr :
2529+
block) {
2530+
if (expr.v) {
2531+
blockValues.push_back(fir::getBase(
2532+
genExprValue(*Fortran::semantics::GetExpr(*expr.v), stmtCtx)));
2533+
} else {
2534+
// TODO: '*'
2535+
}
2536+
}
25242537
mlir::Value streamValue;
25252538
if (stream)
25262539
streamValue = fir::getBase(

‎flang/lib/Parser/executable-parsers.cpp

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -542,19 +542,19 @@ TYPE_CONTEXT_PARSER("UNLOCK statement"_en_US,
542542
// CUF-kernel-do-directive ->
543543
// !$CUF KERNEL DO [ (scalar-int-constant-expr) ] <<< grid, block [, stream]
544544
// >>> do-construct
545-
// grid -> * | scalar-int-expr | ( scalar-int-expr-list )
546-
// block -> * | scalar-int-expr | ( scalar-int-expr-list )
545+
// star-or-expr -> * | scalar-int-expr
546+
// grid -> * | scalar-int-expr | ( star-or-expr-list )
547+
// block -> * | scalar-int-expr | ( star-or-expr-list )
547548
// stream -> ( 0, | STREAM = ) scalar-int-expr
549+
constexpr auto starOrExpr{construct<CUFKernelDoConstruct::StarOrExpr>(
550+
"*" >> pure<std::optional<ScalarIntExpr>>() ||
551+
applyFunction(presentOptional<ScalarIntExpr>, scalarIntExpr))};
552+
constexpr auto gridOrBlock{parenthesized(nonemptyList(starOrExpr)) ||
553+
applyFunction(singletonList<CUFKernelDoConstruct::StarOrExpr>, starOrExpr)};
548554
TYPE_PARSER(sourced(beginDirective >> "$CUF KERNEL DO"_tok >>
549555
construct<CUFKernelDoConstruct::Directive>(
550-
maybe(parenthesized(scalarIntConstantExpr)),
551-
"<<<" >>
552-
("*" >> pure<std::list<ScalarIntExpr>>() ||
553-
parenthesized(nonemptyList(scalarIntExpr)) ||
554-
applyFunction(singletonList<ScalarIntExpr>, scalarIntExpr)),
555-
"," >> ("*" >> pure<std::list<ScalarIntExpr>>() ||
556-
parenthesized(nonemptyList(scalarIntExpr)) ||
557-
applyFunction(singletonList<ScalarIntExpr>, scalarIntExpr)),
556+
maybe(parenthesized(scalarIntConstantExpr)), "<<<" >> gridOrBlock,
557+
"," >> gridOrBlock,
558558
maybe((", 0 ,"_tok || ", STREAM ="_tok) >> scalarIntExpr) / ">>>" /
559559
endDirective)))
560560
TYPE_CONTEXT_PARSER("!$CUF KERNEL DO construct"_en_US,

‎flang/lib/Parser/misc-parsers.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -57,5 +57,10 @@ template <typename A> common::IfNoLvalue<std::list<A>, A> singletonList(A &&x) {
5757
result.emplace_back(std::move(x));
5858
return result;
5959
}
60+
61+
template <typename A>
62+
common::IfNoLvalue<std::optional<A>, A> presentOptional(A &&x) {
63+
return std::make_optional(std::move(x));
64+
}
6065
} // namespace Fortran::parser
6166
#endif

‎flang/lib/Parser/unparse.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2729,6 +2729,13 @@ class UnparseVisitor {
27292729
WALK_NESTED_ENUM(OmpOrderModifier, Kind) // OMP order-modifier
27302730
#undef WALK_NESTED_ENUM
27312731

2732+
void Unparse(const CUFKernelDoConstruct::StarOrExpr &x) {
2733+
if (x.v) {
2734+
Walk(*x.v);
2735+
} else {
2736+
Word("*");
2737+
}
2738+
}
27322739
void Unparse(const CUFKernelDoConstruct::Directive &x) {
27332740
Word("!$CUF KERNEL DO");
27342741
Walk(" (", std::get<std::optional<ScalarIntConstantExpr>>(x.t), ")");

‎flang/test/Lower/CUDA/cuda-kernel-loop-directive.cuf

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -42,10 +42,7 @@ subroutine sub1()
4242
! CHECK: fir.cuda_kernel<<<%c1{{.*}}, (%c256{{.*}}, %c1{{.*}})>>> (%{{.*}} : index, %{{.*}} : index) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index) step (%{{.*}}, %{{.*}} : index, index)
4343
! CHECK: {n = 2 : i64}
4444

45-
! TODO: currently these trigger error in the parser
45+
! TODO: lowering for these cases
4646
! !$cuf kernel do(2) <<< (1,*), (256,1) >>>
4747
! !$cuf kernel do(2) <<< (*,*), (32,4) >>>
4848
end
49-
50-
51-

‎flang/test/Parser/cuf-sanity-tree.CUF

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -144,11 +144,11 @@ include "cuf-sanity-common"
144144
!CHECK: | | | | | | EndDoStmt ->
145145
!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> CUFKernelDoConstruct
146146
!CHECK: | | | | | Directive
147-
!CHECK: | | | | | | Scalar -> Integer -> Expr = '1_4'
147+
!CHECK: | | | | | | StarOrExpr -> Scalar -> Integer -> Expr = '1_4'
148148
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
149-
!CHECK: | | | | | | Scalar -> Integer -> Expr = '2_4'
149+
!CHECK: | | | | | | StarOrExpr -> Scalar -> Integer -> Expr = '2_4'
150150
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
151-
!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4'
151+
!CHECK: | | | | | | StarOrExpr -> Scalar -> Integer -> Expr = '3_4'
152152
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3'
153153
!CHECK: | | | | | | Scalar -> Integer -> Expr = '1_4'
154154
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'

‎flang/test/Semantics/cuf09.cuf

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,15 @@ module m
1010
end
1111

1212
program main
13+
!$cuf kernel do <<< *, * >>> ! ok
14+
do j = 1, 0
15+
end do
16+
!$cuf kernel do <<< (*), (*) >>> ! ok
17+
do j = 1, 0
18+
end do
19+
!$cuf kernel do <<< (1,*), (2,*) >>> ! ok
20+
do j = 1, 0
21+
end do
1322
!ERROR: !$CUF KERNEL DO (1) must be followed by a DO construct with tightly nested outer levels of counted DO loops
1423
!$cuf kernel do <<< 1, 2 >>>
1524
do while (.false.)

0 commit comments

Comments
 (0)