Skip to content

Commit e68153d

Browse files
saiislamsaiislam
authored
merge main into amd-staging (llvm#1215)
2 parents 02df29a + 95adcf2 commit e68153d

26 files changed

+1550
-556
lines changed

‎clang/lib/Format/FormatToken.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -746,6 +746,10 @@ struct FormatToken {
746746
return isOneOf(tok::star, tok::amp, tok::ampamp);
747747
}
748748

749+
bool isPlacementOperator() const {
750+
return isOneOf(tok::kw_new, tok::kw_delete);
751+
}
752+
749753
bool isUnaryOperator() const {
750754
switch (Tok.getKind()) {
751755
case tok::plus:

‎clang/lib/Format/FormatTokenLexer.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -610,9 +610,9 @@ bool FormatTokenLexer::precedesOperand(FormatToken *Tok) {
610610
tok::r_brace, tok::l_square, tok::semi, tok::exclaim,
611611
tok::colon, tok::question, tok::tilde) ||
612612
Tok->isOneOf(tok::kw_return, tok::kw_do, tok::kw_case, tok::kw_throw,
613-
tok::kw_else, tok::kw_new, tok::kw_delete, tok::kw_void,
614-
tok::kw_typeof, Keywords.kw_instanceof, Keywords.kw_in) ||
615-
Tok->isBinaryOperator();
613+
tok::kw_else, tok::kw_void, tok::kw_typeof,
614+
Keywords.kw_instanceof, Keywords.kw_in) ||
615+
Tok->isPlacementOperator() || Tok->isBinaryOperator();
616616
}
617617

618618
bool FormatTokenLexer::canPrecedeRegexLiteral(FormatToken *Prev) {

‎clang/lib/Format/TokenAnnotator.cpp

Lines changed: 32 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1639,6 +1639,31 @@ class AnnotatingParser {
16391639
case tok::kw_operator:
16401640
if (Style.isProto())
16411641
break;
1642+
// Handle C++ user-defined conversion function.
1643+
if (IsCpp && CurrentToken) {
1644+
const auto *Info = CurrentToken->Tok.getIdentifierInfo();
1645+
// What follows Tok is an identifier or a non-operator keyword.
1646+
if (Info && !(CurrentToken->isPlacementOperator() ||
1647+
CurrentToken->is(tok::kw_co_await) ||
1648+
Info->isCPlusPlusOperatorKeyword())) {
1649+
FormatToken *LParen;
1650+
if (CurrentToken->startsSequence(tok::kw_decltype, tok::l_paren,
1651+
tok::kw_auto, tok::r_paren)) {
1652+
// Skip `decltype(auto)`.
1653+
LParen = CurrentToken->Next->Next->Next->Next;
1654+
} else {
1655+
// Skip to l_paren.
1656+
for (LParen = CurrentToken->Next;
1657+
LParen && LParen->isNot(tok::l_paren); LParen = LParen->Next) {
1658+
}
1659+
}
1660+
if (LParen && LParen->is(tok::l_paren)) {
1661+
Tok->setFinalizedType(TT_FunctionDeclarationName);
1662+
LParen->setFinalizedType(TT_FunctionDeclarationLParen);
1663+
break;
1664+
}
1665+
}
1666+
}
16421667
while (CurrentToken &&
16431668
!CurrentToken->isOneOf(tok::l_paren, tok::semi, tok::r_paren)) {
16441669
if (CurrentToken->isOneOf(tok::star, tok::amp))
@@ -2999,7 +3024,7 @@ class AnnotatingParser {
29993024
return TT_UnaryOperator;
30003025
if (PrevToken->is(TT_TypeName))
30013026
return TT_PointerOrReference;
3002-
if (PrevToken->isOneOf(tok::kw_new, tok::kw_delete) && Tok.is(tok::ampamp))
3027+
if (PrevToken->isPlacementOperator() && Tok.is(tok::ampamp))
30033028
return TT_BinaryOperator;
30043029

30053030
const FormatToken *NextToken = Tok.getNextNonComment();
@@ -3071,12 +3096,10 @@ class AnnotatingParser {
30713096
if (InTemplateArgument && NextToken->Tok.isAnyIdentifier())
30723097
return TT_BinaryOperator;
30733098

3074-
// "&&" followed by "(", "*", or "&" is quite unlikely to be two successive
3075-
// unary "&".
3076-
if (Tok.is(tok::ampamp) &&
3077-
NextToken->isOneOf(tok::l_paren, tok::star, tok::amp)) {
3099+
// "&&" followed by "*" or "&" is quite unlikely to be two successive unary
3100+
// "&".
3101+
if (Tok.is(tok::ampamp) && NextToken->isOneOf(tok::star, tok::amp))
30783102
return TT_BinaryOperator;
3079-
}
30803103

30813104
// This catches some cases where evaluation order is used as control flow:
30823105
// aaa && aaa->f();
@@ -3791,7 +3814,7 @@ static bool isFunctionDeclarationName(const LangOptions &LangOpts,
37913814
return Next;
37923815
if (Next->is(TT_OverloadedOperator))
37933816
continue;
3794-
if (Next->isOneOf(tok::kw_new, tok::kw_delete, tok::kw_co_await)) {
3817+
if (Next->isPlacementOperator() || Next->is(tok::kw_co_await)) {
37953818
// For 'new[]' and 'delete[]'.
37963819
if (Next->Next &&
37973820
Next->Next->startsSequence(tok::l_square, tok::r_square)) {
@@ -4802,7 +4825,7 @@ bool TokenAnnotator::spaceRequiredBetween(const AnnotatedLine &Line,
48024825
spaceRequiredBeforeParens(Right);
48034826
}
48044827
if (Style.SpaceBeforeParens == FormatStyle::SBPO_Custom &&
4805-
Left.isOneOf(tok::kw_new, tok::kw_delete) &&
4828+
Left.isPlacementOperator() &&
48064829
Right.isNot(TT_OverloadedOperatorLParen) &&
48074830
!(Line.MightBeFunctionDecl && Left.is(TT_FunctionDeclarationName))) {
48084831
const auto *RParen = Right.MatchingParen;
@@ -4845,7 +4868,7 @@ bool TokenAnnotator::spaceRequiredBetween(const AnnotatedLine &Line,
48454868
return Style.SpaceBeforeParensOptions.AfterControlStatements ||
48464869
spaceRequiredBeforeParens(Right);
48474870
}
4848-
if (Left.isOneOf(tok::kw_new, tok::kw_delete) ||
4871+
if (Left.isPlacementOperator() ||
48494872
(Left.is(tok::r_square) && Left.MatchingParen &&
48504873
Left.MatchingParen->Previous &&
48514874
Left.MatchingParen->Previous->is(tok::kw_delete))) {

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2923,10 +2923,16 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
29232923
if (!S.checkUInt32Argument(AL, E, WGSize[i], i,
29242924
/*StrictlyUnsigned=*/true))
29252925
return;
2926-
if (WGSize[i] == 0) {
2927-
S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero)
2928-
<< AL << E->getSourceRange();
2929-
return;
2926+
}
2927+
2928+
if (!llvm::all_of(WGSize, [](uint32_t Size) { return Size == 0; })) {
2929+
for (unsigned i = 0; i < 3; ++i) {
2930+
const Expr *E = AL.getArgAsExpr(i);
2931+
if (WGSize[i] == 0) {
2932+
S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero)
2933+
<< AL << E->getSourceRange();
2934+
return;
2935+
}
29302936
}
29312937
}
29322938

‎clang/test/CodeGenCUDASPIRV/spirv-attrs.cu

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,9 @@
33

44
#define __global__ __attribute__((global))
55

6+
__attribute__((reqd_work_group_size(0, 0, 0)))
7+
__global__ void reqd_work_group_size_0_0_0() {}
8+
69
__attribute__((reqd_work_group_size(128, 1, 1)))
710
__global__ void reqd_work_group_size_128_1_1() {}
811

@@ -15,13 +18,16 @@ __global__ void vec_type_hint_int() {}
1518
__attribute__((intel_reqd_sub_group_size(64)))
1619
__global__ void intel_reqd_sub_group_size_64() {}
1720

21+
22+
// CHECK: define spir_kernel void @_Z26reqd_work_group_size_0_0_0v() #[[ATTR:[0-9]+]] !reqd_work_group_size ![[WG_SIZE_ZEROS:[0-9]+]]
1823
// CHECK: define spir_kernel void @_Z28reqd_work_group_size_128_1_1v() #[[ATTR:[0-9]+]] !reqd_work_group_size ![[WG_SIZE:[0-9]+]]
1924
// CHECK: define spir_kernel void @_Z26work_group_size_hint_2_2_2v() #[[ATTR]] !work_group_size_hint ![[WG_HINT:[0-9]+]]
2025
// CHECK: define spir_kernel void @_Z17vec_type_hint_intv() #[[ATTR]] !vec_type_hint ![[VEC_HINT:[0-9]+]]
2126
// CHECK: define spir_kernel void @_Z28intel_reqd_sub_group_size_64v() #[[ATTR]] !intel_reqd_sub_group_size ![[SUB_GRP:[0-9]+]]
2227

2328
// CHECK: attributes #[[ATTR]] = { {{.*}} }
2429

30+
// CHECK: ![[WG_SIZE_ZEROS]] = !{i32 0, i32 0, i32 0}
2531
// CHECK: ![[WG_SIZE]] = !{i32 128, i32 1, i32 1}
2632
// CHECK: ![[WG_HINT]] = !{i32 2, i32 2, i32 2}
2733
// CHECK: ![[VEC_HINT]] = !{i32 undef, i32 1}

‎clang/test/SemaOpenCL/invalid-kernel-attrs.cl

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -44,3 +44,5 @@ __kernel __attribute__((reqd_work_group_size(8,16,-32))) void neg2(){} // expect
4444
// 4294967294 is a negative integer if treated as signed.
4545
// Should compile successfully, since we expect an unsigned.
4646
__kernel __attribute__((reqd_work_group_size(8,16,4294967294))) void ok1(){}
47+
48+
__kernel __attribute__((reqd_work_group_size(0,0,0))) void ok_zeros(){}

‎clang/tools/clang-format/git-clang-format

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@ clang-format on the changes in current files or a specific commit.
1919
For further details, run:
2020
git clang-format -h
2121
22-
Requires Python 2.7 or Python 3
22+
Requires Python version >=3.8
2323
"""
2424

2525
from __future__ import absolute_import, division, print_function

‎clang/unittests/Format/FormatTest.cpp

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -10443,27 +10443,17 @@ TEST_F(FormatTest, ReturnTypeBreakingStyle) {
1044310443
"void\n"
1044410444
"A::operator->() {}\n"
1044510445
"void\n"
10446-
"A::operator void *() {}\n"
10446+
"A::operator&() {}\n"
1044710447
"void\n"
10448-
"A::operator void &() {}\n"
10449-
"void\n"
10450-
"A::operator void &&() {}\n"
10451-
"void\n"
10452-
"A::operator char *() {}\n"
10448+
"A::operator&&() {}\n"
1045310449
"void\n"
1045410450
"A::operator[]() {}\n"
1045510451
"void\n"
1045610452
"A::operator!() {}\n"
1045710453
"void\n"
10458-
"A::operator**() {}\n"
10459-
"void\n"
1046010454
"A::operator<Foo> *() {}\n"
1046110455
"void\n"
10462-
"A::operator<Foo> **() {}\n"
10463-
"void\n"
10464-
"A::operator<Foo> &() {}\n"
10465-
"void\n"
10466-
"A::operator void **() {}",
10456+
"A::operator<Foo> &() {}\n",
1046710457
Style);
1046810458
verifyFormat("constexpr auto\n"
1046910459
"operator()() const -> reference {}\n"
@@ -10486,7 +10476,7 @@ TEST_F(FormatTest, ReturnTypeBreakingStyle) {
1048610476
"constexpr auto\n"
1048710477
"operator void &() const -> reference {}\n"
1048810478
"constexpr auto\n"
10489-
"operator void &&() const -> reference {}\n"
10479+
"operator&&() const -> reference {}\n"
1049010480
"constexpr auto\n"
1049110481
"operator char *() const -> reference {}\n"
1049210482
"constexpr auto\n"
@@ -28032,6 +28022,16 @@ TEST_F(FormatTest, BreakAfterAttributes) {
2803228022
" --d;",
2803328023
CtrlStmtCode, Style);
2803428024

28025+
verifyFormat("[[nodiscard]]\n"
28026+
"operator bool();\n"
28027+
"[[nodiscard]]\n"
28028+
"operator bool() {\n"
28029+
" return true;\n"
28030+
"}",
28031+
"[[nodiscard]] operator bool();\n"
28032+
"[[nodiscard]] operator bool() { return true; }",
28033+
Style);
28034+
2803528035
constexpr StringRef CtorDtorCode("struct Foo {\n"
2803628036
" [[deprecated]] Foo();\n"
2803728037
" [[deprecated]] Foo() {}\n"

‎clang/unittests/Format/TokenAnnotatorTest.cpp

Lines changed: 50 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3856,6 +3856,56 @@ TEST_F(TokenAnnotatorTest, AfterPPDirective) {
38563856
EXPECT_TOKEN(Tokens[2], tok::minusminus, TT_AfterPPDirective);
38573857
}
38583858

3859+
TEST_F(TokenAnnotatorTest, UserDefinedConversionFunction) {
3860+
auto Tokens = annotate("operator int(void);");
3861+
ASSERT_EQ(Tokens.size(), 7u) << Tokens;
3862+
EXPECT_TOKEN(Tokens[0], tok::kw_operator, TT_FunctionDeclarationName);
3863+
EXPECT_TOKEN(Tokens[2], tok::l_paren, TT_FunctionDeclarationLParen);
3864+
3865+
Tokens = annotate("explicit operator int *();");
3866+
ASSERT_EQ(Tokens.size(), 8u) << Tokens;
3867+
EXPECT_TOKEN(Tokens[1], tok::kw_operator, TT_FunctionDeclarationName);
3868+
EXPECT_TOKEN(Tokens[3], tok::star, TT_PointerOrReference);
3869+
EXPECT_TOKEN(Tokens[4], tok::l_paren, TT_FunctionDeclarationLParen);
3870+
3871+
Tokens = annotate("operator int &();");
3872+
ASSERT_EQ(Tokens.size(), 7u) << Tokens;
3873+
EXPECT_TOKEN(Tokens[0], tok::kw_operator, TT_FunctionDeclarationName);
3874+
EXPECT_TOKEN(Tokens[2], tok::amp, TT_PointerOrReference);
3875+
EXPECT_TOKEN(Tokens[3], tok::l_paren, TT_FunctionDeclarationLParen);
3876+
3877+
Tokens = annotate("operator auto() const { return 2; }");
3878+
ASSERT_EQ(Tokens.size(), 11u) << Tokens;
3879+
EXPECT_TOKEN(Tokens[0], tok::kw_operator, TT_FunctionDeclarationName);
3880+
EXPECT_TOKEN(Tokens[2], tok::l_paren, TT_FunctionDeclarationLParen);
3881+
EXPECT_TOKEN(Tokens[4], tok::kw_const, TT_TrailingAnnotation);
3882+
EXPECT_TOKEN(Tokens[5], tok::l_brace, TT_FunctionLBrace);
3883+
3884+
Tokens = annotate("operator decltype(auto)() const;");
3885+
ASSERT_EQ(Tokens.size(), 10u) << Tokens;
3886+
EXPECT_TOKEN(Tokens[0], tok::kw_operator, TT_FunctionDeclarationName);
3887+
EXPECT_TOKEN(Tokens[2], tok::l_paren, TT_TypeDeclarationParen);
3888+
EXPECT_TOKEN(Tokens[4], tok::r_paren, TT_TypeDeclarationParen);
3889+
EXPECT_TOKEN(Tokens[5], tok::l_paren, TT_FunctionDeclarationLParen);
3890+
EXPECT_TOKEN(Tokens[7], tok::kw_const, TT_TrailingAnnotation);
3891+
3892+
Tokens = annotate("virtual operator Foo() = 0;");
3893+
ASSERT_EQ(Tokens.size(), 9u) << Tokens;
3894+
EXPECT_TOKEN(Tokens[1], tok::kw_operator, TT_FunctionDeclarationName);
3895+
EXPECT_TOKEN(Tokens[3], tok::l_paren, TT_FunctionDeclarationLParen);
3896+
3897+
Tokens = annotate("operator Foo() override { return Foo(); }");
3898+
ASSERT_EQ(Tokens.size(), 13u) << Tokens;
3899+
EXPECT_TOKEN(Tokens[0], tok::kw_operator, TT_FunctionDeclarationName);
3900+
EXPECT_TOKEN(Tokens[2], tok::l_paren, TT_FunctionDeclarationLParen);
3901+
EXPECT_TOKEN(Tokens[5], tok::l_brace, TT_FunctionLBrace);
3902+
3903+
Tokens = annotate("friend Bar::operator Foo();");
3904+
ASSERT_EQ(Tokens.size(), 9u) << Tokens;
3905+
EXPECT_TOKEN(Tokens[3], tok::kw_operator, TT_FunctionDeclarationName);
3906+
EXPECT_TOKEN(Tokens[5], tok::l_paren, TT_FunctionDeclarationLParen);
3907+
}
3908+
38593909
} // namespace
38603910
} // namespace format
38613911
} // namespace clang

‎flang/lib/Lower/ConvertVariable.cpp

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -738,9 +738,11 @@ static mlir::Value createNewLocal(Fortran::lower::AbstractConverter &converter,
738738
auto idxTy = builder.getIndexType();
739739
for (mlir::Value sh : elidedShape)
740740
indices.push_back(builder.createConvert(loc, idxTy, sh));
741-
mlir::Value alloc = builder.create<cuf::AllocOp>(
742-
loc, ty, nm, symNm, dataAttr, lenParams, indices);
743-
return alloc;
741+
if (dataAttr.getValue() == cuf::DataAttribute::Shared)
742+
return builder.create<cuf::SharedMemoryOp>(loc, ty, nm, symNm, lenParams,
743+
indices);
744+
return builder.create<cuf::AllocOp>(loc, ty, nm, symNm, dataAttr, lenParams,
745+
indices);
744746
}
745747

746748
// Let the builder do all the heavy lifting.
@@ -1032,12 +1034,16 @@ static void instantiateLocal(Fortran::lower::AbstractConverter &converter,
10321034
symMap);
10331035
if (Fortran::semantics::NeedCUDAAlloc(var.getSymbol())) {
10341036
auto *builder = &converter.getFirOpBuilder();
1037+
cuf::DataAttributeAttr dataAttr =
1038+
Fortran::lower::translateSymbolCUFDataAttribute(builder->getContext(),
1039+
var.getSymbol());
10351040
mlir::Location loc = converter.getCurrentLocation();
10361041
fir::ExtendedValue exv =
10371042
converter.getSymbolExtendedValue(var.getSymbol(), &symMap);
10381043
auto *sym = &var.getSymbol();
10391044
const Fortran::semantics::Scope &owner = sym->owner();
1040-
if (owner.kind() != Fortran::semantics::Scope::Kind::MainProgram) {
1045+
if (owner.kind() != Fortran::semantics::Scope::Kind::MainProgram &&
1046+
dataAttr.getValue() != cuf::DataAttribute::Shared) {
10411047
converter.getFctCtx().attachCleanup([builder, loc, exv, sym]() {
10421048
cuf::DataAttributeAttr dataAttr =
10431049
Fortran::lower::translateSymbolCUFDataAttribute(

0 commit comments

Comments
 (0)