Skip to content

Commit e192cd4

Browse files
Merge branch 'llvm:main' into gh-101657
2 parents 30cf01d + 90f5c8b commit e192cd4

File tree

306 files changed

+19877
-5251
lines changed

Some content is hidden

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

306 files changed

+19877
-5251
lines changed

bolt/include/bolt/Profile/DataAggregator.h

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -170,6 +170,9 @@ class DataAggregator : public DataReader {
170170
std::string BuildIDBinaryName;
171171

172172
/// Memory map info for a single file as recorded in perf.data
173+
/// When a binary has multiple text segments, the Size is computed as the
174+
/// difference of the last address of these segments from the BaseAddress.
175+
/// The base addresses of all text segments must be the same.
173176
struct MMapInfo {
174177
uint64_t BaseAddress{0}; /// Base address of the mapped binary.
175178
uint64_t MMapAddress{0}; /// Address of the executable segment.
@@ -493,6 +496,11 @@ class DataAggregator : public DataReader {
493496
/// and return a file name matching a given \p FileBuildID.
494497
std::optional<StringRef> getFileNameForBuildID(StringRef FileBuildID);
495498

499+
/// Get a constant reference to the parsed binary mmap entries.
500+
const std::unordered_map<uint64_t, MMapInfo> &getBinaryMMapInfo() {
501+
return BinaryMMapInfo;
502+
}
503+
496504
friend class YAMLProfileWriter;
497505
};
498506
} // namespace bolt

bolt/lib/Profile/DataAggregator.cpp

Lines changed: 29 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -95,6 +95,12 @@ cl::opt<bool> ReadPreAggregated(
9595
"pa", cl::desc("skip perf and read data from a pre-aggregated file format"),
9696
cl::cat(AggregatorCategory));
9797

98+
cl::opt<std::string>
99+
ReadPerfEvents("perf-script-events",
100+
cl::desc("skip perf event collection by supplying a "
101+
"perf-script output in a textual format"),
102+
cl::ReallyHidden, cl::init(""), cl::cat(AggregatorCategory));
103+
98104
static cl::opt<bool>
99105
TimeAggregator("time-aggr",
100106
cl::desc("time BOLT aggregator"),
@@ -167,8 +173,9 @@ void DataAggregator::findPerfExecutable() {
167173
void DataAggregator::start() {
168174
outs() << "PERF2BOLT: Starting data aggregation job for " << Filename << "\n";
169175

170-
// Don't launch perf for pre-aggregated files
171-
if (opts::ReadPreAggregated)
176+
// Don't launch perf for pre-aggregated files or when perf input is specified
177+
// by the user.
178+
if (opts::ReadPreAggregated || !opts::ReadPerfEvents.empty())
172179
return;
173180

174181
findPerfExecutable();
@@ -464,6 +471,13 @@ void DataAggregator::filterBinaryMMapInfo() {
464471

465472
int DataAggregator::prepareToParse(StringRef Name, PerfProcessInfo &Process,
466473
PerfProcessErrorCallbackTy Callback) {
474+
if (!opts::ReadPerfEvents.empty()) {
475+
outs() << "PERF2BOLT: using pre-processed perf events for '" << Name
476+
<< "' (perf-script-events)\n";
477+
ParsingBuf = opts::ReadPerfEvents;
478+
return 0;
479+
}
480+
467481
std::string Error;
468482
outs() << "PERF2BOLT: waiting for perf " << Name
469483
<< " collection to finish...\n";
@@ -2056,15 +2070,6 @@ std::error_code DataAggregator::parseMMapEvents() {
20562070
if (FileMMapInfo.first == "(deleted)")
20572071
continue;
20582072

2059-
// Consider only the first mapping of the file for any given PID
2060-
auto Range = GlobalMMapInfo.equal_range(FileMMapInfo.first);
2061-
bool PIDExists = llvm::any_of(make_range(Range), [&](const auto &MI) {
2062-
return MI.second.PID == FileMMapInfo.second.PID;
2063-
});
2064-
2065-
if (PIDExists)
2066-
continue;
2067-
20682073
GlobalMMapInfo.insert(FileMMapInfo);
20692074
}
20702075

@@ -2116,12 +2121,22 @@ std::error_code DataAggregator::parseMMapEvents() {
21162121
<< " using file offset 0x" << Twine::utohexstr(MMapInfo.Offset)
21172122
<< ". Ignoring profile data for this mapping\n";
21182123
continue;
2119-
} else {
2120-
MMapInfo.BaseAddress = *BaseAddress;
21212124
}
2125+
MMapInfo.BaseAddress = *BaseAddress;
21222126
}
21232127

2124-
BinaryMMapInfo.insert(std::make_pair(MMapInfo.PID, MMapInfo));
2128+
// Try to add MMapInfo to the map and update its size. Large binaries may
2129+
// span to multiple text segments, so the mapping is inserted only on the
2130+
// first occurrence.
2131+
if (!BinaryMMapInfo.insert(std::make_pair(MMapInfo.PID, MMapInfo)).second)
2132+
assert(MMapInfo.BaseAddress == BinaryMMapInfo[MMapInfo.PID].BaseAddress &&
2133+
"Base address on multiple segment mappings should match");
2134+
2135+
// Update mapping size.
2136+
const uint64_t EndAddress = MMapInfo.MMapAddress + MMapInfo.Size;
2137+
const uint64_t Size = EndAddress - BinaryMMapInfo[MMapInfo.PID].BaseAddress;
2138+
if (Size > BinaryMMapInfo[MMapInfo.PID].Size)
2139+
BinaryMMapInfo[MMapInfo.PID].Size = Size;
21252140
}
21262141

21272142
if (BinaryMMapInfo.empty()) {

bolt/unittests/Core/CMakeLists.txt

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@ set(LLVM_LINK_COMPONENTS
88
add_bolt_unittest(CoreTests
99
BinaryContext.cpp
1010
MCPlusBuilder.cpp
11+
MemoryMaps.cpp
1112
DynoStats.cpp
1213

1314
DISABLE_LLVM_LINK_LLVM_DYLIB
@@ -17,6 +18,8 @@ target_link_libraries(CoreTests
1718
PRIVATE
1819
LLVMBOLTCore
1920
LLVMBOLTRewrite
21+
LLVMBOLTProfile
22+
LLVMTestingSupport
2023
)
2124

2225
foreach (tgt ${BOLT_TARGETS_TO_BUILD})

bolt/unittests/Core/MemoryMaps.cpp

Lines changed: 142 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,142 @@
1+
//===- bolt/unittest/Core/MemoryMaps.cpp ----------------------------------===//
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/BinaryContext.h"
10+
#include "bolt/Profile/DataAggregator.h"
11+
#include "llvm/BinaryFormat/ELF.h"
12+
#include "llvm/DebugInfo/DWARF/DWARFContext.h"
13+
#include "llvm/Support/CommandLine.h"
14+
#include "llvm/Support/TargetSelect.h"
15+
#include "llvm/Testing/Support/Error.h"
16+
#include "gtest/gtest.h"
17+
18+
using namespace llvm;
19+
using namespace llvm::object;
20+
using namespace llvm::ELF;
21+
using namespace bolt;
22+
23+
namespace opts {
24+
extern cl::opt<std::string> ReadPerfEvents;
25+
} // namespace opts
26+
27+
namespace {
28+
29+
/// Perform checks on memory map events normally captured in perf. Tests use
30+
/// the 'opts::ReadPerfEvents' flag to emulate these events, passing a custom
31+
/// 'perf script' output to DataAggregator.
32+
struct MemoryMapsTester : public testing::TestWithParam<Triple::ArchType> {
33+
void SetUp() override {
34+
initalizeLLVM();
35+
prepareElf();
36+
initializeBOLT();
37+
}
38+
39+
protected:
40+
void initalizeLLVM() {
41+
llvm::InitializeAllTargetInfos();
42+
llvm::InitializeAllTargetMCs();
43+
llvm::InitializeAllAsmParsers();
44+
llvm::InitializeAllDisassemblers();
45+
llvm::InitializeAllTargets();
46+
llvm::InitializeAllAsmPrinters();
47+
}
48+
49+
void prepareElf() {
50+
memcpy(ElfBuf, "\177ELF", 4);
51+
ELF64LE::Ehdr *EHdr = reinterpret_cast<typename ELF64LE::Ehdr *>(ElfBuf);
52+
EHdr->e_ident[llvm::ELF::EI_CLASS] = llvm::ELF::ELFCLASS64;
53+
EHdr->e_ident[llvm::ELF::EI_DATA] = llvm::ELF::ELFDATA2LSB;
54+
EHdr->e_machine = GetParam() == Triple::aarch64 ? EM_AARCH64 : EM_X86_64;
55+
MemoryBufferRef Source(StringRef(ElfBuf, sizeof(ElfBuf)), "ELF");
56+
ObjFile = cantFail(ObjectFile::createObjectFile(Source));
57+
}
58+
59+
void initializeBOLT() {
60+
Relocation::Arch = ObjFile->makeTriple().getArch();
61+
BC = cantFail(BinaryContext::createBinaryContext(
62+
ObjFile->makeTriple(), ObjFile->getFileName(), nullptr, true,
63+
DWARFContext::create(*ObjFile.get()), {llvm::outs(), llvm::errs()}));
64+
ASSERT_FALSE(!BC);
65+
}
66+
67+
char ElfBuf[sizeof(typename ELF64LE::Ehdr)] = {};
68+
std::unique_ptr<ObjectFile> ObjFile;
69+
std::unique_ptr<BinaryContext> BC;
70+
};
71+
} // namespace
72+
73+
#ifdef X86_AVAILABLE
74+
75+
INSTANTIATE_TEST_SUITE_P(X86, MemoryMapsTester,
76+
::testing::Values(Triple::x86_64));
77+
78+
#endif
79+
80+
#ifdef AARCH64_AVAILABLE
81+
82+
INSTANTIATE_TEST_SUITE_P(AArch64, MemoryMapsTester,
83+
::testing::Values(Triple::aarch64));
84+
85+
#endif
86+
87+
/// Check that the correct mmap size is computed when we have multiple text
88+
/// segment mappings.
89+
TEST_P(MemoryMapsTester, ParseMultipleSegments) {
90+
const int Pid = 1234;
91+
StringRef Filename = "BINARY";
92+
opts::ReadPerfEvents = formatv(
93+
"name 0 [000] 0.000000: PERF_RECORD_MMAP2 {0}/{0}: "
94+
"[0xabc0000000(0x1000000) @ 0x11c0000 103:01 1573523 0]: r-xp {1}\n"
95+
"name 0 [000] 0.000000: PERF_RECORD_MMAP2 {0}/{0}: "
96+
"[0xabc2000000(0x8000000) @ 0x31d0000 103:01 1573523 0]: r-xp {1}\n",
97+
Pid, Filename);
98+
99+
BC->SegmentMapInfo[0x11da000] =
100+
SegmentInfo{0x11da000, 0x10da000, 0x11ca000, 0x10da000, 0x10000, true};
101+
BC->SegmentMapInfo[0x31d0000] =
102+
SegmentInfo{0x31d0000, 0x51ac82c, 0x31d0000, 0x3000000, 0x200000, true};
103+
104+
DataAggregator DA("");
105+
BC->setFilename(Filename);
106+
Error Err = DA.preprocessProfile(*BC);
107+
108+
// Ignore errors from perf2bolt when parsing memory events later on.
109+
ASSERT_THAT_ERROR(std::move(Err), Succeeded());
110+
111+
auto &BinaryMMapInfo = DA.getBinaryMMapInfo();
112+
auto El = BinaryMMapInfo.find(Pid);
113+
// Check that memory mapping is present and has the expected size.
114+
ASSERT_NE(El, BinaryMMapInfo.end());
115+
ASSERT_EQ(El->second.Size, static_cast<uint64_t>(0xb1d0000));
116+
}
117+
118+
/// Check that DataAggregator aborts when pre-processing an input binary
119+
/// with multiple text segments that have different base addresses.
120+
TEST_P(MemoryMapsTester, MultipleSegmentsMismatchedBaseAddress) {
121+
const int Pid = 1234;
122+
StringRef Filename = "BINARY";
123+
opts::ReadPerfEvents = formatv(
124+
"name 0 [000] 0.000000: PERF_RECORD_MMAP2 {0}/{0}: "
125+
"[0xabc0000000(0x1000000) @ 0x11c0000 103:01 1573523 0]: r-xp {1}\n"
126+
"name 0 [000] 0.000000: PERF_RECORD_MMAP2 {0}/{0}: "
127+
"[0xabc2000000(0x8000000) @ 0x31d0000 103:01 1573523 0]: r-xp {1}\n",
128+
Pid, Filename);
129+
130+
BC->SegmentMapInfo[0x11da000] =
131+
SegmentInfo{0x11da000, 0x10da000, 0x11ca000, 0x10da000, 0x10000, true};
132+
// Using '0x31d0fff' FileOffset which triggers a different base address
133+
// for this second text segment.
134+
BC->SegmentMapInfo[0x31d0000] =
135+
SegmentInfo{0x31d0000, 0x51ac82c, 0x31d0fff, 0x3000000, 0x200000, true};
136+
137+
DataAggregator DA("");
138+
BC->setFilename(Filename);
139+
ASSERT_DEATH(
140+
{ Error Err = DA.preprocessProfile(*BC); },
141+
"Base address on multiple segment mappings should match");
142+
}

bolt/utils/bughunter.sh

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -131,7 +131,7 @@ if [[ $FAIL -eq "0" ]]; then
131131
fi
132132
else
133133
echo "Did it pass? Type the return code [0 = pass, 1 = fail]"
134-
read -n1 PASS
134+
read -n1 FAIL
135135
fi
136136
if [[ $FAIL -eq "0" ]] ; then
137137
echo " Warning: optimized binary passes."
@@ -205,7 +205,7 @@ while [[ "$CONTINUE" -ne "0" ]] ; do
205205
echo " OPTIMIZED_BINARY failure=$FAIL"
206206
else
207207
echo "Did it pass? Type the return code [0 = pass, 1 = fail]"
208-
read -n1 PASS
208+
read -n1 FAIL
209209
fi
210210
else
211211
FAIL=1

clang/docs/ReleaseNotes.rst

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -332,6 +332,8 @@ C23 Feature Support
332332

333333
- Clang now supports `N3029 <https://www.open-std.org/jtc1/sc22/wg14/www/docs/n3029.htm>`_ Improved Normal Enumerations.
334334
- Clang now officially supports `N3030 <https://www.open-std.org/jtc1/sc22/wg14/www/docs/n3030.htm>`_ Enhancements to Enumerations. Clang already supported it as an extension, so there were no changes to compiler behavior.
335+
- Fixed the value of ``BOOL_WIDTH`` in ``<limits.h>`` to return ``1``
336+
explicitly, as mandated by the standard. Fixes #GH117348
335337

336338
Non-comprehensive list of changes in this release
337339
-------------------------------------------------

clang/docs/SanitizerCoverage.rst

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -385,6 +385,20 @@ Users need to implement a single function to capture the CF table at startup:
385385
// the collected control flow.
386386
}
387387
388+
Gated Trace Callbacks
389+
=====================
390+
391+
Gate the invocation of the tracing callbacks with
392+
``-sanitizer-coverage-gated-trace-callbacks``.
393+
394+
When this option is enabled, the instrumentation will not call into the
395+
runtime-provided callbacks for tracing, thus only incurring in a trivial
396+
branch without going through a function call.
397+
398+
It is up to the runtime to toggle the value of the global variable in order to
399+
enable tracing.
400+
401+
This option is only supported for trace-pc-guard and trace-cmp.
388402

389403
Disabling instrumentation with ``__attribute__((no_sanitize("coverage")))``
390404
===========================================================================

clang/include/clang/Basic/Builtins.td

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4882,7 +4882,6 @@ def HLSLSaturate : LangBuiltin<"HLSL_LANG"> {
48824882
let Prototype = "void(...)";
48834883
}
48844884

4885-
48864885
def HLSLSelect : LangBuiltin<"HLSL_LANG"> {
48874886
let Spellings = ["__builtin_hlsl_select"];
48884887
let Attributes = [NoThrow, Const];
@@ -4907,6 +4906,12 @@ def HLSLRadians : LangBuiltin<"HLSL_LANG"> {
49074906
let Prototype = "void(...)";
49084907
}
49094908

4909+
def HLSLBufferUpdateCounter : LangBuiltin<"HLSL_LANG"> {
4910+
let Spellings = ["__builtin_hlsl_buffer_update_counter"];
4911+
let Attributes = [NoThrow];
4912+
let Prototype = "uint32_t(...)";
4913+
}
4914+
49104915
def HLSLSplitDouble: LangBuiltin<"HLSL_LANG"> {
49114916
let Spellings = ["__builtin_hlsl_elementwise_splitdouble"];
49124917
let Attributes = [NoThrow, Const];

clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 17 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -263,7 +263,7 @@ TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "gfx940
263263
TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hfIb", "nc", "dot10-insts")
264264
TARGET_BUILTIN(__builtin_amdgcn_fdot2_f16_f16, "hV2hV2hh", "nc", "dot9-insts")
265265
TARGET_BUILTIN(__builtin_amdgcn_fdot2_bf16_bf16, "sV2sV2ss", "nc", "dot9-insts")
266-
TARGET_BUILTIN(__builtin_amdgcn_fdot2_f32_bf16, "fV2sV2sfIb", "nc", "dot9-insts")
266+
TARGET_BUILTIN(__builtin_amdgcn_fdot2_f32_bf16, "fV2sV2sfIb", "nc", "dot12-insts")
267267
TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSiIb", "nc", "dot2-insts")
268268
TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUiIb", "nc", "dot2-insts")
269269
TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSiIb", "nc", "dot1-insts")
@@ -276,6 +276,7 @@ TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_bf8, "fUiUif", "nc", "dot11-insts")
276276
TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_fp8, "fUiUif", "nc", "dot11-insts")
277277
TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_fp8, "fUiUif", "nc", "dot11-insts")
278278
TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_bf8, "fUiUif", "nc", "dot11-insts")
279+
TARGET_BUILTIN(__builtin_amdgcn_fdot2c_f32_bf16, "fV2yV2yfIb", "nc", "dot13-insts")
279280

280281
//===----------------------------------------------------------------------===//
281282
// GFX10+ only builtins.
@@ -462,6 +463,17 @@ TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8, "V16fV4iV8iV16fiIiI
462463
TARGET_BUILTIN(__builtin_amdgcn_permlane16_swap, "V2UiUiUiIbIb", "nc", "permlane16-swap")
463464
TARGET_BUILTIN(__builtin_amdgcn_permlane32_swap, "V2UiUiUiIbIb", "nc", "permlane32-swap")
464465

466+
TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr4_b64_v2i32, "V2iV2i*3", "nc", "gfx950-insts")
467+
TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr6_b96_v3i32, "V3iV3i*3", "nc", "gfx950-insts")
468+
TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr8_b64_v2i32, "V2iV2i*3", "nc", "gfx950-insts")
469+
TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr16_b64_v4i16, "V4sV4s*3", "nc", "gfx950-insts")
470+
471+
TARGET_BUILTIN(__builtin_amdgcn_ashr_pk_i8_i32, "UsUiUiUi", "nc", "ashr-pk-insts")
472+
TARGET_BUILTIN(__builtin_amdgcn_ashr_pk_u8_i32, "UsUiUiUi", "nc", "ashr-pk-insts")
473+
474+
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_2xpk16_fp6_f32, "V6UiV16fV16ff", "nc", "gfx950-insts")
475+
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32, "V6UiV16fV16ff", "nc", "gfx950-insts")
476+
465477
//===----------------------------------------------------------------------===//
466478
// GFX12+ only builtins.
467479
//===----------------------------------------------------------------------===//
@@ -554,6 +566,10 @@ TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64, "V4fiV2iV4fs",
554566
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64, "V4fiV2iV4fs", "nc", "gfx12-insts,wavefrontsize64")
555567

556568
TARGET_BUILTIN(__builtin_amdgcn_prng_b32, "UiUi", "nc", "prng-inst")
569+
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_fp6_f16, "V6UiV32hf", "nc", "f16bf16-to-fp6bf6-cvt-scale-insts")
570+
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_bf6_f16, "V6UiV32hf", "nc", "f16bf16-to-fp6bf6-cvt-scale-insts")
571+
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_fp6_bf16, "V6UiV32yf", "nc", "f16bf16-to-fp6bf6-cvt-scale-insts")
572+
TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_bf6_bf16, "V6UiV32yf", "nc", "f16bf16-to-fp6bf6-cvt-scale-insts")
557573

558574
#undef BUILTIN
559575
#undef TARGET_BUILTIN

clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7287,6 +7287,8 @@ def err_typecheck_illegal_increment_decrement : Error<
72877287
"cannot %select{decrement|increment}1 value of type %0">;
72887288
def err_typecheck_expect_int : Error<
72897289
"used type %0 where integer is required">;
7290+
def err_typecheck_expect_hlsl_resource : Error<
7291+
"used type %0 where __hlsl_resource_t is required">;
72907292
def err_typecheck_arithmetic_incomplete_or_sizeless_type : Error<
72917293
"arithmetic on a pointer to %select{an incomplete|sizeless}0 type %1">;
72927294
def err_typecheck_pointer_arith_function_type : Error<
@@ -12528,6 +12530,10 @@ def warn_attr_min_eq_max: Warning<
1252812530

1252912531
def err_hlsl_attribute_number_arguments_insufficient_shader_model: Error<
1253012532
"attribute %0 with %1 arguments requires shader model %2 or greater">;
12533+
def err_hlsl_expect_arg_const_int_one_or_neg_one: Error<
12534+
"argument %0 must be constant integer 1 or -1">;
12535+
def err_invalid_hlsl_resource_type: Error<
12536+
"invalid __hlsl_resource_t type attributes">;
1253112537

1253212538
// Layout randomization diagnostics.
1253312539
def err_non_designated_init_used : Error<

0 commit comments

Comments
 (0)