[flang-commits] [flang] fe5a64d - [fang][cuda] Allow * in call chevron syntax (#115381)
via flang-commits
flang-commits at lists.llvm.org
Fri Nov 8 10:33:12 PST 2024
Author: Valentin Clement (バレンタイン クレメン)
Date: 2024-11-08T10:33:07-08:00
New Revision: fe5a64d1160209f22624b112b2629b0d6c4bb264
URL: https://github.com/llvm/llvm-project/commit/fe5a64d1160209f22624b112b2629b0d6c4bb264
DIFF: https://github.com/llvm/llvm-project/commit/fe5a64d1160209f22624b112b2629b0d6c4bb264.diff
LOG: [fang][cuda] Allow * in call chevron syntax (#115381)
Using `*` in call chevron syntax should be allowed. This patch updates
the parser to allow this usage.
```
call sub<<<*,nbBlock>>>()
```
Added:
Modified:
flang/include/flang/Parser/dump-parse-tree.h
flang/include/flang/Parser/parse-tree.h
flang/lib/Parser/program-parsers.cpp
flang/lib/Parser/unparse.cpp
flang/lib/Semantics/expression.cpp
flang/test/Parser/cuf-sanity-common
flang/test/Parser/cuf-sanity-tree.CUF
flang/test/Parser/cuf-sanity-unparse.CUF
Removed:
################################################################################
diff --git a/flang/include/flang/Parser/dump-parse-tree.h b/flang/include/flang/Parser/dump-parse-tree.h
index 456d53389cddfd..03d756eda858b9 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 e85187479380df..a0f8a5d22848fe 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<ScalarExpr>);
struct Chevrons {
TUPLE_CLASS_BOILERPLATE(Chevrons);
- std::tuple<ScalarExpr, ScalarExpr, std::optional<ScalarIntExpr>,
+ std::tuple<StarOrExpr, ScalarExpr, std::optional<ScalarIntExpr>,
std::optional<ScalarIntExpr>>
t;
};
diff --git a/flang/lib/Parser/program-parsers.cpp b/flang/lib/Parser/program-parsers.cpp
index 2b7da18a09bb30..e365cd24a6aed0 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<CallStmt::StarOrExpr>("*" >> pure<std::optional<ScalarExpr>>() ||
+ applyFunction(presentOptional<ScalarExpr>, scalarExpr))};
TYPE_PARSER(extension<LanguageFeature::CUDA>(
- "<<<" >> construct<CallStmt::Chevrons>(scalarExpr, "," >> scalarExpr,
+ "<<<" >> construct<CallStmt::Chevrons>(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 bbb126dcdb6d5e..5d70f3433b4453 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 c70c8a8aecc2f8..ead99821126787 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -3066,11 +3066,17 @@ std::optional<Chevrons> 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<evaluate::CInteger>{-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 9341f054d79d46..7005ef07b22650 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 2820441d5b5f0a..a8b2f93913ca9e 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 d4be347dd044ea..2e2df9ac6646a8 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
More information about the flang-commits
mailing list