Skip to content

Commit 6dbc0b2

Browse files
SC llvm teamSC llvm team
SC llvm team
authored and
SC llvm team
committed
Merge upstream LLVM into amd-gfx12
2 parents a0d7e89 + 59fdd97 commit 6dbc0b2

File tree

8 files changed

+258
-35
lines changed

8 files changed

+258
-35
lines changed

‎clang-tools-extra/clang-tidy/modernize/UseUsingCheck.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -92,7 +92,7 @@ void UseUsingCheck::check(const MatchFinder::MatchResult &Result) {
9292
const LangOptions &LO = getLangOpts();
9393

9494
// Match CXXRecordDecl only to store the range of the last non-implicit full
95-
// declaration, to later check whether it's within the typdef itself.
95+
// declaration, to later check whether it's within the typedef itself.
9696
const auto *MatchedTagDecl = Result.Nodes.getNodeAs<TagDecl>(TagDeclName);
9797
if (MatchedTagDecl) {
9898
// It is not sufficient to just track the last TagDecl that we've seen,
@@ -152,7 +152,7 @@ void UseUsingCheck::check(const MatchFinder::MatchResult &Result) {
152152
StringRef ExtraReference = "";
153153
if (MainTypeEndLoc.isValid() && TypeRange.fullyContains(MainTypeEndLoc)) {
154154
// Each type introduced in a typedef can specify being a reference or
155-
// pointer type seperately, so we need to sigure out if the new using-decl
155+
// pointer type separately, so we need to figure out if the new using-decl
156156
// needs to be to a reference or pointer as well.
157157
const SourceLocation Tok = utils::lexer::findPreviousAnyTokenKind(
158158
MatchedDecl->getLocation(), SM, LO, tok::TokenKind::star,

‎llvm/include/llvm/Support/ThreadPool.h‎

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -149,10 +149,6 @@ class LLVM_ABI StdThreadPool : public ThreadPoolInterface {
149149
/// number of threads!
150150
unsigned getMaxConcurrency() const override { return MaxThreadCount; }
151151

152-
// TODO: Remove, misleading legacy name warning!
153-
LLVM_DEPRECATED("Use getMaxConcurrency instead", "getMaxConcurrency")
154-
unsigned getThreadCount() const { return MaxThreadCount; }
155-
156152
/// Returns true if the current thread is a worker thread of this thread pool.
157153
bool isWorkerThread() const;
158154

@@ -233,10 +229,6 @@ class LLVM_ABI SingleThreadExecutor : public ThreadPoolInterface {
233229
/// Returns always 1: there is no concurrency.
234230
unsigned getMaxConcurrency() const override { return 1; }
235231

236-
// TODO: Remove, misleading legacy name warning!
237-
LLVM_DEPRECATED("Use getMaxConcurrency instead", "getMaxConcurrency")
238-
unsigned getThreadCount() const { return 1; }
239-
240232
/// Returns true if the current thread is a worker thread of this thread pool.
241233
bool isWorkerThread() const;
242234

‎llvm/lib/ExecutionEngine/RuntimeDyld/RuntimeDyld.cpp‎

Lines changed: 10 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -183,7 +183,7 @@ RuntimeDyldImpl::loadObjectImpl(const object::ObjectFile &Obj) {
183183
std::lock_guard<sys::Mutex> locked(lock);
184184

185185
// Save information about our target
186-
Arch = (Triple::ArchType)Obj.getArch();
186+
Arch = Obj.getArch();
187187
IsTargetLittleEndian = Obj.isLittleEndian();
188188
setMipsABI(Obj);
189189

@@ -1361,18 +1361,17 @@ std::unique_ptr<RuntimeDyld::LoadedObjectInfo>
13611361
RuntimeDyld::loadObject(const ObjectFile &Obj) {
13621362
if (!Dyld) {
13631363
if (Obj.isELF())
1364-
Dyld =
1365-
createRuntimeDyldELF(static_cast<Triple::ArchType>(Obj.getArch()),
1366-
MemMgr, Resolver, ProcessAllSections,
1367-
std::move(NotifyStubEmitted));
1364+
Dyld = createRuntimeDyldELF(Obj.getArch(), MemMgr, Resolver,
1365+
ProcessAllSections,
1366+
std::move(NotifyStubEmitted));
13681367
else if (Obj.isMachO())
1369-
Dyld = createRuntimeDyldMachO(
1370-
static_cast<Triple::ArchType>(Obj.getArch()), MemMgr, Resolver,
1371-
ProcessAllSections, std::move(NotifyStubEmitted));
1368+
Dyld = createRuntimeDyldMachO(Obj.getArch(), MemMgr, Resolver,
1369+
ProcessAllSections,
1370+
std::move(NotifyStubEmitted));
13721371
else if (Obj.isCOFF())
1373-
Dyld = createRuntimeDyldCOFF(
1374-
static_cast<Triple::ArchType>(Obj.getArch()), MemMgr, Resolver,
1375-
ProcessAllSections, std::move(NotifyStubEmitted));
1372+
Dyld = createRuntimeDyldCOFF(Obj.getArch(), MemMgr, Resolver,
1373+
ProcessAllSections,
1374+
std::move(NotifyStubEmitted));
13761375
else
13771376
report_fatal_error("Incompatible object format!");
13781377
}

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.xf32.gfx942.ll‎

Lines changed: 243 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,100 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2-
; RUN: llc -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck --check-prefixes=GCN,GFX942 %s
3-
; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -global-isel < %s | FileCheck --check-prefixes=GCN,GISEL %s
4-
; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 < %s | FileCheck --check-prefixes=GCN,GFX942 %s
5-
; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 -global-isel < %s | FileCheck --check-prefixes=GCN,GISEL %s
2+
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck --check-prefixes=GFX942,GFX942-SDAG %s
3+
; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck --check-prefixes=GFX942,GFX942-GISEL %s
4+
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 < %s | FileCheck --check-prefixes=GFX942-STRESS,GFX942-SDAG-STRESS %s
5+
; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 < %s | FileCheck --check-prefixes=GFX942-STRESS,GFX942-GISEL-STRESS %s
66

77
declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x8.xf32(<2 x float>, <2 x float>, <4 x float>, i32, i32, i32)
88
declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x4.xf32(<2 x float>, <2 x float>, <16 x float>, i32, i32, i32)
99

1010
define amdgpu_kernel void @test_mfma_f32_16x16x8xf32(ptr addrspace(1) %arg) #0 {
11+
; GFX942-SDAG-LABEL: test_mfma_f32_16x16x8xf32:
12+
; GFX942-SDAG: ; %bb.0: ; %bb
13+
; GFX942-SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
14+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v4, 1.0
15+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v5, 2.0
16+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v0, 0x40400000
17+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v1, 4.0
18+
; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0)
19+
; GFX942-SDAG-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0
20+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v2, 0
21+
; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0)
22+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a0, s0
23+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a1, s1
24+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a2, s2
25+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a3, s3
26+
; GFX942-SDAG-NEXT: s_nop 1
27+
; GFX942-SDAG-NEXT: v_mfma_f32_16x16x8_xf32 a[0:3], v[4:5], v[0:1], a[0:3] cbsz:1 abid:2 blgp:3
28+
; GFX942-SDAG-NEXT: s_nop 6
29+
; GFX942-SDAG-NEXT: global_store_dwordx4 v2, a[0:3], s[6:7]
30+
; GFX942-SDAG-NEXT: s_endpgm
31+
;
32+
; GFX942-GISEL-LABEL: test_mfma_f32_16x16x8xf32:
33+
; GFX942-GISEL: ; %bb.0: ; %bb
34+
; GFX942-GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
35+
; GFX942-GISEL-NEXT: s_mov_b32 s4, 1.0
36+
; GFX942-GISEL-NEXT: s_mov_b32 s5, 2.0
37+
; GFX942-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[4:5]
38+
; GFX942-GISEL-NEXT: s_mov_b32 s4, 0x40400000
39+
; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0)
40+
; GFX942-GISEL-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0
41+
; GFX942-GISEL-NEXT: s_mov_b32 s5, 4.0
42+
; GFX942-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[4:5]
43+
; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0)
44+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a0, s0
45+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a1, s1
46+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a2, s2
47+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a3, s3
48+
; GFX942-GISEL-NEXT: s_nop 1
49+
; GFX942-GISEL-NEXT: v_mfma_f32_16x16x8_xf32 a[0:3], v[0:1], v[2:3], a[0:3] cbsz:1 abid:2 blgp:3
50+
; GFX942-GISEL-NEXT: v_mov_b32_e32 v0, 0
51+
; GFX942-GISEL-NEXT: s_nop 5
52+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[0:3], s[6:7]
53+
; GFX942-GISEL-NEXT: s_endpgm
54+
;
55+
; GFX942-SDAG-STRESS-LABEL: test_mfma_f32_16x16x8xf32:
56+
; GFX942-SDAG-STRESS: ; %bb.0: ; %bb
57+
; GFX942-SDAG-STRESS-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
58+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v0, 1.0
59+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v1, 2.0
60+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v2, 0x40400000
61+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v3, 4.0
62+
; GFX942-SDAG-STRESS-NEXT: s_waitcnt lgkmcnt(0)
63+
; GFX942-SDAG-STRESS-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0
64+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v4, 0
65+
; GFX942-SDAG-STRESS-NEXT: s_waitcnt lgkmcnt(0)
66+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a0, s0
67+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a1, s1
68+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a2, s2
69+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a3, s3
70+
; GFX942-SDAG-STRESS-NEXT: s_nop 1
71+
; GFX942-SDAG-STRESS-NEXT: v_mfma_f32_16x16x8_xf32 a[0:3], v[0:1], v[2:3], a[0:3] cbsz:1 abid:2 blgp:3
72+
; GFX942-SDAG-STRESS-NEXT: s_nop 6
73+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v4, a[0:3], s[6:7]
74+
; GFX942-SDAG-STRESS-NEXT: s_endpgm
75+
;
76+
; GFX942-GISEL-STRESS-LABEL: test_mfma_f32_16x16x8xf32:
77+
; GFX942-GISEL-STRESS: ; %bb.0: ; %bb
78+
; GFX942-GISEL-STRESS-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
79+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s0, 1.0
80+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s2, 0x40400000
81+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s1, 2.0
82+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s3, 4.0
83+
; GFX942-GISEL-STRESS-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
84+
; GFX942-GISEL-STRESS-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
85+
; GFX942-GISEL-STRESS-NEXT: s_waitcnt lgkmcnt(0)
86+
; GFX942-GISEL-STRESS-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0
87+
; GFX942-GISEL-STRESS-NEXT: s_waitcnt lgkmcnt(0)
88+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a0, s0
89+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a1, s1
90+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a2, s2
91+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a3, s3
92+
; GFX942-GISEL-STRESS-NEXT: s_nop 1
93+
; GFX942-GISEL-STRESS-NEXT: v_mfma_f32_16x16x8_xf32 a[0:3], v[0:1], v[2:3], a[0:3] cbsz:1 abid:2 blgp:3
94+
; GFX942-GISEL-STRESS-NEXT: v_mov_b32_e32 v0, 0
95+
; GFX942-GISEL-STRESS-NEXT: s_nop 5
96+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[0:3], s[6:7]
97+
; GFX942-GISEL-STRESS-NEXT: s_endpgm
1198
bb:
1299
%in.1 = load <4 x float>, ptr addrspace(1) %arg
13100
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8.xf32(<2 x float> <float 1.0, float 2.0>, <2 x float> <float 3.0, float 4.0>, <4 x float> %in.1, i32 1, i32 2, i32 3)
@@ -16,6 +103,157 @@ bb:
16103
}
17104

18105
define amdgpu_kernel void @test_mfma_f32_32x32x4xf32(ptr addrspace(1) %arg) #0 {
106+
; GFX942-SDAG-LABEL: test_mfma_f32_32x32x4xf32:
107+
; GFX942-SDAG: ; %bb.0: ; %bb
108+
; GFX942-SDAG-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24
109+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v2, 1.0
110+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v3, 2.0
111+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v0, 0x40400000
112+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v1, 4.0
113+
; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0)
114+
; GFX942-SDAG-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0
115+
; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0)
116+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a0, s0
117+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a1, s1
118+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a2, s2
119+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a3, s3
120+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a4, s4
121+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a5, s5
122+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a6, s6
123+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a7, s7
124+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a8, s8
125+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a9, s9
126+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a10, s10
127+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a11, s11
128+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a12, s12
129+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a13, s13
130+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a14, s14
131+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a15, s15
132+
; GFX942-SDAG-NEXT: s_nop 1
133+
; GFX942-SDAG-NEXT: v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[0:1], a[0:15] cbsz:1 abid:2 blgp:3
134+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v0, 0
135+
; GFX942-SDAG-NEXT: s_nop 7
136+
; GFX942-SDAG-NEXT: s_nop 1
137+
; GFX942-SDAG-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
138+
; GFX942-SDAG-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
139+
; GFX942-SDAG-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
140+
; GFX942-SDAG-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]
141+
; GFX942-SDAG-NEXT: s_endpgm
142+
;
143+
; GFX942-GISEL-LABEL: test_mfma_f32_32x32x4xf32:
144+
; GFX942-GISEL: ; %bb.0: ; %bb
145+
; GFX942-GISEL-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24
146+
; GFX942-GISEL-NEXT: s_mov_b32 s18, 1.0
147+
; GFX942-GISEL-NEXT: s_mov_b32 s19, 2.0
148+
; GFX942-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[18:19]
149+
; GFX942-GISEL-NEXT: s_mov_b32 s18, 0x40400000
150+
; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0)
151+
; GFX942-GISEL-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0
152+
; GFX942-GISEL-NEXT: s_mov_b32 s19, 4.0
153+
; GFX942-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[18:19]
154+
; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0)
155+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a0, s0
156+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a1, s1
157+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a2, s2
158+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a3, s3
159+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a4, s4
160+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a5, s5
161+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a6, s6
162+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a7, s7
163+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a8, s8
164+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a9, s9
165+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a10, s10
166+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a11, s11
167+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a12, s12
168+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a13, s13
169+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a14, s14
170+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a15, s15
171+
; GFX942-GISEL-NEXT: s_nop 1
172+
; GFX942-GISEL-NEXT: v_mfma_f32_32x32x4_xf32 a[0:15], v[0:1], v[2:3], a[0:15] cbsz:1 abid:2 blgp:3
173+
; GFX942-GISEL-NEXT: v_mov_b32_e32 v0, 0
174+
; GFX942-GISEL-NEXT: s_nop 7
175+
; GFX942-GISEL-NEXT: s_nop 1
176+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]
177+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
178+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
179+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
180+
; GFX942-GISEL-NEXT: s_endpgm
181+
;
182+
; GFX942-SDAG-STRESS-LABEL: test_mfma_f32_32x32x4xf32:
183+
; GFX942-SDAG-STRESS: ; %bb.0: ; %bb
184+
; GFX942-SDAG-STRESS-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24
185+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v0, 1.0
186+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v1, 2.0
187+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v2, 0x40400000
188+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v3, 4.0
189+
; GFX942-SDAG-STRESS-NEXT: s_waitcnt lgkmcnt(0)
190+
; GFX942-SDAG-STRESS-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0
191+
; GFX942-SDAG-STRESS-NEXT: s_waitcnt lgkmcnt(0)
192+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a0, s0
193+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a1, s1
194+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a2, s2
195+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a3, s3
196+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a4, s4
197+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a5, s5
198+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a6, s6
199+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a7, s7
200+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a8, s8
201+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a9, s9
202+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a10, s10
203+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a11, s11
204+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a12, s12
205+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a13, s13
206+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a14, s14
207+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a15, s15
208+
; GFX942-SDAG-STRESS-NEXT: s_nop 1
209+
; GFX942-SDAG-STRESS-NEXT: v_mfma_f32_32x32x4_xf32 a[0:15], v[0:1], v[2:3], a[0:15] cbsz:1 abid:2 blgp:3
210+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v0, 0
211+
; GFX942-SDAG-STRESS-NEXT: s_nop 7
212+
; GFX942-SDAG-STRESS-NEXT: s_nop 1
213+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
214+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
215+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
216+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]
217+
; GFX942-SDAG-STRESS-NEXT: s_endpgm
218+
;
219+
; GFX942-GISEL-STRESS-LABEL: test_mfma_f32_32x32x4xf32:
220+
; GFX942-GISEL-STRESS: ; %bb.0: ; %bb
221+
; GFX942-GISEL-STRESS-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24
222+
; GFX942-GISEL-STRESS-NEXT: s_waitcnt lgkmcnt(0)
223+
; GFX942-GISEL-STRESS-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0
224+
; GFX942-GISEL-STRESS-NEXT: s_waitcnt lgkmcnt(0)
225+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a0, s0
226+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a1, s1
227+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a2, s2
228+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a3, s3
229+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a4, s4
230+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a5, s5
231+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a6, s6
232+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a7, s7
233+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a8, s8
234+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a9, s9
235+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a10, s10
236+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a11, s11
237+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a12, s12
238+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a13, s13
239+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a14, s14
240+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a15, s15
241+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s0, 1.0
242+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s1, 2.0
243+
; GFX942-GISEL-STRESS-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
244+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s0, 0x40400000
245+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s1, 4.0
246+
; GFX942-GISEL-STRESS-NEXT: v_mov_b64_e32 v[2:3], s[0:1]
247+
; GFX942-GISEL-STRESS-NEXT: s_nop 1
248+
; GFX942-GISEL-STRESS-NEXT: v_mfma_f32_32x32x4_xf32 a[0:15], v[0:1], v[2:3], a[0:15] cbsz:1 abid:2 blgp:3
249+
; GFX942-GISEL-STRESS-NEXT: v_mov_b32_e32 v0, 0
250+
; GFX942-GISEL-STRESS-NEXT: s_nop 7
251+
; GFX942-GISEL-STRESS-NEXT: s_nop 1
252+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]
253+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
254+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
255+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
256+
; GFX942-GISEL-STRESS-NEXT: s_endpgm
19257
bb:
20258
%in.1 = load <16 x float>, ptr addrspace(1) %arg
21259
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4.xf32(<2 x float> <float 1.0, float 2.0>, <2 x float> <float 3.0, float 4.0>, <16 x float> %in.1, i32 1, i32 2, i32 3)
@@ -25,6 +263,5 @@ bb:
25263

26264
attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
27265
;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
28-
; GCN: {{.*}}
29266
; GFX942: {{.*}}
30-
; GISEL: {{.*}}
267+
; GFX942-STRESS: {{.*}}

‎llvm/unittests/CodeGen/DroppedVariableStatsMIRTest.cpp‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -39,8 +39,7 @@ createTargetMachine(std::string TargetStr, StringRef CPU, StringRef FS) {
3939
return nullptr;
4040
TargetOptions Options;
4141
return std::unique_ptr<TargetMachine>(
42-
static_cast<TargetMachine *>(T->createTargetMachine(
43-
TT, CPU, FS, Options, std::nullopt, std::nullopt)));
42+
T->createTargetMachine(TT, CPU, FS, Options, std::nullopt, std::nullopt));
4443
}
4544

4645
std::unique_ptr<Module> parseMIR(const TargetMachine &TM, StringRef MIRCode,

‎llvm/utils/gn/secondary/llvm/include/llvm/Config/BUILD.gn‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -296,6 +296,7 @@ write_cmake_config("llvm-config") {
296296
input = "llvm-config.h.cmake"
297297
output = "$target_gen_dir/llvm-config.h"
298298
values = [
299+
"LLVM_ENABLE_PROFCHECK=",
299300
"LLVM_BUILD_LLVM_DYLIB=",
300301
"LLVM_BUILD_SHARED_LIBS=",
301302
"LLVM_ENABLE_LLVM_C_EXPORT_ANNOTATIONS=",

‎llvm/utils/gn/secondary/llvm/lib/MC/BUILD.gn‎

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -57,13 +57,7 @@ static_library("MC") {
5757
"MCSPIRVStreamer.cpp",
5858
"MCSchedule.cpp",
5959
"MCSection.cpp",
60-
"MCSectionCOFF.cpp",
61-
"MCSectionDXContainer.cpp",
62-
"MCSectionELF.cpp",
63-
"MCSectionGOFF.cpp",
6460
"MCSectionMachO.cpp",
65-
"MCSectionWasm.cpp",
66-
"MCSectionXCOFF.cpp",
6761
"MCStreamer.cpp",
6862
"MCSubtargetInfo.cpp",
6963
"MCSymbol.cpp",

‎llvm/utils/gn/secondary/llvm/test/BUILD.gn‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -64,6 +64,7 @@ write_lit_config("lit_site_cfg") {
6464
"LLVM_APPEND_VC_REV=0",
6565
"LLVM_ENABLE_FFI=0",
6666
"LLVM_ENABLE_HTTPLIB=0",
67+
"LLVM_ENABLE_PROFCHECK=0",
6768
"LLVM_EXPERIMENTAL_KEY_INSTRUCTIONS=0",
6869
"LLVM_FORCE_VC_REVISION=",
6970
"LLVM_HAS_LOGF128=0",

0 commit comments

Comments
 (0)