Skip to content

Commit 8c53036

Browse files
arsenmarsenm
authored
AMDGPU: Add v_smfmac_i32_16x16x128_i8 for gfx950 (#117213)
1 parent eaa0a21 commit 8c53036

File tree

12 files changed

+299
-1
lines changed

12 files changed

+299
-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
@@ -3156,6 +3156,7 @@ def int_amdgcn_smfmac_f32_16x16x64_f16 : AMDGPUMSmfmacIntrinsic<llvm_v4f32_ty, l
31563156
def int_amdgcn_smfmac_f32_32x32x32_f16 : AMDGPUMSmfmacIntrinsic<llvm_v16f32_ty, llvm_v8f16_ty, llvm_v16f16_ty>;
31573157
def int_amdgcn_smfmac_f32_16x16x64_bf16 : AMDGPUMSmfmacIntrinsic<llvm_v4f32_ty, llvm_v8bf16_ty, llvm_v16bf16_ty>;
31583158
def int_amdgcn_smfmac_f32_32x32x32_bf16 : AMDGPUMSmfmacIntrinsic<llvm_v16f32_ty, llvm_v8bf16_ty, llvm_v16bf16_ty>;
3159+
def int_amdgcn_smfmac_i32_16x16x128_i8 : AMDGPUMSmfmacIntrinsic<llvm_v4i32_ty, llvm_v4i32_ty, llvm_v8i32_ty>;
31593160
}
31603161

31613162
//===----------------------------------------------------------------------===//

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1094,6 +1094,7 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC(MachineInstr &I) const {
10941094
case Intrinsic::amdgcn_smfmac_f32_32x32x32_f16:
10951095
case Intrinsic::amdgcn_smfmac_f32_16x16x64_bf16:
10961096
case Intrinsic::amdgcn_smfmac_f32_32x32x32_bf16:
1097+
case Intrinsic::amdgcn_smfmac_i32_16x16x128_i8:
10971098
return selectSMFMACIntrin(I);
10981099
default:
10991100
return selectImpl(I, *CoverageInfo);
@@ -3498,6 +3499,9 @@ bool AMDGPUInstructionSelector::selectSMFMACIntrin(MachineInstr &MI) const {
34983499
case Intrinsic::amdgcn_smfmac_f32_32x32x32_bf16:
34993500
Opc = AMDGPU::V_SMFMAC_F32_32X32X32_BF16_e64;
35003501
break;
3502+
case Intrinsic::amdgcn_smfmac_i32_16x16x128_i8:
3503+
Opc = AMDGPU::V_SMFMAC_I32_16X16X128_I8_e64;
3504+
break;
35013505
default:
35023506
llvm_unreachable("unhandled smfmac intrinsic");
35033507
}

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4808,7 +4808,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
48084808
case Intrinsic::amdgcn_smfmac_f32_16x16x64_f16:
48094809
case Intrinsic::amdgcn_smfmac_f32_32x32x32_f16:
48104810
case Intrinsic::amdgcn_smfmac_f32_16x16x64_bf16:
4811-
case Intrinsic::amdgcn_smfmac_f32_32x32x32_bf16: {
4811+
case Intrinsic::amdgcn_smfmac_f32_32x32x32_bf16:
4812+
case Intrinsic::amdgcn_smfmac_i32_16x16x128_i8: {
48124813
// vdst, srcA, srcB, srcC, idx
48134814
OpdsMapping[0] = getAGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI);
48144815
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
@@ -2880,6 +2880,7 @@ def VOP_V4I32_V2I32_V4I32_I32 : VOPProfile <[v4i32, v2i32, v4i32, i32]>;
28802880
def VOP_V16I32_V2I32_V4I32_I32 : VOPProfile <[v16i32, v2i32, v4i32, i32]>;
28812881
def VOP_V4F32_V2I32_V4I32_I32 : VOPProfile <[v4f32, v2i32, v4i32, i32]>;
28822882
def VOP_V16F32_V2I32_V4I32_I32 : VOPProfile <[v16f32, v2i32, v4i32, i32]>;
2883+
def VOP_V4I32_V4I32_V8I32_I32 : VOPProfile <[v4i32, v4i32, v8i32, i32]>;
28832884

28842885
def VOP_V4F32_V8F16_V8F16_V4F32 : VOPProfile <[v4f32, v8f16, v8f16, v4f32]>;
28852886
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
@@ -640,6 +640,7 @@ def VOPProfileSMFMAC_I32_16X16X64_I8 : VOPProfileSMFMAC<VOP_V4I32_V2I32_V4I32_I
640640
def VOPProfileSMFMAC_I32_32X32X32_I8 : VOPProfileSMFMAC<VOP_V16I32_V2I32_V4I32_I32, AVDst_512, AVSrc_64, AVSrc_128>;
641641
def VOPProfileSMFMAC_F32_16X16X64_F8 : VOPProfileSMFMAC<VOP_V4F32_V2I32_V4I32_I32, AVDst_128, AVSrc_64, AVSrc_128>;
642642
def VOPProfileSMFMAC_F32_32X32X32_F8 : VOPProfileSMFMAC<VOP_V16F32_V2I32_V4I32_I32, AVDst_512, AVSrc_64, AVSrc_128>;
643+
def VOPProfileSMFMAC_I32_16X16X128_I8 : VOPProfileSMFMAC<VOP_V4I32_V4I32_V8I32_I32, AVDst_128, AVSrc_128, AVSrc_256>;
643644

644645
def VOPProfileMAI_F32_V8F16_X32 : VOPProfileMAI<VOP_V4F32_V8F16_V8F16_V4F32, AISrc_128_f32, ADst_128, AVSrc_128>;
645646
def VOPProfileMAI_F32_V8F16_X32_VCD : VOPProfileMAI<VOP_V4F32_V8F16_V8F16_V4F32, VISrc_128_f32, VDst_128, AVSrc_128>;
@@ -1051,6 +1052,7 @@ defm V_SMFMAC_F32_16X16X64_F16 : SMFMACInst<"v_smfmac_f32_16x16x64_f16",
10511052
defm V_SMFMAC_F32_32X32X32_F16 : SMFMACInst<"v_smfmac_f32_32x32x32_f16", "F32_32X32X32_F16", int_amdgcn_smfmac_f32_32x32x32_f16>;
10521053
defm V_SMFMAC_F32_16X16X64_BF16 : SMFMACInst<"v_smfmac_f32_16x16x64_bf16", "F32_16X16X64_BF16", int_amdgcn_smfmac_f32_16x16x64_bf16>;
10531054
defm V_SMFMAC_F32_32X32X32_BF16 : SMFMACInst<"v_smfmac_f32_32x32x32_bf16", "F32_32X32X32_BF16", int_amdgcn_smfmac_f32_32x32x32_bf16>;
1055+
defm V_SMFMAC_I32_16X16X128_I8 : SMFMACInst<"v_smfmac_i32_16x16x128_i8", "I32_16X16X128_I8", int_amdgcn_smfmac_i32_16x16x128_i8>;
10541056
}
10551057

10561058
def MAIInstInfoTable : GenericTable {
@@ -2146,6 +2148,7 @@ defm V_SMFMAC_F32_16X16X64_F16 : VOP3P_Real_SMFMAC <0x5a, "v_smfmac_f32_16x1
21462148
defm V_SMFMAC_F32_32X32X32_F16 : VOP3P_Real_SMFMAC <0x5b, "v_smfmac_f32_32x32x32f16">;
21472149
defm V_SMFMAC_F32_16X16X64_BF16 : VOP3P_Real_SMFMAC <0x39, "v_smfmac_f32_16x16x64bf16">;
21482150
defm V_SMFMAC_F32_32X32X32_BF16 : VOP3P_Real_SMFMAC <0x46, "v_smfmac_f32_32x32x32bf16">;
2151+
defm V_SMFMAC_I32_16X16X128_I8 : VOP3P_Real_SMFMAC <0x3a, "v_smfmac_i32_16x16x128i8">;
21492152

21502153
defm V_PK_FMA_F32 : VOP3P_Real_vi <0x30>;
21512154
defm V_PK_MUL_F32 : VOP3P_Real_vi <0x31>;

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

Lines changed: 215 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1481,6 +1481,221 @@ define <16 x float> @test_smfmac_f32_32x32x32_bf16__sgpr(<8 x bfloat> inreg %arg
14811481
ret <16 x float> %result
14821482
}
14831483

1484+
; --------------------------------------------------------------------
1485+
; llvm.amdgcn.smfmac.i32.16x16x128.i8
1486+
; --------------------------------------------------------------------
1487+
1488+
declare <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x128.i8(<4 x i32>, <8 x i32>, <4 x i32>, i32, i32, i32)
1489+
1490+
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 {
1491+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8__vgpr:
1492+
; SDAG: ; %bb.0: ; %bb
1493+
; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
1494+
; SDAG-NEXT: v_and_b32_e32 v0, 0x3ff, v0
1495+
; SDAG-NEXT: v_lshlrev_b32_e32 v0, 4, v0
1496+
; SDAG-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x34
1497+
; SDAG-NEXT: v_mov_b32_e32 v16, 0
1498+
; SDAG-NEXT: s_waitcnt lgkmcnt(0)
1499+
; SDAG-NEXT: global_load_dwordx4 v[8:11], v0, s[6:7]
1500+
; SDAG-NEXT: s_load_dword s16, s[4:5], 0x64
1501+
; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x54
1502+
; SDAG-NEXT: v_mov_b32_e32 v12, s8
1503+
; SDAG-NEXT: v_mov_b32_e32 v13, s9
1504+
; SDAG-NEXT: v_mov_b32_e32 v14, s10
1505+
; SDAG-NEXT: v_mov_b32_e32 v15, s11
1506+
; SDAG-NEXT: v_mov_b32_e32 v0, s12
1507+
; SDAG-NEXT: v_mov_b32_e32 v1, s13
1508+
; SDAG-NEXT: v_mov_b32_e32 v2, s14
1509+
; SDAG-NEXT: v_mov_b32_e32 v3, s15
1510+
; SDAG-NEXT: s_waitcnt lgkmcnt(0)
1511+
; SDAG-NEXT: v_mov_b32_e32 v4, s0
1512+
; SDAG-NEXT: v_mov_b32_e32 v5, s1
1513+
; SDAG-NEXT: v_mov_b32_e32 v6, s2
1514+
; SDAG-NEXT: v_mov_b32_e32 v7, s3
1515+
; SDAG-NEXT: v_mov_b32_e32 v17, s16
1516+
; SDAG-NEXT: s_waitcnt vmcnt(0)
1517+
; SDAG-NEXT: s_nop 0
1518+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 v[8:11], v[12:15], v[0:7], v17 cbsz:1 abid:2
1519+
; SDAG-NEXT: s_nop 6
1520+
; SDAG-NEXT: global_store_dwordx4 v16, v[8:11], s[6:7]
1521+
; SDAG-NEXT: s_endpgm
1522+
;
1523+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8__vgpr:
1524+
; GISEL: ; %bb.0: ; %bb
1525+
; GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
1526+
; GISEL-NEXT: v_and_b32_e32 v0, 0x3ff, v0
1527+
; GISEL-NEXT: v_lshlrev_b32_e32 v0, 4, v0
1528+
; GISEL-NEXT: s_waitcnt lgkmcnt(0)
1529+
; GISEL-NEXT: global_load_dwordx4 v[8:11], v0, s[0:1]
1530+
; GISEL-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x34
1531+
; GISEL-NEXT: s_load_dwordx4 s[16:19], s[4:5], 0x54
1532+
; GISEL-NEXT: s_load_dword s2, s[4:5], 0x64
1533+
; GISEL-NEXT: s_waitcnt lgkmcnt(0)
1534+
; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[10:11]
1535+
; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[8:9]
1536+
; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[12:13]
1537+
; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[14:15]
1538+
; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[16:17]
1539+
; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[18:19]
1540+
; GISEL-NEXT: v_mov_b32_e32 v16, s2
1541+
; GISEL-NEXT: s_waitcnt vmcnt(0)
1542+
; GISEL-NEXT: s_nop 0
1543+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[8:11], v[12:15], v[0:7], v16 cbsz:1 abid:2
1544+
; GISEL-NEXT: v_mov_b32_e32 v0, 0
1545+
; GISEL-NEXT: s_nop 5
1546+
; GISEL-NEXT: global_store_dwordx4 v0, v[8:11], s[0:1]
1547+
; GISEL-NEXT: s_endpgm
1548+
bb:
1549+
%id = call i32 @llvm.amdgcn.workitem.id.x()
1550+
%gep = getelementptr <4 x i32>, ptr addrspace(1) %arg, i32 %id
1551+
%in.1 = load <4 x i32>, ptr addrspace(1) %gep
1552+
%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)
1553+
store <4 x i32> %mai.1, ptr addrspace(1) %arg
1554+
ret void
1555+
}
1556+
1557+
define <4 x i32> @test_smfmac_i32_16x16x128_i8(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3) {
1558+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8:
1559+
; SDAG: ; %bb.0:
1560+
; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1561+
; SDAG-NEXT: v_accvgpr_write_b32 a0, v12
1562+
; SDAG-NEXT: v_accvgpr_write_b32 a1, v13
1563+
; SDAG-NEXT: v_accvgpr_write_b32 a2, v14
1564+
; SDAG-NEXT: v_accvgpr_write_b32 a3, v15
1565+
; SDAG-NEXT: s_nop 1
1566+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 a[0:3], v[0:3], v[4:11], v16
1567+
; SDAG-NEXT: s_nop 6
1568+
; SDAG-NEXT: v_accvgpr_read_b32 v0, a0
1569+
; SDAG-NEXT: v_accvgpr_read_b32 v1, a1
1570+
; SDAG-NEXT: v_accvgpr_read_b32 v2, a2
1571+
; SDAG-NEXT: v_accvgpr_read_b32 v3, a3
1572+
; SDAG-NEXT: s_setpc_b64 s[30:31]
1573+
;
1574+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8:
1575+
; GISEL: ; %bb.0:
1576+
; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1577+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[12:15], v[0:3], v[4:11], v16
1578+
; GISEL-NEXT: s_nop 6
1579+
; GISEL-NEXT: v_mov_b32_e32 v0, v12
1580+
; GISEL-NEXT: v_mov_b32_e32 v1, v13
1581+
; GISEL-NEXT: v_mov_b32_e32 v2, v14
1582+
; GISEL-NEXT: v_mov_b32_e32 v3, v15
1583+
; GISEL-NEXT: s_setpc_b64 s[30:31]
1584+
%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)
1585+
ret <4 x i32> %result
1586+
}
1587+
1588+
define <4 x i32> @test_smfmac_i32_16x16x128_i8__flags0(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3) {
1589+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8__flags0:
1590+
; SDAG: ; %bb.0:
1591+
; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1592+
; SDAG-NEXT: v_accvgpr_write_b32 a0, v12
1593+
; SDAG-NEXT: v_accvgpr_write_b32 a1, v13
1594+
; SDAG-NEXT: v_accvgpr_write_b32 a2, v14
1595+
; SDAG-NEXT: v_accvgpr_write_b32 a3, v15
1596+
; SDAG-NEXT: s_nop 1
1597+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 a[0:3], v[0:3], v[4:11], v16 cbsz:1 abid:3
1598+
; SDAG-NEXT: s_nop 6
1599+
; SDAG-NEXT: v_accvgpr_read_b32 v0, a0
1600+
; SDAG-NEXT: v_accvgpr_read_b32 v1, a1
1601+
; SDAG-NEXT: v_accvgpr_read_b32 v2, a2
1602+
; SDAG-NEXT: v_accvgpr_read_b32 v3, a3
1603+
; SDAG-NEXT: s_setpc_b64 s[30:31]
1604+
;
1605+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8__flags0:
1606+
; GISEL: ; %bb.0:
1607+
; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1608+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[12:15], v[0:3], v[4:11], v16 cbsz:1 abid:3
1609+
; GISEL-NEXT: s_nop 6
1610+
; GISEL-NEXT: v_mov_b32_e32 v0, v12
1611+
; GISEL-NEXT: v_mov_b32_e32 v1, v13
1612+
; GISEL-NEXT: v_mov_b32_e32 v2, v14
1613+
; GISEL-NEXT: v_mov_b32_e32 v3, v15
1614+
; GISEL-NEXT: s_setpc_b64 s[30:31]
1615+
%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)
1616+
ret <4 x i32> %result
1617+
}
1618+
1619+
define <4 x i32> @test_smfmac_i32_16x16x128_i8__flags1(<4 x i32> %arg0, <8 x i32> %arg1, <4 x i32> %arg2, i32 %arg3) {
1620+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8__flags1:
1621+
; SDAG: ; %bb.0:
1622+
; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1623+
; SDAG-NEXT: v_accvgpr_write_b32 a0, v12
1624+
; SDAG-NEXT: v_accvgpr_write_b32 a1, v13
1625+
; SDAG-NEXT: v_accvgpr_write_b32 a2, v14
1626+
; SDAG-NEXT: v_accvgpr_write_b32 a3, v15
1627+
; SDAG-NEXT: s_nop 1
1628+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 a[0:3], v[0:3], v[4:11], v16 cbsz:3 abid:1
1629+
; SDAG-NEXT: s_nop 6
1630+
; SDAG-NEXT: v_accvgpr_read_b32 v0, a0
1631+
; SDAG-NEXT: v_accvgpr_read_b32 v1, a1
1632+
; SDAG-NEXT: v_accvgpr_read_b32 v2, a2
1633+
; SDAG-NEXT: v_accvgpr_read_b32 v3, a3
1634+
; SDAG-NEXT: s_setpc_b64 s[30:31]
1635+
;
1636+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8__flags1:
1637+
; GISEL: ; %bb.0:
1638+
; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1639+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[12:15], v[0:3], v[4:11], v16 cbsz:3 abid:1
1640+
; GISEL-NEXT: s_nop 6
1641+
; GISEL-NEXT: v_mov_b32_e32 v0, v12
1642+
; GISEL-NEXT: v_mov_b32_e32 v1, v13
1643+
; GISEL-NEXT: v_mov_b32_e32 v2, v14
1644+
; GISEL-NEXT: v_mov_b32_e32 v3, v15
1645+
; GISEL-NEXT: s_setpc_b64 s[30:31]
1646+
%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)
1647+
ret <4 x i32> %result
1648+
}
1649+
1650+
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) {
1651+
; SDAG-LABEL: test_smfmac_i32_16x16x128_i8__sgpr:
1652+
; SDAG: ; %bb.0:
1653+
; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1654+
; SDAG-NEXT: v_mov_b32_e32 v8, s0
1655+
; SDAG-NEXT: v_mov_b32_e32 v9, s1
1656+
; SDAG-NEXT: v_mov_b32_e32 v10, s2
1657+
; SDAG-NEXT: v_mov_b32_e32 v11, s3
1658+
; SDAG-NEXT: v_mov_b32_e32 v0, s16
1659+
; SDAG-NEXT: v_mov_b32_e32 v1, s17
1660+
; SDAG-NEXT: v_mov_b32_e32 v2, s18
1661+
; SDAG-NEXT: v_mov_b32_e32 v3, s19
1662+
; SDAG-NEXT: v_mov_b32_e32 v4, s20
1663+
; SDAG-NEXT: v_mov_b32_e32 v5, s21
1664+
; SDAG-NEXT: v_mov_b32_e32 v6, s22
1665+
; SDAG-NEXT: v_mov_b32_e32 v7, s23
1666+
; SDAG-NEXT: v_accvgpr_write_b32 a0, s24
1667+
; SDAG-NEXT: v_accvgpr_write_b32 a1, s25
1668+
; SDAG-NEXT: v_accvgpr_write_b32 a2, s26
1669+
; SDAG-NEXT: v_accvgpr_write_b32 a3, s27
1670+
; SDAG-NEXT: v_mov_b32_e32 v12, s28
1671+
; SDAG-NEXT: s_nop 1
1672+
; SDAG-NEXT: v_smfmac_i32_16x16x128_i8 a[0:3], v[8:11], v[0:7], v12
1673+
; SDAG-NEXT: s_nop 6
1674+
; SDAG-NEXT: v_accvgpr_read_b32 v0, a0
1675+
; SDAG-NEXT: v_accvgpr_read_b32 v1, a1
1676+
; SDAG-NEXT: v_accvgpr_read_b32 v2, a2
1677+
; SDAG-NEXT: v_accvgpr_read_b32 v3, a3
1678+
; SDAG-NEXT: s_setpc_b64 s[30:31]
1679+
;
1680+
; GISEL-LABEL: test_smfmac_i32_16x16x128_i8__sgpr:
1681+
; GISEL: ; %bb.0:
1682+
; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1683+
; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[2:3]
1684+
; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[0:1]
1685+
; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[16:17]
1686+
; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[24:25]
1687+
; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[18:19]
1688+
; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[20:21]
1689+
; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[22:23]
1690+
; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[26:27]
1691+
; GISEL-NEXT: v_mov_b32_e32 v16, s28
1692+
; GISEL-NEXT: s_nop 1
1693+
; GISEL-NEXT: v_smfmac_i32_16x16x128_i8 v[0:3], v[12:15], v[4:11], v16
1694+
; GISEL-NEXT: s_setpc_b64 s[30:31]
1695+
%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)
1696+
ret <4 x i32> %result
1697+
}
1698+
14841699
attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
14851700
;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
14861701
; GCN: {{.*}}

0 commit comments

Comments
 (0)