Skip to content

Commit 65b9d83

Browse files
arsenmPravin Jagtap
arsenm
authored and
Pravin Jagtap
committed
AMDGPU: Add v_smfmac_i32_16x16x128_i8 for gfx950 (llvm#117213)
1 parent f8163ae commit 65b9d83

File tree

12 files changed

+297
-1
lines changed

12 files changed

+297
-1
lines changed

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -448,6 +448,7 @@ TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_f16, "V4fV8hV16hV4fiIiIi", "
448448
TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_f16, "V16fV8hV16hV16fiIiIi", "nc", "gfx950-insts")
449449
TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_bf16, "V4fV8yV16yV4fiIiIi", "nc", "gfx950-insts")
450450
TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_bf16, "V16fV8yV16yV16fiIiIi", "nc", "gfx950-insts")
451+
TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_16x16x128_i8, "V4iV4iV8iV4iiIiIi", "nc", "gfx950-insts")
451452

452453
//===----------------------------------------------------------------------===//
453454
// GFX12+ only builtins.

‎clang/test/CodeGenOpenCL/builtins-amdgcn-mfma.cl‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -496,4 +496,11 @@ void test_smfmac_f32_32x32x32_bf16(global v16f* out, v8bf16 a, v16bf16 b, v16f c
496496
*out = __builtin_amdgcn_smfmac_f32_32x32x32_bf16(a, b, c, idx, 0, 0);
497497
}
498498

499+
// CHECK-GFX950-LABEL: @test_smfmac_i32_16x16x128_i8
500+
// CHECK-GFX950: call <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x128.i8(<4 x i32> %a, <8 x i32> %b, <4 x i32> %c, i32 %idx, i32 0, i32 0)
501+
void test_smfmac_i32_16x16x128_i8(global v4i* out, v4i a, v8i b, v4i c, int idx)
502+
{
503+
*out = __builtin_amdgcn_smfmac_i32_16x16x128_i8(a, b, c, idx, 0, 0);
504+
}
505+
499506
#endif

‎clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950-param.cl‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -88,3 +88,9 @@ void test_smfmac_f32_32x32x32_bf16(global float16* out, bfloat8 a, bfloat16 b, f
8888
*out = __builtin_amdgcn_smfmac_f32_32x32x32_bf16(a, b, c, idx, d, 0); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_32x32x32_bf16' must be a constant integer}}
8989
*out = __builtin_amdgcn_smfmac_f32_32x32x32_bf16(a, b, c, idx, 0, d); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_32x32x32_bf16' must be a constant integer}}
9090
}
91+
92+
void test_smfmac_i32_16x16x128_i8(global int4* out, int4 a, int8 b, int4 c, int idx, int d)
93+
{
94+
*out = __builtin_amdgcn_smfmac_i32_16x16x128_i8(a, b, c, idx, d, 0); // expected-error{{argument to '__builtin_amdgcn_smfmac_i32_16x16x128_i8' must be a constant integer}}
95+
*out = __builtin_amdgcn_smfmac_i32_16x16x128_i8(a, b, c, idx, 0, d); // expected-error{{argument to '__builtin_amdgcn_smfmac_i32_16x16x128_i8' must be a constant integer}}
96+
}

‎clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950.cl‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,7 @@ void test(__global float4* out0, half8 a0, half8 b0, float4 c0,
3838
*out7 = __builtin_amdgcn_smfmac_f32_32x32x32_f16(a7, b7, c7, 0, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_32x32x32_f16' needs target feature gfx950-insts}}
3939
*out8 = __builtin_amdgcn_smfmac_f32_16x16x64_bf16(a8, b8, c8, 0, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_16x16x64_bf16' needs target feature gfx950-insts}}
4040
*out9 = __builtin_amdgcn_smfmac_f32_32x32x32_bf16(a9, b9, c9, 0, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_32x32x32_bf16' needs target feature gfx950-insts}}
41+
*out10 = __builtin_amdgcn_smfmac_i32_16x16x128_i8(a10, b10, c10, 0, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_i32_16x16x128_i8' needs target feature gfx950-insts}}
4142
*out14 = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a14, b14, c14, 0, 0, 0, d14, 0, e14); // expected-error{{'__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' needs target feature gfx950-insts}}
4243
*out15 = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a15, b15, c15, 0, 0, 0, d15, 0, e15); // expected-error{{'__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' needs target feature gfx950-insts}}
4344
}

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3183,6 +3183,7 @@ def int_amdgcn_smfmac_f32_16x16x64_f16 : AMDGPUMSmfmacIntrinsic<llvm_v4f32_ty, l
31833183
def int_amdgcn_smfmac_f32_32x32x32_f16 : AMDGPUMSmfmacIntrinsic<llvm_v16f32_ty, llvm_v8f16_ty, llvm_v16f16_ty>;
31843184
def int_amdgcn_smfmac_f32_16x16x64_bf16 : AMDGPUMSmfmacIntrinsic<llvm_v4f32_ty, llvm_v8bf16_ty, llvm_v16bf16_ty>;
31853185
def int_amdgcn_smfmac_f32_32x32x32_bf16 : AMDGPUMSmfmacIntrinsic<llvm_v16f32_ty, llvm_v8bf16_ty, llvm_v16bf16_ty>;
3186+
def int_amdgcn_smfmac_i32_16x16x128_i8 : AMDGPUMSmfmacIntrinsic<llvm_v4i32_ty, llvm_v4i32_ty, llvm_v8i32_ty>;
31863187
}
31873188

31883189
//===----------------------------------------------------------------------===//

‎llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1079,6 +1079,7 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC(MachineInstr &I) const {
10791079
case Intrinsic::amdgcn_smfmac_f32_32x32x32_f16:
10801080
case Intrinsic::amdgcn_smfmac_f32_16x16x64_bf16:
10811081
case Intrinsic::amdgcn_smfmac_f32_32x32x32_bf16:
1082+
case Intrinsic::amdgcn_smfmac_i32_16x16x128_i8:
10821083
return selectSMFMACIntrin(I);
10831084
default:
10841085
return selectImpl(I, *CoverageInfo);
@@ -3509,6 +3510,9 @@ bool AMDGPUInstructionSelector::selectSMFMACIntrin(MachineInstr &MI) const {
35093510
case Intrinsic::amdgcn_smfmac_f32_32x32x32_bf16:
35103511
Opc = AMDGPU::V_SMFMAC_F32_32X32X32_BF16_e64;
35113512
break;
3513+
case Intrinsic::amdgcn_smfmac_i32_16x16x128_i8:
3514+
Opc = AMDGPU::V_SMFMAC_I32_16X16X128_I8_e64;
3515+
break;
35123516
default:
35133517
llvm_unreachable("unhandled smfmac intrinsic");
35143518
}

‎llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4785,7 +4785,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
47854785
case Intrinsic::amdgcn_smfmac_f32_16x16x64_f16:
47864786
case Intrinsic::amdgcn_smfmac_f32_32x32x32_f16:
47874787
case Intrinsic::amdgcn_smfmac_f32_16x16x64_bf16:
4788-
case Intrinsic::amdgcn_smfmac_f32_32x32x32_bf16: {
4788+
case Intrinsic::amdgcn_smfmac_f32_32x32x32_bf16:
4789+
case Intrinsic::amdgcn_smfmac_i32_16x16x128_i8: {
47894790
// vdst, srcA, srcB, srcC, idx
47904791
OpdsMapping[0] = getAGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI);
47914792
OpdsMapping[2] = getVGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI);

‎llvm/lib/Target/AMDGPU/SIInstrInfo.td‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2688,6 +2688,7 @@ def VOP_V4I32_V2I32_V4I32_I32 : VOPProfile <[v4i32, v2i32, v4i32, i32]>;
26882688
def VOP_V16I32_V2I32_V4I32_I32 : VOPProfile <[v16i32, v2i32, v4i32, i32]>;
26892689
def VOP_V4F32_V2I32_V4I32_I32 : VOPProfile <[v4f32, v2i32, v4i32, i32]>;
26902690
def VOP_V16F32_V2I32_V4I32_I32 : VOPProfile <[v16f32, v2i32, v4i32, i32]>;
2691+
def VOP_V4I32_V4I32_V8I32_I32 : VOPProfile <[v4i32, v4i32, v8i32, i32]>;
26912692

26922693
def VOP_V4F32_V8F16_V8F16_V4F32 : VOPProfile <[v4f32, v8f16, v8f16, v4f32]>;
26932694
def VOP_V16F32_V8F16_V8F16_V16F32 : VOPProfile <[v16f32, v8f16, v8f16, v16f32]>;

‎llvm/lib/Target/AMDGPU/VOP3PInstructions.td‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -644,6 +644,7 @@ def VOPProfileSMFMAC_I32_16X16X64_I8 : VOPProfileSMFMAC<VOP_V4I32_V2I32_V4I32_I
644644
def VOPProfileSMFMAC_I32_32X32X32_I8 : VOPProfileSMFMAC<VOP_V16I32_V2I32_V4I32_I32, AVDst_512, AVSrc_64, AVSrc_128>;
645645
def VOPProfileSMFMAC_F32_16X16X64_F8 : VOPProfileSMFMAC<VOP_V4F32_V2I32_V4I32_I32, AVDst_128, AVSrc_64, AVSrc_128>;
646646
def VOPProfileSMFMAC_F32_32X32X32_F8 : VOPProfileSMFMAC<VOP_V16F32_V2I32_V4I32_I32, AVDst_512, AVSrc_64, AVSrc_128>;
647+
def VOPProfileSMFMAC_I32_16X16X128_I8 : VOPProfileSMFMAC<VOP_V4I32_V4I32_V8I32_I32, AVDst_128, AVSrc_128, AVSrc_256>;
647648

648649
def VOPProfileMAI_F32_V8F16_X32 : VOPProfileMAI<VOP_V4F32_V8F16_V8F16_V4F32, AISrc_128_f32, ADst_128, AVSrc_128>;
649650
def VOPProfileMAI_F32_V8F16_X32_VCD : VOPProfileMAI<VOP_V4F32_V8F16_V8F16_V4F32, VISrc_128_f32, VDst_128, AVSrc_128>;
@@ -1059,6 +1060,7 @@ defm V_SMFMAC_F32_16X16X64_F16 : SMFMACInst<"v_smfmac_f32_16x16x64_f16",
10591060
defm V_SMFMAC_F32_32X32X32_F16 : SMFMACInst<"v_smfmac_f32_32x32x32_f16", "F32_32X32X32_F16", int_amdgcn_smfmac_f32_32x32x32_f16>;
10601061
defm V_SMFMAC_F32_16X16X64_BF16 : SMFMACInst<"v_smfmac_f32_16x16x64_bf16", "F32_16X16X64_BF16", int_amdgcn_smfmac_f32_16x16x64_bf16>;
10611062
defm V_SMFMAC_F32_32X32X32_BF16 : SMFMACInst<"v_smfmac_f32_32x32x32_bf16", "F32_32X32X32_BF16", int_amdgcn_smfmac_f32_32x32x32_bf16>;
1063+
defm V_SMFMAC_I32_16X16X128_I8 : SMFMACInst<"v_smfmac_i32_16x16x128_i8", "I32_16X16X128_I8", int_amdgcn_smfmac_i32_16x16x128_i8>;
10621064
}
10631065

10641066
def MAIInstInfoTable : GenericTable {
@@ -2161,6 +2163,7 @@ defm V_SMFMAC_F32_16X16X64_F16 : VOP3P_Real_SMFMAC <0x5a, "v_smfmac_f32_16x1
21612163
defm V_SMFMAC_F32_32X32X32_F16 : VOP3P_Real_SMFMAC <0x5b, "v_smfmac_f32_32x32x32f16">;
21622164
defm V_SMFMAC_F32_16X16X64_BF16 : VOP3P_Real_SMFMAC <0x39, "v_smfmac_f32_16x16x64bf16">;
21632165
defm V_SMFMAC_F32_32X32X32_BF16 : VOP3P_Real_SMFMAC <0x46, "v_smfmac_f32_32x32x32bf16">;
2166+
defm V_SMFMAC_I32_16X16X128_I8 : VOP3P_Real_SMFMAC <0x3a, "v_smfmac_i32_16x16x128i8">;
21642167

21652168
defm V_PK_FMA_F32 : VOP3P_Real_vi <0x30>;
21662169
defm V_PK_MUL_F32 : VOP3P_Real_vi <0x31>;

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smfmac.gfx950.ll‎

Lines changed: 213 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1430,4 +1430,217 @@ define <16 x float> @test_smfmac_f32_32x32x32_bf16__sgpr(<8 x bfloat> inreg %arg
14301430
ret <16 x float> %result
14311431
}
14321432

1433+
; --------------------------------------------------------------------
1434+
; llvm.amdgcn.smfmac.i32.16x16x128.i8
1435+
; --------------------------------------------------------------------
1436+
1437+
declare <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x128.i8(<4 x i32>, <8 x i32>, <4 x i32>, i32, i32, i32)
1438+
1439+
define amdgpu_kernel void @test_smfmac_i32_16x16x128_i8__vgpr(ptr addrspace(1) %arg, <4 x i32> %a, <8 x i32> %b, i32 %idx) #0 {
1440+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8__vgpr:
1441+
; SDAG: ; %bb.0: ; %bb
1442+
; SDAG-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x24
1443+
; SDAG-NEXT: v_lshlrev_b32_e32 v0, 4, v0
1444+
; SDAG-NEXT: s_load_dwordx8 s[4:11], s[0:1], 0x34
1445+
; SDAG-NEXT: v_mov_b32_e32 v16, 0
1446+
; SDAG-NEXT: s_waitcnt lgkmcnt(0)
1447+
; SDAG-NEXT: global_load_dwordx4 v[8:11], v0, s[2:3]
1448+
; SDAG-NEXT: s_load_dword s16, s[0:1], 0x64
1449+
; SDAG-NEXT: s_load_dwordx4 s[12:15], s[0:1], 0x54
1450+
; SDAG-NEXT: v_mov_b32_e32 v12, s4
1451+
; SDAG-NEXT: v_mov_b32_e32 v13, s5
1452+
; SDAG-NEXT: v_mov_b32_e32 v14, s6
1453+
; SDAG-NEXT: v_mov_b32_e32 v15, s7
1454+
; SDAG-NEXT: v_mov_b32_e32 v0, s8
1455+
; SDAG-NEXT: v_mov_b32_e32 v1, s9
1456+
; SDAG-NEXT: v_mov_b32_e32 v2, s10
1457+
; SDAG-NEXT: v_mov_b32_e32 v3, s11
1458+
; SDAG-NEXT: s_waitcnt lgkmcnt(0)
1459+
; SDAG-NEXT: v_mov_b32_e32 v4, s12
1460+
; SDAG-NEXT: v_mov_b32_e32 v5, s13
1461+
; SDAG-NEXT: v_mov_b32_e32 v6, s14
1462+
; SDAG-NEXT: v_mov_b32_e32 v7, s15
1463+
; SDAG-NEXT: v_mov_b32_e32 v17, s16
1464+
; SDAG-NEXT: s_waitcnt vmcnt(0)
1465+
; SDAG-NEXT: s_nop 0
1466+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 v[8:11], v[12:15], v[0:7], v17 cbsz:1 abid:2
1467+
; SDAG-NEXT: s_nop 6
1468+
; SDAG-NEXT: global_store_dwordx4 v16, v[8:11], s[2:3]
1469+
; SDAG-NEXT: s_endpgm
1470+
;
1471+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8__vgpr:
1472+
; GISEL: ; %bb.0: ; %bb
1473+
; GISEL-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x24
1474+
; GISEL-NEXT: v_lshlrev_b32_e32 v0, 4, v0
1475+
; GISEL-NEXT: s_waitcnt lgkmcnt(0)
1476+
; GISEL-NEXT: global_load_dwordx4 v[8:11], v0, s[2:3]
1477+
; GISEL-NEXT: s_load_dwordx8 s[4:11], s[0:1], 0x34
1478+
; GISEL-NEXT: s_load_dwordx4 s[12:15], s[0:1], 0x54
1479+
; GISEL-NEXT: s_load_dword s16, s[0:1], 0x64
1480+
; GISEL-NEXT: s_waitcnt lgkmcnt(0)
1481+
; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[6:7]
1482+
; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[4:5]
1483+
; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
1484+
; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
1485+
; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
1486+
; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
1487+
; GISEL-NEXT: v_mov_b32_e32 v16, s16
1488+
; GISEL-NEXT: s_waitcnt vmcnt(0)
1489+
; GISEL-NEXT: s_nop 0
1490+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[8:11], v[12:15], v[0:7], v16 cbsz:1 abid:2
1491+
; GISEL-NEXT: v_mov_b32_e32 v0, 0
1492+
; GISEL-NEXT: s_nop 5
1493+
; GISEL-NEXT: global_store_dwordx4 v0, v[8:11], s[2:3]
1494+
; GISEL-NEXT: s_endpgm
1495+
bb:
1496+
%id = call i32 @llvm.amdgcn.workitem.id.x()
1497+
%gep = getelementptr <4 x i32>, ptr addrspace(1) %arg, i32 %id
1498+
%in.1 = load <4 x i32>, ptr addrspace(1) %gep
1499+
%mai.1 = tail call <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x128.i8(<4 x i32> %a, <8 x i32> %b, <4 x i32> %in.1, i32 %idx, i32 1, i32 2)
1500+
store <4 x i32> %mai.1, ptr addrspace(1) %arg
1501+
ret void
1502+
}
1503+
1504+
define <4 x i32> @test_smfmac_i32_16x16x128_i8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3) {
1505+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8:
1506+
; SDAG: ; %bb.0:
1507+
; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1508+
; SDAG-NEXT: v_accvgpr_write_b32 a0, v12
1509+
; SDAG-NEXT: v_accvgpr_write_b32 a1, v13
1510+
; SDAG-NEXT: v_accvgpr_write_b32 a2, v14
1511+
; SDAG-NEXT: v_accvgpr_write_b32 a3, v15
1512+
; SDAG-NEXT: s_nop 1
1513+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 a[0:3], v[0:3], v[4:11], v16
1514+
; SDAG-NEXT: s_nop 6
1515+
; SDAG-NEXT: v_accvgpr_read_b32 v0, a0
1516+
; SDAG-NEXT: v_accvgpr_read_b32 v1, a1
1517+
; SDAG-NEXT: v_accvgpr_read_b32 v2, a2
1518+
; SDAG-NEXT: v_accvgpr_read_b32 v3, a3
1519+
; SDAG-NEXT: s_setpc_b64 s[30:31]
1520+
;
1521+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8:
1522+
; GISEL: ; %bb.0:
1523+
; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1524+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[12:15], v[0:3], v[4:11], v16
1525+
; GISEL-NEXT: s_nop 6
1526+
; GISEL-NEXT: v_mov_b32_e32 v0, v12
1527+
; GISEL-NEXT: v_mov_b32_e32 v1, v13
1528+
; GISEL-NEXT: v_mov_b32_e32 v2, v14
1529+
; GISEL-NEXT: v_mov_b32_e32 v3, v15
1530+
; GISEL-NEXT: s_setpc_b64 s[30:31]
1531+
%result = call <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x128.i8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3, i32 immarg 0, i32 immarg 0)
1532+
ret <4 x i32> %result
1533+
}
1534+
1535+
define <4 x i32> @test_smfmac_i32_16x16x128_i8__flags0(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3) {
1536+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8__flags0:
1537+
; SDAG: ; %bb.0:
1538+
; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1539+
; SDAG-NEXT: v_accvgpr_write_b32 a0, v12
1540+
; SDAG-NEXT: v_accvgpr_write_b32 a1, v13
1541+
; SDAG-NEXT: v_accvgpr_write_b32 a2, v14
1542+
; SDAG-NEXT: v_accvgpr_write_b32 a3, v15
1543+
; SDAG-NEXT: s_nop 1
1544+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 a[0:3], v[0:3], v[4:11], v16 cbsz:1 abid:3
1545+
; SDAG-NEXT: s_nop 6
1546+
; SDAG-NEXT: v_accvgpr_read_b32 v0, a0
1547+
; SDAG-NEXT: v_accvgpr_read_b32 v1, a1
1548+
; SDAG-NEXT: v_accvgpr_read_b32 v2, a2
1549+
; SDAG-NEXT: v_accvgpr_read_b32 v3, a3
1550+
; SDAG-NEXT: s_setpc_b64 s[30:31]
1551+
;
1552+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8__flags0:
1553+
; GISEL: ; %bb.0:
1554+
; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1555+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[12:15], v[0:3], v[4:11], v16 cbsz:1 abid:3
1556+
; GISEL-NEXT: s_nop 6
1557+
; GISEL-NEXT: v_mov_b32_e32 v0, v12
1558+
; GISEL-NEXT: v_mov_b32_e32 v1, v13
1559+
; GISEL-NEXT: v_mov_b32_e32 v2, v14
1560+
; GISEL-NEXT: v_mov_b32_e32 v3, v15
1561+
; GISEL-NEXT: s_setpc_b64 s[30:31]
1562+
%result = call <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x128.i8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3, i32 immarg 1, i32 immarg 3)
1563+
ret <4 x i32> %result
1564+
}
1565+
1566+
define <4 x i32> @test_smfmac_i32_16x16x128_i8__flags1(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3) {
1567+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8__flags1:
1568+
; SDAG: ; %bb.0:
1569+
; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1570+
; SDAG-NEXT: v_accvgpr_write_b32 a0, v12
1571+
; SDAG-NEXT: v_accvgpr_write_b32 a1, v13
1572+
; SDAG-NEXT: v_accvgpr_write_b32 a2, v14
1573+
; SDAG-NEXT: v_accvgpr_write_b32 a3, v15
1574+
; SDAG-NEXT: s_nop 1
1575+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 a[0:3], v[0:3], v[4:11], v16 cbsz:3 abid:1
1576+
; SDAG-NEXT: s_nop 6
1577+
; SDAG-NEXT: v_accvgpr_read_b32 v0, a0
1578+
; SDAG-NEXT: v_accvgpr_read_b32 v1, a1
1579+
; SDAG-NEXT: v_accvgpr_read_b32 v2, a2
1580+
; SDAG-NEXT: v_accvgpr_read_b32 v3, a3
1581+
; SDAG-NEXT: s_setpc_b64 s[30:31]
1582+
;
1583+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8__flags1:
1584+
; GISEL: ; %bb.0:
1585+
; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1586+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[12:15], v[0:3], v[4:11], v16 cbsz:3 abid:1
1587+
; GISEL-NEXT: s_nop 6
1588+
; GISEL-NEXT: v_mov_b32_e32 v0, v12
1589+
; GISEL-NEXT: v_mov_b32_e32 v1, v13
1590+
; GISEL-NEXT: v_mov_b32_e32 v2, v14
1591+
; GISEL-NEXT: v_mov_b32_e32 v3, v15
1592+
; GISEL-NEXT: s_setpc_b64 s[30:31]
1593+
%result = call <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x128.i8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3, i32 immarg 3, i32 immarg 1)
1594+
ret <4 x i32> %result
1595+
}
1596+
1597+
define <4 x i32> @test_smfmac_i32_16x16x128_i8__sgpr(<4 x i32> inreg %arg0, <8 x i32> inreg %arg1, <4 x i32> inreg %arg2, i32 inreg %arg3) {
1598+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8__sgpr:
1599+
; SDAG: ; %bb.0:
1600+
; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1601+
; SDAG-NEXT: v_mov_b32_e32 v8, s0
1602+
; SDAG-NEXT: v_mov_b32_e32 v9, s1
1603+
; SDAG-NEXT: v_mov_b32_e32 v10, s2
1604+
; SDAG-NEXT: v_mov_b32_e32 v11, s3
1605+
; SDAG-NEXT: v_mov_b32_e32 v0, s4
1606+
; SDAG-NEXT: v_mov_b32_e32 v1, s5
1607+
; SDAG-NEXT: v_mov_b32_e32 v2, s6
1608+
; SDAG-NEXT: v_mov_b32_e32 v3, s7
1609+
; SDAG-NEXT: v_mov_b32_e32 v4, s8
1610+
; SDAG-NEXT: v_mov_b32_e32 v5, s9
1611+
; SDAG-NEXT: v_mov_b32_e32 v6, s10
1612+
; SDAG-NEXT: v_mov_b32_e32 v7, s11
1613+
; SDAG-NEXT: v_accvgpr_write_b32 a0, s12
1614+
; SDAG-NEXT: v_accvgpr_write_b32 a1, s13
1615+
; SDAG-NEXT: v_accvgpr_write_b32 a2, s14
1616+
; SDAG-NEXT: v_accvgpr_write_b32 a3, s15
1617+
; SDAG-NEXT: v_mov_b32_e32 v12, s16
1618+
; SDAG-NEXT: s_nop 1
1619+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 a[0:3], v[8:11], v[0:7], v12
1620+
; SDAG-NEXT: s_nop 6
1621+
; SDAG-NEXT: v_accvgpr_read_b32 v0, a0
1622+
; SDAG-NEXT: v_accvgpr_read_b32 v1, a1
1623+
; SDAG-NEXT: v_accvgpr_read_b32 v2, a2
1624+
; SDAG-NEXT: v_accvgpr_read_b32 v3, a3
1625+
; SDAG-NEXT: s_setpc_b64 s[30:31]
1626+
;
1627+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8__sgpr:
1628+
; GISEL: ; %bb.0:
1629+
; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1630+
; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[2:3]
1631+
; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[0:1]
1632+
; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[4:5]
1633+
; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[12:13]
1634+
; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[6:7]
1635+
; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[8:9]
1636+
; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[10:11]
1637+
; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[14:15]
1638+
; GISEL-NEXT: v_mov_b32_e32 v16, s16
1639+
; GISEL-NEXT: s_nop 1
1640+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[0:3], v[12:15], v[4:11], v16
1641+
; GISEL-NEXT: s_setpc_b64 s[30:31]
1642+
%result = call <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x128.i8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3, i32 immarg 0, i32 immarg 0)
1643+
ret <4 x i32> %result
1644+
}
1645+
14331646
attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }

0 commit comments

Comments
 (0)