[flang-commits] [flang] [fang][cuda] Allow * in call chevron syntax (PR #115381)
Valentin Clement バレンタイン クレメン via flang-commits
flang-commits at lists.llvm.org
Fri Nov 8 10:31:42 PST 2024
https://github.com/clementval updated https://github.com/llvm/llvm-project/pull/115381
>From 7f99ff74fe20f3423906c1a2a5326f7271d5fa41 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 7 Nov 2024 14:26:00 -0800
Subject: [PATCH 1/6] [fang][cuda] Allow * in call chevron syntax
---
flang/include/flang/Parser/dump-parse-tree.h | 1 +
flang/include/flang/Parser/parse-tree.h | 5 ++--
flang/lib/Parser/program-parsers.cpp | 7 +++--
flang/lib/Parser/unparse.cpp | 7 +++++
flang/lib/Semantics/expression.cpp | 28 ++++++++++++++------
flang/test/Parser/cuf-sanity-common | 3 +++
flang/test/Parser/cuf-sanity-tree.CUF | 12 ++++-----
flang/test/Parser/cuf-sanity-unparse.CUF | 3 +++
8 files changed, 48 insertions(+), 18 deletions(-)
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
>From c9dabfae6dc10f0d0e2a417ef2e7436b2b327041 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 7 Nov 2024 16:35:50 -0800
Subject: [PATCH 2/6] Allow only one *
---
flang/lib/Semantics/expression.cpp | 8 ++++++++
flang/test/Parser/cuf-sanity-common | 1 -
flang/test/Parser/cuf-sanity-unparse.CUF | 1 -
flang/test/Semantics/cuf04.cuf | 2 ++
4 files changed, 10 insertions(+), 2 deletions(-)
diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp
index e380d9532ee181..b492fe1291b80b 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -3065,7 +3065,10 @@ std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
which);
return false;
}};
+
if (const auto &chevrons{call.chevrons}) {
+ bool gridIsStar{false};
+ bool blockIsStar{false};
auto &starOrExpr0{std::get<0>(chevrons->t)};
if (starOrExpr0.v) {
if (auto expr{Analyze(*starOrExpr0.v)};
@@ -3075,6 +3078,7 @@ std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
return std::nullopt;
}
} else {
+ gridIsStar = true;
result.emplace_back(
AsGenericExpr(evaluate::Constant<evaluate::SubscriptInteger>{-1}));
}
@@ -3087,9 +3091,13 @@ std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
return std::nullopt;
}
} else {
+ blockIsStar = true;
result.emplace_back(
AsGenericExpr(evaluate::Constant<evaluate::SubscriptInteger>{-1}));
}
+ if (gridIsStar && blockIsStar) {
+ Say("Grid and block can not be * in kernel launch parameter"_err_en_US);
+ }
if (const auto &maybeExpr{std::get<2>(chevrons->t)}) {
if (auto expr{Analyze(*maybeExpr)}) {
result.emplace_back(*expr);
diff --git a/flang/test/Parser/cuf-sanity-common b/flang/test/Parser/cuf-sanity-common
index d08048058adbec..ed8ffd09768257 100644
--- a/flang/test/Parser/cuf-sanity-common
+++ b/flang/test/Parser/cuf-sanity-common
@@ -40,7 +40,6 @@ 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)
diff --git a/flang/test/Parser/cuf-sanity-unparse.CUF b/flang/test/Parser/cuf-sanity-unparse.CUF
index 938caa5982c6e5..9345e837d9e184 100644
--- a/flang/test/Parser/cuf-sanity-unparse.CUF
+++ b/flang/test/Parser/cuf-sanity-unparse.CUF
@@ -43,7 +43,6 @@ 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)
diff --git a/flang/test/Semantics/cuf04.cuf b/flang/test/Semantics/cuf04.cuf
index 2e2faa90b490db..32b2102ec43072 100644
--- a/flang/test/Semantics/cuf04.cuf
+++ b/flang/test/Semantics/cuf04.cuf
@@ -20,5 +20,7 @@ module m
call globsubr
!ERROR: Kernel launch parameters in chevrons may not be used unless calling a kernel subroutine
call boring<<<1,2>>>
+ !ERROR: Grid and block can not be * in kernel launch parameter
+ call globsubr<<<*, *>>>
end subroutine
end module
>From c5a48ca2a9f27ad2d3000c3567b439d6bf8c3575 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 7 Nov 2024 20:47:46 -0800
Subject: [PATCH 3/6] Only allow * for grid
---
flang/include/flang/Parser/parse-tree.h | 2 +-
flang/lib/Parser/program-parsers.cpp | 2 +-
flang/lib/Semantics/expression.cpp | 21 ++++-----------------
flang/test/Parser/cuf-sanity-common | 1 -
flang/test/Parser/cuf-sanity-tree.CUF | 6 +++---
flang/test/Parser/cuf-sanity-unparse.CUF | 1 -
flang/test/Semantics/cuf04.cuf | 2 --
7 files changed, 9 insertions(+), 26 deletions(-)
diff --git a/flang/include/flang/Parser/parse-tree.h b/flang/include/flang/Parser/parse-tree.h
index f84fd7565b006d..8898e30a336940 100644
--- a/flang/include/flang/Parser/parse-tree.h
+++ b/flang/include/flang/Parser/parse-tree.h
@@ -3254,7 +3254,7 @@ struct CallStmt {
WRAPPER_CLASS(StarOrExpr, std::optional<ScalarExpr>);
struct Chevrons {
TUPLE_CLASS_BOILERPLATE(Chevrons);
- std::tuple<StarOrExpr, StarOrExpr, 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 a11b9b87765f81..29272da7f899bb 100644
--- a/flang/lib/Parser/program-parsers.cpp
+++ b/flang/lib/Parser/program-parsers.cpp
@@ -480,7 +480,7 @@ constexpr auto starOrExpr{
construct<CallStmt::StarOrExpr>("*" >> pure<std::optional<ScalarExpr>>() ||
applyFunction(presentOptional<ScalarExpr>, scalarExpr))};
TYPE_PARSER(extension<LanguageFeature::CUDA>(
- "<<<" >> construct<CallStmt::Chevrons>(starOrExpr, ", " >> starOrExpr,
+ "<<<" >> construct<CallStmt::Chevrons>(starOrExpr, ", " >> scalarExpr,
maybe("," >> scalarIntExpr), maybe("," >> scalarIntExpr)) /
">>>"))
constexpr auto actualArgSpecList{optionalList(actualArgSpec)};
diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp
index b492fe1291b80b..27da42bfade0c5 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -3067,8 +3067,6 @@ std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
}};
if (const auto &chevrons{call.chevrons}) {
- bool gridIsStar{false};
- bool blockIsStar{false};
auto &starOrExpr0{std::get<0>(chevrons->t)};
if (starOrExpr0.v) {
if (auto expr{Analyze(*starOrExpr0.v)};
@@ -3078,25 +3076,14 @@ std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
return std::nullopt;
}
} else {
- gridIsStar = true;
result.emplace_back(
AsGenericExpr(evaluate::Constant<evaluate::SubscriptInteger>{-1}));
}
- 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;
- }
+ if (auto expr{Analyze(std::get<1>(chevrons->t))};
+ expr && checkLaunchArg(*expr, "block")) {
+ result.emplace_back(*expr);
} else {
- blockIsStar = true;
- result.emplace_back(
- AsGenericExpr(evaluate::Constant<evaluate::SubscriptInteger>{-1}));
- }
- if (gridIsStar && blockIsStar) {
- Say("Grid and block can not be * in kernel launch parameter"_err_en_US);
+ return std::nullopt;
}
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 ed8ffd09768257..7005ef07b22650 100644
--- a/flang/test/Parser/cuf-sanity-common
+++ b/flang/test/Parser/cuf-sanity-common
@@ -41,7 +41,6 @@ module m
call globalsub<<<1, 2, 3>>>
call globalsub<<<1, 2, 3, 4>>>
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 7f097ab6c9c659..a8b2f93913ca9e 100644
--- a/flang/test/Parser/cuf-sanity-tree.CUF
+++ b/flang/test/Parser/cuf-sanity-tree.CUF
@@ -168,7 +168,7 @@ include "cuf-sanity-common"
!CHECK: | | | | | Chevrons
!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '1_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
-!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '2_4'
+!CHECK: | | | | | | Scalar -> Expr = '2_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4,3_4>>>()'
!CHECK: | | | | | Call
@@ -176,7 +176,7 @@ include "cuf-sanity-common"
!CHECK: | | | | | Chevrons
!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '1_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
-!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '2_4'
+!CHECK: | | | | | | Scalar -> Expr = '2_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '2'
!CHECK: | | | | | | Scalar -> Integer -> Expr = '3_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '3'
@@ -186,7 +186,7 @@ include "cuf-sanity-common"
!CHECK: | | | | | Chevrons
!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '1_4'
!CHECK: | | | | | | | LiteralConstant -> IntLiteralConstant = '1'
-!CHECK: | | | | | | StarOrExpr -> Scalar -> Expr = '2_4'
+!CHECK: | | | | | | 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 9345e837d9e184..2c6d2e3528e83e 100644
--- a/flang/test/Parser/cuf-sanity-unparse.CUF
+++ b/flang/test/Parser/cuf-sanity-unparse.CUF
@@ -44,7 +44,6 @@ include "cuf-sanity-common"
!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,5_4>>>()
-!CHECK: CALL globalsub<<<1_4,-1_8>>>()
!CHECK: ALLOCATE(pa(32_4), PINNED=ispinned)
!CHECK: END SUBROUTINE
!CHECK: END MODULE
diff --git a/flang/test/Semantics/cuf04.cuf b/flang/test/Semantics/cuf04.cuf
index 32b2102ec43072..2e2faa90b490db 100644
--- a/flang/test/Semantics/cuf04.cuf
+++ b/flang/test/Semantics/cuf04.cuf
@@ -20,7 +20,5 @@ module m
call globsubr
!ERROR: Kernel launch parameters in chevrons may not be used unless calling a kernel subroutine
call boring<<<1,2>>>
- !ERROR: Grid and block can not be * in kernel launch parameter
- call globsubr<<<*, *>>>
end subroutine
end module
>From 0ae853b472fda6f5bd9ae5d260bd1f5a9713c7b1 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Fri, 8 Nov 2024 10:12:29 -0800
Subject: [PATCH 4/6] Update syntax
---
flang/lib/Parser/program-parsers.cpp | 2 +-
flang/lib/Semantics/expression.cpp | 7 +++----
2 files changed, 4 insertions(+), 5 deletions(-)
diff --git a/flang/lib/Parser/program-parsers.cpp b/flang/lib/Parser/program-parsers.cpp
index 29272da7f899bb..e365cd24a6aed0 100644
--- a/flang/lib/Parser/program-parsers.cpp
+++ b/flang/lib/Parser/program-parsers.cpp
@@ -474,7 +474,7 @@ 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>>() ||
diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp
index 27da42bfade0c5..d75d9604ceebfa 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -3065,11 +3065,10 @@ std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
which);
return false;
}};
-
if (const auto &chevrons{call.chevrons}) {
- auto &starOrExpr0{std::get<0>(chevrons->t)};
- if (starOrExpr0.v) {
- if (auto expr{Analyze(*starOrExpr0.v)};
+ 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 {
>From 66080d694d42f9709083648829462cc1db5b69fc Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Fri, 8 Nov 2024 10:22:37 -0800
Subject: [PATCH 5/6] Fix comment
---
flang/include/flang/Parser/parse-tree.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/flang/include/flang/Parser/parse-tree.h b/flang/include/flang/Parser/parse-tree.h
index 8898e30a336940..37991739026b5a 100644
--- a/flang/include/flang/Parser/parse-tree.h
+++ b/flang/include/flang/Parser/parse-tree.h
@@ -3247,7 +3247,7 @@ 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);
>From 1db337bb5b389056f4eb89d31a808b5f8b83f4ac Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Fri, 8 Nov 2024 10:31:26 -0800
Subject: [PATCH 6/6] Use kind 4
---
flang/lib/Semantics/expression.cpp | 2 +-
flang/test/Parser/cuf-sanity-unparse.CUF | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp
index d75d9604ceebfa..ead99821126787 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -3076,7 +3076,7 @@ std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
}
} else {
result.emplace_back(
- AsGenericExpr(evaluate::Constant<evaluate::SubscriptInteger>{-1}));
+ 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-unparse.CUF b/flang/test/Parser/cuf-sanity-unparse.CUF
index 2c6d2e3528e83e..2e2df9ac6646a8 100644
--- a/flang/test/Parser/cuf-sanity-unparse.CUF
+++ b/flang/test/Parser/cuf-sanity-unparse.CUF
@@ -43,7 +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_8,5_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