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 1 commit
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
8 changes: 8 additions & 0 deletions flang/lib/Semantics/expression.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)};
Expand All @@ -3075,6 +3078,7 @@ std::optional<Chevrons> ExpressionAnalyzer::AnalyzeChevrons(
return std::nullopt;
}
} else {
gridIsStar = true;
result.emplace_back(
AsGenericExpr(evaluate::Constant<evaluate::SubscriptInteger>{-1}));
}
Expand All @@ -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);
Expand Down
1 change: 0 additions & 1 deletion flang/test/Parser/cuf-sanity-common
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
1 change: 0 additions & 1 deletion 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,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)
Expand Down
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