Skip to content

Commit 220b312

Browse files
Merge upstream LLVM into amd-gfx12
2 parents 1a3149a + 9987573 commit 220b312

File tree

92 files changed

+2265
-274
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

92 files changed

+2265
-274
lines changed

bolt/lib/Core/BinaryFunctionCallGraph.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -200,7 +200,8 @@ buildCallGraph(BinaryContext &BC, CgFilterFunction Filter, bool CgFromPerfData,
200200
if (CSI.Symbol)
201201
Counts.emplace_back(CSI.Symbol, CSI.Count);
202202
} else {
203-
const uint64_t Count = BB->getExecutionCount();
203+
const uint64_t Count = BC.MIB->getAnnotationWithDefault(
204+
Inst, "Count", BB->getExecutionCount());
204205
Counts.emplace_back(DstSym, Count);
205206
}
206207

clang-tools-extra/clang-tidy/cppcoreguidelines/InterfacesGlobalInitCheck.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@ void InterfacesGlobalInitCheck::registerMatchers(MatchFinder *Finder) {
1919
hasDeclContext(anyOf(translationUnitDecl(), // Global scope.
2020
namespaceDecl(), // Namespace scope.
2121
recordDecl())), // Class scope.
22-
unless(isConstexpr()));
22+
unless(isConstexpr()), unless(isConstinit()));
2323

2424
const auto ReferencesUndefinedGlobalVar = declRefExpr(hasDeclaration(
2525
varDecl(GlobalVarDecl, unless(isDefinition())).bind("referencee")));

clang-tools-extra/docs/ReleaseNotes.rst

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -232,6 +232,10 @@ Changes in existing checks
232232
<clang-tidy/checks/cppcoreguidelines/avoid-goto>` check by adding the option
233233
`IgnoreMacros` to ignore ``goto`` labels defined in macros.
234234

235+
- Improved :doc:`cppcoreguidelines-interfaces-global-init
236+
<clang-tidy/checks/cppcoreguidelines/interfaces-global-init>` check by
237+
fixing false positives on uses of ``constinit`` variables.
238+
235239
- Improved :doc:`cppcoreguidelines-missing-std-forward
236240
<clang-tidy/checks/cppcoreguidelines/missing-std-forward>` check by adding a
237241
flag to specify the function used for forwarding instead of ``std::forward``.

clang-tools-extra/test/clang-tidy/checkers/cppcoreguidelines/interfaces-global-init.cpp

Lines changed: 33 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -14,11 +14,23 @@ static int GlobalScopeBadInit3 = takesIntPtr(&ExternGlobal);
1414
static int GlobalScopeBadInit4 = 3 * (ExternGlobal + 2);
1515
// CHECK-MESSAGES: [[@LINE-1]]:12: warning: initializing non-local variable with non-const expression depending on uninitialized non-local variable 'ExternGlobal'
1616

17+
#if __cplusplus >= 202002L
18+
extern constinit int ExternConstinitGlobal;
19+
static int GlobalScopeConstinit1 = ExternConstinitGlobal;
20+
static int GlobalScopeConstinit2 = takesInt(ExternConstinitGlobal);
21+
static int GlobalScopeConstinit3 = takesIntPtr(&ExternConstinitGlobal);
22+
static int GlobalScopeConstinit4 = 3 * (ExternConstinitGlobal + 2);
23+
#endif
24+
1725
namespace ns {
1826
static int NamespaceScope = makesInt();
1927
static int NamespaceScopeBadInit = takesInt(ExternGlobal);
2028
// CHECK-MESSAGES: [[@LINE-1]]:12: warning: initializing non-local variable with non-const expression depending on uninitialized non-local variable 'ExternGlobal'
2129

30+
#if __cplusplus >= 202002L
31+
static int NamespaceScopeConstinit = takesInt(ExternConstinitGlobal);
32+
#endif
33+
2234
struct A {
2335
static int ClassScope;
2436
static int ClassScopeBadInit;
@@ -29,6 +41,17 @@ int A::ClassScopeBadInit = takesInt(ExternGlobal);
2941

3042
static int FromClassBadInit = takesInt(A::ClassScope);
3143
// CHECK-MESSAGES: [[@LINE-1]]:12: warning: initializing non-local variable with non-const expression depending on uninitialized non-local variable 'ClassScope'
44+
45+
#if __cplusplus >= 202002L
46+
struct B {
47+
static constinit int ClassScopeConstinit;
48+
static int ClassScopeFromConstinit;
49+
};
50+
51+
int B::ClassScopeFromConstinit = takesInt(ExternConstinitGlobal);
52+
static int FromClassScopeConstinit = takesInt(B::ClassScopeConstinit);
53+
#endif
54+
3255
} // namespace ns
3356

3457
// "const int B::I;" is fine, it just ODR-defines B::I. See [9.4.3] Static
@@ -42,6 +65,16 @@ const int B1::J;
4265
// CHECK-MESSAGES: [[@LINE-1]]:15: warning: initializing non-local variable with non-const expression depending on uninitialized non-local variable 'I'
4366
const int B1::I;
4467

68+
#if __cplusplus >= 202002L
69+
class D {
70+
static const constinit int I = 0;
71+
static const int J = I;
72+
};
73+
74+
const int D::J;
75+
const int D::I;
76+
#endif
77+
4578
void f() {
4679
// This is fine, it's executed after dynamic initialization occurs.
4780
static int G = takesInt(ExternGlobal);
@@ -81,4 +114,3 @@ class B2 {
81114
};
82115
const int B2::I;
83116
const int B2::J;
84-

clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -672,6 +672,8 @@ TARGET_BUILTIN(__builtin_amdgcn_s_monitor_sleep, "vIs", "n", "gfx1250-insts")
672672
TARGET_BUILTIN(__builtin_amdgcn_s_wait_asynccnt, "vIUs", "n", "gfx1250-insts")
673673
TARGET_BUILTIN(__builtin_amdgcn_s_wait_tensorcnt, "vIUs", "n", "gfx1250-insts")
674674

675+
TARGET_BUILTIN(__builtin_amdgcn_tanh_bf16, "yy", "nc", "bf16-trans-insts")
676+
675677
TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_fp8, "hiIi", "nc", "gfx1250-insts")
676678
TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_bf8, "hiIi", "nc", "gfx1250-insts")
677679
TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_fp8, "V2hs", "nc", "gfx1250-insts")

clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -497,6 +497,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
497497
Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, { ResultType });
498498
return Builder.CreateCall(F, { Src });
499499
}
500+
case AMDGPU::BI__builtin_amdgcn_tanh_bf16:
501+
return emitBuiltinWithOneOverloadedType<1>(*this, E,
502+
Intrinsic::amdgcn_tanh);
500503
case AMDGPU::BI__builtin_amdgcn_uicmp:
501504
case AMDGPU::BI__builtin_amdgcn_uicmpl:
502505
case AMDGPU::BI__builtin_amdgcn_sicmp:

clang/lib/Sema/AnalysisBasedWarnings.cpp

Lines changed: 48 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -987,10 +987,11 @@ static void DiagUninitUse(Sema &S, const VarDecl *VD, const UninitUse &Use,
987987
}
988988

989989
/// Diagnose uninitialized const reference usages.
990-
static void DiagnoseUninitializedConstRefUse(Sema &S, const VarDecl *VD,
990+
static bool DiagnoseUninitializedConstRefUse(Sema &S, const VarDecl *VD,
991991
const UninitUse &Use) {
992992
S.Diag(Use.getUser()->getBeginLoc(), diag::warn_uninit_const_reference)
993993
<< VD->getDeclName() << Use.getUser()->getSourceRange();
994+
return !S.getDiagnostics().isLastDiagnosticIgnored();
994995
}
995996

996997
/// DiagnoseUninitializedUse -- Helper function for diagnosing uses of an
@@ -1022,7 +1023,7 @@ static bool DiagnoseUninitializedUse(Sema &S, const VarDecl *VD,
10221023
if (CR.doesContainReference()) {
10231024
S.Diag(DRE->getBeginLoc(), diag::warn_uninit_self_reference_in_init)
10241025
<< VD->getDeclName() << VD->getLocation() << DRE->getSourceRange();
1025-
return true;
1026+
return !S.getDiagnostics().isLastDiagnosticIgnored();
10261027
}
10271028
}
10281029

@@ -1045,7 +1046,7 @@ static bool DiagnoseUninitializedUse(Sema &S, const VarDecl *VD,
10451046
S.Diag(VD->getBeginLoc(), diag::note_var_declared_here)
10461047
<< VD->getDeclName();
10471048

1048-
return true;
1049+
return !S.getDiagnostics().isLastDiagnosticIgnored();
10491050
}
10501051

10511052
namespace {
@@ -1559,43 +1560,7 @@ class UninitValsDiagReporter : public UninitVariablesHandler {
15591560
UsesVec *vec = V.getPointer();
15601561
bool hasSelfInit = V.getInt();
15611562

1562-
// Specially handle the case where we have uses of an uninitialized
1563-
// variable, but the root cause is an idiomatic self-init. We want
1564-
// to report the diagnostic at the self-init since that is the root cause.
1565-
if (!vec->empty() && hasSelfInit && hasAlwaysUninitializedUse(vec))
1566-
DiagnoseUninitializedUse(S, vd,
1567-
UninitUse(vd->getInit()->IgnoreParenCasts(),
1568-
/* isAlwaysUninit */ true),
1569-
/* alwaysReportSelfInit */ true);
1570-
else {
1571-
// Sort the uses by their SourceLocations. While not strictly
1572-
// guaranteed to produce them in line/column order, this will provide
1573-
// a stable ordering.
1574-
llvm::sort(*vec, [](const UninitUse &a, const UninitUse &b) {
1575-
// Move ConstRef uses to the back.
1576-
if (a.isConstRefUse() != b.isConstRefUse())
1577-
return b.isConstRefUse();
1578-
// Prefer a more confident report over a less confident one.
1579-
if (a.getKind() != b.getKind())
1580-
return a.getKind() > b.getKind();
1581-
return a.getUser()->getBeginLoc() < b.getUser()->getBeginLoc();
1582-
});
1583-
1584-
for (const auto &U : *vec) {
1585-
if (U.isConstRefUse()) {
1586-
DiagnoseUninitializedConstRefUse(S, vd, U);
1587-
break;
1588-
}
1589-
1590-
// If we have self-init, downgrade all uses to 'may be uninitialized'.
1591-
UninitUse Use = hasSelfInit ? UninitUse(U.getUser(), false) : U;
1592-
1593-
if (DiagnoseUninitializedUse(S, vd, Use))
1594-
// Skip further diagnostics for this variable. We try to warn only
1595-
// on the first point at which a variable is used uninitialized.
1596-
break;
1597-
}
1598-
}
1563+
diagnoseUnitializedVar(vd, hasSelfInit, vec);
15991564

16001565
// Release the uses vector.
16011566
delete vec;
@@ -1612,6 +1577,49 @@ class UninitValsDiagReporter : public UninitVariablesHandler {
16121577
U.getKind() == UninitUse::AfterDecl;
16131578
});
16141579
}
1580+
1581+
// Print the diagnostic for the variable. We try to warn only on the first
1582+
// point at which a variable is used uninitialized. After the first
1583+
// diagnostic is printed, further diagnostics for this variable are skipped.
1584+
void diagnoseUnitializedVar(const VarDecl *vd, bool hasSelfInit,
1585+
UsesVec *vec) {
1586+
// Specially handle the case where we have uses of an uninitialized
1587+
// variable, but the root cause is an idiomatic self-init. We want
1588+
// to report the diagnostic at the self-init since that is the root cause.
1589+
if (hasSelfInit && hasAlwaysUninitializedUse(vec)) {
1590+
if (DiagnoseUninitializedUse(S, vd,
1591+
UninitUse(vd->getInit()->IgnoreParenCasts(),
1592+
/*isAlwaysUninit=*/true),
1593+
/*alwaysReportSelfInit=*/true))
1594+
return;
1595+
}
1596+
1597+
// Sort the uses by their SourceLocations. While not strictly
1598+
// guaranteed to produce them in line/column order, this will provide
1599+
// a stable ordering.
1600+
llvm::sort(*vec, [](const UninitUse &a, const UninitUse &b) {
1601+
// Prefer the direct use of an uninitialized variable over its use via
1602+
// constant reference.
1603+
if (a.isConstRefUse() != b.isConstRefUse())
1604+
return b.isConstRefUse();
1605+
// Prefer a more confident report over a less confident one.
1606+
if (a.getKind() != b.getKind())
1607+
return a.getKind() > b.getKind();
1608+
return a.getUser()->getBeginLoc() < b.getUser()->getBeginLoc();
1609+
});
1610+
1611+
for (const auto &U : *vec) {
1612+
if (U.isConstRefUse()) {
1613+
if (DiagnoseUninitializedConstRefUse(S, vd, U))
1614+
return;
1615+
} else {
1616+
// If we have self-init, downgrade all uses to 'may be uninitialized'.
1617+
UninitUse Use = hasSelfInit ? UninitUse(U.getUser(), false) : U;
1618+
if (DiagnoseUninitializedUse(S, vd, Use))
1619+
return;
1620+
}
1621+
}
1622+
}
16151623
};
16161624

16171625
/// Inter-procedural data for the called-once checker.

clang/test/CodeGenOpenCL/amdgpu-features.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -110,7 +110,7 @@
110110
// GFX1170: "target-features"="+16-bit-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32,+wmma-128b-insts"
111111
// GFX1200: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32,+wmma-128b-insts"
112112
// GFX1201: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32,+wmma-128b-insts"
113-
// GFX1250: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bitop3-insts,+ci-insts,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+transpose-load-f4f6-insts,+wavefrontsize32"
113+
// GFX1250: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf16-trans-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+transpose-load-f4f6-insts,+wavefrontsize32"
114114

115115
// GFX1103-W64: "target-features"="+16-bit-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot10-insts,+dot12-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize64,+wmma-256b-insts"
116116

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,25 @@ void test_s_wait_tensorcnt() {
4242
__builtin_amdgcn_s_wait_tensorcnt(0);
4343
}
4444

45+
// CHECK-LABEL: @test_tanh_bf16(
46+
// CHECK-NEXT: entry:
47+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
48+
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca bfloat, align 2, addrspace(5)
49+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
50+
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
51+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
52+
// CHECK-NEXT: store bfloat [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 2
53+
// CHECK-NEXT: [[TMP0:%.*]] = load bfloat, ptr [[A_ADDR_ASCAST]], align 2
54+
// CHECK-NEXT: [[TMP1:%.*]] = call bfloat @llvm.amdgcn.tanh.bf16(bfloat [[TMP0]])
55+
// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
56+
// CHECK-NEXT: store bfloat [[TMP1]], ptr addrspace(1) [[TMP2]], align 2
57+
// CHECK-NEXT: ret void
58+
//
59+
void test_tanh_bf16(global __bf16* out, __bf16 a)
60+
{
61+
*out = __builtin_amdgcn_tanh_bf16(a);
62+
}
63+
4564
// CHECK-LABEL: @test_cvt_f16_fp8(
4665
// CHECK-NEXT: entry:
4766
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
// RUN: %clang_cc1 -fsyntax-only -Wuninitialized -Wno-sometimes-uninitialized -verify %s
2+
3+
void foo(const int &);
4+
5+
int f(bool a) {
6+
int v;
7+
if (a) {
8+
foo(v); // expected-warning {{variable 'v' is uninitialized when passed as a const reference argument here}}
9+
v = 5;
10+
}
11+
return v;
12+
}

0 commit comments

Comments
 (0)