[fang][cuda] Allow * in call chevron syntax #115381
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
flang
Member
|
@llvm/pr-subscribers-flang-semantics @llvm/pr-subscribers-flang-parser Author: Valentin Clement (バレンタイン クレメン) (clementval) ChangesUsing Full diff: https://github.com/llvm/llvm-project/pull/115381.diff 8 Files Affected:
diff --git a/flang/include/flang/Parser/dump-parse-tree.h b/flang/include/flang/Parser/dump-parse-tree.h
index bfeb23de535392..675faeb33668f3 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 d2c5b45d995813..f84fd7565b006d 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, StarOrExpr, 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..a11b9b87765f81 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, ", " >> starOrExpr,
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..e380d9532ee181 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -3066,17 +3066,29 @@ 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 &starOrExpr0{std::get<0>(chevrons->t)};
+ if (starOrExpr0.v) {
+ if (auto expr{Analyze(*starOrExpr0.v)};
+ expr && checkLaunchArg(*expr, "grid")) {
+ result.emplace_back(*expr);
+ } else {
+ return std::nullopt;
+ }
} else {
- return std::nullopt;
+ result.emplace_back(
+ AsGenericExpr(evaluate::Constant<evaluate::SubscriptInteger>{-1}));
}
- if (auto expr{Analyze(std::get<1>(chevrons->t))};
- expr && checkLaunchArg(*expr, "block")) {
- result.emplace_back(*expr);
+ auto &starOrExpr1{std::get<1>(chevrons->t)};
+ if (starOrExpr1.v) {
+ if (auto expr{Analyze(*starOrExpr1.v)};
+ expr && checkLaunchArg(*expr, "block")) {
+ result.emplace_back(*expr);
+ } else {
+ return std::nullopt;
+ }
} else {
- return std::nullopt;
+ result.emplace_back(
+ AsGenericExpr(evaluate::Constant<evaluate::SubscriptInteger>{-1}));
}
if (const auto &maybeExpr{std::get<2>(chevrons->t)}) {
if (auto expr{Analyze(*maybeExpr)}) {
diff --git a/flang/test/Parser/cuf-sanity-common b/flang/test/Parser/cuf-sanity-common
index 9341f054d79d46..d08048058adbec 100644
--- a/flang/test/Parser/cuf-sanity-common
+++ b/flang/test/Parser/cuf-sanity-common
@@ -40,6 +40,9 @@ module m
call globalsub<<<1, 2>>>
call globalsub<<<1, 2, 3>>>
call globalsub<<<1, 2, 3, 4>>>
+ call globalsub<<<*,*>>>
+ call globalsub<<<*,5>>>
+ call globalsub<<<1,*>>>
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..7f097ab6c9c659 100644
--- a/flang/test/Parser/cuf-sanity-tree.CUF
+++ b/flang/test/Parser/cuf-sanity-tree.CUF
@@ -166,17 +166,17 @@ 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: | | | | | | StarOrExpr -> Scalar -> Expr = '2_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4,3_4>>>()'
!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: | | | | | | StarOrExpr -> Scalar -> Expr = '2_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3'
@@ -184,9 +184,9 @@ 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: | | | | | | StarOrExpr -> Scalar -> Expr = '2_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3'
diff --git a/flang/test/Parser/cuf-sanity-unparse.CUF b/flang/test/Parser/cuf-sanity-unparse.CUF
index d4be347dd044ea..938caa5982c6e5 100644
--- a/flang/test/Parser/cuf-sanity-unparse.CUF
+++ b/flang/test/Parser/cuf-sanity-unparse.CUF
@@ -43,6 +43,9 @@ 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_8,-1_8>>>()
+!CHECK: CALL globalsub<<<-1_8,5_4>>>()
+!CHECK: CALL globalsub<<<1_4,-1_8>>>()
!CHECK: ALLOCATE(pa(32_4), PINNED=ispinned)
!CHECK: END SUBROUTINE
!CHECK: END MODULE
|
wangzpgi
reviewed
Nov 8, 2024
clementval
commented
Nov 8, 2024
clementval
force-pushed
the
cuf_chevron_star
branch
from
d02922c to
c5a48ca
Compare
wangzpgi
approved these changes
Nov 8, 2024
Groverkss
pushed a commit
to iree-org/llvm-project
that referenced
this pull request
Nov 15, 2024
Using `*` in call chevron syntax should be allowed. This patch updates the parser to allow this usage. ``` call sub<<<*,nbBlock>>>() ```
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Using
*in call chevron syntax should be allowed. This patch updates the parser to allow this usage.