1803 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -verify-machineinstrs -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX90A %s3; RUN: llc -verify-machineinstrs -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX942 %s4; RUN: llc -verify-machineinstrs -mtriple=amdgcn -mcpu=gfx90a -amdgpu-mfma-vgpr-form < %s | FileCheck -enable-var-scope --check-prefixes=VGPR,GFX90A-VGPR %s5; RUN: llc -verify-machineinstrs -mtriple=amdgcn -mcpu=gfx942 -amdgpu-mfma-vgpr-form < %s | FileCheck -enable-var-scope --check-prefixes=VGPR,GFX942-VGPR %s6 7declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16.1k(<4 x i16>, <4 x i16>, <32 x float>, i32, i32, i32)8declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x4bf16.1k(<4 x i16>, <4 x i16>, <16 x float>, i32, i32, i32)9declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x4bf16.1k(<4 x i16>, <4 x i16>, <4 x float>, i32, i32, i32)10declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x8bf16.1k(<4 x i16>, <4 x i16>, <16 x float>, i32, i32, i32)11declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16>, <4 x i16>, <4 x float>, i32, i32, i32)12declare <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double, double, <4 x double>, i32, i32, i32)13declare double @llvm.amdgcn.mfma.f64.4x4x4f64(double, double, double, i32, i32, i32)14declare i32 @llvm.amdgcn.workitem.id.x()15 16define amdgpu_kernel void @test_mfma_f32_32x32x4bf16_1k(ptr addrspace(1) %arg) #0 {17; GFX90A-LABEL: test_mfma_f32_32x32x4bf16_1k:18; GFX90A: ; %bb.0: ; %bb19; GFX90A-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x2420; GFX90A-NEXT: v_mov_b32_e32 v1, 021; GFX90A-NEXT: v_mov_b32_e32 v2, 122; GFX90A-NEXT: v_mov_b32_e32 v3, v123; GFX90A-NEXT: v_mov_b32_e32 v0, 224; GFX90A-NEXT: s_waitcnt lgkmcnt(0)25; GFX90A-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x026; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x4027; GFX90A-NEXT: s_waitcnt lgkmcnt(0)28; GFX90A-NEXT: v_accvgpr_write_b32 a0, s1629; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1730; GFX90A-NEXT: v_accvgpr_write_b32 a2, s1831; GFX90A-NEXT: v_accvgpr_write_b32 a3, s1932; GFX90A-NEXT: v_accvgpr_write_b32 a4, s2033; GFX90A-NEXT: v_accvgpr_write_b32 a5, s2134; GFX90A-NEXT: v_accvgpr_write_b32 a6, s2235; GFX90A-NEXT: v_accvgpr_write_b32 a7, s2336; GFX90A-NEXT: v_accvgpr_write_b32 a8, s2437; GFX90A-NEXT: v_accvgpr_write_b32 a9, s2538; GFX90A-NEXT: v_accvgpr_write_b32 a10, s2639; GFX90A-NEXT: v_accvgpr_write_b32 a11, s2740; GFX90A-NEXT: v_accvgpr_write_b32 a12, s2841; GFX90A-NEXT: v_accvgpr_write_b32 a13, s2942; GFX90A-NEXT: v_accvgpr_write_b32 a14, s3043; GFX90A-NEXT: v_accvgpr_write_b32 a15, s3144; GFX90A-NEXT: v_accvgpr_write_b32 a16, s045; GFX90A-NEXT: v_accvgpr_write_b32 a17, s146; GFX90A-NEXT: v_accvgpr_write_b32 a18, s247; GFX90A-NEXT: v_accvgpr_write_b32 a19, s348; GFX90A-NEXT: v_accvgpr_write_b32 a20, s449; GFX90A-NEXT: v_accvgpr_write_b32 a21, s550; GFX90A-NEXT: v_accvgpr_write_b32 a22, s651; GFX90A-NEXT: v_accvgpr_write_b32 a23, s752; GFX90A-NEXT: v_accvgpr_write_b32 a24, s853; GFX90A-NEXT: v_accvgpr_write_b32 a25, s954; GFX90A-NEXT: v_accvgpr_write_b32 a26, s1055; GFX90A-NEXT: v_accvgpr_write_b32 a27, s1156; GFX90A-NEXT: v_accvgpr_write_b32 a28, s1257; GFX90A-NEXT: v_accvgpr_write_b32 a29, s1358; GFX90A-NEXT: v_accvgpr_write_b32 a30, s1459; GFX90A-NEXT: v_accvgpr_write_b32 a31, s1560; GFX90A-NEXT: s_nop 161; GFX90A-NEXT: v_mfma_f32_32x32x4bf16_1k a[0:31], v[2:3], v[0:1], a[0:31] cbsz:1 abid:2 blgp:362; GFX90A-NEXT: s_nop 1563; GFX90A-NEXT: s_nop 264; GFX90A-NEXT: global_store_dwordx4 v1, a[24:27], s[34:35] offset:9665; GFX90A-NEXT: global_store_dwordx4 v1, a[28:31], s[34:35] offset:11266; GFX90A-NEXT: global_store_dwordx4 v1, a[16:19], s[34:35] offset:6467; GFX90A-NEXT: global_store_dwordx4 v1, a[20:23], s[34:35] offset:8068; GFX90A-NEXT: global_store_dwordx4 v1, a[8:11], s[34:35] offset:3269; GFX90A-NEXT: global_store_dwordx4 v1, a[12:15], s[34:35] offset:4870; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[34:35]71; GFX90A-NEXT: global_store_dwordx4 v1, a[4:7], s[34:35] offset:1672; GFX90A-NEXT: s_endpgm73;74; GFX942-LABEL: test_mfma_f32_32x32x4bf16_1k:75; GFX942: ; %bb.0: ; %bb76; GFX942-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x2477; GFX942-NEXT: v_mov_b32_e32 v1, 078; GFX942-NEXT: v_mov_b32_e32 v2, 179; GFX942-NEXT: v_mov_b32_e32 v3, v180; GFX942-NEXT: v_mov_b32_e32 v0, 281; GFX942-NEXT: s_waitcnt lgkmcnt(0)82; GFX942-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x083; GFX942-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x4084; GFX942-NEXT: s_waitcnt lgkmcnt(0)85; GFX942-NEXT: v_accvgpr_write_b32 a0, s1686; GFX942-NEXT: v_accvgpr_write_b32 a1, s1787; GFX942-NEXT: v_accvgpr_write_b32 a2, s1888; GFX942-NEXT: v_accvgpr_write_b32 a3, s1989; GFX942-NEXT: v_accvgpr_write_b32 a4, s2090; GFX942-NEXT: v_accvgpr_write_b32 a5, s2191; GFX942-NEXT: v_accvgpr_write_b32 a6, s2292; GFX942-NEXT: v_accvgpr_write_b32 a7, s2393; GFX942-NEXT: v_accvgpr_write_b32 a8, s2494; GFX942-NEXT: v_accvgpr_write_b32 a9, s2595; GFX942-NEXT: v_accvgpr_write_b32 a10, s2696; GFX942-NEXT: v_accvgpr_write_b32 a11, s2797; GFX942-NEXT: v_accvgpr_write_b32 a12, s2898; GFX942-NEXT: v_accvgpr_write_b32 a13, s2999; GFX942-NEXT: v_accvgpr_write_b32 a14, s30100; GFX942-NEXT: v_accvgpr_write_b32 a15, s31101; GFX942-NEXT: v_accvgpr_write_b32 a16, s0102; GFX942-NEXT: v_accvgpr_write_b32 a17, s1103; GFX942-NEXT: v_accvgpr_write_b32 a18, s2104; GFX942-NEXT: v_accvgpr_write_b32 a19, s3105; GFX942-NEXT: v_accvgpr_write_b32 a20, s4106; GFX942-NEXT: v_accvgpr_write_b32 a21, s5107; GFX942-NEXT: v_accvgpr_write_b32 a22, s6108; GFX942-NEXT: v_accvgpr_write_b32 a23, s7109; GFX942-NEXT: v_accvgpr_write_b32 a24, s8110; GFX942-NEXT: v_accvgpr_write_b32 a25, s9111; GFX942-NEXT: v_accvgpr_write_b32 a26, s10112; GFX942-NEXT: v_accvgpr_write_b32 a27, s11113; GFX942-NEXT: v_accvgpr_write_b32 a28, s12114; GFX942-NEXT: v_accvgpr_write_b32 a29, s13115; GFX942-NEXT: v_accvgpr_write_b32 a30, s14116; GFX942-NEXT: v_accvgpr_write_b32 a31, s15117; GFX942-NEXT: s_nop 1118; GFX942-NEXT: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[0:1], a[0:31] cbsz:1 abid:2 blgp:3119; GFX942-NEXT: s_nop 15120; GFX942-NEXT: s_nop 2121; GFX942-NEXT: global_store_dwordx4 v1, a[24:27], s[34:35] offset:96122; GFX942-NEXT: global_store_dwordx4 v1, a[28:31], s[34:35] offset:112123; GFX942-NEXT: global_store_dwordx4 v1, a[16:19], s[34:35] offset:64124; GFX942-NEXT: global_store_dwordx4 v1, a[20:23], s[34:35] offset:80125; GFX942-NEXT: global_store_dwordx4 v1, a[8:11], s[34:35] offset:32126; GFX942-NEXT: global_store_dwordx4 v1, a[12:15], s[34:35] offset:48127; GFX942-NEXT: global_store_dwordx4 v1, a[0:3], s[34:35]128; GFX942-NEXT: global_store_dwordx4 v1, a[4:7], s[34:35] offset:16129; GFX942-NEXT: s_endpgm130;131; GFX90A-VGPR-LABEL: test_mfma_f32_32x32x4bf16_1k:132; GFX90A-VGPR: ; %bb.0: ; %bb133; GFX90A-VGPR-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24134; GFX90A-VGPR-NEXT: v_mov_b32_e32 v33, 0135; GFX90A-VGPR-NEXT: v_mov_b32_e32 v34, 1136; GFX90A-VGPR-NEXT: v_mov_b32_e32 v35, v33137; GFX90A-VGPR-NEXT: v_mov_b32_e32 v32, 2138; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)139; GFX90A-VGPR-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0140; GFX90A-VGPR-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40141; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)142; GFX90A-VGPR-NEXT: v_mov_b32_e32 v0, s16143; GFX90A-VGPR-NEXT: v_mov_b32_e32 v1, s17144; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, s18145; GFX90A-VGPR-NEXT: v_mov_b32_e32 v3, s19146; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, s20147; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, s21148; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, s22149; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, s23150; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, s24151; GFX90A-VGPR-NEXT: v_mov_b32_e32 v9, s25152; GFX90A-VGPR-NEXT: v_mov_b32_e32 v10, s26153; GFX90A-VGPR-NEXT: v_mov_b32_e32 v11, s27154; GFX90A-VGPR-NEXT: v_mov_b32_e32 v12, s28155; GFX90A-VGPR-NEXT: v_mov_b32_e32 v13, s29156; GFX90A-VGPR-NEXT: v_mov_b32_e32 v14, s30157; GFX90A-VGPR-NEXT: v_mov_b32_e32 v15, s31158; GFX90A-VGPR-NEXT: v_mov_b32_e32 v16, s0159; GFX90A-VGPR-NEXT: v_mov_b32_e32 v17, s1160; GFX90A-VGPR-NEXT: v_mov_b32_e32 v18, s2161; GFX90A-VGPR-NEXT: v_mov_b32_e32 v19, s3162; GFX90A-VGPR-NEXT: v_mov_b32_e32 v20, s4163; GFX90A-VGPR-NEXT: v_mov_b32_e32 v21, s5164; GFX90A-VGPR-NEXT: v_mov_b32_e32 v22, s6165; GFX90A-VGPR-NEXT: v_mov_b32_e32 v23, s7166; GFX90A-VGPR-NEXT: v_mov_b32_e32 v24, s8167; GFX90A-VGPR-NEXT: v_mov_b32_e32 v25, s9168; GFX90A-VGPR-NEXT: v_mov_b32_e32 v26, s10169; GFX90A-VGPR-NEXT: v_mov_b32_e32 v27, s11170; GFX90A-VGPR-NEXT: v_mov_b32_e32 v28, s12171; GFX90A-VGPR-NEXT: v_mov_b32_e32 v29, s13172; GFX90A-VGPR-NEXT: v_mov_b32_e32 v30, s14173; GFX90A-VGPR-NEXT: v_mov_b32_e32 v31, s15174; GFX90A-VGPR-NEXT: s_nop 1175; GFX90A-VGPR-NEXT: v_mfma_f32_32x32x4bf16_1k v[0:31], v[34:35], v[32:33], v[0:31] cbsz:1 abid:2 blgp:3176; GFX90A-VGPR-NEXT: s_nop 15177; GFX90A-VGPR-NEXT: s_nop 2178; GFX90A-VGPR-NEXT: global_store_dwordx4 v33, v[24:27], s[34:35] offset:96179; GFX90A-VGPR-NEXT: global_store_dwordx4 v33, v[28:31], s[34:35] offset:112180; GFX90A-VGPR-NEXT: global_store_dwordx4 v33, v[16:19], s[34:35] offset:64181; GFX90A-VGPR-NEXT: global_store_dwordx4 v33, v[20:23], s[34:35] offset:80182; GFX90A-VGPR-NEXT: global_store_dwordx4 v33, v[8:11], s[34:35] offset:32183; GFX90A-VGPR-NEXT: global_store_dwordx4 v33, v[12:15], s[34:35] offset:48184; GFX90A-VGPR-NEXT: global_store_dwordx4 v33, v[0:3], s[34:35]185; GFX90A-VGPR-NEXT: global_store_dwordx4 v33, v[4:7], s[34:35] offset:16186; GFX90A-VGPR-NEXT: s_endpgm187;188; GFX942-VGPR-LABEL: test_mfma_f32_32x32x4bf16_1k:189; GFX942-VGPR: ; %bb.0: ; %bb190; GFX942-VGPR-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24191; GFX942-VGPR-NEXT: v_mov_b32_e32 v33, 0192; GFX942-VGPR-NEXT: v_mov_b32_e32 v34, 1193; GFX942-VGPR-NEXT: v_mov_b32_e32 v35, v33194; GFX942-VGPR-NEXT: v_mov_b32_e32 v32, 2195; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)196; GFX942-VGPR-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0197; GFX942-VGPR-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40198; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)199; GFX942-VGPR-NEXT: v_mov_b32_e32 v0, s16200; GFX942-VGPR-NEXT: v_mov_b32_e32 v1, s17201; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, s18202; GFX942-VGPR-NEXT: v_mov_b32_e32 v3, s19203; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, s20204; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, s21205; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, s22206; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, s23207; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, s24208; GFX942-VGPR-NEXT: v_mov_b32_e32 v9, s25209; GFX942-VGPR-NEXT: v_mov_b32_e32 v10, s26210; GFX942-VGPR-NEXT: v_mov_b32_e32 v11, s27211; GFX942-VGPR-NEXT: v_mov_b32_e32 v12, s28212; GFX942-VGPR-NEXT: v_mov_b32_e32 v13, s29213; GFX942-VGPR-NEXT: v_mov_b32_e32 v14, s30214; GFX942-VGPR-NEXT: v_mov_b32_e32 v15, s31215; GFX942-VGPR-NEXT: v_mov_b32_e32 v16, s0216; GFX942-VGPR-NEXT: v_mov_b32_e32 v17, s1217; GFX942-VGPR-NEXT: v_mov_b32_e32 v18, s2218; GFX942-VGPR-NEXT: v_mov_b32_e32 v19, s3219; GFX942-VGPR-NEXT: v_mov_b32_e32 v20, s4220; GFX942-VGPR-NEXT: v_mov_b32_e32 v21, s5221; GFX942-VGPR-NEXT: v_mov_b32_e32 v22, s6222; GFX942-VGPR-NEXT: v_mov_b32_e32 v23, s7223; GFX942-VGPR-NEXT: v_mov_b32_e32 v24, s8224; GFX942-VGPR-NEXT: v_mov_b32_e32 v25, s9225; GFX942-VGPR-NEXT: v_mov_b32_e32 v26, s10226; GFX942-VGPR-NEXT: v_mov_b32_e32 v27, s11227; GFX942-VGPR-NEXT: v_mov_b32_e32 v28, s12228; GFX942-VGPR-NEXT: v_mov_b32_e32 v29, s13229; GFX942-VGPR-NEXT: v_mov_b32_e32 v30, s14230; GFX942-VGPR-NEXT: v_mov_b32_e32 v31, s15231; GFX942-VGPR-NEXT: s_nop 1232; GFX942-VGPR-NEXT: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[34:35], v[32:33], v[0:31] cbsz:1 abid:2 blgp:3233; GFX942-VGPR-NEXT: s_nop 15234; GFX942-VGPR-NEXT: s_nop 2235; GFX942-VGPR-NEXT: global_store_dwordx4 v33, v[24:27], s[34:35] offset:96236; GFX942-VGPR-NEXT: global_store_dwordx4 v33, v[28:31], s[34:35] offset:112237; GFX942-VGPR-NEXT: global_store_dwordx4 v33, v[16:19], s[34:35] offset:64238; GFX942-VGPR-NEXT: global_store_dwordx4 v33, v[20:23], s[34:35] offset:80239; GFX942-VGPR-NEXT: global_store_dwordx4 v33, v[8:11], s[34:35] offset:32240; GFX942-VGPR-NEXT: global_store_dwordx4 v33, v[12:15], s[34:35] offset:48241; GFX942-VGPR-NEXT: global_store_dwordx4 v33, v[0:3], s[34:35]242; GFX942-VGPR-NEXT: global_store_dwordx4 v33, v[4:7], s[34:35] offset:16243; GFX942-VGPR-NEXT: s_endpgm244bb:245 %in.1 = load <32 x float>, ptr addrspace(1) %arg246 %a = bitcast i64 1 to <4 x i16>247 %b = bitcast i64 2 to <4 x i16>248 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16.1k(<4 x i16> %a, <4 x i16> %b, <32 x float> %in.1, i32 1, i32 2, i32 3)249 store <32 x float> %mai.1, ptr addrspace(1) %arg250 ret void251}252 253define amdgpu_kernel void @test_mfma_f32_16x16x4bf16_1k(ptr addrspace(1) %arg) #0 {254; GFX90A-LABEL: test_mfma_f32_16x16x4bf16_1k:255; GFX90A: ; %bb.0: ; %bb256; GFX90A-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24257; GFX90A-NEXT: v_mov_b32_e32 v1, 0258; GFX90A-NEXT: v_mov_b32_e32 v2, 1259; GFX90A-NEXT: v_mov_b32_e32 v3, v1260; GFX90A-NEXT: v_mov_b32_e32 v0, 2261; GFX90A-NEXT: s_waitcnt lgkmcnt(0)262; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0263; GFX90A-NEXT: s_waitcnt lgkmcnt(0)264; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0265; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1266; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2267; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3268; GFX90A-NEXT: v_accvgpr_write_b32 a4, s4269; GFX90A-NEXT: v_accvgpr_write_b32 a5, s5270; GFX90A-NEXT: v_accvgpr_write_b32 a6, s6271; GFX90A-NEXT: v_accvgpr_write_b32 a7, s7272; GFX90A-NEXT: v_accvgpr_write_b32 a8, s8273; GFX90A-NEXT: v_accvgpr_write_b32 a9, s9274; GFX90A-NEXT: v_accvgpr_write_b32 a10, s10275; GFX90A-NEXT: v_accvgpr_write_b32 a11, s11276; GFX90A-NEXT: v_accvgpr_write_b32 a12, s12277; GFX90A-NEXT: v_accvgpr_write_b32 a13, s13278; GFX90A-NEXT: v_accvgpr_write_b32 a14, s14279; GFX90A-NEXT: v_accvgpr_write_b32 a15, s15280; GFX90A-NEXT: s_nop 1281; GFX90A-NEXT: v_mfma_f32_16x16x4bf16_1k a[0:15], v[2:3], v[0:1], a[0:15] cbsz:1 abid:2 blgp:3282; GFX90A-NEXT: s_nop 10283; GFX90A-NEXT: global_store_dwordx4 v1, a[12:15], s[16:17] offset:48284; GFX90A-NEXT: global_store_dwordx4 v1, a[8:11], s[16:17] offset:32285; GFX90A-NEXT: global_store_dwordx4 v1, a[4:7], s[16:17] offset:16286; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[16:17]287; GFX90A-NEXT: s_endpgm288;289; GFX942-LABEL: test_mfma_f32_16x16x4bf16_1k:290; GFX942: ; %bb.0: ; %bb291; GFX942-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24292; GFX942-NEXT: v_mov_b32_e32 v1, 0293; GFX942-NEXT: v_mov_b32_e32 v2, 1294; GFX942-NEXT: v_mov_b32_e32 v3, v1295; GFX942-NEXT: v_mov_b32_e32 v0, 2296; GFX942-NEXT: s_waitcnt lgkmcnt(0)297; GFX942-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0298; GFX942-NEXT: s_waitcnt lgkmcnt(0)299; GFX942-NEXT: v_accvgpr_write_b32 a0, s0300; GFX942-NEXT: v_accvgpr_write_b32 a1, s1301; GFX942-NEXT: v_accvgpr_write_b32 a2, s2302; GFX942-NEXT: v_accvgpr_write_b32 a3, s3303; GFX942-NEXT: v_accvgpr_write_b32 a4, s4304; GFX942-NEXT: v_accvgpr_write_b32 a5, s5305; GFX942-NEXT: v_accvgpr_write_b32 a6, s6306; GFX942-NEXT: v_accvgpr_write_b32 a7, s7307; GFX942-NEXT: v_accvgpr_write_b32 a8, s8308; GFX942-NEXT: v_accvgpr_write_b32 a9, s9309; GFX942-NEXT: v_accvgpr_write_b32 a10, s10310; GFX942-NEXT: v_accvgpr_write_b32 a11, s11311; GFX942-NEXT: v_accvgpr_write_b32 a12, s12312; GFX942-NEXT: v_accvgpr_write_b32 a13, s13313; GFX942-NEXT: v_accvgpr_write_b32 a14, s14314; GFX942-NEXT: v_accvgpr_write_b32 a15, s15315; GFX942-NEXT: s_nop 1316; GFX942-NEXT: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[0:1], a[0:15] cbsz:1 abid:2 blgp:3317; GFX942-NEXT: s_nop 10318; GFX942-NEXT: global_store_dwordx4 v1, a[12:15], s[16:17] offset:48319; GFX942-NEXT: global_store_dwordx4 v1, a[8:11], s[16:17] offset:32320; GFX942-NEXT: global_store_dwordx4 v1, a[4:7], s[16:17] offset:16321; GFX942-NEXT: global_store_dwordx4 v1, a[0:3], s[16:17]322; GFX942-NEXT: s_endpgm323;324; GFX90A-VGPR-LABEL: test_mfma_f32_16x16x4bf16_1k:325; GFX90A-VGPR: ; %bb.0: ; %bb326; GFX90A-VGPR-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24327; GFX90A-VGPR-NEXT: v_mov_b32_e32 v17, 0328; GFX90A-VGPR-NEXT: v_mov_b32_e32 v18, 1329; GFX90A-VGPR-NEXT: v_mov_b32_e32 v19, v17330; GFX90A-VGPR-NEXT: v_mov_b32_e32 v16, 2331; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)332; GFX90A-VGPR-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0333; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)334; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]335; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]336; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], s[4:5], s[4:5] op_sel:[0,1]337; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[6:7], s[6:7], s[6:7] op_sel:[0,1]338; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[8:9], s[8:9] op_sel:[0,1]339; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[10:11], s[10:11] op_sel:[0,1]340; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[12:13], s[12:13], s[12:13] op_sel:[0,1]341; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[14:15], s[14:15], s[14:15] op_sel:[0,1]342; GFX90A-VGPR-NEXT: s_nop 1343; GFX90A-VGPR-NEXT: v_mfma_f32_16x16x4bf16_1k v[0:15], v[18:19], v[16:17], v[0:15] cbsz:1 abid:2 blgp:3344; GFX90A-VGPR-NEXT: s_nop 10345; GFX90A-VGPR-NEXT: global_store_dwordx4 v17, v[12:15], s[16:17] offset:48346; GFX90A-VGPR-NEXT: global_store_dwordx4 v17, v[8:11], s[16:17] offset:32347; GFX90A-VGPR-NEXT: global_store_dwordx4 v17, v[4:7], s[16:17] offset:16348; GFX90A-VGPR-NEXT: global_store_dwordx4 v17, v[0:3], s[16:17]349; GFX90A-VGPR-NEXT: s_endpgm350;351; GFX942-VGPR-LABEL: test_mfma_f32_16x16x4bf16_1k:352; GFX942-VGPR: ; %bb.0: ; %bb353; GFX942-VGPR-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24354; GFX942-VGPR-NEXT: v_mov_b32_e32 v17, 0355; GFX942-VGPR-NEXT: v_mov_b32_e32 v18, 1356; GFX942-VGPR-NEXT: v_mov_b32_e32 v19, v17357; GFX942-VGPR-NEXT: v_mov_b32_e32 v16, 2358; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)359; GFX942-VGPR-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0360; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)361; GFX942-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[0:1]362; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[2:3]363; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[4:5]364; GFX942-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[6:7]365; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[8:9]366; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[10:11]367; GFX942-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[12:13]368; GFX942-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[14:15]369; GFX942-VGPR-NEXT: s_nop 1370; GFX942-VGPR-NEXT: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[18:19], v[16:17], v[0:15] cbsz:1 abid:2 blgp:3371; GFX942-VGPR-NEXT: s_nop 10372; GFX942-VGPR-NEXT: global_store_dwordx4 v17, v[12:15], s[16:17] offset:48373; GFX942-VGPR-NEXT: global_store_dwordx4 v17, v[8:11], s[16:17] offset:32374; GFX942-VGPR-NEXT: global_store_dwordx4 v17, v[4:7], s[16:17] offset:16375; GFX942-VGPR-NEXT: global_store_dwordx4 v17, v[0:3], s[16:17]376; GFX942-VGPR-NEXT: s_endpgm377bb:378 %in.1 = load <16 x float>, ptr addrspace(1) %arg379 %a = bitcast i64 1 to <4 x i16>380 %b = bitcast i64 2 to <4 x i16>381 %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4bf16.1k(<4 x i16> %a, <4 x i16> %b, <16 x float> %in.1, i32 1, i32 2, i32 3)382 store <16 x float> %mai.1, ptr addrspace(1) %arg383 ret void384}385 386define amdgpu_kernel void @test_mfma_f32_4x4x4bf16_1k(ptr addrspace(1) %arg) #0 {387; GFX90A-LABEL: test_mfma_f32_4x4x4bf16_1k:388; GFX90A: ; %bb.0: ; %bb389; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24390; GFX90A-NEXT: v_mov_b32_e32 v1, 0391; GFX90A-NEXT: v_mov_b32_e32 v2, 1392; GFX90A-NEXT: v_mov_b32_e32 v3, v1393; GFX90A-NEXT: v_mov_b32_e32 v0, 2394; GFX90A-NEXT: s_waitcnt lgkmcnt(0)395; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0396; GFX90A-NEXT: s_waitcnt lgkmcnt(0)397; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0398; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1399; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2400; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3401; GFX90A-NEXT: s_nop 1402; GFX90A-NEXT: v_mfma_f32_4x4x4bf16_1k a[0:3], v[2:3], v[0:1], a[0:3] cbsz:1 abid:2 blgp:3403; GFX90A-NEXT: s_nop 4404; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[6:7]405; GFX90A-NEXT: s_endpgm406;407; GFX942-LABEL: test_mfma_f32_4x4x4bf16_1k:408; GFX942: ; %bb.0: ; %bb409; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24410; GFX942-NEXT: v_mov_b32_e32 v1, 0411; GFX942-NEXT: v_mov_b32_e32 v2, 1412; GFX942-NEXT: v_mov_b32_e32 v3, v1413; GFX942-NEXT: v_mov_b32_e32 v0, 2414; GFX942-NEXT: s_waitcnt lgkmcnt(0)415; GFX942-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0416; GFX942-NEXT: s_waitcnt lgkmcnt(0)417; GFX942-NEXT: v_accvgpr_write_b32 a0, s0418; GFX942-NEXT: v_accvgpr_write_b32 a1, s1419; GFX942-NEXT: v_accvgpr_write_b32 a2, s2420; GFX942-NEXT: v_accvgpr_write_b32 a3, s3421; GFX942-NEXT: s_nop 1422; GFX942-NEXT: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[0:1], a[0:3] cbsz:1 abid:2 blgp:3423; GFX942-NEXT: s_nop 4424; GFX942-NEXT: global_store_dwordx4 v1, a[0:3], s[6:7]425; GFX942-NEXT: s_endpgm426;427; GFX90A-VGPR-LABEL: test_mfma_f32_4x4x4bf16_1k:428; GFX90A-VGPR: ; %bb.0: ; %bb429; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24430; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, 0431; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, 1432; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, v5433; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, 2434; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)435; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0436; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)437; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]438; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]439; GFX90A-VGPR-NEXT: s_nop 1440; GFX90A-VGPR-NEXT: v_mfma_f32_4x4x4bf16_1k v[0:3], v[6:7], v[4:5], v[0:3] cbsz:1 abid:2 blgp:3441; GFX90A-VGPR-NEXT: s_nop 4442; GFX90A-VGPR-NEXT: global_store_dwordx4 v5, v[0:3], s[6:7]443; GFX90A-VGPR-NEXT: s_endpgm444;445; GFX942-VGPR-LABEL: test_mfma_f32_4x4x4bf16_1k:446; GFX942-VGPR: ; %bb.0: ; %bb447; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24448; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, 0449; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, 1450; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, v5451; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, 2452; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)453; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0454; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)455; GFX942-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[0:1]456; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[2:3]457; GFX942-VGPR-NEXT: s_nop 1458; GFX942-VGPR-NEXT: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[6:7], v[4:5], v[0:3] cbsz:1 abid:2 blgp:3459; GFX942-VGPR-NEXT: s_nop 4460; GFX942-VGPR-NEXT: global_store_dwordx4 v5, v[0:3], s[6:7]461; GFX942-VGPR-NEXT: s_endpgm462bb:463 %in.1 = load <4 x float>, ptr addrspace(1) %arg464 %a = bitcast i64 1 to <4 x i16>465 %b = bitcast i64 2 to <4 x i16>466 %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4bf16.1k(<4 x i16> %a, <4 x i16> %b, <4 x float> %in.1, i32 1, i32 2, i32 3)467 store <4 x float> %mai.1, ptr addrspace(1) %arg468 ret void469}470 471define amdgpu_kernel void @test_mfma_f32_32x32x8bf16_1k(ptr addrspace(1) %arg) #0 {472; GFX90A-LABEL: test_mfma_f32_32x32x8bf16_1k:473; GFX90A: ; %bb.0: ; %bb474; GFX90A-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24475; GFX90A-NEXT: v_mov_b32_e32 v1, 0476; GFX90A-NEXT: v_mov_b32_e32 v2, 1477; GFX90A-NEXT: v_mov_b32_e32 v3, v1478; GFX90A-NEXT: v_mov_b32_e32 v0, 2479; GFX90A-NEXT: s_waitcnt lgkmcnt(0)480; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0481; GFX90A-NEXT: s_waitcnt lgkmcnt(0)482; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0483; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1484; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2485; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3486; GFX90A-NEXT: v_accvgpr_write_b32 a4, s4487; GFX90A-NEXT: v_accvgpr_write_b32 a5, s5488; GFX90A-NEXT: v_accvgpr_write_b32 a6, s6489; GFX90A-NEXT: v_accvgpr_write_b32 a7, s7490; GFX90A-NEXT: v_accvgpr_write_b32 a8, s8491; GFX90A-NEXT: v_accvgpr_write_b32 a9, s9492; GFX90A-NEXT: v_accvgpr_write_b32 a10, s10493; GFX90A-NEXT: v_accvgpr_write_b32 a11, s11494; GFX90A-NEXT: v_accvgpr_write_b32 a12, s12495; GFX90A-NEXT: v_accvgpr_write_b32 a13, s13496; GFX90A-NEXT: v_accvgpr_write_b32 a14, s14497; GFX90A-NEXT: v_accvgpr_write_b32 a15, s15498; GFX90A-NEXT: s_nop 1499; GFX90A-NEXT: v_mfma_f32_32x32x8bf16_1k a[0:15], v[2:3], v[0:1], a[0:15] cbsz:1 abid:2 blgp:3500; GFX90A-NEXT: s_nop 15501; GFX90A-NEXT: s_nop 2502; GFX90A-NEXT: global_store_dwordx4 v1, a[12:15], s[16:17] offset:48503; GFX90A-NEXT: global_store_dwordx4 v1, a[8:11], s[16:17] offset:32504; GFX90A-NEXT: global_store_dwordx4 v1, a[4:7], s[16:17] offset:16505; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[16:17]506; GFX90A-NEXT: s_endpgm507;508; GFX942-LABEL: test_mfma_f32_32x32x8bf16_1k:509; GFX942: ; %bb.0: ; %bb510; GFX942-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24511; GFX942-NEXT: v_mov_b32_e32 v1, 0512; GFX942-NEXT: v_mov_b32_e32 v2, 1513; GFX942-NEXT: v_mov_b32_e32 v3, v1514; GFX942-NEXT: v_mov_b32_e32 v0, 2515; GFX942-NEXT: s_waitcnt lgkmcnt(0)516; GFX942-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0517; GFX942-NEXT: s_waitcnt lgkmcnt(0)518; GFX942-NEXT: v_accvgpr_write_b32 a0, s0519; GFX942-NEXT: v_accvgpr_write_b32 a1, s1520; GFX942-NEXT: v_accvgpr_write_b32 a2, s2521; GFX942-NEXT: v_accvgpr_write_b32 a3, s3522; GFX942-NEXT: v_accvgpr_write_b32 a4, s4523; GFX942-NEXT: v_accvgpr_write_b32 a5, s5524; GFX942-NEXT: v_accvgpr_write_b32 a6, s6525; GFX942-NEXT: v_accvgpr_write_b32 a7, s7526; GFX942-NEXT: v_accvgpr_write_b32 a8, s8527; GFX942-NEXT: v_accvgpr_write_b32 a9, s9528; GFX942-NEXT: v_accvgpr_write_b32 a10, s10529; GFX942-NEXT: v_accvgpr_write_b32 a11, s11530; GFX942-NEXT: v_accvgpr_write_b32 a12, s12531; GFX942-NEXT: v_accvgpr_write_b32 a13, s13532; GFX942-NEXT: v_accvgpr_write_b32 a14, s14533; GFX942-NEXT: v_accvgpr_write_b32 a15, s15534; GFX942-NEXT: s_nop 1535; GFX942-NEXT: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[0:1], a[0:15] cbsz:1 abid:2 blgp:3536; GFX942-NEXT: s_nop 10537; GFX942-NEXT: global_store_dwordx4 v1, a[12:15], s[16:17] offset:48538; GFX942-NEXT: global_store_dwordx4 v1, a[8:11], s[16:17] offset:32539; GFX942-NEXT: global_store_dwordx4 v1, a[4:7], s[16:17] offset:16540; GFX942-NEXT: global_store_dwordx4 v1, a[0:3], s[16:17]541; GFX942-NEXT: s_endpgm542;543; GFX90A-VGPR-LABEL: test_mfma_f32_32x32x8bf16_1k:544; GFX90A-VGPR: ; %bb.0: ; %bb545; GFX90A-VGPR-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24546; GFX90A-VGPR-NEXT: v_mov_b32_e32 v17, 0547; GFX90A-VGPR-NEXT: v_mov_b32_e32 v18, 1548; GFX90A-VGPR-NEXT: v_mov_b32_e32 v19, v17549; GFX90A-VGPR-NEXT: v_mov_b32_e32 v16, 2550; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)551; GFX90A-VGPR-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0552; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)553; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]554; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]555; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], s[4:5], s[4:5] op_sel:[0,1]556; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[6:7], s[6:7], s[6:7] op_sel:[0,1]557; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[8:9], s[8:9] op_sel:[0,1]558; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[10:11], s[10:11] op_sel:[0,1]559; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[12:13], s[12:13], s[12:13] op_sel:[0,1]560; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[14:15], s[14:15], s[14:15] op_sel:[0,1]561; GFX90A-VGPR-NEXT: s_nop 1562; GFX90A-VGPR-NEXT: v_mfma_f32_32x32x8bf16_1k v[0:15], v[18:19], v[16:17], v[0:15] cbsz:1 abid:2 blgp:3563; GFX90A-VGPR-NEXT: s_nop 15564; GFX90A-VGPR-NEXT: s_nop 2565; GFX90A-VGPR-NEXT: global_store_dwordx4 v17, v[12:15], s[16:17] offset:48566; GFX90A-VGPR-NEXT: global_store_dwordx4 v17, v[8:11], s[16:17] offset:32567; GFX90A-VGPR-NEXT: global_store_dwordx4 v17, v[4:7], s[16:17] offset:16568; GFX90A-VGPR-NEXT: global_store_dwordx4 v17, v[0:3], s[16:17]569; GFX90A-VGPR-NEXT: s_endpgm570;571; GFX942-VGPR-LABEL: test_mfma_f32_32x32x8bf16_1k:572; GFX942-VGPR: ; %bb.0: ; %bb573; GFX942-VGPR-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24574; GFX942-VGPR-NEXT: v_mov_b32_e32 v17, 0575; GFX942-VGPR-NEXT: v_mov_b32_e32 v18, 1576; GFX942-VGPR-NEXT: v_mov_b32_e32 v19, v17577; GFX942-VGPR-NEXT: v_mov_b32_e32 v16, 2578; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)579; GFX942-VGPR-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0580; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)581; GFX942-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[0:1]582; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[2:3]583; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[4:5]584; GFX942-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[6:7]585; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[8:9]586; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[10:11]587; GFX942-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[12:13]588; GFX942-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[14:15]589; GFX942-VGPR-NEXT: s_nop 1590; GFX942-VGPR-NEXT: v_mfma_f32_32x32x8_bf16 v[0:15], v[18:19], v[16:17], v[0:15] cbsz:1 abid:2 blgp:3591; GFX942-VGPR-NEXT: s_nop 10592; GFX942-VGPR-NEXT: global_store_dwordx4 v17, v[12:15], s[16:17] offset:48593; GFX942-VGPR-NEXT: global_store_dwordx4 v17, v[8:11], s[16:17] offset:32594; GFX942-VGPR-NEXT: global_store_dwordx4 v17, v[4:7], s[16:17] offset:16595; GFX942-VGPR-NEXT: global_store_dwordx4 v17, v[0:3], s[16:17]596; GFX942-VGPR-NEXT: s_endpgm597bb:598 %in.1 = load <16 x float>, ptr addrspace(1) %arg599 %a = bitcast i64 1 to <4 x i16>600 %b = bitcast i64 2 to <4 x i16>601 %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8bf16.1k(<4 x i16> %a, <4 x i16> %b, <16 x float> %in.1, i32 1, i32 2, i32 3)602 store <16 x float> %mai.1, ptr addrspace(1) %arg603 ret void604}605 606define amdgpu_kernel void @test_mfma_f32_16x16x16bf16_1k(ptr addrspace(1) %arg) #0 {607; GFX90A-LABEL: test_mfma_f32_16x16x16bf16_1k:608; GFX90A: ; %bb.0: ; %bb609; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24610; GFX90A-NEXT: v_mov_b32_e32 v1, 0611; GFX90A-NEXT: v_mov_b32_e32 v2, 1612; GFX90A-NEXT: v_mov_b32_e32 v3, v1613; GFX90A-NEXT: v_mov_b32_e32 v0, 2614; GFX90A-NEXT: s_waitcnt lgkmcnt(0)615; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0616; GFX90A-NEXT: s_waitcnt lgkmcnt(0)617; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0618; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1619; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2620; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3621; GFX90A-NEXT: s_nop 1622; GFX90A-NEXT: v_mfma_f32_16x16x16bf16_1k a[0:3], v[2:3], v[0:1], a[0:3] cbsz:1 abid:2 blgp:3623; GFX90A-NEXT: s_nop 10624; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[6:7]625; GFX90A-NEXT: s_endpgm626;627; GFX942-LABEL: test_mfma_f32_16x16x16bf16_1k:628; GFX942: ; %bb.0: ; %bb629; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24630; GFX942-NEXT: v_mov_b32_e32 v1, 0631; GFX942-NEXT: v_mov_b32_e32 v2, 1632; GFX942-NEXT: v_mov_b32_e32 v3, v1633; GFX942-NEXT: v_mov_b32_e32 v0, 2634; GFX942-NEXT: s_waitcnt lgkmcnt(0)635; GFX942-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0636; GFX942-NEXT: s_waitcnt lgkmcnt(0)637; GFX942-NEXT: v_accvgpr_write_b32 a0, s0638; GFX942-NEXT: v_accvgpr_write_b32 a1, s1639; GFX942-NEXT: v_accvgpr_write_b32 a2, s2640; GFX942-NEXT: v_accvgpr_write_b32 a3, s3641; GFX942-NEXT: s_nop 1642; GFX942-NEXT: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[0:1], a[0:3] cbsz:1 abid:2 blgp:3643; GFX942-NEXT: s_nop 6644; GFX942-NEXT: global_store_dwordx4 v1, a[0:3], s[6:7]645; GFX942-NEXT: s_endpgm646;647; GFX90A-VGPR-LABEL: test_mfma_f32_16x16x16bf16_1k:648; GFX90A-VGPR: ; %bb.0: ; %bb649; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24650; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, 0651; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, 1652; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, v5653; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, 2654; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)655; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0656; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)657; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]658; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]659; GFX90A-VGPR-NEXT: s_nop 1660; GFX90A-VGPR-NEXT: v_mfma_f32_16x16x16bf16_1k v[0:3], v[6:7], v[4:5], v[0:3] cbsz:1 abid:2 blgp:3661; GFX90A-VGPR-NEXT: s_nop 10662; GFX90A-VGPR-NEXT: global_store_dwordx4 v5, v[0:3], s[6:7]663; GFX90A-VGPR-NEXT: s_endpgm664;665; GFX942-VGPR-LABEL: test_mfma_f32_16x16x16bf16_1k:666; GFX942-VGPR: ; %bb.0: ; %bb667; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24668; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, 0669; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, 1670; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, v5671; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, 2672; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)673; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0674; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)675; GFX942-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[0:1]676; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[2:3]677; GFX942-VGPR-NEXT: s_nop 1678; GFX942-VGPR-NEXT: v_mfma_f32_16x16x16_bf16 v[0:3], v[6:7], v[4:5], v[0:3] cbsz:1 abid:2 blgp:3679; GFX942-VGPR-NEXT: s_nop 6680; GFX942-VGPR-NEXT: global_store_dwordx4 v5, v[0:3], s[6:7]681; GFX942-VGPR-NEXT: s_endpgm682bb:683 %in.1 = load <4 x float>, ptr addrspace(1) %arg684 %a = bitcast i64 1 to <4 x i16>685 %b = bitcast i64 2 to <4 x i16>686 %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %a, <4 x i16> %b, <4 x float> %in.1, i32 1, i32 2, i32 3)687 store <4 x float> %mai.1, ptr addrspace(1) %arg688 ret void689}690 691define amdgpu_kernel void @test_mfma_f64_4x4x4f64(ptr addrspace(1) %arg, double %a, double %b) #0 {692; GFX90A-LABEL: test_mfma_f64_4x4x4f64:693; GFX90A: ; %bb.0: ; %bb694; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24695; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34696; GFX90A-NEXT: s_waitcnt lgkmcnt(0)697; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]698; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]699; GFX90A-NEXT: s_nop 1700; GFX90A-NEXT: v_mfma_f64_4x4x4f64 a[0:1], v[0:1], v[2:3], 0701; GFX90A-NEXT: s_nop 3702; GFX90A-NEXT: v_mfma_f64_4x4x4f64 a[0:1], v[0:1], v[2:3], a[0:1] cbsz:1 abid:2 blgp:3703; GFX90A-NEXT: v_mov_b32_e32 v0, 0704; GFX90A-NEXT: s_nop 7705; GFX90A-NEXT: global_store_dwordx2 v0, a[0:1], s[0:1]706; GFX90A-NEXT: s_endpgm707;708; GFX942-LABEL: test_mfma_f64_4x4x4f64:709; GFX942: ; %bb.0: ; %bb710; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24711; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34712; GFX942-NEXT: s_waitcnt lgkmcnt(0)713; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[2:3]714; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]715; GFX942-NEXT: s_nop 1716; GFX942-NEXT: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], v[2:3], 0717; GFX942-NEXT: s_nop 3718; GFX942-NEXT: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], v[2:3], a[0:1] cbsz:1 abid:2 neg:[1,1,0]719; GFX942-NEXT: v_mov_b32_e32 v0, 0720; GFX942-NEXT: s_nop 7721; GFX942-NEXT: global_store_dwordx2 v0, a[0:1], s[0:1]722; GFX942-NEXT: s_endpgm723;724; GFX90A-VGPR-LABEL: test_mfma_f64_4x4x4f64:725; GFX90A-VGPR: ; %bb.0: ; %bb726; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24727; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34728; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)729; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]730; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], s[6:7], s[6:7] op_sel:[0,1]731; GFX90A-VGPR-NEXT: s_nop 1732; GFX90A-VGPR-NEXT: v_mfma_f64_4x4x4f64 v[0:1], v[2:3], v[4:5], 0733; GFX90A-VGPR-NEXT: s_nop 3734; GFX90A-VGPR-NEXT: v_mfma_f64_4x4x4f64 v[0:1], v[2:3], v[4:5], v[0:1] cbsz:1 abid:2 blgp:3735; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, 0736; GFX90A-VGPR-NEXT: s_nop 7737; GFX90A-VGPR-NEXT: global_store_dwordx2 v2, v[0:1], s[0:1]738; GFX90A-VGPR-NEXT: s_endpgm739;740; GFX942-VGPR-LABEL: test_mfma_f64_4x4x4f64:741; GFX942-VGPR: ; %bb.0: ; %bb742; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24743; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34744; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)745; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[2:3]746; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[6:7]747; GFX942-VGPR-NEXT: s_nop 1748; GFX942-VGPR-NEXT: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[2:3], v[4:5], 0749; GFX942-VGPR-NEXT: s_nop 3750; GFX942-VGPR-NEXT: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[2:3], v[4:5], v[0:1] cbsz:1 abid:2 neg:[1,1,0]751; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, 0752; GFX942-VGPR-NEXT: s_nop 7753; GFX942-VGPR-NEXT: global_store_dwordx2 v2, v[0:1], s[0:1]754; GFX942-VGPR-NEXT: s_endpgm755bb:756 %mai.1 = tail call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %a, double %b, double 0.0, i32 0, i32 0, i32 0)757 %mai.2 = tail call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %a, double %b, double %mai.1, i32 1, i32 2, i32 3)758 store double %mai.2, ptr addrspace(1) %arg759 ret void760}761 762define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, double %a, double %b) #0 {763; GFX90A-LABEL: test_mfma_f64_16x16x4f64:764; GFX90A: ; %bb.0: ; %bb765; GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24766; GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34767; GFX90A-NEXT: s_waitcnt lgkmcnt(0)768; GFX90A-NEXT: v_mov_b32_e32 v0, s10769; GFX90A-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0770; GFX90A-NEXT: v_mov_b32_e32 v1, s11771; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[12:13], s[12:13] op_sel:[0,1]772; GFX90A-NEXT: s_waitcnt lgkmcnt(0)773; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0774; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1775; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2776; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3777; GFX90A-NEXT: v_accvgpr_write_b32 a4, s4778; GFX90A-NEXT: v_accvgpr_write_b32 a5, s5779; GFX90A-NEXT: v_accvgpr_write_b32 a6, s6780; GFX90A-NEXT: v_accvgpr_write_b32 a7, s7781; GFX90A-NEXT: s_nop 1782; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:3783; GFX90A-NEXT: v_mov_b32_e32 v0, 0784; GFX90A-NEXT: s_nop 15785; GFX90A-NEXT: s_nop 0786; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[8:9] offset:16787; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[8:9]788; GFX90A-NEXT: s_endpgm789;790; GFX942-LABEL: test_mfma_f64_16x16x4f64:791; GFX942: ; %bb.0: ; %bb792; GFX942-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24793; GFX942-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34794; GFX942-NEXT: s_waitcnt lgkmcnt(0)795; GFX942-NEXT: v_mov_b32_e32 v0, s10796; GFX942-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0797; GFX942-NEXT: v_mov_b32_e32 v1, s11798; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[12:13]799; GFX942-NEXT: s_waitcnt lgkmcnt(0)800; GFX942-NEXT: v_accvgpr_write_b32 a0, s0801; GFX942-NEXT: v_accvgpr_write_b32 a1, s1802; GFX942-NEXT: v_accvgpr_write_b32 a2, s2803; GFX942-NEXT: v_accvgpr_write_b32 a3, s3804; GFX942-NEXT: v_accvgpr_write_b32 a4, s4805; GFX942-NEXT: v_accvgpr_write_b32 a5, s5806; GFX942-NEXT: v_accvgpr_write_b32 a6, s6807; GFX942-NEXT: v_accvgpr_write_b32 a7, s7808; GFX942-NEXT: s_nop 1809; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]810; GFX942-NEXT: v_mov_b32_e32 v0, 0811; GFX942-NEXT: s_nop 15812; GFX942-NEXT: s_nop 0813; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[8:9] offset:16814; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[8:9]815; GFX942-NEXT: s_endpgm816;817; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64:818; GFX90A-VGPR: ; %bb.0: ; %bb819; GFX90A-VGPR-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24820; GFX90A-VGPR-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34821; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)822; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, s10823; GFX90A-VGPR-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0824; GFX90A-VGPR-NEXT: v_mov_b32_e32 v9, s11825; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[12:13], s[12:13] op_sel:[0,1]826; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)827; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]828; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]829; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], s[4:5], s[4:5] op_sel:[0,1]830; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[6:7], s[6:7], s[6:7] op_sel:[0,1]831; GFX90A-VGPR-NEXT: s_nop 1832; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 blgp:3833; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, 0834; GFX90A-VGPR-NEXT: s_nop 15835; GFX90A-VGPR-NEXT: s_nop 0836; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[8:9] offset:16837; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[8:9]838; GFX90A-VGPR-NEXT: s_endpgm839;840; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64:841; GFX942-VGPR: ; %bb.0: ; %bb842; GFX942-VGPR-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24843; GFX942-VGPR-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34844; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)845; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, s10846; GFX942-VGPR-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0847; GFX942-VGPR-NEXT: v_mov_b32_e32 v9, s11848; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[12:13]849; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)850; GFX942-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[0:1]851; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[2:3]852; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[4:5]853; GFX942-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[6:7]854; GFX942-VGPR-NEXT: s_nop 1855; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 neg:[1,1,0]856; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, 0857; GFX942-VGPR-NEXT: s_nop 15858; GFX942-VGPR-NEXT: s_nop 0859; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[8:9] offset:16860; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[8:9]861; GFX942-VGPR-NEXT: s_endpgm862bb:863 %in.1 = load <4 x double>, ptr addrspace(1) %arg864 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> %in.1, i32 1, i32 2, i32 3)865 store <4 x double> %mai.1, ptr addrspace(1) %arg866 ret void867}868 869define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_imm_0(ptr addrspace(1) %arg, double %a, double %b) #0 {870; GFX90A-LABEL: test_mfma_f64_16x16x4f64_splat_imm_0:871; GFX90A: ; %bb.0: ; %bb872; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24873; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34874; GFX90A-NEXT: s_waitcnt lgkmcnt(0)875; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]876; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]877; GFX90A-NEXT: s_nop 1878; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], 0879; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:3880; GFX90A-NEXT: v_mov_b32_e32 v0, 0881; GFX90A-NEXT: s_nop 15882; GFX90A-NEXT: s_nop 0883; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:16884; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]885; GFX90A-NEXT: s_endpgm886;887; GFX942-LABEL: test_mfma_f64_16x16x4f64_splat_imm_0:888; GFX942: ; %bb.0: ; %bb889; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24890; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34891; GFX942-NEXT: s_waitcnt lgkmcnt(0)892; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[2:3]893; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]894; GFX942-NEXT: s_nop 1895; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], 0896; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]897; GFX942-NEXT: v_mov_b32_e32 v0, 0898; GFX942-NEXT: s_nop 15899; GFX942-NEXT: s_nop 0900; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:16901; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]902; GFX942-NEXT: s_endpgm903;904; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_0:905; GFX90A-VGPR: ; %bb.0: ; %bb906; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24907; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34908; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)909; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[2:3], s[2:3] op_sel:[0,1]910; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]911; GFX90A-VGPR-NEXT: s_nop 1912; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], 0913; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 blgp:3914; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, 0915; GFX90A-VGPR-NEXT: s_nop 15916; GFX90A-VGPR-NEXT: s_nop 0917; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:16918; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]919; GFX90A-VGPR-NEXT: s_endpgm920;921; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_0:922; GFX942-VGPR: ; %bb.0: ; %bb923; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24924; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34925; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)926; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[2:3]927; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[6:7]928; GFX942-VGPR-NEXT: s_nop 1929; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], 0930; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 neg:[1,1,0]931; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, 0932; GFX942-VGPR-NEXT: s_nop 15933; GFX942-VGPR-NEXT: s_nop 0934; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:16935; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]936; GFX942-VGPR-NEXT: s_endpgm937bb:938 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> zeroinitializer, i32 0, i32 0, i32 0)939 %mai.2 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> %mai.1, i32 1, i32 2, i32 3)940 store <4 x double> %mai.2, ptr addrspace(1) %arg941 ret void942}943 944define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_imm_int_neg1(ptr addrspace(1) %arg, double %a, double %b) #0 {945; GFX90A-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_neg1:946; GFX90A: ; %bb.0: ; %bb947; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24948; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34949; GFX90A-NEXT: s_waitcnt lgkmcnt(0)950; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]951; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]952; GFX90A-NEXT: s_nop 1953; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], -1954; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:3955; GFX90A-NEXT: v_mov_b32_e32 v0, 0956; GFX90A-NEXT: s_nop 15957; GFX90A-NEXT: s_nop 0958; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:16959; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]960; GFX90A-NEXT: s_endpgm961;962; GFX942-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_neg1:963; GFX942: ; %bb.0: ; %bb964; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24965; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34966; GFX942-NEXT: s_waitcnt lgkmcnt(0)967; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[2:3]968; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]969; GFX942-NEXT: s_nop 1970; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], -1971; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]972; GFX942-NEXT: v_mov_b32_e32 v0, 0973; GFX942-NEXT: s_nop 15974; GFX942-NEXT: s_nop 0975; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:16976; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]977; GFX942-NEXT: s_endpgm978;979; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_neg1:980; GFX90A-VGPR: ; %bb.0: ; %bb981; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24982; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34983; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)984; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[2:3], s[2:3] op_sel:[0,1]985; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]986; GFX90A-VGPR-NEXT: s_nop 1987; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], -1988; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 blgp:3989; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, 0990; GFX90A-VGPR-NEXT: s_nop 15991; GFX90A-VGPR-NEXT: s_nop 0992; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:16993; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]994; GFX90A-VGPR-NEXT: s_endpgm995;996; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_neg1:997; GFX942-VGPR: ; %bb.0: ; %bb998; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24999; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341000; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)1001; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[2:3]1002; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[6:7]1003; GFX942-VGPR-NEXT: s_nop 11004; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], -11005; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 neg:[1,1,0]1006; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, 01007; GFX942-VGPR-NEXT: s_nop 151008; GFX942-VGPR-NEXT: s_nop 01009; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161010; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1011; GFX942-VGPR-NEXT: s_endpgm1012bb:1013 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> splat (double bitcast (i64 -1 to double)), i32 0, i32 0, i32 0)1014 %mai.2 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> %mai.1, i32 1, i32 2, i32 3)1015 store <4 x double> %mai.2, ptr addrspace(1) %arg1016 ret void1017}1018 1019define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_imm_1(ptr addrspace(1) %arg, double %a, double %b) #0 {1020; GFX90A-LABEL: test_mfma_f64_16x16x4f64_splat_imm_1:1021; GFX90A: ; %bb.0: ; %bb1022; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241023; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341024; GFX90A-NEXT: s_waitcnt lgkmcnt(0)1025; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]1026; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]1027; GFX90A-NEXT: s_nop 11028; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], 1.01029; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:31030; GFX90A-NEXT: v_mov_b32_e32 v0, 01031; GFX90A-NEXT: s_nop 151032; GFX90A-NEXT: s_nop 01033; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161034; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1035; GFX90A-NEXT: s_endpgm1036;1037; GFX942-LABEL: test_mfma_f64_16x16x4f64_splat_imm_1:1038; GFX942: ; %bb.0: ; %bb1039; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241040; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341041; GFX942-NEXT: s_waitcnt lgkmcnt(0)1042; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[2:3]1043; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]1044; GFX942-NEXT: s_nop 11045; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], 1.01046; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]1047; GFX942-NEXT: v_mov_b32_e32 v0, 01048; GFX942-NEXT: s_nop 151049; GFX942-NEXT: s_nop 01050; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161051; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1052; GFX942-NEXT: s_endpgm1053;1054; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_1:1055; GFX90A-VGPR: ; %bb.0: ; %bb1056; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241057; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341058; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)1059; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[2:3], s[2:3] op_sel:[0,1]1060; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]1061; GFX90A-VGPR-NEXT: s_nop 11062; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], 1.01063; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 blgp:31064; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, 01065; GFX90A-VGPR-NEXT: s_nop 151066; GFX90A-VGPR-NEXT: s_nop 01067; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161068; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1069; GFX90A-VGPR-NEXT: s_endpgm1070;1071; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_1:1072; GFX942-VGPR: ; %bb.0: ; %bb1073; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241074; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341075; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)1076; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[2:3]1077; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[6:7]1078; GFX942-VGPR-NEXT: s_nop 11079; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], 1.01080; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 neg:[1,1,0]1081; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, 01082; GFX942-VGPR-NEXT: s_nop 151083; GFX942-VGPR-NEXT: s_nop 01084; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161085; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1086; GFX942-VGPR-NEXT: s_endpgm1087bb:1088 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> splat (double 1.0), i32 0, i32 0, i32 0)1089 %mai.2 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> %mai.1, i32 1, i32 2, i32 3)1090 store <4 x double> %mai.2, ptr addrspace(1) %arg1091 ret void1092}1093 1094define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_imm_neg1(ptr addrspace(1) %arg, double %a, double %b) #0 {1095; GFX90A-LABEL: test_mfma_f64_16x16x4f64_splat_imm_neg1:1096; GFX90A: ; %bb.0: ; %bb1097; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241098; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341099; GFX90A-NEXT: s_waitcnt lgkmcnt(0)1100; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]1101; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]1102; GFX90A-NEXT: s_nop 11103; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], -1.01104; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:31105; GFX90A-NEXT: v_mov_b32_e32 v0, 01106; GFX90A-NEXT: s_nop 151107; GFX90A-NEXT: s_nop 01108; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161109; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1110; GFX90A-NEXT: s_endpgm1111;1112; GFX942-LABEL: test_mfma_f64_16x16x4f64_splat_imm_neg1:1113; GFX942: ; %bb.0: ; %bb1114; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241115; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341116; GFX942-NEXT: s_waitcnt lgkmcnt(0)1117; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[2:3]1118; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]1119; GFX942-NEXT: s_nop 11120; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], -1.01121; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]1122; GFX942-NEXT: v_mov_b32_e32 v0, 01123; GFX942-NEXT: s_nop 151124; GFX942-NEXT: s_nop 01125; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161126; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1127; GFX942-NEXT: s_endpgm1128;1129; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_neg1:1130; GFX90A-VGPR: ; %bb.0: ; %bb1131; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241132; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341133; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)1134; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[2:3], s[2:3] op_sel:[0,1]1135; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]1136; GFX90A-VGPR-NEXT: s_nop 11137; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], -1.01138; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 blgp:31139; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, 01140; GFX90A-VGPR-NEXT: s_nop 151141; GFX90A-VGPR-NEXT: s_nop 01142; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161143; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1144; GFX90A-VGPR-NEXT: s_endpgm1145;1146; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_neg1:1147; GFX942-VGPR: ; %bb.0: ; %bb1148; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241149; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341150; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)1151; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[2:3]1152; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[6:7]1153; GFX942-VGPR-NEXT: s_nop 11154; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], -1.01155; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 neg:[1,1,0]1156; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, 01157; GFX942-VGPR-NEXT: s_nop 151158; GFX942-VGPR-NEXT: s_nop 01159; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161160; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1161; GFX942-VGPR-NEXT: s_endpgm1162bb:1163 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> splat (double -1.0), i32 0, i32 0, i32 0)1164 %mai.2 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> %mai.1, i32 1, i32 2, i32 3)1165 store <4 x double> %mai.2, ptr addrspace(1) %arg1166 ret void1167}1168 1169define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_imm_int_64(ptr addrspace(1) %arg, double %a, double %b) #0 {1170; GFX90A-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64:1171; GFX90A: ; %bb.0: ; %bb1172; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241173; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341174; GFX90A-NEXT: s_waitcnt lgkmcnt(0)1175; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]1176; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]1177; GFX90A-NEXT: s_nop 11178; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], 641179; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:31180; GFX90A-NEXT: v_mov_b32_e32 v0, 01181; GFX90A-NEXT: s_nop 151182; GFX90A-NEXT: s_nop 01183; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161184; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1185; GFX90A-NEXT: s_endpgm1186;1187; GFX942-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64:1188; GFX942: ; %bb.0: ; %bb1189; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241190; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341191; GFX942-NEXT: s_waitcnt lgkmcnt(0)1192; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[2:3]1193; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]1194; GFX942-NEXT: s_nop 11195; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], 641196; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]1197; GFX942-NEXT: v_mov_b32_e32 v0, 01198; GFX942-NEXT: s_nop 151199; GFX942-NEXT: s_nop 01200; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161201; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1202; GFX942-NEXT: s_endpgm1203;1204; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64:1205; GFX90A-VGPR: ; %bb.0: ; %bb1206; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241207; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341208; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)1209; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[2:3], s[2:3] op_sel:[0,1]1210; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]1211; GFX90A-VGPR-NEXT: s_nop 11212; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], 641213; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 blgp:31214; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, 01215; GFX90A-VGPR-NEXT: s_nop 151216; GFX90A-VGPR-NEXT: s_nop 01217; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161218; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1219; GFX90A-VGPR-NEXT: s_endpgm1220;1221; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64:1222; GFX942-VGPR: ; %bb.0: ; %bb1223; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241224; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341225; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)1226; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[2:3]1227; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[6:7]1228; GFX942-VGPR-NEXT: s_nop 11229; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], 641230; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 neg:[1,1,0]1231; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, 01232; GFX942-VGPR-NEXT: s_nop 151233; GFX942-VGPR-NEXT: s_nop 01234; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161235; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1236; GFX942-VGPR-NEXT: s_endpgm1237bb:1238 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> splat (double bitcast (i64 64 to double)), i32 0, i32 0, i32 0)1239 %mai.2 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> %mai.1, i32 1, i32 2, i32 3)1240 store <4 x double> %mai.2, ptr addrspace(1) %arg1241 ret void1242}1243 1244define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_bits(ptr addrspace(1) %arg, double %a, double %b) #0 {1245; GFX90A-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_bits:1246; GFX90A: ; %bb.0: ; %bb1247; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241248; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341249; GFX90A-NEXT: v_accvgpr_write_b32 a0, 01250; GFX90A-NEXT: v_accvgpr_write_b32 a1, 641251; GFX90A-NEXT: v_accvgpr_mov_b32 a2, a01252; GFX90A-NEXT: s_waitcnt lgkmcnt(0)1253; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]1254; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a11255; GFX90A-NEXT: v_accvgpr_mov_b32 a4, a01256; GFX90A-NEXT: v_accvgpr_mov_b32 a5, a11257; GFX90A-NEXT: v_accvgpr_mov_b32 a6, a01258; GFX90A-NEXT: v_accvgpr_mov_b32 a7, a11259; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]1260; GFX90A-NEXT: s_nop 11261; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]1262; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:31263; GFX90A-NEXT: v_mov_b32_e32 v0, 01264; GFX90A-NEXT: s_nop 151265; GFX90A-NEXT: s_nop 01266; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161267; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1268; GFX90A-NEXT: s_endpgm1269;1270; GFX942-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_bits:1271; GFX942: ; %bb.0: ; %bb1272; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241273; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341274; GFX942-NEXT: v_accvgpr_write_b32 a0, 01275; GFX942-NEXT: v_accvgpr_write_b32 a1, 641276; GFX942-NEXT: v_accvgpr_mov_b32 a2, a01277; GFX942-NEXT: s_waitcnt lgkmcnt(0)1278; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[2:3]1279; GFX942-NEXT: v_accvgpr_mov_b32 a3, a11280; GFX942-NEXT: v_accvgpr_mov_b32 a4, a01281; GFX942-NEXT: v_accvgpr_mov_b32 a5, a11282; GFX942-NEXT: v_accvgpr_mov_b32 a6, a01283; GFX942-NEXT: v_accvgpr_mov_b32 a7, a11284; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]1285; GFX942-NEXT: s_nop 11286; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]1287; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]1288; GFX942-NEXT: v_mov_b32_e32 v0, 01289; GFX942-NEXT: s_nop 151290; GFX942-NEXT: s_nop 01291; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161292; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1293; GFX942-NEXT: s_endpgm1294;1295; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_bits:1296; GFX90A-VGPR: ; %bb.0: ; %bb1297; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241298; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341299; GFX90A-VGPR-NEXT: v_mov_b32_e32 v0, 01300; GFX90A-VGPR-NEXT: v_mov_b32_e32 v1, 641301; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, v01302; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)1303; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[2:3], s[2:3] op_sel:[0,1]1304; GFX90A-VGPR-NEXT: v_mov_b32_e32 v3, v11305; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, v01306; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, v11307; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, v01308; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, v11309; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], v[6:7], v[6:7] op_sel:[0,1]1310; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[12:13], s[6:7], s[6:7] op_sel:[0,1]1311; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[6:7], v[4:5], v[4:5] op_sel:[0,1]1312; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], v[2:3], v[2:3] op_sel:[0,1]1313; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], v[0:1], v[0:1] op_sel:[0,1]1314; GFX90A-VGPR-NEXT: s_nop 11315; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[2:9], v[10:11], v[12:13], v[2:9]1316; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[2:9], v[10:11], v[12:13], v[2:9] cbsz:1 abid:2 blgp:31317; GFX90A-VGPR-NEXT: s_nop 151318; GFX90A-VGPR-NEXT: s_nop 11319; GFX90A-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:161320; GFX90A-VGPR-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1]1321; GFX90A-VGPR-NEXT: s_endpgm1322;1323; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_bits:1324; GFX942-VGPR: ; %bb.0: ; %bb1325; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241326; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341327; GFX942-VGPR-NEXT: v_mov_b32_e32 v0, 01328; GFX942-VGPR-NEXT: v_mov_b32_e32 v1, 641329; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, v01330; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)1331; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[2:3]1332; GFX942-VGPR-NEXT: v_mov_b32_e32 v3, v11333; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, v01334; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, v11335; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, v01336; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, v11337; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], v[6:7]1338; GFX942-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[6:7]1339; GFX942-VGPR-NEXT: v_mov_b64_e32 v[6:7], v[4:5]1340; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], v[2:3]1341; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], v[0:1]1342; GFX942-VGPR-NEXT: s_nop 11343; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[2:9], v[10:11], v[12:13], v[2:9]1344; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[2:9], v[10:11], v[12:13], v[2:9] cbsz:1 abid:2 neg:[1,1,0]1345; GFX942-VGPR-NEXT: s_nop 151346; GFX942-VGPR-NEXT: s_nop 11347; GFX942-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:161348; GFX942-VGPR-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1]1349; GFX942-VGPR-NEXT: s_endpgm1350bb:1351 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> splat (double bitcast (i64 274877906944 to double)), i32 0, i32 0, i32 0)1352 %mai.2 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> %mai.1, i32 1, i32 2, i32 3)1353 store <4 x double> %mai.2, ptr addrspace(1) %arg1354 ret void1355}1356 1357define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_and_low(ptr addrspace(1) %arg, double %a, double %b) #0 {1358; GFX90A-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_and_low:1359; GFX90A: ; %bb.0: ; %bb1360; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241361; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341362; GFX90A-NEXT: v_accvgpr_write_b32 a0, 641363; GFX90A-NEXT: v_accvgpr_mov_b32 a1, a01364; GFX90A-NEXT: v_accvgpr_mov_b32 a2, a01365; GFX90A-NEXT: s_waitcnt lgkmcnt(0)1366; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]1367; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a01368; GFX90A-NEXT: v_accvgpr_mov_b32 a4, a01369; GFX90A-NEXT: v_accvgpr_mov_b32 a5, a01370; GFX90A-NEXT: v_accvgpr_mov_b32 a6, a01371; GFX90A-NEXT: v_accvgpr_mov_b32 a7, a01372; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]1373; GFX90A-NEXT: s_nop 11374; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]1375; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:31376; GFX90A-NEXT: v_mov_b32_e32 v0, 01377; GFX90A-NEXT: s_nop 151378; GFX90A-NEXT: s_nop 01379; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161380; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1381; GFX90A-NEXT: s_endpgm1382;1383; GFX942-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_and_low:1384; GFX942: ; %bb.0: ; %bb1385; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241386; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341387; GFX942-NEXT: v_accvgpr_write_b32 a0, 641388; GFX942-NEXT: v_accvgpr_mov_b32 a1, a01389; GFX942-NEXT: v_accvgpr_mov_b32 a2, a01390; GFX942-NEXT: s_waitcnt lgkmcnt(0)1391; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[2:3]1392; GFX942-NEXT: v_accvgpr_mov_b32 a3, a01393; GFX942-NEXT: v_accvgpr_mov_b32 a4, a01394; GFX942-NEXT: v_accvgpr_mov_b32 a5, a01395; GFX942-NEXT: v_accvgpr_mov_b32 a6, a01396; GFX942-NEXT: v_accvgpr_mov_b32 a7, a01397; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]1398; GFX942-NEXT: s_nop 11399; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]1400; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]1401; GFX942-NEXT: v_mov_b32_e32 v0, 01402; GFX942-NEXT: s_nop 151403; GFX942-NEXT: s_nop 01404; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161405; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1406; GFX942-NEXT: s_endpgm1407;1408; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_and_low:1409; GFX90A-VGPR: ; %bb.0: ; %bb1410; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241411; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341412; GFX90A-VGPR-NEXT: v_mov_b32_e32 v0, 641413; GFX90A-VGPR-NEXT: v_mov_b32_e32 v1, v01414; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, v01415; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)1416; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[2:3], s[2:3] op_sel:[0,1]1417; GFX90A-VGPR-NEXT: v_mov_b32_e32 v3, v01418; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, v01419; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, v01420; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, v01421; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, v01422; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]1423; GFX90A-VGPR-NEXT: s_nop 11424; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7]1425; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 blgp:31426; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, 01427; GFX90A-VGPR-NEXT: s_nop 151428; GFX90A-VGPR-NEXT: s_nop 01429; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161430; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1431; GFX90A-VGPR-NEXT: s_endpgm1432;1433; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_int_64_in_high_and_low:1434; GFX942-VGPR: ; %bb.0: ; %bb1435; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241436; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341437; GFX942-VGPR-NEXT: v_mov_b32_e32 v0, 641438; GFX942-VGPR-NEXT: v_mov_b32_e32 v1, v01439; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, v01440; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)1441; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[2:3]1442; GFX942-VGPR-NEXT: v_mov_b32_e32 v3, v01443; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, v01444; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, v01445; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, v01446; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, v01447; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[6:7]1448; GFX942-VGPR-NEXT: s_nop 11449; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7]1450; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 neg:[1,1,0]1451; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, 01452; GFX942-VGPR-NEXT: s_nop 151453; GFX942-VGPR-NEXT: s_nop 01454; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161455; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1456; GFX942-VGPR-NEXT: s_endpgm1457bb:1458 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> splat (double bitcast (i64 274877907008 to double)), i32 0, i32 0, i32 0)1459 %mai.2 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> %mai.1, i32 1, i32 2, i32 3)1460 store <4 x double> %mai.2, ptr addrspace(1) %arg1461 ret void1462}1463 1464define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_imm_f32_1_in_high_and_low(ptr addrspace(1) %arg, double %a, double %b) #0 {1465; GFX90A-LABEL: test_mfma_f64_16x16x4f64_splat_imm_f32_1_in_high_and_low:1466; GFX90A: ; %bb.0: ; %bb1467; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241468; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341469; GFX90A-NEXT: v_accvgpr_write_b32 a0, 1.01470; GFX90A-NEXT: v_accvgpr_mov_b32 a1, a01471; GFX90A-NEXT: v_accvgpr_mov_b32 a2, a01472; GFX90A-NEXT: s_waitcnt lgkmcnt(0)1473; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]1474; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a01475; GFX90A-NEXT: v_accvgpr_mov_b32 a4, a01476; GFX90A-NEXT: v_accvgpr_mov_b32 a5, a01477; GFX90A-NEXT: v_accvgpr_mov_b32 a6, a01478; GFX90A-NEXT: v_accvgpr_mov_b32 a7, a01479; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]1480; GFX90A-NEXT: s_nop 11481; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]1482; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:31483; GFX90A-NEXT: v_mov_b32_e32 v0, 01484; GFX90A-NEXT: s_nop 151485; GFX90A-NEXT: s_nop 01486; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161487; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1488; GFX90A-NEXT: s_endpgm1489;1490; GFX942-LABEL: test_mfma_f64_16x16x4f64_splat_imm_f32_1_in_high_and_low:1491; GFX942: ; %bb.0: ; %bb1492; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241493; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341494; GFX942-NEXT: v_accvgpr_write_b32 a0, 1.01495; GFX942-NEXT: v_accvgpr_mov_b32 a1, a01496; GFX942-NEXT: v_accvgpr_mov_b32 a2, a01497; GFX942-NEXT: s_waitcnt lgkmcnt(0)1498; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[2:3]1499; GFX942-NEXT: v_accvgpr_mov_b32 a3, a01500; GFX942-NEXT: v_accvgpr_mov_b32 a4, a01501; GFX942-NEXT: v_accvgpr_mov_b32 a5, a01502; GFX942-NEXT: v_accvgpr_mov_b32 a6, a01503; GFX942-NEXT: v_accvgpr_mov_b32 a7, a01504; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]1505; GFX942-NEXT: s_nop 11506; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]1507; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]1508; GFX942-NEXT: v_mov_b32_e32 v0, 01509; GFX942-NEXT: s_nop 151510; GFX942-NEXT: s_nop 01511; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161512; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1513; GFX942-NEXT: s_endpgm1514;1515; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_f32_1_in_high_and_low:1516; GFX90A-VGPR: ; %bb.0: ; %bb1517; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241518; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341519; GFX90A-VGPR-NEXT: v_mov_b32_e32 v0, 1.01520; GFX90A-VGPR-NEXT: v_mov_b32_e32 v1, v01521; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, v01522; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)1523; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], s[2:3], s[2:3] op_sel:[0,1]1524; GFX90A-VGPR-NEXT: v_mov_b32_e32 v3, v01525; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, v01526; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, v01527; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, v01528; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, v01529; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]1530; GFX90A-VGPR-NEXT: s_nop 11531; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7]1532; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 blgp:31533; GFX90A-VGPR-NEXT: v_mov_b32_e32 v8, 01534; GFX90A-VGPR-NEXT: s_nop 151535; GFX90A-VGPR-NEXT: s_nop 01536; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161537; GFX90A-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1538; GFX90A-VGPR-NEXT: s_endpgm1539;1540; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_imm_f32_1_in_high_and_low:1541; GFX942-VGPR: ; %bb.0: ; %bb1542; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241543; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341544; GFX942-VGPR-NEXT: v_mov_b32_e32 v0, 1.01545; GFX942-VGPR-NEXT: v_mov_b32_e32 v1, v01546; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, v01547; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)1548; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[2:3]1549; GFX942-VGPR-NEXT: v_mov_b32_e32 v3, v01550; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, v01551; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, v01552; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, v01553; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, v01554; GFX942-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[6:7]1555; GFX942-VGPR-NEXT: s_nop 11556; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7]1557; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[0:7], v[8:9], v[10:11], v[0:7] cbsz:1 abid:2 neg:[1,1,0]1558; GFX942-VGPR-NEXT: v_mov_b32_e32 v8, 01559; GFX942-VGPR-NEXT: s_nop 151560; GFX942-VGPR-NEXT: s_nop 01561; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[4:7], s[0:1] offset:161562; GFX942-VGPR-NEXT: global_store_dwordx4 v8, v[0:3], s[0:1]1563; GFX942-VGPR-NEXT: s_endpgm1564bb:1565 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> splat (double bitcast (<2 x float> splat (float 1.0) to double)), i32 0, i32 0, i32 0)1566 %mai.2 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> %mai.1, i32 1, i32 2, i32 3)1567 store <4 x double> %mai.2, ptr addrspace(1) %arg1568 ret void1569}1570 1571define amdgpu_kernel void @test_mfma_f64_16x16x4f64_imm(ptr addrspace(1) %arg, double %a, double %b) #0 {1572; GFX90A-LABEL: test_mfma_f64_16x16x4f64_imm:1573; GFX90A: ; %bb.0: ; %bb1574; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241575; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341576; GFX90A-NEXT: v_accvgpr_write_b32 a0, 01577; GFX90A-NEXT: v_mov_b32_e32 v0, 0x3ff000001578; GFX90A-NEXT: v_accvgpr_write_b32 a7, v01579; GFX90A-NEXT: s_waitcnt lgkmcnt(0)1580; GFX90A-NEXT: v_mov_b32_e32 v2, s21581; GFX90A-NEXT: v_mov_b32_e32 v3, s31582; GFX90A-NEXT: v_accvgpr_mov_b32 a1, a01583; GFX90A-NEXT: v_accvgpr_mov_b32 a2, a01584; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a01585; GFX90A-NEXT: v_accvgpr_mov_b32 a4, a01586; GFX90A-NEXT: v_accvgpr_mov_b32 a5, a01587; GFX90A-NEXT: v_accvgpr_mov_b32 a6, a01588; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[6:7], s[6:7] op_sel:[0,1]1589; GFX90A-NEXT: s_nop 11590; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[2:3], v[0:1], a[0:7]1591; GFX90A-NEXT: v_mov_b32_e32 v0, 01592; GFX90A-NEXT: s_nop 151593; GFX90A-NEXT: s_nop 01594; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161595; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1596; GFX90A-NEXT: s_endpgm1597;1598; GFX942-LABEL: test_mfma_f64_16x16x4f64_imm:1599; GFX942: ; %bb.0: ; %bb1600; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241601; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341602; GFX942-NEXT: v_accvgpr_write_b32 a0, 01603; GFX942-NEXT: v_mov_b32_e32 v0, 0x3ff000001604; GFX942-NEXT: v_accvgpr_write_b32 a7, v01605; GFX942-NEXT: s_waitcnt lgkmcnt(0)1606; GFX942-NEXT: v_mov_b32_e32 v2, s21607; GFX942-NEXT: v_mov_b32_e32 v3, s31608; GFX942-NEXT: v_accvgpr_mov_b32 a1, a01609; GFX942-NEXT: v_accvgpr_mov_b32 a2, a01610; GFX942-NEXT: v_accvgpr_mov_b32 a3, a01611; GFX942-NEXT: v_accvgpr_mov_b32 a4, a01612; GFX942-NEXT: v_accvgpr_mov_b32 a5, a01613; GFX942-NEXT: v_accvgpr_mov_b32 a6, a01614; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[6:7]1615; GFX942-NEXT: s_nop 11616; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[2:3], v[0:1], a[0:7]1617; GFX942-NEXT: v_mov_b32_e32 v0, 01618; GFX942-NEXT: s_nop 151619; GFX942-NEXT: s_nop 01620; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161621; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1622; GFX942-NEXT: s_endpgm1623;1624; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_imm:1625; GFX90A-VGPR: ; %bb.0: ; %bb1626; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241627; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341628; GFX90A-VGPR-NEXT: v_mov_b32_e32 v0, 01629; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, 0x3ff000001630; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, v01631; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)1632; GFX90A-VGPR-NEXT: v_mov_b32_e32 v10, s21633; GFX90A-VGPR-NEXT: v_mov_b32_e32 v11, s31634; GFX90A-VGPR-NEXT: v_mov_b32_e32 v3, v01635; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, v01636; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, v01637; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, v01638; GFX90A-VGPR-NEXT: v_mov_b32_e32 v1, v01639; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], v[6:7], v[6:7] op_sel:[0,1]1640; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[12:13], s[6:7], s[6:7] op_sel:[0,1]1641; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[6:7], v[4:5], v[4:5] op_sel:[0,1]1642; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], v[2:3], v[2:3] op_sel:[0,1]1643; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], v[0:1], v[0:1] op_sel:[0,1]1644; GFX90A-VGPR-NEXT: s_nop 11645; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[2:9], v[10:11], v[12:13], v[2:9]1646; GFX90A-VGPR-NEXT: s_nop 151647; GFX90A-VGPR-NEXT: s_nop 11648; GFX90A-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:161649; GFX90A-VGPR-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1]1650; GFX90A-VGPR-NEXT: s_endpgm1651;1652; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_imm:1653; GFX942-VGPR: ; %bb.0: ; %bb1654; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241655; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341656; GFX942-VGPR-NEXT: v_mov_b32_e32 v0, 01657; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, 0x3ff000001658; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, v01659; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)1660; GFX942-VGPR-NEXT: v_mov_b32_e32 v10, s21661; GFX942-VGPR-NEXT: v_mov_b32_e32 v11, s31662; GFX942-VGPR-NEXT: v_mov_b32_e32 v3, v01663; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, v01664; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, v01665; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, v01666; GFX942-VGPR-NEXT: v_mov_b32_e32 v1, v01667; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], v[6:7]1668; GFX942-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[6:7]1669; GFX942-VGPR-NEXT: v_mov_b64_e32 v[6:7], v[4:5]1670; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], v[2:3]1671; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], v[0:1]1672; GFX942-VGPR-NEXT: s_nop 11673; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[2:9], v[10:11], v[12:13], v[2:9]1674; GFX942-VGPR-NEXT: s_nop 151675; GFX942-VGPR-NEXT: s_nop 11676; GFX942-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:161677; GFX942-VGPR-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1]1678; GFX942-VGPR-NEXT: s_endpgm1679bb:1680 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> <double 0.0, double 0.0, double 0.0, double 1.0>, i32 0, i32 0, i32 0)1681 store <4 x double> %mai.1, ptr addrspace(1) %arg1682 ret void1683}1684 1685define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_lit(ptr addrspace(1) %arg, double %a, double %b) #0 {1686; GFX90A-LABEL: test_mfma_f64_16x16x4f64_splat_lit:1687; GFX90A: ; %bb.0: ; %bb1688; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241689; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341690; GFX90A-NEXT: v_mov_b32_e32 v0, 0x405ec0001691; GFX90A-NEXT: v_accvgpr_write_b32 a0, 01692; GFX90A-NEXT: v_accvgpr_write_b32 a1, v01693; GFX90A-NEXT: s_waitcnt lgkmcnt(0)1694; GFX90A-NEXT: v_mov_b32_e32 v2, s21695; GFX90A-NEXT: v_mov_b32_e32 v3, s31696; GFX90A-NEXT: v_accvgpr_mov_b32 a2, a01697; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a11698; GFX90A-NEXT: v_accvgpr_mov_b32 a4, a01699; GFX90A-NEXT: v_accvgpr_mov_b32 a5, a11700; GFX90A-NEXT: v_accvgpr_mov_b32 a6, a01701; GFX90A-NEXT: v_accvgpr_mov_b32 a7, a11702; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[6:7], s[6:7] op_sel:[0,1]1703; GFX90A-NEXT: s_nop 11704; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[2:3], v[0:1], a[0:7]1705; GFX90A-NEXT: v_mov_b32_e32 v0, 01706; GFX90A-NEXT: s_nop 151707; GFX90A-NEXT: s_nop 01708; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161709; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1710; GFX90A-NEXT: s_endpgm1711;1712; GFX942-LABEL: test_mfma_f64_16x16x4f64_splat_lit:1713; GFX942: ; %bb.0: ; %bb1714; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241715; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341716; GFX942-NEXT: v_mov_b32_e32 v0, 0x405ec0001717; GFX942-NEXT: v_accvgpr_write_b32 a0, 01718; GFX942-NEXT: v_accvgpr_write_b32 a1, v01719; GFX942-NEXT: s_waitcnt lgkmcnt(0)1720; GFX942-NEXT: v_mov_b32_e32 v2, s21721; GFX942-NEXT: v_mov_b32_e32 v3, s31722; GFX942-NEXT: v_accvgpr_mov_b32 a2, a01723; GFX942-NEXT: v_accvgpr_mov_b32 a3, a11724; GFX942-NEXT: v_accvgpr_mov_b32 a4, a01725; GFX942-NEXT: v_accvgpr_mov_b32 a5, a11726; GFX942-NEXT: v_accvgpr_mov_b32 a6, a01727; GFX942-NEXT: v_accvgpr_mov_b32 a7, a11728; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[6:7]1729; GFX942-NEXT: s_nop 11730; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[2:3], v[0:1], a[0:7]1731; GFX942-NEXT: v_mov_b32_e32 v0, 01732; GFX942-NEXT: s_nop 151733; GFX942-NEXT: s_nop 01734; GFX942-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:161735; GFX942-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1736; GFX942-NEXT: s_endpgm1737;1738; GFX90A-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_lit:1739; GFX90A-VGPR: ; %bb.0: ; %bb1740; GFX90A-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241741; GFX90A-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341742; GFX90A-VGPR-NEXT: v_mov_b32_e32 v0, 01743; GFX90A-VGPR-NEXT: v_mov_b32_e32 v1, 0x405ec0001744; GFX90A-VGPR-NEXT: v_mov_b32_e32 v2, v01745; GFX90A-VGPR-NEXT: s_waitcnt lgkmcnt(0)1746; GFX90A-VGPR-NEXT: v_mov_b32_e32 v10, s21747; GFX90A-VGPR-NEXT: v_mov_b32_e32 v11, s31748; GFX90A-VGPR-NEXT: v_mov_b32_e32 v3, v11749; GFX90A-VGPR-NEXT: v_mov_b32_e32 v4, v01750; GFX90A-VGPR-NEXT: v_mov_b32_e32 v5, v11751; GFX90A-VGPR-NEXT: v_mov_b32_e32 v6, v01752; GFX90A-VGPR-NEXT: v_mov_b32_e32 v7, v11753; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[8:9], v[6:7], v[6:7] op_sel:[0,1]1754; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[12:13], s[6:7], s[6:7] op_sel:[0,1]1755; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[6:7], v[4:5], v[4:5] op_sel:[0,1]1756; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[4:5], v[2:3], v[2:3] op_sel:[0,1]1757; GFX90A-VGPR-NEXT: v_pk_mov_b32 v[2:3], v[0:1], v[0:1] op_sel:[0,1]1758; GFX90A-VGPR-NEXT: s_nop 11759; GFX90A-VGPR-NEXT: v_mfma_f64_16x16x4f64 v[2:9], v[10:11], v[12:13], v[2:9]1760; GFX90A-VGPR-NEXT: s_nop 151761; GFX90A-VGPR-NEXT: s_nop 11762; GFX90A-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:161763; GFX90A-VGPR-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1]1764; GFX90A-VGPR-NEXT: s_endpgm1765;1766; GFX942-VGPR-LABEL: test_mfma_f64_16x16x4f64_splat_lit:1767; GFX942-VGPR: ; %bb.0: ; %bb1768; GFX942-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x241769; GFX942-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x341770; GFX942-VGPR-NEXT: v_mov_b32_e32 v0, 01771; GFX942-VGPR-NEXT: v_mov_b32_e32 v1, 0x405ec0001772; GFX942-VGPR-NEXT: v_mov_b32_e32 v2, v01773; GFX942-VGPR-NEXT: s_waitcnt lgkmcnt(0)1774; GFX942-VGPR-NEXT: v_mov_b32_e32 v10, s21775; GFX942-VGPR-NEXT: v_mov_b32_e32 v11, s31776; GFX942-VGPR-NEXT: v_mov_b32_e32 v3, v11777; GFX942-VGPR-NEXT: v_mov_b32_e32 v4, v01778; GFX942-VGPR-NEXT: v_mov_b32_e32 v5, v11779; GFX942-VGPR-NEXT: v_mov_b32_e32 v6, v01780; GFX942-VGPR-NEXT: v_mov_b32_e32 v7, v11781; GFX942-VGPR-NEXT: v_mov_b64_e32 v[8:9], v[6:7]1782; GFX942-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[6:7]1783; GFX942-VGPR-NEXT: v_mov_b64_e32 v[6:7], v[4:5]1784; GFX942-VGPR-NEXT: v_mov_b64_e32 v[4:5], v[2:3]1785; GFX942-VGPR-NEXT: v_mov_b64_e32 v[2:3], v[0:1]1786; GFX942-VGPR-NEXT: s_nop 11787; GFX942-VGPR-NEXT: v_mfma_f64_16x16x4_f64 v[2:9], v[10:11], v[12:13], v[2:9]1788; GFX942-VGPR-NEXT: s_nop 151789; GFX942-VGPR-NEXT: s_nop 11790; GFX942-VGPR-NEXT: global_store_dwordx4 v0, v[6:9], s[0:1] offset:161791; GFX942-VGPR-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1]1792; GFX942-VGPR-NEXT: s_endpgm1793bb:1794 %mai.1 = tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %a, double %b, <4 x double> <double 123.0, double 123.0, double 123.0, double 123.0>, i32 0, i32 0, i32 0)1795 store <4 x double> %mai.1, ptr addrspace(1) %arg1796 ret void1797}1798 1799attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }1800;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:1801; GCN: {{.*}}1802; VGPR: {{.*}}1803