Skip to content

[fang][cuda] Allow * in call chevron syntax #115381

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 6 commits into from
Nov 8, 2024
Merged

[fang][cuda] Allow * in call chevron syntax #115381

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
7f99ff7
[fang][cuda] Allow * in call chevron syntax
clementval Nov 7, 2024
c9dabfa
Allow only one *
clementval Nov 8, 2024
c5a48ca
Only allow * for grid
clementval Nov 8, 2024
0ae853b
Update syntax
clementval Nov 8, 2024
66080d6
Fix comment
clementval Nov 8, 2024
1db337b
Use kind 4
clementval Nov 8, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions flang/include/flang/Parser/dump-parse-tree.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
5 changes: 3 additions & 2 deletions flang/include/flang/Parser/parse-tree.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
};
Expand Down
7 changes: 5 additions & 2 deletions flang/lib/Parser/program-parsers.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)};
Expand Down
7 changes: 7 additions & 0 deletions flang/lib/Parser/unparse.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
38 changes: 29 additions & 9 deletions flang/lib/Semantics/expression.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3065,18 +3065,38 @@ std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
which);
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);
bool gridIsStar{false};
bool blockIsStar{false};
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;
}
if (auto expr{Analyze(std::get<1>(chevrons->t))};
expr && checkLaunchArg(*expr, "block")) {
result.emplace_back(*expr);
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;
}
} else {
return std::nullopt;
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)}) {
Expand Down
2 changes: 2 additions & 0 deletions flang/test/Parser/cuf-sanity-common
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,8 @@ module m
call globalsub<<<1, 2>>>
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
12 changes: 6 additions & 6 deletions flang/test/Parser/cuf-sanity-tree.CUF
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -166,27 +166,27 @@ 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'
!CHECK: | | | | ExecutionPartConstruct -> ExecutableConstruct -> ActionStmt -> CallStmt = 'CALL globalsub<<<1_4,2_4,3_4,4_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'
Expand Down
2 changes: 2 additions & 0 deletions flang/test/Parser/cuf-sanity-unparse.CUF
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,8 @@ 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,-1_8>>>()
!CHECK: ALLOCATE(pa(32_4), PINNED=ispinned)
!CHECK: END SUBROUTINE
!CHECK: END MODULE
2 changes: 2 additions & 0 deletions flang/test/Semantics/cuf04.cuf
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Loading