Skip to content

Commit ce419d8

Browse files
jrbyrnesjrbyrnes
committed
Finish porting changes
Change-Id: I7fc0bdd7d74464f5f95b9c2fb977cdd759d81215
1 parent 7bc4fce commit ce419d8

File tree

3 files changed

+13
-52
lines changed

3 files changed

+13
-52
lines changed

‎llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8734,7 +8734,7 @@ void AMDGPUAsmParser::cvtScaledMFMA(MCInst &Inst,
87348734
if (OpselHiIdx != OptionalIdx.end())
87358735
OpSelHi = static_cast<const AMDGPUOperand &>(*Operands[OpselHiIdx->second])
87368736
.getImm();
8737-
static const AMDGPU::OpName ModOps[] = {AMDGPU::OpName::src0_modifiers,
8737+
static const int ModOps[] = {AMDGPU::OpName::src0_modifiers,
87388738
AMDGPU::OpName::src1_modifiers};
87398739

87408740
for (unsigned J = 0; J < 2; ++J) {

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.ll‎

Lines changed: 4 additions & 43 deletions
Original file line numberDiff line numberDiff line change
@@ -1899,11 +1899,7 @@ define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd(<8 x i32
18991899
; SDAG-NEXT: v_mov_b32_e32 v19, s3
19001900
; SDAG-NEXT: v_mov_b32_e32 v21, s5
19011901
; SDAG-NEXT: s_nop 1
1902-
<<<<<<< HEAD
1903-
; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], s4, v21 op_sel_hi:[0,0,0] blgp:2
1904-
=======
1905-
; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], s12, v17 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
1906-
>>>>>>> f33f71291d5d... [AMDGPU] Fix opsel for scaled MFMA operations
1902+
; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], s4, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
19071903
; SDAG-NEXT: s_nop 7
19081904
; SDAG-NEXT: s_nop 3
19091905
; SDAG-NEXT: global_store_dwordx4 v20, v[0:3], s[6:7]
@@ -1926,13 +1922,8 @@ define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd(<8 x i32
19261922
; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[22:23]
19271923
; GISEL-NEXT: v_mov_b32_e32 v20, s25
19281924
; GISEL-NEXT: s_nop 1
1929-
<<<<<<< HEAD
1930-
; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], s24, v20 op_sel_hi:[0,0,0] blgp:2
1925+
; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], s24, v20 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
19311926
; GISEL-NEXT: v_mov_b32_e32 v4, 0
1932-
=======
1933-
; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], s28, v16 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
1934-
; GISEL-NEXT: v_mov_b32_e32 v0, 0
1935-
>>>>>>> f33f71291d5d... [AMDGPU] Fix opsel for scaled MFMA operations
19361927
; GISEL-NEXT: s_nop 7
19371928
; GISEL-NEXT: s_nop 2
19381929
; GISEL-NEXT: global_store_dwordx4 v4, v[0:3], s[26:27]
@@ -1971,33 +1962,8 @@ define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA
19711962
; SDAG-NEXT: v_mov_b64_e32 v[16:17], s[4:5]
19721963
; SDAG-NEXT: s_load_dwordx2 s[0:1], s[0:1], 0x50
19731964
; SDAG-NEXT: s_nop 0
1974-
; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], s2, -2 op_sel_hi:[0,0,0]
1965+
; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], s2, -2 op_sel:[1,1,0] op_sel_hi:[1,0,0]
19751966
; SDAG-NEXT: s_waitcnt lgkmcnt(0)
1976-
<<<<<<< HEAD
1977-
=======
1978-
; SDAG-NEXT: v_mov_b32_e32 v0, s8
1979-
; SDAG-NEXT: v_mov_b32_e32 v1, s9
1980-
; SDAG-NEXT: v_mov_b32_e32 v2, s10
1981-
; SDAG-NEXT: v_mov_b32_e32 v3, s11
1982-
; SDAG-NEXT: v_mov_b32_e32 v4, s12
1983-
; SDAG-NEXT: v_mov_b32_e32 v5, s13
1984-
; SDAG-NEXT: v_mov_b32_e32 v6, s14
1985-
; SDAG-NEXT: v_mov_b32_e32 v7, s15
1986-
; SDAG-NEXT: v_accvgpr_write_b32 a0, s0
1987-
; SDAG-NEXT: v_mov_b32_e32 v8, s16
1988-
; SDAG-NEXT: v_mov_b32_e32 v9, s17
1989-
; SDAG-NEXT: v_mov_b32_e32 v10, s18
1990-
; SDAG-NEXT: v_mov_b32_e32 v11, s19
1991-
; SDAG-NEXT: v_mov_b32_e32 v12, s20
1992-
; SDAG-NEXT: v_mov_b32_e32 v13, s21
1993-
; SDAG-NEXT: v_mov_b32_e32 v14, s22
1994-
; SDAG-NEXT: v_mov_b32_e32 v15, s23
1995-
; SDAG-NEXT: v_accvgpr_write_b32 a1, s1
1996-
; SDAG-NEXT: v_accvgpr_write_b32 a2, s2
1997-
; SDAG-NEXT: v_accvgpr_write_b32 a3, s3
1998-
; SDAG-NEXT: s_nop 1
1999-
; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], s6, -2 op_sel:[1,1,0] op_sel_hi:[1,0,0]
2000-
>>>>>>> f33f71291d5d... [AMDGPU] Fix opsel for scaled MFMA operations
20011967
; SDAG-NEXT: s_nop 7
20021968
; SDAG-NEXT: s_nop 2
20031969
; SDAG-NEXT: global_store_dwordx4 v20, v[0:3], s[0:1]
@@ -2021,13 +1987,8 @@ define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA
20211987
; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[18:19]
20221988
; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[22:23]
20231989
; GISEL-NEXT: s_nop 1
2024-
<<<<<<< HEAD
2025-
; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v20, -2 op_sel_hi:[0,0,0]
1990+
; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v20, -2 op_sel:[1,1,0] op_sel_hi:[1,0,0]
20261991
; GISEL-NEXT: v_mov_b32_e32 v4, 0
2027-
=======
2028-
; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, -2 op_sel:[1,1,0] op_sel_hi:[1,0,0]
2029-
; GISEL-NEXT: v_mov_b32_e32 v0, 0
2030-
>>>>>>> f33f71291d5d... [AMDGPU] Fix opsel for scaled MFMA operations
20311992
; GISEL-NEXT: s_nop 7
20321993
; GISEL-NEXT: s_nop 2
20331994
; GISEL-NEXT: global_store_dwordx4 v4, v[0:3], s[0:1]

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.ll‎

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -4370,8 +4370,8 @@ define amdgpu_kernel void @test_mfma_scale_f32_32x32x64_f8f6f4__vgprcd(<8 x i32>
43704370
; SDAG-NEXT: v_mov_b64_e32 v[14:15], s[50:51]
43714371
; SDAG-NEXT: v_mov_b32_e32 v32, s5
43724372
; SDAG-NEXT: s_nop 1
4373-
; SDAG-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[0:7], v[8:15], a[0:15], s0, v16 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
4374-
; SDAG-NEXT: v_mov_b32_e32 v0, 0
4373+
; SDAG-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15], s4, v32 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
4374+
; SDAG-NEXT: v_mov_b32_e32 v16, 0
43754375
; SDAG-NEXT: s_nop 7
43764376
; SDAG-NEXT: s_nop 7
43774377
; SDAG-NEXT: s_nop 2
@@ -4405,8 +4405,8 @@ define amdgpu_kernel void @test_mfma_scale_f32_32x32x64_f8f6f4__vgprcd(<8 x i32>
44054405
; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[50:51]
44064406
; GISEL-NEXT: v_mov_b32_e32 v32, s21
44074407
; GISEL-NEXT: s_nop 1
4408-
; GISEL-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[0:7], v[8:15], a[0:15], s0, v16 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
4409-
; GISEL-NEXT: v_mov_b32_e32 v0, 0
4408+
; GISEL-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15], s20, v32 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
4409+
; GISEL-NEXT: v_mov_b32_e32 v16, 0
44104410
; GISEL-NEXT: s_nop 7
44114411
; GISEL-NEXT: s_nop 7
44124412
; GISEL-NEXT: s_nop 2
@@ -4453,8 +4453,8 @@ define amdgpu_kernel void @test_mfma_scale_f32_32x32x64_f8f6f4__vgprcd___scaleA_
44534453
; SDAG-NEXT: v_mov_b64_e32 v[12:13], s[48:49]
44544454
; SDAG-NEXT: v_mov_b64_e32 v[14:15], s[50:51]
44554455
; SDAG-NEXT: s_nop 1
4456-
; SDAG-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[0:7], v[8:15], a[0:15], s2, -2 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
4457-
; SDAG-NEXT: v_mov_b32_e32 v0, 0
4456+
; SDAG-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15], s2, -2 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
4457+
; SDAG-NEXT: v_mov_b32_e32 v16, 0
44584458
; SDAG-NEXT: s_nop 7
44594459
; SDAG-NEXT: s_nop 7
44604460
; SDAG-NEXT: s_nop 2
@@ -4488,8 +4488,8 @@ define amdgpu_kernel void @test_mfma_scale_f32_32x32x64_f8f6f4__vgprcd___scaleA_
44884488
; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[48:49]
44894489
; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[50:51]
44904490
; GISEL-NEXT: s_nop 1
4491-
; GISEL-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[0:7], v[8:15], a[0:15], v16, -2 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
4492-
; GISEL-NEXT: v_mov_b32_e32 v0, 0
4491+
; GISEL-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15], v32, -2 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:2
4492+
; GISEL-NEXT: v_mov_b32_e32 v16, 0
44934493
; GISEL-NEXT: s_nop 7
44944494
; GISEL-NEXT: s_nop 7
44954495
; GISEL-NEXT: s_nop 2

0 commit comments

Comments
 (0)