diff --git a/flang/include/flang/Parser/dump-parse-tree.h b/flang/include/flang/Parser/dump-parse-tree.h index bfeb23de53539..675faeb33668f 100644 --- a/flang/include/flang/Parser/dump-parse-tree.h +++ b/flang/include/flang/Parser/dump-parse-tree.h @@ -177,6 +177,7 @@ class ParseTreeDumper { NODE(parser, Call) NODE(parser, CallStmt) NODE(CallStmt, Chevrons) + NODE(CallStmt, StarOrExpr) NODE(parser, CaseConstruct) NODE(CaseConstruct, Case) NODE(parser, CaseSelector) diff --git a/flang/include/flang/Parser/parse-tree.h b/flang/include/flang/Parser/parse-tree.h index d2c5b45d99581..37991739026b5 100644 --- a/flang/include/flang/Parser/parse-tree.h +++ b/flang/include/flang/Parser/parse-tree.h @@ -3247,13 +3247,14 @@ struct FunctionReference { // R1521 call-stmt -> CALL procedure-designator [ chevrons ] // [( [actual-arg-spec-list] )] -// (CUDA) chevrons -> <<< scalar-expr, scalar-expr [, +// (CUDA) chevrons -> <<< * | scalar-expr, scalar-expr [, // scalar-int-expr [, scalar-int-expr ] ] >>> struct CallStmt { BOILERPLATE(CallStmt); + WRAPPER_CLASS(StarOrExpr, std::optional); struct Chevrons { TUPLE_CLASS_BOILERPLATE(Chevrons); - std::tuple, + std::tuple, std::optional> t; }; diff --git a/flang/lib/Parser/program-parsers.cpp b/flang/lib/Parser/program-parsers.cpp index 2b7da18a09bb3..e365cd24a6aed 100644 --- a/flang/lib/Parser/program-parsers.cpp +++ b/flang/lib/Parser/program-parsers.cpp @@ -474,10 +474,13 @@ TYPE_CONTEXT_PARSER("function reference"_en_US, // R1521 call-stmt -> CALL procedure-designator [chevrons] /// [( [actual-arg-spec-list] )] -// (CUDA) chevrons -> <<< scalar-expr, scalar-expr [, scalar-int-expr +// (CUDA) chevrons -> <<< * | scalar-expr, scalar-expr [, scalar-int-expr // [, scalar-int-expr ] ] >>> +constexpr auto starOrExpr{ + construct("*" >> pure>() || + applyFunction(presentOptional, scalarExpr))}; TYPE_PARSER(extension( - "<<<" >> construct(scalarExpr, "," >> scalarExpr, + "<<<" >> construct(starOrExpr, ", " >> scalarExpr, maybe("," >> scalarIntExpr), maybe("," >> scalarIntExpr)) / ">>>")) constexpr auto actualArgSpecList{optionalList(actualArgSpec)}; diff --git a/flang/lib/Parser/unparse.cpp b/flang/lib/Parser/unparse.cpp index bbb126dcdb6d5..5d70f3433b445 100644 --- a/flang/lib/Parser/unparse.cpp +++ b/flang/lib/Parser/unparse.cpp @@ -1703,6 +1703,13 @@ class UnparseVisitor { void Unparse(const IntrinsicStmt &x) { // R1519 Word("INTRINSIC :: "), Walk(x.v, ", "); } + void Unparse(const CallStmt::StarOrExpr &x) { + if (x.v) { + Walk(*x.v); + } else { + Word("*"); + } + } void Unparse(const CallStmt::Chevrons &x) { // CUDA Walk(std::get<0>(x.t)); // grid Word(","), Walk(std::get<1>(x.t)); // block diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp index c70c8a8aecc2f..ead9982112678 100644 --- a/flang/lib/Semantics/expression.cpp +++ b/flang/lib/Semantics/expression.cpp @@ -3066,11 +3066,17 @@ std::optional ExpressionAnalyzer::AnalyzeChevrons( return false; }}; if (const auto &chevrons{call.chevrons}) { - if (auto expr{Analyze(std::get<0>(chevrons->t))}; - expr && checkLaunchArg(*expr, "grid")) { - result.emplace_back(*expr); + auto &starOrExpr{std::get<0>(chevrons->t)}; + if (starOrExpr.v) { + if (auto expr{Analyze(*starOrExpr.v)}; + expr && checkLaunchArg(*expr, "grid")) { + result.emplace_back(*expr); + } else { + return std::nullopt; + } } else { - return std::nullopt; + result.emplace_back( + AsGenericExpr(evaluate::Constant{-1})); } if (auto expr{Analyze(std::get<1>(chevrons->t))}; expr && checkLaunchArg(*expr, "block")) { diff --git a/flang/test/Parser/cuf-sanity-common b/flang/test/Parser/cuf-sanity-common index 9341f054d79d4..7005ef07b2265 100644 --- a/flang/test/Parser/cuf-sanity-common +++ b/flang/test/Parser/cuf-sanity-common @@ -40,6 +40,7 @@ module m call globalsub<<<1, 2>>> call globalsub<<<1, 2, 3>>> call globalsub<<<1, 2, 3, 4>>> + call globalsub<<<*,5>>> allocate(pa(32), pinned = isPinned) end subroutine end module diff --git a/flang/test/Parser/cuf-sanity-tree.CUF b/flang/test/Parser/cuf-sanity-tree.CUF index 2820441d5b5f0..a8b2f93913ca9 100644 --- a/flang/test/Parser/cuf-sanity-tree.CUF +++ b/flang/test/Parser/cuf-sanity-tree.CUF @@ -166,7 +166,7 @@ include "cuf-sanity-common" !CHECK: | | | | | Call !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' !CHECK: | | | | | Chevrons -!CHECK: | | | | | | Scalar -> Expr = '1_4' +!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '1_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1' !CHECK: | | | | | | Scalar -> Expr = '2_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2' @@ -174,7 +174,7 @@ include "cuf-sanity-common" !CHECK: | | | | | Call !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' !CHECK: | | | | | Chevrons -!CHECK: | | | | | | Scalar -> Expr = '1_4' +!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '1_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1' !CHECK: | | | | | | Scalar -> Expr = '2_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2' @@ -184,7 +184,7 @@ include "cuf-sanity-common" !CHECK: | | | | | Call !CHECK: | | | | | | ProcedureDesignator -> Name = 'globalsub' !CHECK: | | | | | Chevrons -!CHECK: | | | | | | Scalar -> Expr = '1_4' +!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '1_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1' !CHECK: | | | | | | Scalar -> Expr = '2_4' !CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2' diff --git a/flang/test/Parser/cuf-sanity-unparse.CUF b/flang/test/Parser/cuf-sanity-unparse.CUF index d4be347dd044e..2e2df9ac6646a 100644 --- a/flang/test/Parser/cuf-sanity-unparse.CUF +++ b/flang/test/Parser/cuf-sanity-unparse.CUF @@ -43,6 +43,7 @@ include "cuf-sanity-common" !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>>>() +!CHECK: CALL globalsub<<<-1_4,5_4>>>() !CHECK: ALLOCATE(pa(32_4), PINNED=ispinned) !CHECK: END SUBROUTINE !CHECK: END MODULE