2554 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -global-isel=0 < %s | FileCheck -check-prefixes=GCN,SDAG %s3; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -global-isel=1 < %s | FileCheck -check-prefixes=GCN,GISEL %s4 5; 0 = fp86; 1 = bf87; 2 = fp68; 3 = bf69; 4 = fp410 11; --------------------------------------------------------------------12; Different format signatures13; --------------------------------------------------------------------14 15; fp8 x fp816define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp0(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {17; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp0:18; GCN: ; %bb.0:19; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)20; GCN-NEXT: v_accvgpr_write_b32 a0, v1621; GCN-NEXT: v_accvgpr_write_b32 a1, v1722; GCN-NEXT: v_accvgpr_write_b32 a2, v1823; GCN-NEXT: v_accvgpr_write_b32 a3, v1924; GCN-NEXT: s_nop 125; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0]26; GCN-NEXT: s_nop 1127; GCN-NEXT: v_accvgpr_read_b32 v0, a028; GCN-NEXT: v_accvgpr_read_b32 v1, a129; GCN-NEXT: v_accvgpr_read_b32 v2, a230; GCN-NEXT: v_accvgpr_read_b32 v3, a331; GCN-NEXT: s_setpc_b64 s[30:31]32 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,33 i32 0, ; cbsz34 i32 0, ; blgp35 i32 0, i32 %scale0, i32 0, i32 %scale1)36 ret <4 x float> %result37}38 39define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_1_1__cbsz1__blgp1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {40; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_1_1__cbsz1__blgp1:41; GCN: ; %bb.0:42; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)43; GCN-NEXT: v_accvgpr_write_b32 a0, v1644; GCN-NEXT: v_accvgpr_write_b32 a1, v1745; GCN-NEXT: v_accvgpr_write_b32 a2, v1846; GCN-NEXT: v_accvgpr_write_b32 a3, v1947; GCN-NEXT: s_nop 148; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel:[1,1,0] op_sel_hi:[0,0,0]49; GCN-NEXT: s_nop 1150; GCN-NEXT: v_accvgpr_read_b32 v0, a051; GCN-NEXT: v_accvgpr_read_b32 v1, a152; GCN-NEXT: v_accvgpr_read_b32 v2, a253; GCN-NEXT: v_accvgpr_read_b32 v3, a354; GCN-NEXT: s_setpc_b64 s[30:31]55 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,56 i32 0, ; cbsz57 i32 0, ; blgp58 i32 1, i32 %scale0, i32 1, i32 %scale1)59 ret <4 x float> %result60}61 62define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_2_2__cbsz1__blgp1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {63; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_2_2__cbsz1__blgp1:64; GCN: ; %bb.0:65; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)66; GCN-NEXT: v_accvgpr_write_b32 a0, v1667; GCN-NEXT: v_accvgpr_write_b32 a1, v1768; GCN-NEXT: v_accvgpr_write_b32 a2, v1869; GCN-NEXT: v_accvgpr_write_b32 a3, v1970; GCN-NEXT: s_nop 171; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[1,1,0]72; GCN-NEXT: s_nop 1173; GCN-NEXT: v_accvgpr_read_b32 v0, a074; GCN-NEXT: v_accvgpr_read_b32 v1, a175; GCN-NEXT: v_accvgpr_read_b32 v2, a276; GCN-NEXT: v_accvgpr_read_b32 v3, a377; GCN-NEXT: s_setpc_b64 s[30:31]78 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,79 i32 0, ; cbsz80 i32 0, ; blgp81 i32 2, i32 %scale0, i32 2, i32 %scale1)82 ret <4 x float> %result83}84 85define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_3_3__cbsz1__blgp1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {86; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_3_3__cbsz1__blgp1:87; GCN: ; %bb.0:88; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)89; GCN-NEXT: v_accvgpr_write_b32 a0, v1690; GCN-NEXT: v_accvgpr_write_b32 a1, v1791; GCN-NEXT: v_accvgpr_write_b32 a2, v1892; GCN-NEXT: v_accvgpr_write_b32 a3, v1993; GCN-NEXT: s_nop 194; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel:[1,1,0] op_sel_hi:[1,1,0]95; GCN-NEXT: s_nop 1196; GCN-NEXT: v_accvgpr_read_b32 v0, a097; GCN-NEXT: v_accvgpr_read_b32 v1, a198; GCN-NEXT: v_accvgpr_read_b32 v2, a299; GCN-NEXT: v_accvgpr_read_b32 v3, a3100; GCN-NEXT: s_setpc_b64 s[30:31]101 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,102 i32 0, ; cbsz103 i32 0, ; blgp104 i32 3, i32 %scale0, i32 3, i32 %scale1)105 ret <4 x float> %result106}107 108define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_3__cbsz1__blgp1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {109; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_3__cbsz1__blgp1:110; GCN: ; %bb.0:111; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)112; GCN-NEXT: v_accvgpr_write_b32 a0, v16113; GCN-NEXT: v_accvgpr_write_b32 a1, v17114; GCN-NEXT: v_accvgpr_write_b32 a2, v18115; GCN-NEXT: v_accvgpr_write_b32 a3, v19116; GCN-NEXT: s_nop 1117; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel:[0,1,0] op_sel_hi:[0,1,0]118; GCN-NEXT: s_nop 11119; GCN-NEXT: v_accvgpr_read_b32 v0, a0120; GCN-NEXT: v_accvgpr_read_b32 v1, a1121; GCN-NEXT: v_accvgpr_read_b32 v2, a2122; GCN-NEXT: v_accvgpr_read_b32 v3, a3123; GCN-NEXT: s_setpc_b64 s[30:31]124 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,125 i32 0, ; cbsz126 i32 0, ; blgp127 i32 0, i32 %scale0, i32 3, i32 %scale1)128 ret <4 x float> %result129}130 131define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_3_0__cbsz1__blgp1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {132; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_3_0__cbsz1__blgp1:133; GCN: ; %bb.0:134; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)135; GCN-NEXT: v_accvgpr_write_b32 a0, v16136; GCN-NEXT: v_accvgpr_write_b32 a1, v17137; GCN-NEXT: v_accvgpr_write_b32 a2, v18138; GCN-NEXT: v_accvgpr_write_b32 a3, v19139; GCN-NEXT: s_nop 1140; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel:[1,0,0] op_sel_hi:[1,0,0]141; GCN-NEXT: s_nop 11142; GCN-NEXT: v_accvgpr_read_b32 v0, a0143; GCN-NEXT: v_accvgpr_read_b32 v1, a1144; GCN-NEXT: v_accvgpr_read_b32 v2, a2145; GCN-NEXT: v_accvgpr_read_b32 v3, a3146; GCN-NEXT: s_setpc_b64 s[30:31]147 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,148 i32 0, ; cbsz149 i32 0, ; blgp150 i32 3, i32 %scale0, i32 0, i32 %scale1)151 ret <4 x float> %result152}153 154define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_2_3__cbsz1__blgp1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {155; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_2_3__cbsz1__blgp1:156; GCN: ; %bb.0:157; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)158; GCN-NEXT: v_accvgpr_write_b32 a0, v16159; GCN-NEXT: v_accvgpr_write_b32 a1, v17160; GCN-NEXT: v_accvgpr_write_b32 a2, v18161; GCN-NEXT: v_accvgpr_write_b32 a3, v19162; GCN-NEXT: s_nop 1163; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel:[0,1,0] op_sel_hi:[1,1,0]164; GCN-NEXT: s_nop 11165; GCN-NEXT: v_accvgpr_read_b32 v0, a0166; GCN-NEXT: v_accvgpr_read_b32 v1, a1167; GCN-NEXT: v_accvgpr_read_b32 v2, a2168; GCN-NEXT: v_accvgpr_read_b32 v3, a3169; GCN-NEXT: s_setpc_b64 s[30:31]170 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,171 i32 0, ; cbsz172 i32 0, ; blgp173 i32 2, i32 %scale0, i32 3, i32 %scale1)174 ret <4 x float> %result175}176 177define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_3_2__cbsz1__blgp1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {178; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_3_2__cbsz1__blgp1:179; GCN: ; %bb.0:180; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)181; GCN-NEXT: v_accvgpr_write_b32 a0, v16182; GCN-NEXT: v_accvgpr_write_b32 a1, v17183; GCN-NEXT: v_accvgpr_write_b32 a2, v18184; GCN-NEXT: v_accvgpr_write_b32 a3, v19185; GCN-NEXT: s_nop 1186; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel:[1,0,0] op_sel_hi:[1,1,0]187; GCN-NEXT: s_nop 11188; GCN-NEXT: v_accvgpr_read_b32 v0, a0189; GCN-NEXT: v_accvgpr_read_b32 v1, a1190; GCN-NEXT: v_accvgpr_read_b32 v2, a2191; GCN-NEXT: v_accvgpr_read_b32 v3, a3192; GCN-NEXT: s_setpc_b64 s[30:31]193 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,194 i32 0, ; cbsz195 i32 0, ; blgp196 i32 3, i32 %scale0, i32 2, i32 %scale1)197 ret <4 x float> %result198}199 200; This should be optimized to avoid the scale201define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp0__constant_scale_0_0(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {202; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp0__constant_scale_0_0:203; GCN: ; %bb.0:204; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)205; GCN-NEXT: v_accvgpr_write_b32 a0, v16206; GCN-NEXT: v_accvgpr_write_b32 a1, v17207; GCN-NEXT: v_accvgpr_write_b32 a2, v18208; GCN-NEXT: v_accvgpr_write_b32 a3, v19209; GCN-NEXT: s_nop 1210; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3]211; GCN-NEXT: s_nop 11212; GCN-NEXT: v_accvgpr_read_b32 v0, a0213; GCN-NEXT: v_accvgpr_read_b32 v1, a1214; GCN-NEXT: v_accvgpr_read_b32 v2, a2215; GCN-NEXT: v_accvgpr_read_b32 v3, a3216; GCN-NEXT: s_setpc_b64 s[30:31]217 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,218 i32 0, ; cbsz219 i32 0, ; blgp220 i32 0, i32 0, i32 0, i32 0)221 ret <4 x float> %result222}223 224; fp8 x bf8225define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {226; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp1:227; GCN: ; %bb.0:228; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)229; GCN-NEXT: v_accvgpr_write_b32 a0, v16230; GCN-NEXT: v_accvgpr_write_b32 a1, v17231; GCN-NEXT: v_accvgpr_write_b32 a2, v18232; GCN-NEXT: v_accvgpr_write_b32 a3, v19233; GCN-NEXT: s_nop 1234; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] blgp:1235; GCN-NEXT: s_nop 11236; GCN-NEXT: v_accvgpr_read_b32 v0, a0237; GCN-NEXT: v_accvgpr_read_b32 v1, a1238; GCN-NEXT: v_accvgpr_read_b32 v2, a2239; GCN-NEXT: v_accvgpr_read_b32 v3, a3240; GCN-NEXT: s_setpc_b64 s[30:31]241 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,242 i32 0, ; cbsz243 i32 1, ; blgp244 i32 0, i32 %scale0, i32 0, i32 %scale1)245 ret <4 x float> %result246}247 248; This should be optimized to avoid the scale249define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp1__constant_scale_0_0(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {250; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp1__constant_scale_0_0:251; GCN: ; %bb.0:252; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)253; GCN-NEXT: v_accvgpr_write_b32 a0, v16254; GCN-NEXT: v_accvgpr_write_b32 a1, v17255; GCN-NEXT: v_accvgpr_write_b32 a2, v18256; GCN-NEXT: v_accvgpr_write_b32 a3, v19257; GCN-NEXT: s_nop 1258; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3] blgp:1259; GCN-NEXT: s_nop 11260; GCN-NEXT: v_accvgpr_read_b32 v0, a0261; GCN-NEXT: v_accvgpr_read_b32 v1, a1262; GCN-NEXT: v_accvgpr_read_b32 v2, a2263; GCN-NEXT: v_accvgpr_read_b32 v3, a3264; GCN-NEXT: s_setpc_b64 s[30:31]265 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,266 i32 0, ; cbsz267 i32 1, ; blgp268 i32 0, i32 0, i32 0, i32 0)269 ret <4 x float> %result270}271 272; fp8 x fp6273define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp2(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {274; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp2:275; GCN: ; %bb.0:276; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)277; GCN-NEXT: v_accvgpr_write_b32 a0, v14278; GCN-NEXT: v_accvgpr_write_b32 a1, v15279; GCN-NEXT: v_accvgpr_write_b32 a2, v16280; GCN-NEXT: v_accvgpr_write_b32 a3, v17281; GCN-NEXT: s_nop 1282; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] blgp:2283; GCN-NEXT: s_nop 11284; GCN-NEXT: v_accvgpr_read_b32 v0, a0285; GCN-NEXT: v_accvgpr_read_b32 v1, a1286; GCN-NEXT: v_accvgpr_read_b32 v2, a2287; GCN-NEXT: v_accvgpr_read_b32 v3, a3288; GCN-NEXT: s_setpc_b64 s[30:31]289 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,290 i32 0, ; cbsz291 i32 2, ; blgp292 i32 0, i32 %scale0, i32 0, i32 %scale1)293 ret <4 x float> %result294}295 296; This should be optimized to avoid the scale297define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp2__constant_scale_0_0(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {298; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp2__constant_scale_0_0:299; GCN: ; %bb.0:300; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)301; GCN-NEXT: v_accvgpr_write_b32 a0, v14302; GCN-NEXT: v_accvgpr_write_b32 a1, v15303; GCN-NEXT: v_accvgpr_write_b32 a2, v16304; GCN-NEXT: v_accvgpr_write_b32 a3, v17305; GCN-NEXT: s_nop 1306; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:13], a[0:3] blgp:2307; GCN-NEXT: s_nop 11308; GCN-NEXT: v_accvgpr_read_b32 v0, a0309; GCN-NEXT: v_accvgpr_read_b32 v1, a1310; GCN-NEXT: v_accvgpr_read_b32 v2, a2311; GCN-NEXT: v_accvgpr_read_b32 v3, a3312; GCN-NEXT: s_setpc_b64 s[30:31]313 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,314 i32 0, ; cbsz315 i32 2, ; blgp316 i32 0, i32 0, i32 0, i32 0)317 ret <4 x float> %result318}319 320; fp8 x bf6321define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp3(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {322; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp3:323; GCN: ; %bb.0:324; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)325; GCN-NEXT: v_accvgpr_write_b32 a0, v14326; GCN-NEXT: v_accvgpr_write_b32 a1, v15327; GCN-NEXT: v_accvgpr_write_b32 a2, v16328; GCN-NEXT: v_accvgpr_write_b32 a3, v17329; GCN-NEXT: s_nop 1330; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] blgp:3331; GCN-NEXT: s_nop 11332; GCN-NEXT: v_accvgpr_read_b32 v0, a0333; GCN-NEXT: v_accvgpr_read_b32 v1, a1334; GCN-NEXT: v_accvgpr_read_b32 v2, a2335; GCN-NEXT: v_accvgpr_read_b32 v3, a3336; GCN-NEXT: s_setpc_b64 s[30:31]337 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,338 i32 0, ; cbsz339 i32 3, ; blgp340 i32 0, i32 %scale0, i32 0, i32 %scale1)341 ret <4 x float> %result342}343 344; This should be optimized to avoid the scale345define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp3__constant_scale_0_0(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {346; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp3__constant_scale_0_0:347; GCN: ; %bb.0:348; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)349; GCN-NEXT: v_accvgpr_write_b32 a0, v14350; GCN-NEXT: v_accvgpr_write_b32 a1, v15351; GCN-NEXT: v_accvgpr_write_b32 a2, v16352; GCN-NEXT: v_accvgpr_write_b32 a3, v17353; GCN-NEXT: s_nop 1354; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:13], a[0:3] blgp:3355; GCN-NEXT: s_nop 11356; GCN-NEXT: v_accvgpr_read_b32 v0, a0357; GCN-NEXT: v_accvgpr_read_b32 v1, a1358; GCN-NEXT: v_accvgpr_read_b32 v2, a2359; GCN-NEXT: v_accvgpr_read_b32 v3, a3360; GCN-NEXT: s_setpc_b64 s[30:31]361 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,362 i32 0, ; cbsz363 i32 3, ; blgp364 i32 0, i32 0, i32 0, i32 0)365 ret <4 x float> %result366}367 368; fp8 x fp4369define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp4(<8 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {370; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp4:371; GCN: ; %bb.0:372; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)373; GCN-NEXT: v_accvgpr_write_b32 a0, v12374; GCN-NEXT: v_accvgpr_write_b32 a1, v13375; GCN-NEXT: v_accvgpr_write_b32 a2, v14376; GCN-NEXT: v_accvgpr_write_b32 a3, v15377; GCN-NEXT: s_nop 1378; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:11], a[0:3], v16, v17 op_sel_hi:[0,0,0] blgp:4379; GCN-NEXT: s_nop 11380; GCN-NEXT: v_accvgpr_read_b32 v0, a0381; GCN-NEXT: v_accvgpr_read_b32 v1, a1382; GCN-NEXT: v_accvgpr_read_b32 v2, a2383; GCN-NEXT: v_accvgpr_read_b32 v3, a3384; GCN-NEXT: s_setpc_b64 s[30:31]385 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v4i32(<8 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,386 i32 0, ; cbsz387 i32 4, ; blgp388 i32 0, i32 %scale0, i32 0, i32 %scale1)389 ret <4 x float> %result390}391 392; This should be optimized to avoid the scale393define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp4__constant_scale_0_0(<8 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2) {394; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz0__blgp4__constant_scale_0_0:395; GCN: ; %bb.0:396; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)397; GCN-NEXT: v_accvgpr_write_b32 a0, v12398; GCN-NEXT: v_accvgpr_write_b32 a1, v13399; GCN-NEXT: v_accvgpr_write_b32 a2, v14400; GCN-NEXT: v_accvgpr_write_b32 a3, v15401; GCN-NEXT: s_nop 1402; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:11], a[0:3] blgp:4403; GCN-NEXT: s_nop 11404; GCN-NEXT: v_accvgpr_read_b32 v0, a0405; GCN-NEXT: v_accvgpr_read_b32 v1, a1406; GCN-NEXT: v_accvgpr_read_b32 v2, a2407; GCN-NEXT: v_accvgpr_read_b32 v3, a3408; GCN-NEXT: s_setpc_b64 s[30:31]409 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v4i32(<8 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,410 i32 0, ; cbsz411 i32 4, ; blgp412 i32 0, i32 0, i32 0, i32 0)413 ret <4 x float> %result414}415 416; bf8 x fp8417define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp0(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {418; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp0:419; GCN: ; %bb.0:420; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)421; GCN-NEXT: v_accvgpr_write_b32 a0, v16422; GCN-NEXT: v_accvgpr_write_b32 a1, v17423; GCN-NEXT: v_accvgpr_write_b32 a2, v18424; GCN-NEXT: v_accvgpr_write_b32 a3, v19425; GCN-NEXT: s_nop 1426; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] cbsz:1427; GCN-NEXT: s_nop 11428; GCN-NEXT: v_accvgpr_read_b32 v0, a0429; GCN-NEXT: v_accvgpr_read_b32 v1, a1430; GCN-NEXT: v_accvgpr_read_b32 v2, a2431; GCN-NEXT: v_accvgpr_read_b32 v3, a3432; GCN-NEXT: s_setpc_b64 s[30:31]433 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,434 i32 1, ; cbsz435 i32 0, ; blgp436 i32 0, i32 %scale0, i32 0, i32 %scale1)437 ret <4 x float> %result438}439 440; This should be optimized to avoid the scale441define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp0__constant_scale_0_0(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {442; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp0__constant_scale_0_0:443; GCN: ; %bb.0:444; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)445; GCN-NEXT: v_accvgpr_write_b32 a0, v16446; GCN-NEXT: v_accvgpr_write_b32 a1, v17447; GCN-NEXT: v_accvgpr_write_b32 a2, v18448; GCN-NEXT: v_accvgpr_write_b32 a3, v19449; GCN-NEXT: s_nop 1450; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3] cbsz:1451; GCN-NEXT: s_nop 11452; GCN-NEXT: v_accvgpr_read_b32 v0, a0453; GCN-NEXT: v_accvgpr_read_b32 v1, a1454; GCN-NEXT: v_accvgpr_read_b32 v2, a2455; GCN-NEXT: v_accvgpr_read_b32 v3, a3456; GCN-NEXT: s_setpc_b64 s[30:31]457 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,458 i32 1, ; cbsz459 i32 0, ; blgp460 i32 0, i32 0, i32 0, i32 0)461 ret <4 x float> %result462}463 464; bf8 x bf8465define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {466; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp1:467; GCN: ; %bb.0:468; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)469; GCN-NEXT: v_accvgpr_write_b32 a0, v16470; GCN-NEXT: v_accvgpr_write_b32 a1, v17471; GCN-NEXT: v_accvgpr_write_b32 a2, v18472; GCN-NEXT: v_accvgpr_write_b32 a3, v19473; GCN-NEXT: s_nop 1474; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] cbsz:1 blgp:1475; GCN-NEXT: s_nop 11476; GCN-NEXT: v_accvgpr_read_b32 v0, a0477; GCN-NEXT: v_accvgpr_read_b32 v1, a1478; GCN-NEXT: v_accvgpr_read_b32 v2, a2479; GCN-NEXT: v_accvgpr_read_b32 v3, a3480; GCN-NEXT: s_setpc_b64 s[30:31]481 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,482 i32 1, ; cbsz483 i32 1, ; blgp484 i32 0, i32 %scale0, i32 0, i32 %scale1)485 ret <4 x float> %result486}487 488 489; This should be optimized to avoid the scale490define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp1__constant_scale_0_0(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {491; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp1__constant_scale_0_0:492; GCN: ; %bb.0:493; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)494; GCN-NEXT: v_accvgpr_write_b32 a0, v16495; GCN-NEXT: v_accvgpr_write_b32 a1, v17496; GCN-NEXT: v_accvgpr_write_b32 a2, v18497; GCN-NEXT: v_accvgpr_write_b32 a3, v19498; GCN-NEXT: s_nop 1499; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3] cbsz:1 blgp:1500; GCN-NEXT: s_nop 11501; GCN-NEXT: v_accvgpr_read_b32 v0, a0502; GCN-NEXT: v_accvgpr_read_b32 v1, a1503; GCN-NEXT: v_accvgpr_read_b32 v2, a2504; GCN-NEXT: v_accvgpr_read_b32 v3, a3505; GCN-NEXT: s_setpc_b64 s[30:31]506 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,507 i32 1, ; cbsz508 i32 1, ; blgp509 i32 0, i32 0, i32 0, i32 0)510 ret <4 x float> %result511}512 513; bf8 x fp6514define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp2(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {515; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp2:516; GCN: ; %bb.0:517; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)518; GCN-NEXT: v_accvgpr_write_b32 a0, v14519; GCN-NEXT: v_accvgpr_write_b32 a1, v15520; GCN-NEXT: v_accvgpr_write_b32 a2, v16521; GCN-NEXT: v_accvgpr_write_b32 a3, v17522; GCN-NEXT: s_nop 1523; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] cbsz:1 blgp:2524; GCN-NEXT: s_nop 11525; GCN-NEXT: v_accvgpr_read_b32 v0, a0526; GCN-NEXT: v_accvgpr_read_b32 v1, a1527; GCN-NEXT: v_accvgpr_read_b32 v2, a2528; GCN-NEXT: v_accvgpr_read_b32 v3, a3529; GCN-NEXT: s_setpc_b64 s[30:31]530 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,531 i32 1, ; cbsz532 i32 2, ; blgp533 i32 0, i32 %scale0, i32 0, i32 %scale1)534 ret <4 x float> %result535}536 537define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp2__constant_scale_0(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {538; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp2__constant_scale_0:539; GCN: ; %bb.0:540; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)541; GCN-NEXT: v_accvgpr_write_b32 a0, v14542; GCN-NEXT: v_accvgpr_write_b32 a1, v15543; GCN-NEXT: v_accvgpr_write_b32 a2, v16544; GCN-NEXT: v_accvgpr_write_b32 a3, v17545; GCN-NEXT: s_nop 1546; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:13], a[0:3] cbsz:1 blgp:2547; GCN-NEXT: s_nop 11548; GCN-NEXT: v_accvgpr_read_b32 v0, a0549; GCN-NEXT: v_accvgpr_read_b32 v1, a1550; GCN-NEXT: v_accvgpr_read_b32 v2, a2551; GCN-NEXT: v_accvgpr_read_b32 v3, a3552; GCN-NEXT: s_setpc_b64 s[30:31]553 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,554 i32 1, ; cbsz555 i32 2, ; blgp556 i32 0, i32 0, i32 0, i32 0)557 ret <4 x float> %result558}559 560; bf8 x bf6561define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp3(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {562; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp3:563; GCN: ; %bb.0:564; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)565; GCN-NEXT: v_accvgpr_write_b32 a0, v14566; GCN-NEXT: v_accvgpr_write_b32 a1, v15567; GCN-NEXT: v_accvgpr_write_b32 a2, v16568; GCN-NEXT: v_accvgpr_write_b32 a3, v17569; GCN-NEXT: s_nop 1570; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] cbsz:1 blgp:3571; GCN-NEXT: s_nop 11572; GCN-NEXT: v_accvgpr_read_b32 v0, a0573; GCN-NEXT: v_accvgpr_read_b32 v1, a1574; GCN-NEXT: v_accvgpr_read_b32 v2, a2575; GCN-NEXT: v_accvgpr_read_b32 v3, a3576; GCN-NEXT: s_setpc_b64 s[30:31]577 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,578 i32 1, ; cbsz579 i32 3, ; blgp580 i32 0, i32 %scale0, i32 0, i32 %scale1)581 ret <4 x float> %result582}583 584; This should be optimized to avoid the scale585define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp3__constant_scale_0_0(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {586; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp3__constant_scale_0_0:587; GCN: ; %bb.0:588; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)589; GCN-NEXT: v_accvgpr_write_b32 a0, v14590; GCN-NEXT: v_accvgpr_write_b32 a1, v15591; GCN-NEXT: v_accvgpr_write_b32 a2, v16592; GCN-NEXT: v_accvgpr_write_b32 a3, v17593; GCN-NEXT: s_nop 1594; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:13], a[0:3] cbsz:1 blgp:3595; GCN-NEXT: s_nop 11596; GCN-NEXT: v_accvgpr_read_b32 v0, a0597; GCN-NEXT: v_accvgpr_read_b32 v1, a1598; GCN-NEXT: v_accvgpr_read_b32 v2, a2599; GCN-NEXT: v_accvgpr_read_b32 v3, a3600; GCN-NEXT: s_setpc_b64 s[30:31]601 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,602 i32 1, ; cbsz603 i32 3, ; blgp604 i32 0, i32 0, i32 0, i32 0)605 ret <4 x float> %result606}607 608; bf8 x fp4609define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp4(<8 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {610; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp4:611; GCN: ; %bb.0:612; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)613; GCN-NEXT: v_accvgpr_write_b32 a0, v12614; GCN-NEXT: v_accvgpr_write_b32 a1, v13615; GCN-NEXT: v_accvgpr_write_b32 a2, v14616; GCN-NEXT: v_accvgpr_write_b32 a3, v15617; GCN-NEXT: s_nop 1618; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:11], a[0:3], v16, v17 op_sel_hi:[0,0,0] cbsz:1 blgp:4619; GCN-NEXT: s_nop 11620; GCN-NEXT: v_accvgpr_read_b32 v0, a0621; GCN-NEXT: v_accvgpr_read_b32 v1, a1622; GCN-NEXT: v_accvgpr_read_b32 v2, a2623; GCN-NEXT: v_accvgpr_read_b32 v3, a3624; GCN-NEXT: s_setpc_b64 s[30:31]625 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v4i32(<8 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,626 i32 1, ; cbsz627 i32 4, ; blgp628 i32 0, i32 %scale0, i32 0, i32 %scale1)629 ret <4 x float> %result630}631 632; This should be optimized to avoid the scale633define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp4__constant_scale_0_0(<8 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2) {634; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz1__blgp4__constant_scale_0_0:635; GCN: ; %bb.0:636; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)637; GCN-NEXT: v_accvgpr_write_b32 a0, v12638; GCN-NEXT: v_accvgpr_write_b32 a1, v13639; GCN-NEXT: v_accvgpr_write_b32 a2, v14640; GCN-NEXT: v_accvgpr_write_b32 a3, v15641; GCN-NEXT: s_nop 1642; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:11], a[0:3] cbsz:1 blgp:4643; GCN-NEXT: s_nop 11644; GCN-NEXT: v_accvgpr_read_b32 v0, a0645; GCN-NEXT: v_accvgpr_read_b32 v1, a1646; GCN-NEXT: v_accvgpr_read_b32 v2, a2647; GCN-NEXT: v_accvgpr_read_b32 v3, a3648; GCN-NEXT: s_setpc_b64 s[30:31]649 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v4i32(<8 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,650 i32 1, ; cbsz651 i32 4, ; blgp652 i32 0, i32 0, i32 0, i32 0)653 ret <4 x float> %result654}655 656; fp6 x fp8657define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp0(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {658; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp0:659; GCN: ; %bb.0:660; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)661; GCN-NEXT: v_accvgpr_write_b32 a0, v14662; GCN-NEXT: v_accvgpr_write_b32 a1, v15663; GCN-NEXT: v_accvgpr_write_b32 a2, v16664; GCN-NEXT: v_accvgpr_write_b32 a3, v17665; GCN-NEXT: s_nop 1666; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] cbsz:2667; GCN-NEXT: s_nop 11668; GCN-NEXT: v_accvgpr_read_b32 v0, a0669; GCN-NEXT: v_accvgpr_read_b32 v1, a1670; GCN-NEXT: v_accvgpr_read_b32 v2, a2671; GCN-NEXT: v_accvgpr_read_b32 v3, a3672; GCN-NEXT: s_setpc_b64 s[30:31]673 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,674 i32 2, ; cbsz675 i32 0, ; blgp676 i32 0, i32 %scale0, i32 0, i32 %scale1)677 ret <4 x float> %result678}679 680; This should be optimized to avoid the scale681define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp0__constant_scale_0_0(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {682; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp0__constant_scale_0_0:683; GCN: ; %bb.0:684; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)685; GCN-NEXT: v_accvgpr_write_b32 a0, v14686; GCN-NEXT: v_accvgpr_write_b32 a1, v15687; GCN-NEXT: v_accvgpr_write_b32 a2, v16688; GCN-NEXT: v_accvgpr_write_b32 a3, v17689; GCN-NEXT: s_nop 1690; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:13], a[0:3] cbsz:2691; GCN-NEXT: s_nop 11692; GCN-NEXT: v_accvgpr_read_b32 v0, a0693; GCN-NEXT: v_accvgpr_read_b32 v1, a1694; GCN-NEXT: v_accvgpr_read_b32 v2, a2695; GCN-NEXT: v_accvgpr_read_b32 v3, a3696; GCN-NEXT: s_setpc_b64 s[30:31]697 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,698 i32 2, ; cbsz699 i32 0, ; blgp700 i32 0, i32 0, i32 0, i32 0)701 ret <4 x float> %result702}703 704; fp6 x bf8705define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp1(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {706; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp1:707; GCN: ; %bb.0:708; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)709; GCN-NEXT: v_accvgpr_write_b32 a0, v14710; GCN-NEXT: v_accvgpr_write_b32 a1, v15711; GCN-NEXT: v_accvgpr_write_b32 a2, v16712; GCN-NEXT: v_accvgpr_write_b32 a3, v17713; GCN-NEXT: s_nop 1714; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] cbsz:2 blgp:1715; GCN-NEXT: s_nop 11716; GCN-NEXT: v_accvgpr_read_b32 v0, a0717; GCN-NEXT: v_accvgpr_read_b32 v1, a1718; GCN-NEXT: v_accvgpr_read_b32 v2, a2719; GCN-NEXT: v_accvgpr_read_b32 v3, a3720; GCN-NEXT: s_setpc_b64 s[30:31]721 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,722 i32 2, ; cbsz723 i32 1, ; blgp724 i32 0, i32 %scale0, i32 0, i32 %scale1)725 ret <4 x float> %result726}727 728; This should be optimized to avoid the scale729define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp1__constant_scale_0_0(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {730; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp1__constant_scale_0_0:731; GCN: ; %bb.0:732; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)733; GCN-NEXT: v_accvgpr_write_b32 a0, v14734; GCN-NEXT: v_accvgpr_write_b32 a1, v15735; GCN-NEXT: v_accvgpr_write_b32 a2, v16736; GCN-NEXT: v_accvgpr_write_b32 a3, v17737; GCN-NEXT: s_nop 1738; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:13], a[0:3] cbsz:2 blgp:1739; GCN-NEXT: s_nop 11740; GCN-NEXT: v_accvgpr_read_b32 v0, a0741; GCN-NEXT: v_accvgpr_read_b32 v1, a1742; GCN-NEXT: v_accvgpr_read_b32 v2, a2743; GCN-NEXT: v_accvgpr_read_b32 v3, a3744; GCN-NEXT: s_setpc_b64 s[30:31]745 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,746 i32 2, ; cbsz747 i32 1, ; blgp748 i32 0, i32 0, i32 0, i32 0)749 ret <4 x float> %result750}751 752; fp6 x fp6753define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp2(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {754; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp2:755; GCN: ; %bb.0:756; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)757; GCN-NEXT: v_accvgpr_write_b32 a0, v12758; GCN-NEXT: v_accvgpr_write_b32 a1, v13759; GCN-NEXT: v_accvgpr_write_b32 a2, v14760; GCN-NEXT: v_accvgpr_write_b32 a3, v15761; GCN-NEXT: s_nop 1762; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3], v16, v17 op_sel_hi:[0,0,0] cbsz:2 blgp:2763; GCN-NEXT: s_nop 7764; GCN-NEXT: v_accvgpr_read_b32 v0, a0765; GCN-NEXT: v_accvgpr_read_b32 v1, a1766; GCN-NEXT: v_accvgpr_read_b32 v2, a2767; GCN-NEXT: v_accvgpr_read_b32 v3, a3768; GCN-NEXT: s_setpc_b64 s[30:31]769 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,770 i32 2, ; cbsz771 i32 2, ; blgp772 i32 0, i32 %scale0, i32 0, i32 %scale1)773 ret <4 x float> %result774}775 776; This should be optimized to avoid the scale777define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp2__constant_scale_0_0(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {778; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp2__constant_scale_0_0:779; GCN: ; %bb.0:780; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)781; GCN-NEXT: v_accvgpr_write_b32 a0, v12782; GCN-NEXT: v_accvgpr_write_b32 a1, v13783; GCN-NEXT: v_accvgpr_write_b32 a2, v14784; GCN-NEXT: v_accvgpr_write_b32 a3, v15785; GCN-NEXT: s_nop 1786; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3] cbsz:2 blgp:2787; GCN-NEXT: s_nop 7788; GCN-NEXT: v_accvgpr_read_b32 v0, a0789; GCN-NEXT: v_accvgpr_read_b32 v1, a1790; GCN-NEXT: v_accvgpr_read_b32 v2, a2791; GCN-NEXT: v_accvgpr_read_b32 v3, a3792; GCN-NEXT: s_setpc_b64 s[30:31]793 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,794 i32 2, ; cbsz795 i32 2, ; blgp796 i32 0, i32 0, i32 0, i32 0)797 ret <4 x float> %result798}799 800; fp6 x bf6801define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp3(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {802; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp3:803; GCN: ; %bb.0:804; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)805; GCN-NEXT: v_accvgpr_write_b32 a0, v12806; GCN-NEXT: v_accvgpr_write_b32 a1, v13807; GCN-NEXT: v_accvgpr_write_b32 a2, v14808; GCN-NEXT: v_accvgpr_write_b32 a3, v15809; GCN-NEXT: s_nop 1810; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3], v16, v17 op_sel_hi:[0,0,0] cbsz:2 blgp:3811; GCN-NEXT: s_nop 7812; GCN-NEXT: v_accvgpr_read_b32 v0, a0813; GCN-NEXT: v_accvgpr_read_b32 v1, a1814; GCN-NEXT: v_accvgpr_read_b32 v2, a2815; GCN-NEXT: v_accvgpr_read_b32 v3, a3816; GCN-NEXT: s_setpc_b64 s[30:31]817 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,818 i32 2, ; cbsz819 i32 3, ; blgp820 i32 0, i32 %scale0, i32 0, i32 %scale1)821 ret <4 x float> %result822}823 824; This should be optimized to avoid the scale825define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp3__constant_scale_0_0(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {826; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp3__constant_scale_0_0:827; GCN: ; %bb.0:828; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)829; GCN-NEXT: v_accvgpr_write_b32 a0, v12830; GCN-NEXT: v_accvgpr_write_b32 a1, v13831; GCN-NEXT: v_accvgpr_write_b32 a2, v14832; GCN-NEXT: v_accvgpr_write_b32 a3, v15833; GCN-NEXT: s_nop 1834; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3] cbsz:2 blgp:3835; GCN-NEXT: s_nop 7836; GCN-NEXT: v_accvgpr_read_b32 v0, a0837; GCN-NEXT: v_accvgpr_read_b32 v1, a1838; GCN-NEXT: v_accvgpr_read_b32 v2, a2839; GCN-NEXT: v_accvgpr_read_b32 v3, a3840; GCN-NEXT: s_setpc_b64 s[30:31]841 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,842 i32 2, ; cbsz843 i32 3, ; blgp844 i32 0, i32 0, i32 0, i32 0)845 ret <4 x float> %result846}847 848 849; bf6 x fp8850define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp0(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {851; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp0:852; GCN: ; %bb.0:853; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)854; GCN-NEXT: v_accvgpr_write_b32 a0, v14855; GCN-NEXT: v_accvgpr_write_b32 a1, v15856; GCN-NEXT: v_accvgpr_write_b32 a2, v16857; GCN-NEXT: v_accvgpr_write_b32 a3, v17858; GCN-NEXT: s_nop 1859; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] cbsz:3860; GCN-NEXT: s_nop 11861; GCN-NEXT: v_accvgpr_read_b32 v0, a0862; GCN-NEXT: v_accvgpr_read_b32 v1, a1863; GCN-NEXT: v_accvgpr_read_b32 v2, a2864; GCN-NEXT: v_accvgpr_read_b32 v3, a3865; GCN-NEXT: s_setpc_b64 s[30:31]866 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,867 i32 3, ; cbsz868 i32 0, ; blgp869 i32 0, i32 %scale0, i32 0, i32 %scale1)870 ret <4 x float> %result871}872 873; This should be optimized to avoid the scale874define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp0__constant_scale_0_0(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {875; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp0__constant_scale_0_0:876; GCN: ; %bb.0:877; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)878; GCN-NEXT: v_accvgpr_write_b32 a0, v14879; GCN-NEXT: v_accvgpr_write_b32 a1, v15880; GCN-NEXT: v_accvgpr_write_b32 a2, v16881; GCN-NEXT: v_accvgpr_write_b32 a3, v17882; GCN-NEXT: s_nop 1883; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:13], a[0:3] cbsz:3884; GCN-NEXT: s_nop 11885; GCN-NEXT: v_accvgpr_read_b32 v0, a0886; GCN-NEXT: v_accvgpr_read_b32 v1, a1887; GCN-NEXT: v_accvgpr_read_b32 v2, a2888; GCN-NEXT: v_accvgpr_read_b32 v3, a3889; GCN-NEXT: s_setpc_b64 s[30:31]890 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,891 i32 3, ; cbsz892 i32 0, ; blgp893 i32 0, i32 0, i32 0, i32 0)894 ret <4 x float> %result895}896 897; bf6 x bf8898define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp1(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {899; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp1:900; GCN: ; %bb.0:901; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)902; GCN-NEXT: v_accvgpr_write_b32 a0, v14903; GCN-NEXT: v_accvgpr_write_b32 a1, v15904; GCN-NEXT: v_accvgpr_write_b32 a2, v16905; GCN-NEXT: v_accvgpr_write_b32 a3, v17906; GCN-NEXT: s_nop 1907; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] cbsz:3 blgp:1908; GCN-NEXT: s_nop 11909; GCN-NEXT: v_accvgpr_read_b32 v0, a0910; GCN-NEXT: v_accvgpr_read_b32 v1, a1911; GCN-NEXT: v_accvgpr_read_b32 v2, a2912; GCN-NEXT: v_accvgpr_read_b32 v3, a3913; GCN-NEXT: s_setpc_b64 s[30:31]914 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,915 i32 3, ; cbsz916 i32 1, ; blgp917 i32 0, i32 %scale0, i32 0, i32 %scale1)918 ret <4 x float> %result919}920 921; This should be optimized to avoid the scale922define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp1__constant_scale_0_0(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {923; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp1__constant_scale_0_0:924; GCN: ; %bb.0:925; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)926; GCN-NEXT: v_accvgpr_write_b32 a0, v14927; GCN-NEXT: v_accvgpr_write_b32 a1, v15928; GCN-NEXT: v_accvgpr_write_b32 a2, v16929; GCN-NEXT: v_accvgpr_write_b32 a3, v17930; GCN-NEXT: s_nop 1931; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:13], a[0:3] cbsz:3 blgp:1932; GCN-NEXT: s_nop 11933; GCN-NEXT: v_accvgpr_read_b32 v0, a0934; GCN-NEXT: v_accvgpr_read_b32 v1, a1935; GCN-NEXT: v_accvgpr_read_b32 v2, a2936; GCN-NEXT: v_accvgpr_read_b32 v3, a3937; GCN-NEXT: s_setpc_b64 s[30:31]938 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,939 i32 3, ; cbsz940 i32 1, ; blgp941 i32 0, i32 0, i32 0, i32 0)942 ret <4 x float> %result943}944 945; bf6 x fp6946define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp2(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {947; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp2:948; GCN: ; %bb.0:949; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)950; GCN-NEXT: v_accvgpr_write_b32 a0, v12951; GCN-NEXT: v_accvgpr_write_b32 a1, v13952; GCN-NEXT: v_accvgpr_write_b32 a2, v14953; GCN-NEXT: v_accvgpr_write_b32 a3, v15954; GCN-NEXT: s_nop 1955; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3], v16, v17 op_sel_hi:[0,0,0] cbsz:3 blgp:2956; GCN-NEXT: s_nop 7957; GCN-NEXT: v_accvgpr_read_b32 v0, a0958; GCN-NEXT: v_accvgpr_read_b32 v1, a1959; GCN-NEXT: v_accvgpr_read_b32 v2, a2960; GCN-NEXT: v_accvgpr_read_b32 v3, a3961; GCN-NEXT: s_setpc_b64 s[30:31]962 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,963 i32 3, ; cbsz964 i32 2, ; blgp965 i32 0, i32 %scale0, i32 0, i32 %scale1)966 ret <4 x float> %result967}968 969; This should be optimized to avoid the scale970define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp2__constant_scale_0_0(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {971; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp2__constant_scale_0_0:972; GCN: ; %bb.0:973; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)974; GCN-NEXT: v_accvgpr_write_b32 a0, v12975; GCN-NEXT: v_accvgpr_write_b32 a1, v13976; GCN-NEXT: v_accvgpr_write_b32 a2, v14977; GCN-NEXT: v_accvgpr_write_b32 a3, v15978; GCN-NEXT: s_nop 1979; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3] cbsz:3 blgp:2980; GCN-NEXT: s_nop 7981; GCN-NEXT: v_accvgpr_read_b32 v0, a0982; GCN-NEXT: v_accvgpr_read_b32 v1, a1983; GCN-NEXT: v_accvgpr_read_b32 v2, a2984; GCN-NEXT: v_accvgpr_read_b32 v3, a3985; GCN-NEXT: s_setpc_b64 s[30:31]986 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,987 i32 3, ; cbsz988 i32 2, ; blgp989 i32 0, i32 0, i32 0, i32 0)990 ret <4 x float> %result991}992 993; bf6 x fp4994define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp4(<6 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {995; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp4:996; GCN: ; %bb.0:997; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)998; GCN-NEXT: v_accvgpr_write_b32 a0, v10999; GCN-NEXT: v_accvgpr_write_b32 a1, v111000; GCN-NEXT: v_accvgpr_write_b32 a2, v121001; GCN-NEXT: v_accvgpr_write_b32 a3, v131002; GCN-NEXT: s_nop 11003; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:9], a[0:3], v14, v15 op_sel_hi:[0,0,0] cbsz:3 blgp:41004; GCN-NEXT: s_nop 71005; GCN-NEXT: v_accvgpr_read_b32 v0, a01006; GCN-NEXT: v_accvgpr_read_b32 v1, a11007; GCN-NEXT: v_accvgpr_read_b32 v2, a21008; GCN-NEXT: v_accvgpr_read_b32 v3, a31009; GCN-NEXT: s_setpc_b64 s[30:31]1010 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v4i32(<6 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,1011 i32 3, ; cbsz1012 i32 4, ; blgp1013 i32 0, i32 %scale0, i32 0, i32 %scale1)1014 ret <4 x float> %result1015}1016 1017; This should be optimized to avoid the scale1018define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp4__constant_scale_0_0(<6 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2) {1019; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp4__constant_scale_0_0:1020; GCN: ; %bb.0:1021; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1022; GCN-NEXT: v_accvgpr_write_b32 a0, v101023; GCN-NEXT: v_accvgpr_write_b32 a1, v111024; GCN-NEXT: v_accvgpr_write_b32 a2, v121025; GCN-NEXT: v_accvgpr_write_b32 a3, v131026; GCN-NEXT: s_nop 11027; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:9], a[0:3] cbsz:3 blgp:41028; GCN-NEXT: s_nop 71029; GCN-NEXT: v_accvgpr_read_b32 v0, a01030; GCN-NEXT: v_accvgpr_read_b32 v1, a11031; GCN-NEXT: v_accvgpr_read_b32 v2, a21032; GCN-NEXT: v_accvgpr_read_b32 v3, a31033; GCN-NEXT: s_setpc_b64 s[30:31]1034 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v4i32(<6 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,1035 i32 3, ; cbsz1036 i32 4, ; blgp1037 i32 0, i32 0, i32 0, i32 0)1038 ret <4 x float> %result1039}1040 1041; bf6 x bf61042define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp3(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1043; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp3:1044; GCN: ; %bb.0:1045; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1046; GCN-NEXT: v_accvgpr_write_b32 a0, v121047; GCN-NEXT: v_accvgpr_write_b32 a1, v131048; GCN-NEXT: v_accvgpr_write_b32 a2, v141049; GCN-NEXT: v_accvgpr_write_b32 a3, v151050; GCN-NEXT: s_nop 11051; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3], v16, v17 op_sel_hi:[0,0,0] cbsz:3 blgp:31052; GCN-NEXT: s_nop 71053; GCN-NEXT: v_accvgpr_read_b32 v0, a01054; GCN-NEXT: v_accvgpr_read_b32 v1, a11055; GCN-NEXT: v_accvgpr_read_b32 v2, a21056; GCN-NEXT: v_accvgpr_read_b32 v3, a31057; GCN-NEXT: s_setpc_b64 s[30:31]1058 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,1059 i32 3, ; cbsz1060 i32 3, ; blgp1061 i32 0, i32 %scale0, i32 0, i32 %scale1)1062 ret <4 x float> %result1063}1064 1065; This should be optimized to avoid the scale1066define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp3__constant_scale_0_0(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {1067; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz3__blgp3__constant_scale_0_0:1068; GCN: ; %bb.0:1069; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1070; GCN-NEXT: v_accvgpr_write_b32 a0, v121071; GCN-NEXT: v_accvgpr_write_b32 a1, v131072; GCN-NEXT: v_accvgpr_write_b32 a2, v141073; GCN-NEXT: v_accvgpr_write_b32 a3, v151074; GCN-NEXT: s_nop 11075; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:11], a[0:3] cbsz:3 blgp:31076; GCN-NEXT: s_nop 71077; GCN-NEXT: v_accvgpr_read_b32 v0, a01078; GCN-NEXT: v_accvgpr_read_b32 v1, a11079; GCN-NEXT: v_accvgpr_read_b32 v2, a21080; GCN-NEXT: v_accvgpr_read_b32 v3, a31081; GCN-NEXT: s_setpc_b64 s[30:31]1082 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,1083 i32 3, ; cbsz1084 i32 3, ; blgp1085 i32 0, i32 0, i32 0, i32 0)1086 ret <4 x float> %result1087}1088 1089; fp6 x fp41090define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp4(<6 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1091; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp4:1092; GCN: ; %bb.0:1093; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1094; GCN-NEXT: v_accvgpr_write_b32 a0, v101095; GCN-NEXT: v_accvgpr_write_b32 a1, v111096; GCN-NEXT: v_accvgpr_write_b32 a2, v121097; GCN-NEXT: v_accvgpr_write_b32 a3, v131098; GCN-NEXT: s_nop 11099; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:9], a[0:3], v14, v15 op_sel_hi:[0,0,0] cbsz:2 blgp:41100; GCN-NEXT: s_nop 71101; GCN-NEXT: v_accvgpr_read_b32 v0, a01102; GCN-NEXT: v_accvgpr_read_b32 v1, a11103; GCN-NEXT: v_accvgpr_read_b32 v2, a21104; GCN-NEXT: v_accvgpr_read_b32 v3, a31105; GCN-NEXT: s_setpc_b64 s[30:31]1106 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v4i32(<6 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,1107 i32 2, ; cbsz1108 i32 4, ; blgp1109 i32 0, i32 %scale0, i32 0, i32 %scale1)1110 ret <4 x float> %result1111}1112 1113; This should be optimized to avoid the scale1114define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp4__constant_scale_0_0(<6 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2) {1115; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz2__blgp4__constant_scale_0_0:1116; GCN: ; %bb.0:1117; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1118; GCN-NEXT: v_accvgpr_write_b32 a0, v101119; GCN-NEXT: v_accvgpr_write_b32 a1, v111120; GCN-NEXT: v_accvgpr_write_b32 a2, v121121; GCN-NEXT: v_accvgpr_write_b32 a3, v131122; GCN-NEXT: s_nop 11123; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:9], a[0:3] cbsz:2 blgp:41124; GCN-NEXT: s_nop 71125; GCN-NEXT: v_accvgpr_read_b32 v0, a01126; GCN-NEXT: v_accvgpr_read_b32 v1, a11127; GCN-NEXT: v_accvgpr_read_b32 v2, a21128; GCN-NEXT: v_accvgpr_read_b32 v3, a31129; GCN-NEXT: s_setpc_b64 s[30:31]1130 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v4i32(<6 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,1131 i32 2, ; cbsz1132 i32 4, ; blgp1133 i32 0, i32 0, i32 0, i32 0)1134 ret <4 x float> %result1135}1136 1137; fp4 x fp81138define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp0(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1139; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp0:1140; GCN: ; %bb.0:1141; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1142; GCN-NEXT: v_accvgpr_write_b32 a0, v121143; GCN-NEXT: v_accvgpr_write_b32 a1, v131144; GCN-NEXT: v_accvgpr_write_b32 a2, v141145; GCN-NEXT: v_accvgpr_write_b32 a3, v151146; GCN-NEXT: s_nop 11147; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:11], a[0:3], v16, v17 op_sel_hi:[0,0,0] cbsz:41148; GCN-NEXT: s_nop 111149; GCN-NEXT: v_accvgpr_read_b32 v0, a01150; GCN-NEXT: v_accvgpr_read_b32 v1, a11151; GCN-NEXT: v_accvgpr_read_b32 v2, a21152; GCN-NEXT: v_accvgpr_read_b32 v3, a31153; GCN-NEXT: s_setpc_b64 s[30:31]1154 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v8i32(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,1155 i32 4, ; cbsz1156 i32 0, ; blgp1157 i32 0, i32 %scale0, i32 0, i32 %scale1)1158 ret <4 x float> %result1159}1160 1161; This should be optimized to avoid the scale1162define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp0__constant_scale_0_0(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {1163; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp0__constant_scale_0_0:1164; GCN: ; %bb.0:1165; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1166; GCN-NEXT: v_accvgpr_write_b32 a0, v121167; GCN-NEXT: v_accvgpr_write_b32 a1, v131168; GCN-NEXT: v_accvgpr_write_b32 a2, v141169; GCN-NEXT: v_accvgpr_write_b32 a3, v151170; GCN-NEXT: s_nop 11171; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:11], a[0:3] cbsz:41172; GCN-NEXT: s_nop 111173; GCN-NEXT: v_accvgpr_read_b32 v0, a01174; GCN-NEXT: v_accvgpr_read_b32 v1, a11175; GCN-NEXT: v_accvgpr_read_b32 v2, a21176; GCN-NEXT: v_accvgpr_read_b32 v3, a31177; GCN-NEXT: s_setpc_b64 s[30:31]1178 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v8i32(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,1179 i32 4, ; cbsz1180 i32 0, ; blgp1181 i32 0, i32 0, i32 0, i32 0)1182 ret <4 x float> %result1183}1184 1185; fp4 x bf81186define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp1(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1187; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp1:1188; GCN: ; %bb.0:1189; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1190; GCN-NEXT: v_accvgpr_write_b32 a0, v121191; GCN-NEXT: v_accvgpr_write_b32 a1, v131192; GCN-NEXT: v_accvgpr_write_b32 a2, v141193; GCN-NEXT: v_accvgpr_write_b32 a3, v151194; GCN-NEXT: s_nop 11195; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:11], a[0:3], v16, v17 op_sel_hi:[0,0,0] cbsz:4 blgp:11196; GCN-NEXT: s_nop 111197; GCN-NEXT: v_accvgpr_read_b32 v0, a01198; GCN-NEXT: v_accvgpr_read_b32 v1, a11199; GCN-NEXT: v_accvgpr_read_b32 v2, a21200; GCN-NEXT: v_accvgpr_read_b32 v3, a31201; GCN-NEXT: s_setpc_b64 s[30:31]1202 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v8i32(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,1203 i32 4, ; cbsz1204 i32 1, ; blgp1205 i32 0, i32 %scale0, i32 0, i32 %scale1)1206 ret <4 x float> %result1207}1208 1209; This should be optimized to avoid the scale1210define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp1__constant_scale_0_0(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {1211; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp1__constant_scale_0_0:1212; GCN: ; %bb.0:1213; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1214; GCN-NEXT: v_accvgpr_write_b32 a0, v121215; GCN-NEXT: v_accvgpr_write_b32 a1, v131216; GCN-NEXT: v_accvgpr_write_b32 a2, v141217; GCN-NEXT: v_accvgpr_write_b32 a3, v151218; GCN-NEXT: s_nop 11219; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:11], a[0:3] cbsz:4 blgp:11220; GCN-NEXT: s_nop 111221; GCN-NEXT: v_accvgpr_read_b32 v0, a01222; GCN-NEXT: v_accvgpr_read_b32 v1, a11223; GCN-NEXT: v_accvgpr_read_b32 v2, a21224; GCN-NEXT: v_accvgpr_read_b32 v3, a31225; GCN-NEXT: s_setpc_b64 s[30:31]1226 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v8i32(<4 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,1227 i32 4, ; cbsz1228 i32 1, ; blgp1229 i32 0, i32 0, i32 0, i32 0)1230 ret <4 x float> %result1231}1232 1233; fp4 x fp61234define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp2(<4 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1235; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp2:1236; GCN: ; %bb.0:1237; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1238; GCN-NEXT: v_accvgpr_write_b32 a0, v101239; GCN-NEXT: v_accvgpr_write_b32 a1, v111240; GCN-NEXT: v_accvgpr_write_b32 a2, v121241; GCN-NEXT: v_accvgpr_write_b32 a3, v131242; GCN-NEXT: s_nop 11243; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:9], a[0:3], v14, v15 op_sel_hi:[0,0,0] cbsz:4 blgp:21244; GCN-NEXT: s_nop 71245; GCN-NEXT: v_accvgpr_read_b32 v0, a01246; GCN-NEXT: v_accvgpr_read_b32 v1, a11247; GCN-NEXT: v_accvgpr_read_b32 v2, a21248; GCN-NEXT: v_accvgpr_read_b32 v3, a31249; GCN-NEXT: s_setpc_b64 s[30:31]1250 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v6i32(<4 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,1251 i32 4, ; cbsz1252 i32 2, ; blgp1253 i32 0, i32 %scale0, i32 0, i32 %scale1)1254 ret <4 x float> %result1255}1256 1257; This should be optimized to avoid the scale1258define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp2__constant_scale_0_0(<4 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {1259; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp2__constant_scale_0_0:1260; GCN: ; %bb.0:1261; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1262; GCN-NEXT: v_accvgpr_write_b32 a0, v101263; GCN-NEXT: v_accvgpr_write_b32 a1, v111264; GCN-NEXT: v_accvgpr_write_b32 a2, v121265; GCN-NEXT: v_accvgpr_write_b32 a3, v131266; GCN-NEXT: s_nop 11267; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:9], a[0:3] cbsz:4 blgp:21268; GCN-NEXT: s_nop 71269; GCN-NEXT: v_accvgpr_read_b32 v0, a01270; GCN-NEXT: v_accvgpr_read_b32 v1, a11271; GCN-NEXT: v_accvgpr_read_b32 v2, a21272; GCN-NEXT: v_accvgpr_read_b32 v3, a31273; GCN-NEXT: s_setpc_b64 s[30:31]1274 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v6i32(<4 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,1275 i32 4, ; cbsz1276 i32 2, ; blgp1277 i32 0, i32 0, i32 0, i32 0)1278 ret <4 x float> %result1279}1280 1281; fp4 x bf61282define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp3(<4 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1283; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp3:1284; GCN: ; %bb.0:1285; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1286; GCN-NEXT: v_accvgpr_write_b32 a0, v101287; GCN-NEXT: v_accvgpr_write_b32 a1, v111288; GCN-NEXT: v_accvgpr_write_b32 a2, v121289; GCN-NEXT: v_accvgpr_write_b32 a3, v131290; GCN-NEXT: s_nop 11291; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:9], a[0:3], v14, v15 op_sel_hi:[0,0,0] cbsz:4 blgp:31292; GCN-NEXT: s_nop 71293; GCN-NEXT: v_accvgpr_read_b32 v0, a01294; GCN-NEXT: v_accvgpr_read_b32 v1, a11295; GCN-NEXT: v_accvgpr_read_b32 v2, a21296; GCN-NEXT: v_accvgpr_read_b32 v3, a31297; GCN-NEXT: s_setpc_b64 s[30:31]1298 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v6i32(<4 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,1299 i32 4, ; cbsz1300 i32 3, ; blgp1301 i32 0, i32 %scale0, i32 0, i32 %scale1)1302 ret <4 x float> %result1303}1304 1305; This should be optimized to avoid the scale1306define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp3__constant_scale_0_0(<4 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2) {1307; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp3__constant_scale_0_0:1308; GCN: ; %bb.0:1309; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1310; GCN-NEXT: v_accvgpr_write_b32 a0, v101311; GCN-NEXT: v_accvgpr_write_b32 a1, v111312; GCN-NEXT: v_accvgpr_write_b32 a2, v121313; GCN-NEXT: v_accvgpr_write_b32 a3, v131314; GCN-NEXT: s_nop 11315; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:9], a[0:3] cbsz:4 blgp:31316; GCN-NEXT: s_nop 71317; GCN-NEXT: v_accvgpr_read_b32 v0, a01318; GCN-NEXT: v_accvgpr_read_b32 v1, a11319; GCN-NEXT: v_accvgpr_read_b32 v2, a21320; GCN-NEXT: v_accvgpr_read_b32 v3, a31321; GCN-NEXT: s_setpc_b64 s[30:31]1322 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v6i32(<4 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,1323 i32 4, ; cbsz1324 i32 3, ; blgp1325 i32 0, i32 0, i32 0, i32 0)1326 ret <4 x float> %result1327}1328 1329; fp4 x fp41330define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp4(<4 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1331; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp4:1332; GCN: ; %bb.0:1333; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1334; GCN-NEXT: v_accvgpr_write_b32 a0, v81335; GCN-NEXT: v_accvgpr_write_b32 a1, v91336; GCN-NEXT: v_accvgpr_write_b32 a2, v101337; GCN-NEXT: v_accvgpr_write_b32 a3, v111338; GCN-NEXT: s_nop 11339; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:7], a[0:3], v12, v13 op_sel_hi:[0,0,0] cbsz:4 blgp:41340; GCN-NEXT: s_nop 71341; GCN-NEXT: v_accvgpr_read_b32 v0, a01342; GCN-NEXT: v_accvgpr_read_b32 v1, a11343; GCN-NEXT: v_accvgpr_read_b32 v2, a21344; GCN-NEXT: v_accvgpr_read_b32 v3, a31345; GCN-NEXT: s_setpc_b64 s[30:31]1346 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v4i32(<4 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,1347 i32 4, ; cbsz1348 i32 4, ; blgp1349 i32 0, i32 %scale0, i32 0, i32 %scale1)1350 ret <4 x float> %result1351}1352 1353; This should be optimized to avoid the scale1354define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp4__constant_scale_0_0(<4 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2) {1355; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__cbsz4__blgp4__constant_scale_0_0:1356; GCN: ; %bb.0:1357; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1358; GCN-NEXT: v_accvgpr_write_b32 a0, v81359; GCN-NEXT: v_accvgpr_write_b32 a1, v91360; GCN-NEXT: v_accvgpr_write_b32 a2, v101361; GCN-NEXT: v_accvgpr_write_b32 a3, v111362; GCN-NEXT: s_nop 11363; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:3], v[4:7], a[0:3] cbsz:4 blgp:41364; GCN-NEXT: s_nop 71365; GCN-NEXT: v_accvgpr_read_b32 v0, a01366; GCN-NEXT: v_accvgpr_read_b32 v1, a11367; GCN-NEXT: v_accvgpr_read_b32 v2, a21368; GCN-NEXT: v_accvgpr_read_b32 v3, a31369; GCN-NEXT: s_setpc_b64 s[30:31]1370 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v4i32(<4 x i32> %arg0, <4 x i32> %arg1, <4 x float> %arg2,1371 i32 4, ; cbsz1372 i32 4, ; blgp1373 i32 0, i32 0, i32 0, i32 0)1374 ret <4 x float> %result1375}1376 1377; --------------------------------------------------------------------1378; Different input parameter classes1379; --------------------------------------------------------------------1380 1381define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__sgpr_scaleA__sgpr_scaleB(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 inreg %scale0, i32 inreg %scale1) {1382; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__sgpr_scaleA__sgpr_scaleB:1383; GCN: ; %bb.0:1384; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1385; GCN-NEXT: v_accvgpr_write_b32 a0, v161386; GCN-NEXT: v_accvgpr_write_b32 a1, v171387; GCN-NEXT: v_accvgpr_write_b32 a2, v181388; GCN-NEXT: v_accvgpr_write_b32 a3, v191389; GCN-NEXT: v_mov_b32_e32 v16, s01390; GCN-NEXT: v_mov_b32_e32 v17, s11391; GCN-NEXT: s_nop 11392; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, v17 op_sel_hi:[0,0,0]1393; GCN-NEXT: s_nop 111394; GCN-NEXT: v_accvgpr_read_b32 v0, a01395; GCN-NEXT: v_accvgpr_read_b32 v1, a11396; GCN-NEXT: v_accvgpr_read_b32 v2, a21397; GCN-NEXT: v_accvgpr_read_b32 v3, a31398; GCN-NEXT: s_setpc_b64 s[30:31]1399 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 %scale0, i32 0, i32 %scale1)1400 ret <4 x float> %result1401}1402 1403define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__sgpr_scaleA__vgpr_scaleB(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 inreg %scale0, i32 %scale1) {1404; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__sgpr_scaleA__vgpr_scaleB:1405; GCN: ; %bb.0:1406; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1407; GCN-NEXT: v_accvgpr_write_b32 a0, v161408; GCN-NEXT: v_accvgpr_write_b32 a1, v171409; GCN-NEXT: v_accvgpr_write_b32 a2, v181410; GCN-NEXT: v_accvgpr_write_b32 a3, v191411; GCN-NEXT: v_mov_b32_e32 v16, s01412; GCN-NEXT: s_nop 11413; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, v20 op_sel_hi:[0,0,0]1414; GCN-NEXT: s_nop 111415; GCN-NEXT: v_accvgpr_read_b32 v0, a01416; GCN-NEXT: v_accvgpr_read_b32 v1, a11417; GCN-NEXT: v_accvgpr_read_b32 v2, a21418; GCN-NEXT: v_accvgpr_read_b32 v3, a31419; GCN-NEXT: s_setpc_b64 s[30:31]1420 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 %scale0, i32 0, i32 %scale1)1421 ret <4 x float> %result1422}1423 1424define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__vgpr_scaleA__sgpr_scaleB(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 inreg %scale1) {1425; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__vgpr_scaleA__sgpr_scaleB:1426; GCN: ; %bb.0:1427; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1428; GCN-NEXT: v_accvgpr_write_b32 a0, v161429; GCN-NEXT: v_accvgpr_write_b32 a1, v171430; GCN-NEXT: v_accvgpr_write_b32 a2, v181431; GCN-NEXT: v_accvgpr_write_b32 a3, v191432; GCN-NEXT: v_mov_b32_e32 v16, s01433; GCN-NEXT: s_nop 11434; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v16 op_sel_hi:[0,0,0]1435; GCN-NEXT: s_nop 111436; GCN-NEXT: v_accvgpr_read_b32 v0, a01437; GCN-NEXT: v_accvgpr_read_b32 v1, a11438; GCN-NEXT: v_accvgpr_read_b32 v2, a21439; GCN-NEXT: v_accvgpr_read_b32 v3, a31440; GCN-NEXT: s_setpc_b64 s[30:31]1441 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 %scale0, i32 0, i32 %scale1)1442 ret <4 x float> %result1443}1444 1445define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgprs(<8 x i32> inreg %arg0, <8 x i32> inreg %arg1, <4 x float> inreg %arg2, i32 %scale0, i32 %scale1) {1446; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgprs:1447; SDAG: ; %bb.0:1448; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1449; SDAG-NEXT: v_mov_b32_e32 v14, s01450; SDAG-NEXT: v_mov_b32_e32 v15, s11451; SDAG-NEXT: v_mov_b32_e32 v16, s21452; SDAG-NEXT: v_mov_b32_e32 v17, s31453; SDAG-NEXT: v_mov_b32_e32 v18, s161454; SDAG-NEXT: v_mov_b32_e32 v19, s171455; SDAG-NEXT: v_mov_b32_e32 v20, s181456; SDAG-NEXT: v_mov_b32_e32 v21, s191457; SDAG-NEXT: v_mov_b32_e32 v4, s281458; SDAG-NEXT: v_mov_b32_e32 v5, s291459; SDAG-NEXT: v_mov_b32_e32 v6, s201460; SDAG-NEXT: v_mov_b32_e32 v7, s211461; SDAG-NEXT: v_mov_b32_e32 v8, s221462; SDAG-NEXT: v_mov_b32_e32 v9, s231463; SDAG-NEXT: v_mov_b32_e32 v10, s241464; SDAG-NEXT: v_mov_b32_e32 v11, s251465; SDAG-NEXT: v_mov_b32_e32 v12, s261466; SDAG-NEXT: v_mov_b32_e32 v13, s271467; SDAG-NEXT: v_accvgpr_write_b32 a0, v41468; SDAG-NEXT: v_accvgpr_write_b32 a1, v51469; SDAG-NEXT: v_accvgpr_write_b32 a2, v01470; SDAG-NEXT: v_accvgpr_write_b32 a3, v11471; SDAG-NEXT: s_nop 11472; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[14:21], v[6:13], a[0:3], v2, v3 op_sel_hi:[0,0,0]1473; SDAG-NEXT: s_nop 111474; SDAG-NEXT: v_accvgpr_read_b32 v0, a01475; SDAG-NEXT: v_accvgpr_read_b32 v1, a11476; SDAG-NEXT: v_accvgpr_read_b32 v2, a21477; SDAG-NEXT: v_accvgpr_read_b32 v3, a31478; SDAG-NEXT: s_setpc_b64 s[30:31]1479;1480; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgprs:1481; GISEL: ; %bb.0:1482; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1483; GISEL-NEXT: s_mov_b32 s12, s01484; GISEL-NEXT: s_mov_b32 s13, s11485; GISEL-NEXT: s_mov_b32 s14, s21486; GISEL-NEXT: s_mov_b32 s15, s31487; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]1488; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]1489; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[16:17]1490; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[18:19]1491; GISEL-NEXT: v_mov_b32_e32 v20, s281492; GISEL-NEXT: v_mov_b32_e32 v21, s291493; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[20:21]1494; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[22:23]1495; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[24:25]1496; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[26:27]1497; GISEL-NEXT: v_accvgpr_write_b32 a0, v201498; GISEL-NEXT: v_accvgpr_write_b32 a1, v211499; GISEL-NEXT: v_accvgpr_write_b32 a2, v01500; GISEL-NEXT: v_accvgpr_write_b32 a3, v11501; GISEL-NEXT: s_nop 11502; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[4:11], v[12:19], a[0:3], v2, v3 op_sel_hi:[0,0,0]1503; GISEL-NEXT: s_nop 111504; GISEL-NEXT: v_accvgpr_read_b32 v0, a01505; GISEL-NEXT: v_accvgpr_read_b32 v1, a11506; GISEL-NEXT: v_accvgpr_read_b32 v2, a21507; GISEL-NEXT: v_accvgpr_read_b32 v3, a31508; GISEL-NEXT: s_setpc_b64 s[30:31]1509 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 %scale0, i32 0, i32 %scale1)1510 ret <4 x float> %result1511}1512 1513define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgpr_vgpr_vgpr__sgpr_vgpr(<8 x i32> inreg %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 inreg %scale0, i32 %scale1) {1514; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgpr_vgpr_vgpr__sgpr_vgpr:1515; SDAG: ; %bb.0:1516; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1517; SDAG-NEXT: v_mov_b32_e32 v14, s01518; SDAG-NEXT: v_mov_b32_e32 v15, s11519; SDAG-NEXT: v_mov_b32_e32 v16, s21520; SDAG-NEXT: v_mov_b32_e32 v17, s31521; SDAG-NEXT: v_mov_b32_e32 v18, s161522; SDAG-NEXT: v_mov_b32_e32 v19, s171523; SDAG-NEXT: v_mov_b32_e32 v20, s181524; SDAG-NEXT: v_mov_b32_e32 v21, s191525; SDAG-NEXT: v_accvgpr_write_b32 a0, v81526; SDAG-NEXT: v_accvgpr_write_b32 a1, v91527; SDAG-NEXT: v_accvgpr_write_b32 a2, v101528; SDAG-NEXT: v_accvgpr_write_b32 a3, v111529; SDAG-NEXT: v_mov_b32_e32 v8, s201530; SDAG-NEXT: s_nop 11531; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[14:21], v[0:7], a[0:3], v8, v12 op_sel_hi:[0,0,0]1532; SDAG-NEXT: s_nop 111533; SDAG-NEXT: v_accvgpr_read_b32 v0, a01534; SDAG-NEXT: v_accvgpr_read_b32 v1, a11535; SDAG-NEXT: v_accvgpr_read_b32 v2, a21536; SDAG-NEXT: v_accvgpr_read_b32 v3, a31537; SDAG-NEXT: s_setpc_b64 s[30:31]1538;1539; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgpr_vgpr_vgpr__sgpr_vgpr:1540; GISEL: ; %bb.0:1541; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1542; GISEL-NEXT: s_mov_b32 s12, s01543; GISEL-NEXT: s_mov_b32 s13, s11544; GISEL-NEXT: s_mov_b32 s14, s21545; GISEL-NEXT: s_mov_b32 s15, s31546; GISEL-NEXT: v_mov_b64_e32 v[20:21], s[18:19]1547; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[16:17]1548; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[14:15]1549; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[12:13]1550; GISEL-NEXT: v_accvgpr_write_b32 a0, v81551; GISEL-NEXT: v_accvgpr_write_b32 a1, v91552; GISEL-NEXT: v_accvgpr_write_b32 a2, v101553; GISEL-NEXT: v_accvgpr_write_b32 a3, v111554; GISEL-NEXT: v_mov_b32_e32 v8, s201555; GISEL-NEXT: s_nop 11556; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[14:21], v[0:7], a[0:3], v8, v12 op_sel_hi:[0,0,0]1557; GISEL-NEXT: s_nop 111558; GISEL-NEXT: v_accvgpr_read_b32 v0, a01559; GISEL-NEXT: v_accvgpr_read_b32 v1, a11560; GISEL-NEXT: v_accvgpr_read_b32 v2, a21561; GISEL-NEXT: v_accvgpr_read_b32 v3, a31562; GISEL-NEXT: s_setpc_b64 s[30:31]1563 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 %scale0, i32 0, i32 %scale1)1564 ret <4 x float> %result1565}1566 1567define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgpr_vgpr_vgpr__vgpr_sgpr(<8 x i32> inreg %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 inreg %scale1) {1568; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgpr_vgpr_vgpr__vgpr_sgpr:1569; SDAG: ; %bb.0:1570; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1571; SDAG-NEXT: v_mov_b32_e32 v14, s01572; SDAG-NEXT: v_mov_b32_e32 v15, s11573; SDAG-NEXT: v_mov_b32_e32 v16, s21574; SDAG-NEXT: v_mov_b32_e32 v17, s31575; SDAG-NEXT: v_mov_b32_e32 v18, s161576; SDAG-NEXT: v_mov_b32_e32 v19, s171577; SDAG-NEXT: v_mov_b32_e32 v20, s181578; SDAG-NEXT: v_mov_b32_e32 v21, s191579; SDAG-NEXT: v_accvgpr_write_b32 a0, v81580; SDAG-NEXT: v_accvgpr_write_b32 a1, v91581; SDAG-NEXT: v_accvgpr_write_b32 a2, v101582; SDAG-NEXT: v_accvgpr_write_b32 a3, v111583; SDAG-NEXT: v_mov_b32_e32 v8, s201584; SDAG-NEXT: s_nop 11585; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[14:21], v[0:7], a[0:3], v12, v8 op_sel_hi:[0,0,0]1586; SDAG-NEXT: s_nop 111587; SDAG-NEXT: v_accvgpr_read_b32 v0, a01588; SDAG-NEXT: v_accvgpr_read_b32 v1, a11589; SDAG-NEXT: v_accvgpr_read_b32 v2, a21590; SDAG-NEXT: v_accvgpr_read_b32 v3, a31591; SDAG-NEXT: s_setpc_b64 s[30:31]1592;1593; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgpr_vgpr_vgpr__vgpr_sgpr:1594; GISEL: ; %bb.0:1595; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1596; GISEL-NEXT: s_mov_b32 s12, s01597; GISEL-NEXT: s_mov_b32 s13, s11598; GISEL-NEXT: s_mov_b32 s14, s21599; GISEL-NEXT: s_mov_b32 s15, s31600; GISEL-NEXT: v_mov_b64_e32 v[20:21], s[18:19]1601; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[16:17]1602; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[14:15]1603; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[12:13]1604; GISEL-NEXT: v_accvgpr_write_b32 a0, v81605; GISEL-NEXT: v_accvgpr_write_b32 a1, v91606; GISEL-NEXT: v_accvgpr_write_b32 a2, v101607; GISEL-NEXT: v_accvgpr_write_b32 a3, v111608; GISEL-NEXT: v_mov_b32_e32 v8, s201609; GISEL-NEXT: s_nop 11610; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[14:21], v[0:7], a[0:3], v12, v8 op_sel_hi:[0,0,0]1611; GISEL-NEXT: s_nop 111612; GISEL-NEXT: v_accvgpr_read_b32 v0, a01613; GISEL-NEXT: v_accvgpr_read_b32 v1, a11614; GISEL-NEXT: v_accvgpr_read_b32 v2, a21615; GISEL-NEXT: v_accvgpr_read_b32 v3, a31616; GISEL-NEXT: s_setpc_b64 s[30:31]1617 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 %scale0, i32 0, i32 %scale1)1618 ret <4 x float> %result1619}1620 1621define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0_vgpr_sgpr_vgpr__vgpr_sgpr(<8 x i32> %arg0, <8 x i32> inreg %arg1, <4 x float> %arg2, i32 %scale0, i32 inreg %scale1) {1622; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_vgpr_sgpr_vgpr__vgpr_sgpr:1623; SDAG: ; %bb.0:1624; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1625; SDAG-NEXT: v_mov_b32_e32 v14, s01626; SDAG-NEXT: v_mov_b32_e32 v15, s11627; SDAG-NEXT: v_mov_b32_e32 v16, s21628; SDAG-NEXT: v_mov_b32_e32 v17, s31629; SDAG-NEXT: v_mov_b32_e32 v18, s161630; SDAG-NEXT: v_mov_b32_e32 v19, s171631; SDAG-NEXT: v_mov_b32_e32 v20, s181632; SDAG-NEXT: v_mov_b32_e32 v21, s191633; SDAG-NEXT: v_accvgpr_write_b32 a0, v81634; SDAG-NEXT: v_accvgpr_write_b32 a1, v91635; SDAG-NEXT: v_accvgpr_write_b32 a2, v101636; SDAG-NEXT: v_accvgpr_write_b32 a3, v111637; SDAG-NEXT: v_mov_b32_e32 v8, s201638; SDAG-NEXT: s_nop 11639; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[14:21], a[0:3], v12, v8 op_sel_hi:[0,0,0]1640; SDAG-NEXT: s_nop 111641; SDAG-NEXT: v_accvgpr_read_b32 v0, a01642; SDAG-NEXT: v_accvgpr_read_b32 v1, a11643; SDAG-NEXT: v_accvgpr_read_b32 v2, a21644; SDAG-NEXT: v_accvgpr_read_b32 v3, a31645; SDAG-NEXT: s_setpc_b64 s[30:31]1646;1647; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_vgpr_sgpr_vgpr__vgpr_sgpr:1648; GISEL: ; %bb.0:1649; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1650; GISEL-NEXT: s_mov_b32 s12, s01651; GISEL-NEXT: s_mov_b32 s13, s11652; GISEL-NEXT: s_mov_b32 s14, s21653; GISEL-NEXT: s_mov_b32 s15, s31654; GISEL-NEXT: v_mov_b64_e32 v[20:21], s[18:19]1655; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[16:17]1656; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[14:15]1657; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[12:13]1658; GISEL-NEXT: v_accvgpr_write_b32 a0, v81659; GISEL-NEXT: v_accvgpr_write_b32 a1, v91660; GISEL-NEXT: v_accvgpr_write_b32 a2, v101661; GISEL-NEXT: v_accvgpr_write_b32 a3, v111662; GISEL-NEXT: v_mov_b32_e32 v8, s201663; GISEL-NEXT: s_nop 11664; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[14:21], a[0:3], v12, v8 op_sel_hi:[0,0,0]1665; GISEL-NEXT: s_nop 111666; GISEL-NEXT: v_accvgpr_read_b32 v0, a01667; GISEL-NEXT: v_accvgpr_read_b32 v1, a11668; GISEL-NEXT: v_accvgpr_read_b32 v2, a21669; GISEL-NEXT: v_accvgpr_read_b32 v3, a31670; GISEL-NEXT: s_setpc_b64 s[30:31]1671 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 %scale0, i32 0, i32 %scale1)1672 ret <4 x float> %result1673}1674 1675define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0_vgpr_vgpr_sgpr__vgpr_sgpr(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> inreg %arg2, i32 %scale0, i32 inreg %scale1) {1676; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_vgpr_vgpr_sgpr__vgpr_sgpr:1677; GCN: ; %bb.0:1678; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1679; GCN-NEXT: v_accvgpr_write_b32 a0, s01680; GCN-NEXT: v_accvgpr_write_b32 a1, s11681; GCN-NEXT: v_accvgpr_write_b32 a2, s21682; GCN-NEXT: v_accvgpr_write_b32 a3, s31683; GCN-NEXT: v_mov_b32_e32 v17, s161684; GCN-NEXT: s_nop 11685; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, v17 op_sel_hi:[0,0,0]1686; GCN-NEXT: s_nop 111687; GCN-NEXT: v_accvgpr_read_b32 v0, a01688; GCN-NEXT: v_accvgpr_read_b32 v1, a11689; GCN-NEXT: v_accvgpr_read_b32 v2, a21690; GCN-NEXT: v_accvgpr_read_b32 v3, a31691; GCN-NEXT: s_setpc_b64 s[30:31]1692 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 %scale0, i32 0, i32 %scale1)1693 ret <4 x float> %result1694}1695 1696define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgpr_vgpr_sgpr__vgpr_sgpr(<8 x i32> inreg %arg0, <8 x i32> %arg1, <4 x float> inreg %arg2, i32 %scale0, i32 inreg %scale1) {1697; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgpr_vgpr_sgpr__vgpr_sgpr:1698; SDAG: ; %bb.0:1699; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1700; SDAG-NEXT: v_mov_b32_e32 v10, s01701; SDAG-NEXT: v_mov_b32_e32 v11, s11702; SDAG-NEXT: v_mov_b32_e32 v12, s21703; SDAG-NEXT: v_mov_b32_e32 v13, s31704; SDAG-NEXT: v_mov_b32_e32 v14, s161705; SDAG-NEXT: v_mov_b32_e32 v15, s171706; SDAG-NEXT: v_mov_b32_e32 v16, s181707; SDAG-NEXT: v_mov_b32_e32 v17, s191708; SDAG-NEXT: v_accvgpr_write_b32 a0, s201709; SDAG-NEXT: v_accvgpr_write_b32 a1, s211710; SDAG-NEXT: v_accvgpr_write_b32 a2, s221711; SDAG-NEXT: v_accvgpr_write_b32 a3, s231712; SDAG-NEXT: v_mov_b32_e32 v9, s241713; SDAG-NEXT: s_nop 11714; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[10:17], v[0:7], a[0:3], v8, v9 op_sel_hi:[0,0,0]1715; SDAG-NEXT: s_nop 111716; SDAG-NEXT: v_accvgpr_read_b32 v0, a01717; SDAG-NEXT: v_accvgpr_read_b32 v1, a11718; SDAG-NEXT: v_accvgpr_read_b32 v2, a21719; SDAG-NEXT: v_accvgpr_read_b32 v3, a31720; SDAG-NEXT: s_setpc_b64 s[30:31]1721;1722; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0_sgpr_vgpr_sgpr__vgpr_sgpr:1723; GISEL: ; %bb.0:1724; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1725; GISEL-NEXT: s_mov_b32 s12, s01726; GISEL-NEXT: s_mov_b32 s13, s11727; GISEL-NEXT: s_mov_b32 s14, s21728; GISEL-NEXT: s_mov_b32 s15, s31729; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[12:13]1730; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[14:15]1731; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[16:17]1732; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[18:19]1733; GISEL-NEXT: v_accvgpr_write_b32 a0, s201734; GISEL-NEXT: v_accvgpr_write_b32 a1, s211735; GISEL-NEXT: v_accvgpr_write_b32 a2, s221736; GISEL-NEXT: v_accvgpr_write_b32 a3, s231737; GISEL-NEXT: v_mov_b32_e32 v9, s241738; GISEL-NEXT: s_nop 11739; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[10:17], v[0:7], a[0:3], v8, v9 op_sel_hi:[0,0,0]1740; GISEL-NEXT: s_nop 111741; GISEL-NEXT: v_accvgpr_read_b32 v0, a01742; GISEL-NEXT: v_accvgpr_read_b32 v1, a11743; GISEL-NEXT: v_accvgpr_read_b32 v2, a21744; GISEL-NEXT: v_accvgpr_read_b32 v3, a31745; GISEL-NEXT: s_setpc_b64 s[30:31]1746 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 %scale0, i32 0, i32 %scale1)1747 ret <4 x float> %result1748}1749 1750define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__scaleA_inlineimm__scaleB_inlineimm(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1751; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__scaleA_inlineimm__scaleB_inlineimm:1752; SDAG: ; %bb.0:1753; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1754; SDAG-NEXT: v_mov_b32_e32 v20, -21755; SDAG-NEXT: v_mov_b32_e32 v21, 331756; SDAG-NEXT: v_accvgpr_write_b32 a0, v161757; SDAG-NEXT: v_accvgpr_write_b32 a1, v171758; SDAG-NEXT: v_accvgpr_write_b32 a2, v181759; SDAG-NEXT: v_accvgpr_write_b32 a3, v191760; SDAG-NEXT: s_nop 11761; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v21, v20 op_sel_hi:[1,1,0]1762; SDAG-NEXT: s_nop 111763; SDAG-NEXT: v_accvgpr_read_b32 v0, a01764; SDAG-NEXT: v_accvgpr_read_b32 v1, a11765; SDAG-NEXT: v_accvgpr_read_b32 v2, a21766; SDAG-NEXT: v_accvgpr_read_b32 v3, a31767; SDAG-NEXT: s_setpc_b64 s[30:31]1768;1769; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__scaleA_inlineimm__scaleB_inlineimm:1770; GISEL: ; %bb.0:1771; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1772; GISEL-NEXT: v_accvgpr_write_b32 a0, v161773; GISEL-NEXT: v_accvgpr_write_b32 a1, v171774; GISEL-NEXT: v_accvgpr_write_b32 a2, v181775; GISEL-NEXT: v_accvgpr_write_b32 a3, v191776; GISEL-NEXT: v_mov_b32_e32 v16, 331777; GISEL-NEXT: v_mov_b32_e32 v17, -21778; GISEL-NEXT: s_nop 11779; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, v17 op_sel_hi:[1,1,0]1780; GISEL-NEXT: s_nop 111781; GISEL-NEXT: v_accvgpr_read_b32 v0, a01782; GISEL-NEXT: v_accvgpr_read_b32 v1, a11783; GISEL-NEXT: v_accvgpr_read_b32 v2, a21784; GISEL-NEXT: v_accvgpr_read_b32 v3, a31785; GISEL-NEXT: s_setpc_b64 s[30:31]1786 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 2, i32 33, i32 2, i32 -2)1787 ret <4 x float> %result1788}1789 1790define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__scaleA_kimm__scaleB_inlineimm(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1791; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__scaleA_kimm__scaleB_inlineimm:1792; SDAG: ; %bb.0:1793; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1794; SDAG-NEXT: v_mov_b32_e32 v20, -21795; SDAG-NEXT: v_mov_b32_e32 v21, 0x411796; SDAG-NEXT: v_accvgpr_write_b32 a0, v161797; SDAG-NEXT: v_accvgpr_write_b32 a1, v171798; SDAG-NEXT: v_accvgpr_write_b32 a2, v181799; SDAG-NEXT: v_accvgpr_write_b32 a3, v191800; SDAG-NEXT: s_nop 11801; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v21, v20 op_sel_hi:[1,1,0]1802; SDAG-NEXT: s_nop 111803; SDAG-NEXT: v_accvgpr_read_b32 v0, a01804; SDAG-NEXT: v_accvgpr_read_b32 v1, a11805; SDAG-NEXT: v_accvgpr_read_b32 v2, a21806; SDAG-NEXT: v_accvgpr_read_b32 v3, a31807; SDAG-NEXT: s_setpc_b64 s[30:31]1808;1809; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__scaleA_kimm__scaleB_inlineimm:1810; GISEL: ; %bb.0:1811; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1812; GISEL-NEXT: v_accvgpr_write_b32 a0, v161813; GISEL-NEXT: v_accvgpr_write_b32 a1, v171814; GISEL-NEXT: v_accvgpr_write_b32 a2, v181815; GISEL-NEXT: v_accvgpr_write_b32 a3, v191816; GISEL-NEXT: v_mov_b32_e32 v16, 0x411817; GISEL-NEXT: v_mov_b32_e32 v17, -21818; GISEL-NEXT: s_nop 11819; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, v17 op_sel_hi:[1,1,0]1820; GISEL-NEXT: s_nop 111821; GISEL-NEXT: v_accvgpr_read_b32 v0, a01822; GISEL-NEXT: v_accvgpr_read_b32 v1, a11823; GISEL-NEXT: v_accvgpr_read_b32 v2, a21824; GISEL-NEXT: v_accvgpr_read_b32 v3, a31825; GISEL-NEXT: s_setpc_b64 s[30:31]1826 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 2, i32 65, i32 2, i32 -2)1827 ret <4 x float> %result1828}1829 1830define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4_0_0__scaleA_kimm__scaleB_kimm(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {1831; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__scaleA_kimm__scaleB_kimm:1832; SDAG: ; %bb.0:1833; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1834; SDAG-NEXT: v_mov_b32_e32 v20, 0x4d1835; SDAG-NEXT: v_mov_b32_e32 v21, 0x411836; SDAG-NEXT: v_accvgpr_write_b32 a0, v161837; SDAG-NEXT: v_accvgpr_write_b32 a1, v171838; SDAG-NEXT: v_accvgpr_write_b32 a2, v181839; SDAG-NEXT: v_accvgpr_write_b32 a3, v191840; SDAG-NEXT: s_nop 11841; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v21, v20 op_sel_hi:[1,1,0]1842; SDAG-NEXT: s_nop 111843; SDAG-NEXT: v_accvgpr_read_b32 v0, a01844; SDAG-NEXT: v_accvgpr_read_b32 v1, a11845; SDAG-NEXT: v_accvgpr_read_b32 v2, a21846; SDAG-NEXT: v_accvgpr_read_b32 v3, a31847; SDAG-NEXT: s_setpc_b64 s[30:31]1848;1849; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4_0_0__scaleA_kimm__scaleB_kimm:1850; GISEL: ; %bb.0:1851; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1852; GISEL-NEXT: v_accvgpr_write_b32 a0, v161853; GISEL-NEXT: v_accvgpr_write_b32 a1, v171854; GISEL-NEXT: v_accvgpr_write_b32 a2, v181855; GISEL-NEXT: v_accvgpr_write_b32 a3, v191856; GISEL-NEXT: v_mov_b32_e32 v16, 0x411857; GISEL-NEXT: v_mov_b32_e32 v17, 0x4d1858; GISEL-NEXT: s_nop 11859; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, v17 op_sel_hi:[1,1,0]1860; GISEL-NEXT: s_nop 111861; GISEL-NEXT: v_accvgpr_read_b32 v0, a01862; GISEL-NEXT: v_accvgpr_read_b32 v1, a11863; GISEL-NEXT: v_accvgpr_read_b32 v2, a21864; GISEL-NEXT: v_accvgpr_read_b32 v3, a31865; GISEL-NEXT: s_setpc_b64 s[30:31]1866 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 2, i32 65, i32 2, i32 77)1867 ret <4 x float> %result1868}1869 1870define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1, ptr addrspace(1) %ptr) #0 {1871; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd:1872; SDAG: ; %bb.0:1873; SDAG-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x01874; SDAG-NEXT: v_mov_b32_e32 v20, 01875; SDAG-NEXT: s_waitcnt lgkmcnt(0)1876; SDAG-NEXT: v_mov_b32_e32 v0, s81877; SDAG-NEXT: v_mov_b32_e32 v1, s91878; SDAG-NEXT: v_mov_b32_e32 v2, s101879; SDAG-NEXT: v_mov_b32_e32 v3, s111880; SDAG-NEXT: v_mov_b32_e32 v4, s121881; SDAG-NEXT: v_mov_b32_e32 v5, s131882; SDAG-NEXT: v_mov_b32_e32 v6, s141883; SDAG-NEXT: v_mov_b32_e32 v7, s151884; SDAG-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x401885; SDAG-NEXT: v_mov_b32_e32 v8, s161886; SDAG-NEXT: v_mov_b32_e32 v9, s171887; SDAG-NEXT: v_mov_b32_e32 v10, s181888; SDAG-NEXT: v_mov_b32_e32 v11, s191889; SDAG-NEXT: v_mov_b32_e32 v12, s201890; SDAG-NEXT: v_mov_b32_e32 v13, s211891; SDAG-NEXT: v_mov_b32_e32 v14, s221892; SDAG-NEXT: v_mov_b32_e32 v15, s231893; SDAG-NEXT: s_waitcnt lgkmcnt(0)1894; SDAG-NEXT: v_mov_b32_e32 v16, s81895; SDAG-NEXT: v_mov_b32_e32 v17, s91896; SDAG-NEXT: v_mov_b32_e32 v18, s101897; SDAG-NEXT: v_mov_b32_e32 v19, s111898; SDAG-NEXT: v_mov_b32_e32 v21, s121899; SDAG-NEXT: v_mov_b32_e32 v22, s131900; SDAG-NEXT: s_nop 11901; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v21, v22 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:21902; SDAG-NEXT: s_nop 111903; SDAG-NEXT: global_store_dwordx4 v20, v[0:3], s[14:15]1904; SDAG-NEXT: s_endpgm1905;1906; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd:1907; GISEL: ; %bb.0:1908; GISEL-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x01909; GISEL-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x401910; GISEL-NEXT: s_waitcnt lgkmcnt(0)1911; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]1912; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]1913; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]1914; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]1915; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[16:17]1916; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[24:25]1917; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[18:19]1918; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[20:21]1919; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[22:23]1920; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[26:27]1921; GISEL-NEXT: v_mov_b32_e32 v20, s281922; GISEL-NEXT: v_mov_b32_e32 v21, s291923; GISEL-NEXT: s_nop 11924; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v20, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0] blgp:21925; GISEL-NEXT: v_mov_b32_e32 v4, 01926; GISEL-NEXT: s_nop 101927; GISEL-NEXT: global_store_dwordx4 v4, v[0:3], s[30:31]1928; GISEL-NEXT: s_endpgm1929 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 2, i32 3, i32 %scale0, i32 1, i32 %scale1)1930 store <4 x float> %result, ptr addrspace(1) %ptr, align 161931 ret void1932}1933 1934define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_kimm__scaleB__inlineimm(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, ptr addrspace(1) %ptr) #0 {1935; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_kimm__scaleB__inlineimm:1936; SDAG: ; %bb.0:1937; SDAG-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x01938; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x401939; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x501940; SDAG-NEXT: v_mov_b32_e32 v21, -21941; SDAG-NEXT: v_mov_b32_e32 v22, 0x411942; SDAG-NEXT: v_mov_b32_e32 v20, 01943; SDAG-NEXT: s_waitcnt lgkmcnt(0)1944; SDAG-NEXT: v_mov_b32_e32 v0, s81945; SDAG-NEXT: v_mov_b32_e32 v1, s91946; SDAG-NEXT: v_mov_b32_e32 v2, s101947; SDAG-NEXT: v_mov_b32_e32 v3, s111948; SDAG-NEXT: v_mov_b32_e32 v4, s121949; SDAG-NEXT: v_mov_b32_e32 v5, s131950; SDAG-NEXT: v_mov_b32_e32 v6, s141951; SDAG-NEXT: v_mov_b32_e32 v7, s151952; SDAG-NEXT: v_mov_b64_e32 v[18:19], s[2:3]1953; SDAG-NEXT: v_mov_b32_e32 v8, s161954; SDAG-NEXT: v_mov_b32_e32 v9, s171955; SDAG-NEXT: v_mov_b32_e32 v10, s181956; SDAG-NEXT: v_mov_b32_e32 v11, s191957; SDAG-NEXT: v_mov_b32_e32 v12, s201958; SDAG-NEXT: v_mov_b32_e32 v13, s211959; SDAG-NEXT: v_mov_b32_e32 v14, s221960; SDAG-NEXT: v_mov_b32_e32 v15, s231961; SDAG-NEXT: v_mov_b64_e32 v[16:17], s[0:1]1962; SDAG-NEXT: s_nop 11963; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v22, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0]1964; SDAG-NEXT: s_nop 111965; SDAG-NEXT: global_store_dwordx4 v20, v[0:3], s[6:7]1966; SDAG-NEXT: s_endpgm1967;1968; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_kimm__scaleB__inlineimm:1969; GISEL: ; %bb.0:1970; GISEL-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x01971; GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x401972; GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x501973; GISEL-NEXT: v_mov_b32_e32 v20, 0x411974; GISEL-NEXT: v_mov_b32_e32 v21, -21975; GISEL-NEXT: s_waitcnt lgkmcnt(0)1976; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]1977; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]1978; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]1979; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]1980; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[16:17]1981; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[2:3]1982; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[18:19]1983; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[20:21]1984; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[22:23]1985; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[0:1]1986; GISEL-NEXT: s_nop 11987; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v20, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0]1988; GISEL-NEXT: v_mov_b32_e32 v4, 01989; GISEL-NEXT: s_nop 101990; GISEL-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]1991; GISEL-NEXT: s_endpgm1992 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 3, i32 65, i32 1, i32 -2)1993 store <4 x float> %result, ptr addrspace(1) %ptr, align 161994 ret void1995}1996 1997define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_kimm__scaleB__FP_literal(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, ptr addrspace(1) %ptr) #0 {1998; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_kimm__scaleB__FP_literal:1999; SDAG: ; %bb.0:2000; SDAG-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x02001; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x402002; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x502003; SDAG-NEXT: v_mov_b32_e32 v21, 1.02004; SDAG-NEXT: v_mov_b32_e32 v22, 0x412005; SDAG-NEXT: v_mov_b32_e32 v20, 02006; SDAG-NEXT: s_waitcnt lgkmcnt(0)2007; SDAG-NEXT: v_mov_b32_e32 v0, s82008; SDAG-NEXT: v_mov_b32_e32 v1, s92009; SDAG-NEXT: v_mov_b32_e32 v2, s102010; SDAG-NEXT: v_mov_b32_e32 v3, s112011; SDAG-NEXT: v_mov_b32_e32 v4, s122012; SDAG-NEXT: v_mov_b32_e32 v5, s132013; SDAG-NEXT: v_mov_b32_e32 v6, s142014; SDAG-NEXT: v_mov_b32_e32 v7, s152015; SDAG-NEXT: v_mov_b64_e32 v[18:19], s[2:3]2016; SDAG-NEXT: v_mov_b32_e32 v8, s162017; SDAG-NEXT: v_mov_b32_e32 v9, s172018; SDAG-NEXT: v_mov_b32_e32 v10, s182019; SDAG-NEXT: v_mov_b32_e32 v11, s192020; SDAG-NEXT: v_mov_b32_e32 v12, s202021; SDAG-NEXT: v_mov_b32_e32 v13, s212022; SDAG-NEXT: v_mov_b32_e32 v14, s222023; SDAG-NEXT: v_mov_b32_e32 v15, s232024; SDAG-NEXT: v_mov_b64_e32 v[16:17], s[0:1]2025; SDAG-NEXT: s_nop 12026; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v22, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0]2027; SDAG-NEXT: s_nop 112028; SDAG-NEXT: global_store_dwordx4 v20, v[0:3], s[6:7]2029; SDAG-NEXT: s_endpgm2030;2031; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_kimm__scaleB__FP_literal:2032; GISEL: ; %bb.0:2033; GISEL-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x02034; GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x402035; GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x502036; GISEL-NEXT: v_mov_b32_e32 v20, 0x412037; GISEL-NEXT: v_mov_b32_e32 v21, 1.02038; GISEL-NEXT: s_waitcnt lgkmcnt(0)2039; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]2040; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]2041; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]2042; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]2043; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[16:17]2044; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[2:3]2045; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[18:19]2046; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[20:21]2047; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[22:23]2048; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[0:1]2049; GISEL-NEXT: s_nop 12050; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v20, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0]2051; GISEL-NEXT: v_mov_b32_e32 v4, 02052; GISEL-NEXT: s_nop 102053; GISEL-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]2054; GISEL-NEXT: s_endpgm2055 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 3, i32 65, i32 1, i32 1065353216)2056 store <4 x float> %result, ptr addrspace(1) %ptr, align 162057 ret void2058}2059 2060define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_FP_literal__scaleB__inline_imm(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, ptr addrspace(1) %ptr) #0 {2061; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_FP_literal__scaleB__inline_imm:2062; SDAG: ; %bb.0:2063; SDAG-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x02064; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x402065; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x502066; SDAG-NEXT: v_mov_b32_e32 v21, -22067; SDAG-NEXT: v_mov_b32_e32 v22, 1.02068; SDAG-NEXT: v_mov_b32_e32 v20, 02069; SDAG-NEXT: s_waitcnt lgkmcnt(0)2070; SDAG-NEXT: v_mov_b32_e32 v0, s82071; SDAG-NEXT: v_mov_b32_e32 v1, s92072; SDAG-NEXT: v_mov_b32_e32 v2, s102073; SDAG-NEXT: v_mov_b32_e32 v3, s112074; SDAG-NEXT: v_mov_b32_e32 v4, s122075; SDAG-NEXT: v_mov_b32_e32 v5, s132076; SDAG-NEXT: v_mov_b32_e32 v6, s142077; SDAG-NEXT: v_mov_b32_e32 v7, s152078; SDAG-NEXT: v_mov_b64_e32 v[18:19], s[2:3]2079; SDAG-NEXT: v_mov_b32_e32 v8, s162080; SDAG-NEXT: v_mov_b32_e32 v9, s172081; SDAG-NEXT: v_mov_b32_e32 v10, s182082; SDAG-NEXT: v_mov_b32_e32 v11, s192083; SDAG-NEXT: v_mov_b32_e32 v12, s202084; SDAG-NEXT: v_mov_b32_e32 v13, s212085; SDAG-NEXT: v_mov_b32_e32 v14, s222086; SDAG-NEXT: v_mov_b32_e32 v15, s232087; SDAG-NEXT: v_mov_b64_e32 v[16:17], s[0:1]2088; SDAG-NEXT: s_nop 12089; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v22, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0]2090; SDAG-NEXT: s_nop 112091; SDAG-NEXT: global_store_dwordx4 v20, v[0:3], s[6:7]2092; SDAG-NEXT: s_endpgm2093;2094; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_FP_literal__scaleB__inline_imm:2095; GISEL: ; %bb.0:2096; GISEL-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x02097; GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x402098; GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x502099; GISEL-NEXT: v_mov_b32_e32 v20, 1.02100; GISEL-NEXT: v_mov_b32_e32 v21, -22101; GISEL-NEXT: s_waitcnt lgkmcnt(0)2102; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]2103; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]2104; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]2105; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]2106; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[16:17]2107; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[2:3]2108; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[18:19]2109; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[20:21]2110; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[22:23]2111; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[0:1]2112; GISEL-NEXT: s_nop 12113; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v20, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0]2114; GISEL-NEXT: v_mov_b32_e32 v4, 02115; GISEL-NEXT: s_nop 102116; GISEL-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]2117; GISEL-NEXT: s_endpgm2118 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 3, i32 1065353216, i32 1, i32 -2)2119 store <4 x float> %result, ptr addrspace(1) %ptr, align 162120 ret void2121}2122 2123define amdgpu_kernel void @test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_FP_literal__scaleB__FP_literal(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, ptr addrspace(1) %ptr) #0 {2124; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_FP_literal__scaleB__FP_literal:2125; SDAG: ; %bb.0:2126; SDAG-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x02127; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x402128; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x502129; SDAG-NEXT: v_mov_b32_e32 v21, 0.159154942130; SDAG-NEXT: v_mov_b32_e32 v22, 1.02131; SDAG-NEXT: v_mov_b32_e32 v20, 02132; SDAG-NEXT: s_waitcnt lgkmcnt(0)2133; SDAG-NEXT: v_mov_b32_e32 v0, s82134; SDAG-NEXT: v_mov_b32_e32 v1, s92135; SDAG-NEXT: v_mov_b32_e32 v2, s102136; SDAG-NEXT: v_mov_b32_e32 v3, s112137; SDAG-NEXT: v_mov_b32_e32 v4, s122138; SDAG-NEXT: v_mov_b32_e32 v5, s132139; SDAG-NEXT: v_mov_b32_e32 v6, s142140; SDAG-NEXT: v_mov_b32_e32 v7, s152141; SDAG-NEXT: v_mov_b64_e32 v[18:19], s[2:3]2142; SDAG-NEXT: v_mov_b32_e32 v8, s162143; SDAG-NEXT: v_mov_b32_e32 v9, s172144; SDAG-NEXT: v_mov_b32_e32 v10, s182145; SDAG-NEXT: v_mov_b32_e32 v11, s192146; SDAG-NEXT: v_mov_b32_e32 v12, s202147; SDAG-NEXT: v_mov_b32_e32 v13, s212148; SDAG-NEXT: v_mov_b32_e32 v14, s222149; SDAG-NEXT: v_mov_b32_e32 v15, s232150; SDAG-NEXT: v_mov_b64_e32 v[16:17], s[0:1]2151; SDAG-NEXT: s_nop 12152; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v22, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0]2153; SDAG-NEXT: s_nop 112154; SDAG-NEXT: global_store_dwordx4 v20, v[0:3], s[6:7]2155; SDAG-NEXT: s_endpgm2156;2157; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4__vgprcd___scaleA_FP_literal__scaleB__FP_literal:2158; GISEL: ; %bb.0:2159; GISEL-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x02160; GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x402161; GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x502162; GISEL-NEXT: v_mov_b32_e32 v20, 1.02163; GISEL-NEXT: v_mov_b32_e32 v21, 0.159154942164; GISEL-NEXT: s_waitcnt lgkmcnt(0)2165; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]2166; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]2167; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]2168; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]2169; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[16:17]2170; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[2:3]2171; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[18:19]2172; GISEL-NEXT: v_mov_b64_e32 v[12:13], s[20:21]2173; GISEL-NEXT: v_mov_b64_e32 v[14:15], s[22:23]2174; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[0:1]2175; GISEL-NEXT: s_nop 12176; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[0:7], v[8:15], v[16:19], v20, v21 op_sel:[1,1,0] op_sel_hi:[1,0,0]2177; GISEL-NEXT: v_mov_b32_e32 v4, 02178; GISEL-NEXT: s_nop 102179; GISEL-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]2180; GISEL-NEXT: s_endpgm2181 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 3, i32 1065353216, i32 1, i32 1042479491)2182 store <4 x float> %result, ptr addrspace(1) %ptr, align 162183 ret void2184}2185 2186; This should be optimized to avoid the scale2187define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_0_0_a(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2188; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_0_0_a:2189; GCN: ; %bb.0:2190; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2191; GCN-NEXT: v_accvgpr_write_b32 a0, v162192; GCN-NEXT: v_accvgpr_write_b32 a1, v172193; GCN-NEXT: v_accvgpr_write_b32 a2, v182194; GCN-NEXT: v_accvgpr_write_b32 a3, v192195; GCN-NEXT: s_nop 12196; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3]2197; GCN-NEXT: s_nop 112198; GCN-NEXT: v_accvgpr_read_b32 v0, a02199; GCN-NEXT: v_accvgpr_read_b32 v1, a12200; GCN-NEXT: v_accvgpr_read_b32 v2, a22201; GCN-NEXT: v_accvgpr_read_b32 v3, a32202; GCN-NEXT: s_setpc_b64 s[30:31]2203 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0)2204 ret <4 x float> %result2205}2206 2207; This should be optimized to avoid the scale, with non-0 op_sel arguments.2208define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_0_0_b(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2209; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_0_0_b:2210; GCN: ; %bb.0:2211; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2212; GCN-NEXT: v_accvgpr_write_b32 a0, v162213; GCN-NEXT: v_accvgpr_write_b32 a1, v172214; GCN-NEXT: v_accvgpr_write_b32 a2, v182215; GCN-NEXT: v_accvgpr_write_b32 a3, v192216; GCN-NEXT: s_nop 12217; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3]2218; GCN-NEXT: s_nop 112219; GCN-NEXT: v_accvgpr_read_b32 v0, a02220; GCN-NEXT: v_accvgpr_read_b32 v1, a12221; GCN-NEXT: v_accvgpr_read_b32 v2, a22222; GCN-NEXT: v_accvgpr_read_b32 v3, a32223; GCN-NEXT: s_setpc_b64 s[30:31]2224 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 3, i32 0, i32 1, i32 0)2225 ret <4 x float> %result2226}2227 2228define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_0_1(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2229; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_0_1:2230; SDAG: ; %bb.0:2231; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2232; SDAG-NEXT: v_mov_b32_e32 v20, 12233; SDAG-NEXT: v_mov_b32_e32 v21, 02234; SDAG-NEXT: v_accvgpr_write_b32 a0, v162235; SDAG-NEXT: v_accvgpr_write_b32 a1, v172236; SDAG-NEXT: v_accvgpr_write_b32 a2, v182237; SDAG-NEXT: v_accvgpr_write_b32 a3, v192238; SDAG-NEXT: s_nop 12239; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v21, v20 op_sel_hi:[0,0,0]2240; SDAG-NEXT: s_nop 112241; SDAG-NEXT: v_accvgpr_read_b32 v0, a02242; SDAG-NEXT: v_accvgpr_read_b32 v1, a12243; SDAG-NEXT: v_accvgpr_read_b32 v2, a22244; SDAG-NEXT: v_accvgpr_read_b32 v3, a32245; SDAG-NEXT: s_setpc_b64 s[30:31]2246;2247; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_0_1:2248; GISEL: ; %bb.0:2249; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2250; GISEL-NEXT: v_accvgpr_write_b32 a0, v162251; GISEL-NEXT: v_accvgpr_write_b32 a1, v172252; GISEL-NEXT: v_accvgpr_write_b32 a2, v182253; GISEL-NEXT: v_accvgpr_write_b32 a3, v192254; GISEL-NEXT: v_mov_b32_e32 v16, 02255; GISEL-NEXT: v_mov_b32_e32 v17, 12256; GISEL-NEXT: s_nop 12257; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, v17 op_sel_hi:[0,0,0]2258; GISEL-NEXT: s_nop 112259; GISEL-NEXT: v_accvgpr_read_b32 v0, a02260; GISEL-NEXT: v_accvgpr_read_b32 v1, a12261; GISEL-NEXT: v_accvgpr_read_b32 v2, a22262; GISEL-NEXT: v_accvgpr_read_b32 v3, a32263; GISEL-NEXT: s_setpc_b64 s[30:31]2264 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 0, i32 0, i32 1)2265 ret <4 x float> %result2266}2267 2268define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_1_0_a(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2269; SDAG-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_1_0_a:2270; SDAG: ; %bb.0:2271; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2272; SDAG-NEXT: v_mov_b32_e32 v20, 02273; SDAG-NEXT: v_mov_b32_e32 v21, 12274; SDAG-NEXT: v_accvgpr_write_b32 a0, v162275; SDAG-NEXT: v_accvgpr_write_b32 a1, v172276; SDAG-NEXT: v_accvgpr_write_b32 a2, v182277; SDAG-NEXT: v_accvgpr_write_b32 a3, v192278; SDAG-NEXT: s_nop 12279; SDAG-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v21, v20 op_sel_hi:[0,0,0]2280; SDAG-NEXT: s_nop 112281; SDAG-NEXT: v_accvgpr_read_b32 v0, a02282; SDAG-NEXT: v_accvgpr_read_b32 v1, a12283; SDAG-NEXT: v_accvgpr_read_b32 v2, a22284; SDAG-NEXT: v_accvgpr_read_b32 v3, a32285; SDAG-NEXT: s_setpc_b64 s[30:31]2286;2287; GISEL-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___constant_scale_1_0_a:2288; GISEL: ; %bb.0:2289; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2290; GISEL-NEXT: v_accvgpr_write_b32 a0, v162291; GISEL-NEXT: v_accvgpr_write_b32 a1, v172292; GISEL-NEXT: v_accvgpr_write_b32 a2, v182293; GISEL-NEXT: v_accvgpr_write_b32 a3, v192294; GISEL-NEXT: v_mov_b32_e32 v16, 12295; GISEL-NEXT: v_mov_b32_e32 v17, 02296; GISEL-NEXT: s_nop 12297; GISEL-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v16, v17 op_sel_hi:[0,0,0]2298; GISEL-NEXT: s_nop 112299; GISEL-NEXT: v_accvgpr_read_b32 v0, a02300; GISEL-NEXT: v_accvgpr_read_b32 v1, a12301; GISEL-NEXT: v_accvgpr_read_b32 v2, a22302; GISEL-NEXT: v_accvgpr_read_b32 v3, a32303; GISEL-NEXT: s_setpc_b64 s[30:31]2304 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0, i32 1, i32 0, i32 0)2305 ret <4 x float> %result2306}2307 2308; --------------------------------------------------------------------2309; Incorrect signature for format cases (IR vector too large)2310; --------------------------------------------------------------------2311 2312define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp8__v8i32_fp6(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2313; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp8__v8i32_fp6:2314; GCN: ; %bb.0:2315; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2316; GCN-NEXT: v_accvgpr_write_b32 a0, v162317; GCN-NEXT: v_accvgpr_write_b32 a1, v172318; GCN-NEXT: v_accvgpr_write_b32 a2, v182319; GCN-NEXT: v_accvgpr_write_b32 a3, v192320; GCN-NEXT: s_nop 12321; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] blgp:22322; GCN-NEXT: s_nop 112323; GCN-NEXT: v_accvgpr_read_b32 v0, a02324; GCN-NEXT: v_accvgpr_read_b32 v1, a12325; GCN-NEXT: v_accvgpr_read_b32 v2, a22326; GCN-NEXT: v_accvgpr_read_b32 v3, a32327; GCN-NEXT: s_setpc_b64 s[30:31]2328 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,2329 i32 0, ; cbsz2330 i32 2, ; blgp2331 i32 0, i32 %scale0, i32 0, i32 %scale1)2332 ret <4 x float> %result2333}2334 2335define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp6__v8i32_fp8(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2336; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp6__v8i32_fp8:2337; GCN: ; %bb.0:2338; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2339; GCN-NEXT: v_accvgpr_write_b32 a0, v162340; GCN-NEXT: v_accvgpr_write_b32 a1, v172341; GCN-NEXT: v_accvgpr_write_b32 a2, v182342; GCN-NEXT: v_accvgpr_write_b32 a3, v192343; GCN-NEXT: s_nop 12344; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] cbsz:22345; GCN-NEXT: s_nop 112346; GCN-NEXT: v_accvgpr_read_b32 v0, a02347; GCN-NEXT: v_accvgpr_read_b32 v1, a12348; GCN-NEXT: v_accvgpr_read_b32 v2, a22349; GCN-NEXT: v_accvgpr_read_b32 v3, a32350; GCN-NEXT: s_setpc_b64 s[30:31]2351 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,2352 i32 2, ; cbsz2353 i32 0, ; blgp2354 i32 0, i32 %scale0, i32 0, i32 %scale1)2355 ret <4 x float> %result2356}2357 2358define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp6__v8i32_fp6(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2359; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp6__v8i32_fp6:2360; GCN: ; %bb.0:2361; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2362; GCN-NEXT: v_accvgpr_write_b32 a0, v162363; GCN-NEXT: v_accvgpr_write_b32 a1, v172364; GCN-NEXT: v_accvgpr_write_b32 a2, v182365; GCN-NEXT: v_accvgpr_write_b32 a3, v192366; GCN-NEXT: s_nop 12367; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] cbsz:2 blgp:22368; GCN-NEXT: s_nop 72369; GCN-NEXT: v_accvgpr_read_b32 v0, a02370; GCN-NEXT: v_accvgpr_read_b32 v1, a12371; GCN-NEXT: v_accvgpr_read_b32 v2, a22372; GCN-NEXT: v_accvgpr_read_b32 v3, a32373; GCN-NEXT: s_setpc_b64 s[30:31]2374 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,2375 i32 2, ; cbsz2376 i32 2, ; blgp2377 i32 0, i32 %scale0, i32 0, i32 %scale1)2378 ret <4 x float> %result2379}2380 2381define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp6__v8i32_fp6__0_scale(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {2382; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp6__v8i32_fp6__0_scale:2383; GCN: ; %bb.0:2384; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2385; GCN-NEXT: v_accvgpr_write_b32 a0, v162386; GCN-NEXT: v_accvgpr_write_b32 a1, v172387; GCN-NEXT: v_accvgpr_write_b32 a2, v182388; GCN-NEXT: v_accvgpr_write_b32 a3, v192389; GCN-NEXT: s_nop 12390; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3] cbsz:2 blgp:22391; GCN-NEXT: s_nop 72392; GCN-NEXT: v_accvgpr_read_b32 v0, a02393; GCN-NEXT: v_accvgpr_read_b32 v1, a12394; GCN-NEXT: v_accvgpr_read_b32 v2, a22395; GCN-NEXT: v_accvgpr_read_b32 v3, a32396; GCN-NEXT: s_setpc_b64 s[30:31]2397 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,2398 i32 2, ; cbsz2399 i32 2, ; blgp2400 i32 0, i32 0, i32 0, i32 0)2401 ret <4 x float> %result2402}2403 2404define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp8__v8i32_fp4(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2405; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp8__v8i32_fp4:2406; GCN: ; %bb.0:2407; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2408; GCN-NEXT: v_accvgpr_write_b32 a0, v162409; GCN-NEXT: v_accvgpr_write_b32 a1, v172410; GCN-NEXT: v_accvgpr_write_b32 a2, v182411; GCN-NEXT: v_accvgpr_write_b32 a3, v192412; GCN-NEXT: s_nop 12413; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] blgp:42414; GCN-NEXT: s_nop 112415; GCN-NEXT: v_accvgpr_read_b32 v0, a02416; GCN-NEXT: v_accvgpr_read_b32 v1, a12417; GCN-NEXT: v_accvgpr_read_b32 v2, a22418; GCN-NEXT: v_accvgpr_read_b32 v3, a32419; GCN-NEXT: s_setpc_b64 s[30:31]2420 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,2421 i32 0, ; cbsz2422 i32 4, ; blgp2423 i32 0, i32 %scale0, i32 0, i32 %scale1)2424 ret <4 x float> %result2425}2426 2427define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp4__v8i32_fp8(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2428; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp4__v8i32_fp8:2429; GCN: ; %bb.0:2430; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2431; GCN-NEXT: v_accvgpr_write_b32 a0, v162432; GCN-NEXT: v_accvgpr_write_b32 a1, v172433; GCN-NEXT: v_accvgpr_write_b32 a2, v182434; GCN-NEXT: v_accvgpr_write_b32 a3, v192435; GCN-NEXT: s_nop 12436; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] cbsz:42437; GCN-NEXT: s_nop 112438; GCN-NEXT: v_accvgpr_read_b32 v0, a02439; GCN-NEXT: v_accvgpr_read_b32 v1, a12440; GCN-NEXT: v_accvgpr_read_b32 v2, a22441; GCN-NEXT: v_accvgpr_read_b32 v3, a32442; GCN-NEXT: s_setpc_b64 s[30:31]2443 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,2444 i32 4, ; cbsz2445 i32 0, ; blgp2446 i32 0, i32 %scale0, i32 0, i32 %scale1)2447 ret <4 x float> %result2448}2449 2450define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp8__v6i32_fp4(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2451; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp8__v6i32_fp4:2452; GCN: ; %bb.0:2453; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2454; GCN-NEXT: v_accvgpr_write_b32 a0, v142455; GCN-NEXT: v_accvgpr_write_b32 a1, v152456; GCN-NEXT: v_accvgpr_write_b32 a2, v162457; GCN-NEXT: v_accvgpr_write_b32 a3, v172458; GCN-NEXT: s_nop 12459; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] blgp:42460; GCN-NEXT: s_nop 112461; GCN-NEXT: v_accvgpr_read_b32 v0, a02462; GCN-NEXT: v_accvgpr_read_b32 v1, a12463; GCN-NEXT: v_accvgpr_read_b32 v2, a22464; GCN-NEXT: v_accvgpr_read_b32 v3, a32465; GCN-NEXT: s_setpc_b64 s[30:31]2466 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32> %arg0, <6 x i32> %arg1, <4 x float> %arg2,2467 i32 0, ; cbsz2468 i32 4, ; blgp2469 i32 0, i32 %scale0, i32 0, i32 %scale1)2470 ret <4 x float> %result2471}2472 2473define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v6i32_fp4__v8i32_fp8(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2474; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v6i32_fp4__v8i32_fp8:2475; GCN: ; %bb.0:2476; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2477; GCN-NEXT: v_accvgpr_write_b32 a0, v142478; GCN-NEXT: v_accvgpr_write_b32 a1, v152479; GCN-NEXT: v_accvgpr_write_b32 a2, v162480; GCN-NEXT: v_accvgpr_write_b32 a3, v172481; GCN-NEXT: s_nop 12482; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:5], v[6:13], a[0:3], v18, v19 op_sel_hi:[0,0,0] cbsz:42483; GCN-NEXT: s_nop 112484; GCN-NEXT: v_accvgpr_read_b32 v0, a02485; GCN-NEXT: v_accvgpr_read_b32 v1, a12486; GCN-NEXT: v_accvgpr_read_b32 v2, a22487; GCN-NEXT: v_accvgpr_read_b32 v3, a32488; GCN-NEXT: s_setpc_b64 s[30:31]2489 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,2490 i32 4, ; cbsz2491 i32 0, ; blgp2492 i32 0, i32 %scale0, i32 0, i32 %scale1)2493 ret <4 x float> %result2494}2495 2496define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp4__v8i32_fp4(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2, i32 %scale0, i32 %scale1) {2497; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp4__v8i32_fp4:2498; GCN: ; %bb.0:2499; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2500; GCN-NEXT: v_accvgpr_write_b32 a0, v162501; GCN-NEXT: v_accvgpr_write_b32 a1, v172502; GCN-NEXT: v_accvgpr_write_b32 a2, v182503; GCN-NEXT: v_accvgpr_write_b32 a3, v192504; GCN-NEXT: s_nop 12505; GCN-NEXT: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] cbsz:4 blgp:42506; GCN-NEXT: s_nop 72507; GCN-NEXT: v_accvgpr_read_b32 v0, a02508; GCN-NEXT: v_accvgpr_read_b32 v1, a12509; GCN-NEXT: v_accvgpr_read_b32 v2, a22510; GCN-NEXT: v_accvgpr_read_b32 v3, a32511; GCN-NEXT: s_setpc_b64 s[30:31]2512 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,2513 i32 4, ; cbsz2514 i32 4, ; blgp2515 i32 0, i32 %scale0, i32 0, i32 %scale1)2516 ret <4 x float> %result2517}2518 2519define <4 x float> @test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp4__v8i32_fp4__0_scale(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2) {2520; GCN-LABEL: test_mfma_scale_f32_16x16x128_f8f6f4___v8i32_fp4__v8i32_fp4__0_scale:2521; GCN: ; %bb.0:2522; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)2523; GCN-NEXT: v_accvgpr_write_b32 a0, v162524; GCN-NEXT: v_accvgpr_write_b32 a1, v172525; GCN-NEXT: v_accvgpr_write_b32 a2, v182526; GCN-NEXT: v_accvgpr_write_b32 a3, v192527; GCN-NEXT: s_nop 12528; GCN-NEXT: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3] cbsz:4 blgp:42529; GCN-NEXT: s_nop 72530; GCN-NEXT: v_accvgpr_read_b32 v0, a02531; GCN-NEXT: v_accvgpr_read_b32 v1, a12532; GCN-NEXT: v_accvgpr_read_b32 v2, a22533; GCN-NEXT: v_accvgpr_read_b32 v3, a32534; GCN-NEXT: s_setpc_b64 s[30:31]2535 %result = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <4 x float> %arg2,2536 i32 4, ; cbsz2537 i32 4, ; blgp2538 i32 0, i32 0, i32 0, i32 0)2539 ret <4 x float> %result2540}2541 2542declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32>, <8 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32) #12543declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v6i32(<6 x i32>, <6 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32) #12544declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v4i32(<4 x i32>, <4 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32) #12545declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v6i32(<4 x i32>, <6 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32) #12546declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v4i32.v8i32(<4 x i32>, <8 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32) #12547declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v4i32(<6 x i32>, <4 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32) #12548declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32>, <8 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32) #12549declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v4i32(<8 x i32>, <4 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32) #12550declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v6i32(<8 x i32>, <6 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32) #12551 2552attributes #0 = { "amdgpu-flat-work-group-size"="512,512" "amdgpu-agpr-alloc"="0,0" }2553attributes #1 = { convergent nocallback nofree nosync nounwind willreturn memory(none) }2554