Skip to content

Commit a595ff6

Browse files
ronliebronlieb
committed
merge main into amd-staging
Change-Id: I0bf1b1a277fe9dfab90fdd9c05abcf2ba7bc7930
2 parents 94eb3ec + f078548 commit a595ff6

File tree

224 files changed

+11642
-18948
lines changed

Some content is hidden

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

224 files changed

+11642
-18948
lines changed

‎bolt/include/bolt/Core/GDBIndex.h

Lines changed: 61 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,61 @@
1+
//===-- bolt/Core/GDBIndex.h - GDB Index support ----------------*- C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
///
9+
/// This file contains declaration of classes required for generation of
10+
/// .gdb_index section.
11+
///
12+
//===----------------------------------------------------------------------===//
13+
14+
#ifndef BOLT_CORE_GDB_INDEX_H
15+
#define BOLT_CORE_GDB_INDEX_H
16+
17+
#include "bolt/Core/BinaryContext.h"
18+
#include <vector>
19+
20+
namespace llvm {
21+
namespace bolt {
22+
23+
class GDBIndex {
24+
public:
25+
/// Contains information about TU so we can write out correct entries in GDB
26+
/// index.
27+
struct GDBIndexTUEntry {
28+
uint64_t UnitOffset;
29+
uint64_t TypeHash;
30+
uint64_t TypeDIERelativeOffset;
31+
};
32+
33+
private:
34+
BinaryContext &BC;
35+
36+
/// Entries for GDB Index Types CU List.
37+
using GDBIndexTUEntryType = std::vector<GDBIndexTUEntry>;
38+
GDBIndexTUEntryType GDBIndexTUEntryVector;
39+
40+
public:
41+
GDBIndex(BinaryContext &BC) : BC(BC) {}
42+
43+
std::mutex GDBIndexMutex;
44+
45+
/// Adds an GDBIndexTUEntry if .gdb_index section exists.
46+
void addGDBTypeUnitEntry(const GDBIndexTUEntry &&Entry);
47+
48+
/// Rewrite .gdb_index section if present.
49+
void updateGdbIndexSection(const CUOffsetMap &CUMap, const uint32_t NumCUs,
50+
DebugARangesSectionWriter &ARangesSectionWriter);
51+
52+
/// Returns all entries needed for Types CU list.
53+
const GDBIndexTUEntryType &getGDBIndexTUEntryVector() const {
54+
return GDBIndexTUEntryVector;
55+
}
56+
};
57+
58+
} // namespace bolt
59+
} // namespace llvm
60+
61+
#endif

‎bolt/lib/Core/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,7 @@ add_llvm_library(LLVMBOLTCore
2525
DynoStats.cpp
2626
Exceptions.cpp
2727
FunctionLayout.cpp
28+
GDBIndex.cpp
2829
HashUtilities.cpp
2930
JumpTable.cpp
3031
MCPlusBuilder.cpp

‎bolt/lib/Core/GDBIndex.cpp

Lines changed: 185 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,185 @@
1+
//===- bolt/Core/GDBIndex.cpp - GDB Index support ------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include "bolt/Core/GDBIndex.h"
10+
11+
using namespace llvm::bolt;
12+
using namespace llvm::support::endian;
13+
14+
void GDBIndex::addGDBTypeUnitEntry(const GDBIndexTUEntry &&Entry) {
15+
std::lock_guard<std::mutex> Lock(GDBIndexMutex);
16+
if (!BC.getGdbIndexSection())
17+
return;
18+
GDBIndexTUEntryVector.emplace_back(Entry);
19+
}
20+
21+
void GDBIndex::updateGdbIndexSection(
22+
const CUOffsetMap &CUMap, const uint32_t NumCUs,
23+
DebugARangesSectionWriter &ARangesSectionWriter) {
24+
if (!BC.getGdbIndexSection())
25+
return;
26+
27+
// See https://sourceware.org/gdb/onlinedocs/gdb/Index-Section-Format.html
28+
// for .gdb_index section format.
29+
30+
StringRef GdbIndexContents = BC.getGdbIndexSection()->getContents();
31+
32+
const char *Data = GdbIndexContents.data();
33+
34+
// Parse the header.
35+
const uint32_t Version = read32le(Data);
36+
if (Version != 7 && Version != 8) {
37+
errs() << "BOLT-ERROR: can only process .gdb_index versions 7 and 8\n";
38+
exit(1);
39+
}
40+
41+
// Some .gdb_index generators use file offsets while others use section
42+
// offsets. Hence we can only rely on offsets relative to each other,
43+
// and ignore their absolute values.
44+
const uint32_t CUListOffset = read32le(Data + 4);
45+
const uint32_t CUTypesOffset = read32le(Data + 8);
46+
const uint32_t AddressTableOffset = read32le(Data + 12);
47+
const uint32_t SymbolTableOffset = read32le(Data + 16);
48+
const uint32_t ConstantPoolOffset = read32le(Data + 20);
49+
Data += 24;
50+
51+
// Map CUs offsets to indices and verify existing index table.
52+
std::map<uint32_t, uint32_t> OffsetToIndexMap;
53+
const uint32_t CUListSize = CUTypesOffset - CUListOffset;
54+
const uint32_t TUListSize = AddressTableOffset - CUTypesOffset;
55+
const unsigned NUmCUsEncoded = CUListSize / 16;
56+
unsigned MaxDWARFVersion = BC.DwCtx->getMaxVersion();
57+
unsigned NumDWARF5TUs =
58+
getGDBIndexTUEntryVector().size() - BC.DwCtx->getNumTypeUnits();
59+
bool SkipTypeUnits = false;
60+
// For DWARF5 Types are in .debug_info.
61+
// LLD doesn't generate Types CU List, and in CU list offset
62+
// only includes CUs.
63+
// GDB 11+ includes only CUs in CU list and generates Types
64+
// list.
65+
// GDB 9 includes CUs and TUs in CU list and generates TYpes
66+
// list. The NumCUs is CUs + TUs, so need to modify the check.
67+
// For split-dwarf
68+
// GDB-11, DWARF5: TU units from dwo are not included.
69+
// GDB-11, DWARF4: TU units from dwo are included.
70+
if (MaxDWARFVersion >= 5)
71+
SkipTypeUnits = !TUListSize ? true
72+
: ((NUmCUsEncoded + NumDWARF5TUs) ==
73+
BC.DwCtx->getNumCompileUnits());
74+
75+
if (!((CUListSize == NumCUs * 16) ||
76+
(CUListSize == (NumCUs + NumDWARF5TUs) * 16))) {
77+
errs() << "BOLT-ERROR: .gdb_index: CU count mismatch\n";
78+
exit(1);
79+
}
80+
DenseSet<uint64_t> OriginalOffsets;
81+
for (unsigned Index = 0, Units = BC.DwCtx->getNumCompileUnits();
82+
Index < Units; ++Index) {
83+
const DWARFUnit *CU = BC.DwCtx->getUnitAtIndex(Index);
84+
if (SkipTypeUnits && CU->isTypeUnit())
85+
continue;
86+
const uint64_t Offset = read64le(Data);
87+
Data += 16;
88+
if (CU->getOffset() != Offset) {
89+
errs() << "BOLT-ERROR: .gdb_index CU offset mismatch\n";
90+
exit(1);
91+
}
92+
93+
OriginalOffsets.insert(Offset);
94+
OffsetToIndexMap[Offset] = Index;
95+
}
96+
97+
// Ignore old address table.
98+
const uint32_t OldAddressTableSize = SymbolTableOffset - AddressTableOffset;
99+
// Move Data to the beginning of symbol table.
100+
Data += SymbolTableOffset - CUTypesOffset;
101+
102+
// Calculate the size of the new address table.
103+
uint32_t NewAddressTableSize = 0;
104+
for (const auto &CURangesPair : ARangesSectionWriter.getCUAddressRanges()) {
105+
const SmallVector<DebugAddressRange, 2> &Ranges = CURangesPair.second;
106+
NewAddressTableSize += Ranges.size() * 20;
107+
}
108+
109+
// Difference between old and new table (and section) sizes.
110+
// Could be negative.
111+
int32_t Delta = NewAddressTableSize - OldAddressTableSize;
112+
113+
size_t NewGdbIndexSize = GdbIndexContents.size() + Delta;
114+
115+
// Free'd by ExecutableFileMemoryManager.
116+
auto *NewGdbIndexContents = new uint8_t[NewGdbIndexSize];
117+
uint8_t *Buffer = NewGdbIndexContents;
118+
119+
write32le(Buffer, Version);
120+
write32le(Buffer + 4, CUListOffset);
121+
write32le(Buffer + 8, CUTypesOffset);
122+
write32le(Buffer + 12, AddressTableOffset);
123+
write32le(Buffer + 16, SymbolTableOffset + Delta);
124+
write32le(Buffer + 20, ConstantPoolOffset + Delta);
125+
Buffer += 24;
126+
127+
using MapEntry = std::pair<uint32_t, CUInfo>;
128+
std::vector<MapEntry> CUVector(CUMap.begin(), CUMap.end());
129+
// Need to sort since we write out all of TUs in .debug_info before CUs.
130+
std::sort(CUVector.begin(), CUVector.end(),
131+
[](const MapEntry &E1, const MapEntry &E2) -> bool {
132+
return E1.second.Offset < E2.second.Offset;
133+
});
134+
// Writing out CU List <Offset, Size>
135+
for (auto &CUInfo : CUVector) {
136+
// Skipping TU for DWARF5 when they are not included in CU list.
137+
if (!OriginalOffsets.count(CUInfo.first))
138+
continue;
139+
write64le(Buffer, CUInfo.second.Offset);
140+
// Length encoded in CU doesn't contain first 4 bytes that encode length.
141+
write64le(Buffer + 8, CUInfo.second.Length + 4);
142+
Buffer += 16;
143+
}
144+
145+
// Rewrite TU CU List, since abbrevs can be different.
146+
// Entry example:
147+
// 0: offset = 0x00000000, type_offset = 0x0000001e, type_signature =
148+
// 0x418503b8111e9a7b Spec says " triplet, the first value is the CU offset,
149+
// the second value is the type offset in the CU, and the third value is the
150+
// type signature" Looking at what is being generated by gdb-add-index. The
151+
// first entry is TU offset, second entry is offset from it, and third entry
152+
// is the type signature.
153+
if (TUListSize)
154+
for (const GDBIndexTUEntry &Entry : getGDBIndexTUEntryVector()) {
155+
write64le(Buffer, Entry.UnitOffset);
156+
write64le(Buffer + 8, Entry.TypeDIERelativeOffset);
157+
write64le(Buffer + 16, Entry.TypeHash);
158+
Buffer += sizeof(GDBIndexTUEntry);
159+
}
160+
161+
// Generate new address table.
162+
for (const std::pair<const uint64_t, DebugAddressRangesVector> &CURangesPair :
163+
ARangesSectionWriter.getCUAddressRanges()) {
164+
const uint32_t CUIndex = OffsetToIndexMap[CURangesPair.first];
165+
const DebugAddressRangesVector &Ranges = CURangesPair.second;
166+
for (const DebugAddressRange &Range : Ranges) {
167+
write64le(Buffer, Range.LowPC);
168+
write64le(Buffer + 8, Range.HighPC);
169+
write32le(Buffer + 16, CUIndex);
170+
Buffer += 20;
171+
}
172+
}
173+
174+
const size_t TrailingSize =
175+
GdbIndexContents.data() + GdbIndexContents.size() - Data;
176+
assert(Buffer + TrailingSize == NewGdbIndexContents + NewGdbIndexSize &&
177+
"size calculation error");
178+
179+
// Copy over the rest of the original data.
180+
memcpy(Buffer, Data, TrailingSize);
181+
182+
// Register the new section.
183+
BC.registerOrUpdateNoteSection(".gdb_index", NewGdbIndexContents,
184+
NewGdbIndexSize);
185+
}

‎clang-tools-extra/clang-tidy/readability/ContainerSizeEmptyCheck.cpp

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -96,9 +96,14 @@ AST_MATCHER(QualType, isIntegralType) {
9696

9797
AST_MATCHER_P(UserDefinedLiteral, hasLiteral,
9898
clang::ast_matchers::internal::Matcher<Expr>, InnerMatcher) {
99-
if (const Expr *CookedLiteral = Node.getCookedLiteral()) {
99+
const UserDefinedLiteral::LiteralOperatorKind LOK =
100+
Node.getLiteralOperatorKind();
101+
if (LOK == UserDefinedLiteral::LOK_Template ||
102+
LOK == UserDefinedLiteral::LOK_Raw)
103+
return false;
104+
105+
if (const Expr *CookedLiteral = Node.getCookedLiteral())
100106
return InnerMatcher.matches(*CookedLiteral, Finder, Builder);
101-
}
102107
return false;
103108
}
104109

‎clang-tools-extra/docs/ReleaseNotes.rst

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -376,6 +376,7 @@ Changes in existing checks
376376
- Improved :doc:`readability-container-size-empty
377377
<clang-tidy/checks/readability/container-size-empty>` check to prevent false
378378
positives when utilizing ``size`` or ``length`` methods that accept parameter.
379+
Fixed crash when facing template user defined literals.
379380

380381
- Improved :doc:`readability-duplicate-include
381382
<clang-tidy/checks/readability/duplicate-include>` check by excluding include

‎clang-tools-extra/test/clang-tidy/checkers/readability/container-size-empty.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -889,3 +889,9 @@ namespace PR88203 {
889889
// CHECK-FIXES: {{^ }}if (s.empty()) {}{{$}}
890890
}
891891
}
892+
893+
namespace PR94454 {
894+
template <char...>
895+
int operator""_ci() { return 0; }
896+
auto eq = 0_ci == 0;
897+
}

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -240,7 +240,7 @@ TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s", "t", "at
240240
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t", "atomic-global-pk-add-bf16-inst")
241241
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "atomic-ds-pk-add-16-insts")
242242
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h", "t", "atomic-ds-pk-add-16-insts")
243-
TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3UiiUi", "t", "gfx940-insts")
243+
TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "gfx940-insts")
244244

245245
//===----------------------------------------------------------------------===//
246246
// Deep learning builtins.

‎clang/include/clang/Basic/Cuda.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -124,6 +124,7 @@ enum class CudaArch {
124124
GFX1103,
125125
GFX1150,
126126
GFX1151,
127+
GFX1152,
127128
GFX12_GENERIC,
128129
GFX1200,
129130
GFX1201,

‎clang/include/clang/Basic/riscv_vector.td

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2637,7 +2637,8 @@ let UnMaskedPolicyScheme = HasPassthruOperand in {
26372637
defm vbrev : RVVOutBuiltinSetZvbb;
26382638
defm vclz : RVVOutBuiltinSetZvbb;
26392639
defm vctz : RVVOutBuiltinSetZvbb;
2640-
defm vcpopv : RVVOutBuiltinSetZvbb;
2640+
let IRName = "vcpopv", MaskedIRName = "vcpopv_mask" in
2641+
defm vcpop : RVVOutBuiltinSetZvbb;
26412642
let OverloadedName = "vwsll" in
26422643
defm vwsll : RVVSignedWidenBinBuiltinSetVwsll;
26432644
}

‎clang/lib/AST/Interp/Function.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -40,7 +40,8 @@ SourceInfo Function::getSource(CodePtr PC) const {
4040
unsigned Offset = PC - getCodeBegin();
4141
using Elem = std::pair<unsigned, SourceInfo>;
4242
auto It = llvm::lower_bound(SrcMap, Elem{Offset, {}}, llvm::less_first());
43-
assert(It != SrcMap.end());
43+
if (It == SrcMap.end())
44+
return SrcMap.back().second;
4445
return It->second;
4546
}
4647

‎clang/lib/AST/Interp/Interp.cpp

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -538,7 +538,7 @@ bool CheckCallable(InterpState &S, CodePtr OpPC, const Function *F) {
538538
return false;
539539
}
540540

541-
if (!F->isConstexpr()) {
541+
if (!F->isConstexpr() || !F->hasBody()) {
542542
const SourceLocation &Loc = S.Current->getLocation(OpPC);
543543
if (S.getLangOpts().CPlusPlus11) {
544544
const FunctionDecl *DiagDecl = F->getDecl();
@@ -572,9 +572,10 @@ bool CheckCallable(InterpState &S, CodePtr OpPC, const Function *F) {
572572
S.checkingPotentialConstantExpression())
573573
return false;
574574

575-
// If the declaration is defined _and_ declared 'constexpr', the below
576-
// diagnostic doesn't add anything useful.
577-
if (DiagDecl->isDefined() && DiagDecl->isConstexpr())
575+
// If the declaration is defined, declared 'constexpr' _and_ has a body,
576+
// the below diagnostic doesn't add anything useful.
577+
if (DiagDecl->isDefined() && DiagDecl->isConstexpr() &&
578+
DiagDecl->hasBody())
578579
return false;
579580

580581
S.FFDiag(Loc, diag::note_constexpr_invalid_function, 1)

‎clang/lib/AST/Interp/Interp.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2194,7 +2194,7 @@ inline bool ArrayDecay(InterpState &S, CodePtr OpPC) {
21942194
if (!CheckRange(S, OpPC, Ptr, CSK_ArrayToPointer))
21952195
return false;
21962196

2197-
if (!Ptr.isUnknownSizeArray() || Ptr.isDummy()) {
2197+
if (Ptr.isRoot() || !Ptr.isUnknownSizeArray() || Ptr.isDummy()) {
21982198
S.Stk.push<Pointer>(Ptr.atIndex(0));
21992199
return true;
22002200
}

‎clang/lib/Basic/Cuda.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -144,6 +144,7 @@ static const CudaArchToStringMap arch_names[] = {
144144
GFX(1103), // gfx1103
145145
GFX(1150), // gfx1150
146146
GFX(1151), // gfx1151
147+
GFX(1152), // gfx1152
147148
{CudaArch::GFX12_GENERIC, "gfx12-generic", "compute_amdgcn"},
148149
GFX(1200), // gfx1200
149150
GFX(1201), // gfx1201

‎clang/lib/Basic/Targets/NVPTX.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -228,6 +228,7 @@ void NVPTXTargetInfo::getTargetDefines(const LangOptions &Opts,
228228
case CudaArch::GFX1103:
229229
case CudaArch::GFX1150:
230230
case CudaArch::GFX1151:
231+
case CudaArch::GFX1152:
231232
case CudaArch::GFX12_GENERIC:
232233
case CudaArch::GFX1200:
233234
case CudaArch::GFX1201:

‎clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3763,6 +3763,7 @@ void CGOpenMPRuntimeGPU::processRequiresDirective(
37633763
case CudaArch::GFX1103:
37643764
case CudaArch::GFX1150:
37653765
case CudaArch::GFX1151:
3766+
case CudaArch::GFX1152:
37663767
case CudaArch::GFX12_GENERIC:
37673768
case CudaArch::GFX1200:
37683769
case CudaArch::GFX1201:

0 commit comments

Comments
 (0)