brintos / llvm-project-archived public Read only

1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -verify-machineinstrs -mcpu=gfx942 -amdgpu-mfma-vgpr-form < %s | FileCheck %s3 4target triple = "amdgcn-amd-amdhsa"5 6define amdgpu_kernel void @respect_optnone(double %arg0, double %arg1, ptr addrspace(1) %ptr) #4 {7; CHECK-LABEL: respect_optnone:8; CHECK:       ; %bb.0: ; %bb9; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x010; CHECK-NEXT:    s_load_dwordx2 s[2:3], s[4:5], 0x811; CHECK-NEXT:    s_nop 012; CHECK-NEXT:    s_load_dwordx2 s[4:5], s[4:5], 0x1013; CHECK-NEXT:    s_mov_b32 s6, 0x3ff14; CHECK-NEXT:    v_and_b32_e64 v0, v0, s615; CHECK-NEXT:    s_mov_b32 s6, 316; CHECK-NEXT:    v_lshlrev_b32_e64 v0, s6, v017; CHECK-NEXT:    s_waitcnt lgkmcnt(0)18; CHECK-NEXT:    global_load_dwordx2 v[0:1], v0, s[4:5]19; CHECK-NEXT:    v_mov_b64_e32 v[2:3], s[0:1]20; CHECK-NEXT:    v_mov_b64_e32 v[4:5], s[2:3]21; CHECK-NEXT:    s_waitcnt vmcnt(0)22; CHECK-NEXT:    s_nop 023; CHECK-NEXT:    v_mfma_f64_4x4x4_4b_f64 v[0:1], v[2:3], v[4:5], v[0:1]24; CHECK-NEXT:    s_nop 525; CHECK-NEXT:    v_accvgpr_write_b32 a0, v026; CHECK-NEXT:    v_accvgpr_write_b32 a1, v127; CHECK-NEXT:    ;;#ASMSTART28; CHECK-NEXT:    ; use a[0:1]29; CHECK-NEXT:    ;;#ASMEND30; CHECK-NEXT:    s_endpgm31bb:32  %id = call i32 @llvm.amdgcn.workitem.id.x()33  %gep = getelementptr double, ptr addrspace(1) %ptr, i32 %id34  %src2 = load double, ptr addrspace(1) %gep35  %mai = call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %arg0, double %arg1, double %src2, i32 0, i32 0, i32 0)36  call void asm sideeffect "; use $0", "a"(double %mai)37  ret void38}39 40define amdgpu_kernel void @test_mfma_f32_32x32x1f32_rewrite_vgpr_mfma(ptr addrspace(1) %arg) #0 {41; CHECK-LABEL: test_mfma_f32_32x32x1f32_rewrite_vgpr_mfma:42; CHECK:       ; %bb.0: ; %bb43; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x044; CHECK-NEXT:    v_and_b32_e32 v0, 0x3ff, v045; CHECK-NEXT:    v_lshlrev_b32_e32 v0, 7, v046; CHECK-NEXT:    v_mov_b32_e32 v32, 1.047; CHECK-NEXT:    v_mov_b32_e32 v33, 2.048; CHECK-NEXT:    s_waitcnt lgkmcnt(0)49; CHECK-NEXT:    global_load_dwordx4 v[28:31], v0, s[0:1] offset:11250; CHECK-NEXT:    global_load_dwordx4 v[24:27], v0, s[0:1] offset:9651; CHECK-NEXT:    global_load_dwordx4 v[20:23], v0, s[0:1] offset:8052; CHECK-NEXT:    global_load_dwordx4 v[16:19], v0, s[0:1] offset:6453; CHECK-NEXT:    global_load_dwordx4 v[12:15], v0, s[0:1] offset:4854; CHECK-NEXT:    global_load_dwordx4 v[8:11], v0, s[0:1] offset:3255; CHECK-NEXT:    global_load_dwordx4 v[4:7], v0, s[0:1] offset:1656; CHECK-NEXT:    s_nop 057; CHECK-NEXT:    global_load_dwordx4 v[0:3], v0, s[0:1]58; CHECK-NEXT:    v_accvgpr_write_b32 a0, 1.059; CHECK-NEXT:    v_accvgpr_write_b32 a1, 2.060; CHECK-NEXT:    s_waitcnt vmcnt(0)61; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, v[0:31]62; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[32:63], a0, a1, v[0:31]63; CHECK-NEXT:    s_nop 1564; CHECK-NEXT:    s_nop 165; CHECK-NEXT:    v_mov_b32_e32 v2, v3266; CHECK-NEXT:    v_mov_b32_e32 v3, v3367; CHECK-NEXT:    v_mov_b32_e32 v4, v3468; CHECK-NEXT:    v_mov_b32_e32 v5, v3569; CHECK-NEXT:    v_mov_b32_e32 v6, v3670; CHECK-NEXT:    v_mov_b32_e32 v7, v3771; CHECK-NEXT:    v_mov_b32_e32 v8, v3872; CHECK-NEXT:    v_mov_b32_e32 v9, v3973; CHECK-NEXT:    v_mov_b32_e32 v10, v4074; CHECK-NEXT:    v_mov_b32_e32 v11, v4175; CHECK-NEXT:    v_mov_b32_e32 v12, v4276; CHECK-NEXT:    v_mov_b32_e32 v13, v4377; CHECK-NEXT:    v_mov_b32_e32 v14, v4478; CHECK-NEXT:    v_mov_b32_e32 v15, v4579; CHECK-NEXT:    v_mov_b32_e32 v16, v4680; CHECK-NEXT:    v_mov_b32_e32 v17, v4781; CHECK-NEXT:    v_mov_b32_e32 v18, v4882; CHECK-NEXT:    v_mov_b32_e32 v19, v4983; CHECK-NEXT:    v_mov_b32_e32 v20, v5084; CHECK-NEXT:    v_mov_b32_e32 v21, v5185; CHECK-NEXT:    v_mov_b32_e32 v22, v5286; CHECK-NEXT:    v_mov_b32_e32 v23, v5387; CHECK-NEXT:    v_mov_b32_e32 v24, v5488; CHECK-NEXT:    v_mov_b32_e32 v25, v5589; CHECK-NEXT:    v_mov_b32_e32 v26, v5690; CHECK-NEXT:    v_mov_b32_e32 v27, v5791; CHECK-NEXT:    v_mov_b32_e32 v28, v5892; CHECK-NEXT:    v_mov_b32_e32 v29, v5993; CHECK-NEXT:    v_mov_b32_e32 v30, v6094; CHECK-NEXT:    v_mov_b32_e32 v31, v6195; CHECK-NEXT:    v_mov_b32_e32 v32, 096; CHECK-NEXT:    s_nop 097; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], a0, a1, v[0:31]98; CHECK-NEXT:    s_nop 1599; CHECK-NEXT:    s_nop 1100; CHECK-NEXT:    global_store_dwordx4 v32, v[24:27], s[0:1] offset:96101; CHECK-NEXT:    global_store_dwordx4 v32, v[28:31], s[0:1] offset:112102; CHECK-NEXT:    global_store_dwordx4 v32, v[16:19], s[0:1] offset:64103; CHECK-NEXT:    global_store_dwordx4 v32, v[20:23], s[0:1] offset:80104; CHECK-NEXT:    global_store_dwordx4 v32, v[8:11], s[0:1] offset:32105; CHECK-NEXT:    global_store_dwordx4 v32, v[12:15], s[0:1] offset:48106; CHECK-NEXT:    global_store_dwordx4 v32, v[0:3], s[0:1]107; CHECK-NEXT:    global_store_dwordx4 v32, v[4:7], s[0:1] offset:16108; CHECK-NEXT:    s_endpgm109bb:110  %id = call i32 @llvm.amdgcn.workitem.id.x()111  %gep = getelementptr <32 x float>, ptr addrspace(1) %arg, i32 %id112  %in.1 = load <32 x float>, ptr addrspace(1) %gep, align 128113  %mai.1 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %in.1, i32 0, i32 0, i32 0)114  %mai.2 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %mai.1, i32 0, i32 0, i32 0)115  %tmp.1 = shufflevector <32 x float> %mai.2, <32 x float> %mai.1, <32 x i32> <i32 32, i32 33, i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15, i32 16, i32 17, i32 18, i32 19, i32 20, i32 21, i32 22, i32 23, i32 24, i32 25, i32 26, i32 27, i32 28, i32 29>116  %mai.3 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %tmp.1, i32 0, i32 0, i32 0)117  store <32 x float> %mai.3, ptr addrspace(1) %arg, align 128118  ret void119}120 121define amdgpu_kernel void @test_mfma_f32_32x32x1f32_rewrite_vgpr_mfma_noshuffle(ptr addrspace(1) %arg) #0 {122; CHECK-LABEL: test_mfma_f32_32x32x1f32_rewrite_vgpr_mfma_noshuffle:123; CHECK:       ; %bb.0: ; %bb124; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0125; CHECK-NEXT:    v_and_b32_e32 v0, 0x3ff, v0126; CHECK-NEXT:    v_lshlrev_b32_e32 v0, 7, v0127; CHECK-NEXT:    v_mov_b32_e32 v32, 1.0128; CHECK-NEXT:    v_mov_b32_e32 v33, 2.0129; CHECK-NEXT:    s_waitcnt lgkmcnt(0)130; CHECK-NEXT:    global_load_dwordx4 v[28:31], v0, s[0:1] offset:112131; CHECK-NEXT:    global_load_dwordx4 v[24:27], v0, s[0:1] offset:96132; CHECK-NEXT:    global_load_dwordx4 v[20:23], v0, s[0:1] offset:80133; CHECK-NEXT:    global_load_dwordx4 v[16:19], v0, s[0:1] offset:64134; CHECK-NEXT:    global_load_dwordx4 v[12:15], v0, s[0:1] offset:48135; CHECK-NEXT:    global_load_dwordx4 v[8:11], v0, s[0:1] offset:32136; CHECK-NEXT:    global_load_dwordx4 v[4:7], v0, s[0:1] offset:16137; CHECK-NEXT:    s_nop 0138; CHECK-NEXT:    global_load_dwordx4 v[0:3], v0, s[0:1]139; CHECK-NEXT:    s_waitcnt vmcnt(0)140; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, v[0:31]141; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, v[0:31]142; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, v[0:31]143; CHECK-NEXT:    v_mov_b32_e32 v32, 0144; CHECK-NEXT:    s_nop 15145; CHECK-NEXT:    s_nop 0146; CHECK-NEXT:    global_store_dwordx4 v32, v[24:27], s[0:1] offset:96147; CHECK-NEXT:    global_store_dwordx4 v32, v[28:31], s[0:1] offset:112148; CHECK-NEXT:    global_store_dwordx4 v32, v[16:19], s[0:1] offset:64149; CHECK-NEXT:    global_store_dwordx4 v32, v[20:23], s[0:1] offset:80150; CHECK-NEXT:    global_store_dwordx4 v32, v[8:11], s[0:1] offset:32151; CHECK-NEXT:    global_store_dwordx4 v32, v[12:15], s[0:1] offset:48152; CHECK-NEXT:    global_store_dwordx4 v32, v[0:3], s[0:1]153; CHECK-NEXT:    global_store_dwordx4 v32, v[4:7], s[0:1] offset:16154; CHECK-NEXT:    s_endpgm155bb:156  %id = call i32 @llvm.amdgcn.workitem.id.x()157  %gep = getelementptr <32 x float>, ptr addrspace(1) %arg, i32 %id158  %in.1 = load <32 x float>, ptr addrspace(1) %gep, align 128159  %mai.1 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %in.1, i32 0, i32 0, i32 0)160  %mai.2 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %mai.1, i32 0, i32 0, i32 0)161  %mai.3 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %mai.2, i32 0, i32 0, i32 0)162  store <32 x float> %mai.3, ptr addrspace(1) %arg, align 128163  ret void164}165 166define amdgpu_kernel void @test_mfma_f32_32x32x1f32_rewrite_vgpr_mfma_imm0_src2(ptr addrspace(1) %arg) #0 {167; CHECK-LABEL: test_mfma_f32_32x32x1f32_rewrite_vgpr_mfma_imm0_src2:168; CHECK:       ; %bb.0: ; %bb169; CHECK-NEXT:    v_mov_b32_e32 v32, 1.0170; CHECK-NEXT:    v_mov_b32_e32 v33, 2.0171; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0172; CHECK-NEXT:    s_nop 0173; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, 0174; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, v[0:31]175; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, v[0:31]176; CHECK-NEXT:    v_mov_b32_e32 v32, 0177; CHECK-NEXT:    s_waitcnt lgkmcnt(0)178; CHECK-NEXT:    s_nop 15179; CHECK-NEXT:    global_store_dwordx4 v32, v[28:31], s[0:1] offset:112180; CHECK-NEXT:    global_store_dwordx4 v32, v[24:27], s[0:1] offset:96181; CHECK-NEXT:    global_store_dwordx4 v32, v[20:23], s[0:1] offset:80182; CHECK-NEXT:    global_store_dwordx4 v32, v[16:19], s[0:1] offset:64183; CHECK-NEXT:    global_store_dwordx4 v32, v[12:15], s[0:1] offset:48184; CHECK-NEXT:    global_store_dwordx4 v32, v[8:11], s[0:1] offset:32185; CHECK-NEXT:    global_store_dwordx4 v32, v[4:7], s[0:1] offset:16186; CHECK-NEXT:    global_store_dwordx4 v32, v[0:3], s[0:1]187; CHECK-NEXT:    s_endpgm188bb:189  %id = call i32 @llvm.amdgcn.workitem.id.x()190  %gep = getelementptr <32 x float>, ptr addrspace(1) %arg, i32 %id191  %in.1 = load <32 x float>, ptr addrspace(1) %gep, align 128192  %mai.1 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> zeroinitializer, i32 0, i32 0, i32 0)193  %mai.2 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %mai.1, i32 0, i32 0, i32 0)194  %mai.3 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %mai.2, i32 0, i32 0, i32 0)195  store <32 x float> %mai.3, ptr addrspace(1) %arg, align 128196  ret void197}198 199define amdgpu_kernel void @test_mfma_f32_32x32x1f32_rewrite_vgpr_mfma_imm1_src2(ptr addrspace(1) %arg) #0 {200; CHECK-LABEL: test_mfma_f32_32x32x1f32_rewrite_vgpr_mfma_imm1_src2:201; CHECK:       ; %bb.0: ; %bb202; CHECK-NEXT:    v_mov_b32_e32 v32, 1.0203; CHECK-NEXT:    v_mov_b32_e32 v33, 2.0204; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0205; CHECK-NEXT:    s_nop 0206; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, 1.0207; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, v[0:31]208; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[0:31], v32, v33, v[0:31]209; CHECK-NEXT:    v_mov_b32_e32 v32, 0210; CHECK-NEXT:    s_waitcnt lgkmcnt(0)211; CHECK-NEXT:    s_nop 15212; CHECK-NEXT:    global_store_dwordx4 v32, v[28:31], s[0:1] offset:112213; CHECK-NEXT:    global_store_dwordx4 v32, v[24:27], s[0:1] offset:96214; CHECK-NEXT:    global_store_dwordx4 v32, v[20:23], s[0:1] offset:80215; CHECK-NEXT:    global_store_dwordx4 v32, v[16:19], s[0:1] offset:64216; CHECK-NEXT:    global_store_dwordx4 v32, v[12:15], s[0:1] offset:48217; CHECK-NEXT:    global_store_dwordx4 v32, v[8:11], s[0:1] offset:32218; CHECK-NEXT:    global_store_dwordx4 v32, v[4:7], s[0:1] offset:16219; CHECK-NEXT:    global_store_dwordx4 v32, v[0:3], s[0:1]220; CHECK-NEXT:    s_endpgm221bb:222  %id = call i32 @llvm.amdgcn.workitem.id.x()223  %gep = getelementptr <32 x float>, ptr addrspace(1) %arg, i32 %id224  %in.1 = load <32 x float>, ptr addrspace(1) %gep, align 128225  %mai.1 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> splat (float 1.000000e+00), i32 0, i32 0, i32 0)226  %mai.2 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %mai.1, i32 0, i32 0, i32 0)227  %mai.3 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %mai.2, i32 0, i32 0, i32 0)228  store <32 x float> %mai.3, ptr addrspace(1) %arg, align 128229  ret void230}231 232; The inline asm requires the value be copied to an AGPR class, not233; the AV_* pseudo we usually expect for register allocator live range234; splits.235define amdgpu_kernel void @test_rewrite_mfma_direct_copy_to_agpr_class(ptr addrspace(1) %arg) #0 {236; CHECK-LABEL: test_rewrite_mfma_direct_copy_to_agpr_class:237; CHECK:       ; %bb.0: ; %bb238; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0239; CHECK-NEXT:    v_and_b32_e32 v0, 0x3ff, v0240; CHECK-NEXT:    v_lshlrev_b32_e32 v0, 7, v0241; CHECK-NEXT:    v_mov_b32_e32 v32, 2.0242; CHECK-NEXT:    v_mov_b32_e32 v33, 4.0243; CHECK-NEXT:    s_waitcnt lgkmcnt(0)244; CHECK-NEXT:    global_load_dwordx4 a[28:31], v0, s[0:1] offset:112245; CHECK-NEXT:    global_load_dwordx4 a[24:27], v0, s[0:1] offset:96246; CHECK-NEXT:    global_load_dwordx4 a[20:23], v0, s[0:1] offset:80247; CHECK-NEXT:    global_load_dwordx4 a[16:19], v0, s[0:1] offset:64248; CHECK-NEXT:    global_load_dwordx4 a[12:15], v0, s[0:1] offset:48249; CHECK-NEXT:    global_load_dwordx4 a[8:11], v0, s[0:1] offset:32250; CHECK-NEXT:    global_load_dwordx4 a[4:7], v0, s[0:1] offset:16251; CHECK-NEXT:    global_load_dwordx4 a[0:3], v0, s[0:1]252; CHECK-NEXT:    s_waitcnt vmcnt(0)253; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v32, v33, a[0:31]254; CHECK-NEXT:    ;;#ASMSTART255; CHECK-NEXT:    ; use a[0:31]256; CHECK-NEXT:    ;;#ASMEND257; CHECK-NEXT:    s_endpgm258bb:259  %id = call i32 @llvm.amdgcn.workitem.id.x()260  %gep = getelementptr <32 x float>, ptr addrspace(1) %arg, i32 %id261  %in = load <32 x float>, ptr addrspace(1) %gep, align 128262  %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 2.0, float 4.0, <32 x float> %in, i32 0, i32 0, i32 0)263  call void asm sideeffect "; use $0", "a"(<32 x float> %mai)264  ret void265}266 267define void @test_rewrite_mfma_imm_src2(float %arg0, float %arg1) #0 {268; CHECK-LABEL: test_rewrite_mfma_imm_src2:269; CHECK:       ; %bb.0: ; %bb270; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)271; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, 2.0272; CHECK-NEXT:    ;;#ASMSTART273; CHECK-NEXT:    ; use a[0:31]274; CHECK-NEXT:    ;;#ASMEND275; CHECK-NEXT:    s_setpc_b64 s[30:31]276bb:277  %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float %arg0, float %arg1, <32 x float> splat (float 2.0), i32 0, i32 0, i32 0)278  call void asm sideeffect "; use $0", "a"(<32 x float> %mai)279  ret void280}281 282define void @test_rewrite_mfma_subreg_extract0(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {283; CHECK-LABEL: test_rewrite_mfma_subreg_extract0:284; CHECK:       ; %bb.0: ; %bb285; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)286; CHECK-NEXT:    global_load_dwordx4 a[28:31], v[2:3], off offset:112287; CHECK-NEXT:    global_load_dwordx4 a[24:27], v[2:3], off offset:96288; CHECK-NEXT:    global_load_dwordx4 a[20:23], v[2:3], off offset:80289; CHECK-NEXT:    global_load_dwordx4 a[16:19], v[2:3], off offset:64290; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[2:3], off offset:48291; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[2:3], off offset:32292; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[2:3], off offset:16293; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off294; CHECK-NEXT:    s_waitcnt vmcnt(0)295; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[0:31]296; CHECK-NEXT:    ;;#ASMSTART297; CHECK-NEXT:    ; use a[0:3]298; CHECK-NEXT:    ;;#ASMEND299; CHECK-NEXT:    s_setpc_b64 s[30:31]300bb:301  %src2 = load <32 x float>, ptr addrspace(1) %ptr302  %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float %arg0, float %arg1, <32 x float> %src2, i32 0, i32 0, i32 0)303  %extract.sub4 = shufflevector <32 x float> %mai, <32 x float> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3>304  call void asm sideeffect "; use $0", "a"(<4 x float> %extract.sub4)305  ret void306}307 308define void @test_rewrite_mfma_subreg_extract1(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {309; CHECK-LABEL: test_rewrite_mfma_subreg_extract1:310; CHECK:       ; %bb.0: ; %bb311; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)312; CHECK-NEXT:    global_load_dwordx4 a[28:31], v[2:3], off offset:112313; CHECK-NEXT:    global_load_dwordx4 a[24:27], v[2:3], off offset:96314; CHECK-NEXT:    global_load_dwordx4 a[20:23], v[2:3], off offset:80315; CHECK-NEXT:    global_load_dwordx4 a[16:19], v[2:3], off offset:64316; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[2:3], off offset:48317; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[2:3], off offset:32318; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[2:3], off offset:16319; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off320; CHECK-NEXT:    s_waitcnt vmcnt(0)321; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[0:31]322; CHECK-NEXT:    ;;#ASMSTART323; CHECK-NEXT:    ; use a[4:7]324; CHECK-NEXT:    ;;#ASMEND325; CHECK-NEXT:    s_setpc_b64 s[30:31]326bb:327  %src2 = load <32 x float>, ptr addrspace(1) %ptr328  %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float %arg0, float %arg1, <32 x float> %src2, i32 0, i32 0, i32 0)329  %extract.sub4 = shufflevector <32 x float> %mai, <32 x float> poison, <4 x i32> <i32 4, i32 5, i32 6, i32 7>330  call void asm sideeffect "; use $0", "a"(<4 x float> %extract.sub4)331  ret void332}333 334; odd offset335define void @test_rewrite_mfma_subreg_extract2(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {336; CHECK-LABEL: test_rewrite_mfma_subreg_extract2:337; CHECK:       ; %bb.0: ; %bb338; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)339; CHECK-NEXT:    global_load_dwordx4 a[28:31], v[2:3], off offset:112340; CHECK-NEXT:    global_load_dwordx4 a[24:27], v[2:3], off offset:96341; CHECK-NEXT:    global_load_dwordx4 a[20:23], v[2:3], off offset:80342; CHECK-NEXT:    global_load_dwordx4 a[16:19], v[2:3], off offset:64343; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[2:3], off offset:48344; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[2:3], off offset:32345; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[2:3], off offset:16346; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off347; CHECK-NEXT:    s_waitcnt vmcnt(0)348; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[0:31]349; CHECK-NEXT:    s_nop 15350; CHECK-NEXT:    s_nop 1351; CHECK-NEXT:    v_accvgpr_mov_b32 a0, a1352; CHECK-NEXT:    v_accvgpr_mov_b32 a1, a2353; CHECK-NEXT:    v_accvgpr_mov_b32 a2, a3354; CHECK-NEXT:    v_accvgpr_mov_b32 a3, a4355; CHECK-NEXT:    ;;#ASMSTART356; CHECK-NEXT:    ; use a[0:3]357; CHECK-NEXT:    ;;#ASMEND358; CHECK-NEXT:    s_setpc_b64 s[30:31]359bb:360  %src2 = load <32 x float>, ptr addrspace(1) %ptr361  %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float %arg0, float %arg1, <32 x float> %src2, i32 0, i32 0, i32 0)362  %extract.sub4 = shufflevector <32 x float> %mai, <32 x float> poison, <4 x i32> <i32 1, i32 2, i32 3, i32 4>363  call void asm sideeffect "; use $0", "a"(<4 x float> %extract.sub4)364  ret void365}366 367define amdgpu_kernel void @illegal_mfma_after_rewrite() #1 {368; CHECK-LABEL: illegal_mfma_after_rewrite:369; CHECK:       ; %bb.0: ; %entry370; CHECK-NEXT:    s_mov_b32 s4, 0371; CHECK-NEXT:    s_mov_b32 s5, s4372; CHECK-NEXT:    v_mov_b64_e32 v[26:27], s[4:5]373; CHECK-NEXT:    ;;#ASMSTART374; CHECK-NEXT:    ; def s[0:3]375; CHECK-NEXT:    ;;#ASMEND376; CHECK-NEXT:    ;;#ASMSTART377; CHECK-NEXT:    ; def v[16:19]378; CHECK-NEXT:    ;;#ASMEND379; CHECK-NEXT:    s_nop 0380; CHECK-NEXT:    v_mov_b64_e32 v[0:1], s[0:1]381; CHECK-NEXT:    v_mov_b64_e32 v[2:3], s[2:3]382; CHECK-NEXT:    s_mov_b32 s0, 0x3c003c00383; CHECK-NEXT:    s_mov_b32 s1, s0384; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[4:7], v[26:27], v[26:27], v[0:3]385; CHECK-NEXT:    v_mov_b64_e32 v[28:29], s[0:1]386; CHECK-NEXT:    s_mov_b32 s0, 0x7e007e00387; CHECK-NEXT:    s_mov_b32 s1, s0388; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[4:7], v[26:27], v[26:27], v[4:7]389; CHECK-NEXT:    v_mov_b64_e32 v[30:31], s[0:1]390; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[6:9], v[26:27], v[28:29], v[0:3]391; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[6:9], v[26:27], v[26:27], v[6:9]392; CHECK-NEXT:    s_nop 3393; CHECK-NEXT:    v_cvt_f16_f32_e32 v24, v4394; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[12:15], v[26:27], v[30:31], v[0:3]395; CHECK-NEXT:    s_nop 0396; CHECK-NEXT:    v_mov_b32_e32 v8, 0x7fc00000397; CHECK-NEXT:    v_mov_b32_e32 v9, v8398; CHECK-NEXT:    v_mov_b32_e32 v10, v8399; CHECK-NEXT:    v_mov_b32_e32 v11, v8400; CHECK-NEXT:    v_cvt_f16_f32_e32 v2, v6401; CHECK-NEXT:    v_mov_b64_e32 v[0:1], 0402; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[8:11], v[26:27], v[26:27], v[8:11]403; CHECK-NEXT:    global_store_short v[0:1], v2, off404; CHECK-NEXT:    buffer_wbl2 sc0 sc1405; CHECK-NEXT:    s_waitcnt vmcnt(0)406; CHECK-NEXT:    buffer_inv sc0 sc1407; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[2:5], v[26:27], v[28:29], v[16:19]408; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[6:9], v[26:27], v[26:27], v[8:11]409; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[20:23], v[26:27], v[26:27], v[16:19]410; CHECK-NEXT:    s_nop 5411; CHECK-NEXT:    v_cvt_f16_f32_e32 v10, v6412; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[6:9], v[26:27], v[26:27], v[12:15]413; CHECK-NEXT:    global_store_short v[0:1], v10, off414; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[2:5], v[26:27], v[26:27], v[2:5]415; CHECK-NEXT:    buffer_wbl2 sc0 sc1416; CHECK-NEXT:    s_waitcnt vmcnt(0)417; CHECK-NEXT:    buffer_inv sc0 sc1418; CHECK-NEXT:    s_nop 1419; CHECK-NEXT:    v_cvt_f16_f32_e32 v6, v6420; CHECK-NEXT:    global_store_short v[0:1], v6, off421; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[16:19], v[26:27], v[26:27], v[20:23]422; CHECK-NEXT:    buffer_wbl2 sc0 sc1423; CHECK-NEXT:    s_waitcnt vmcnt(0)424; CHECK-NEXT:    buffer_inv sc0 sc1425; CHECK-NEXT:    global_store_short v[0:1], v24, off426; CHECK-NEXT:    buffer_wbl2 sc0 sc1427; CHECK-NEXT:    s_waitcnt vmcnt(0)428; CHECK-NEXT:    buffer_inv sc0 sc1429; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[2:5], v[28:29], v[26:27], v[2:5]430; CHECK-NEXT:    s_nop 6431; CHECK-NEXT:    v_cvt_f16_f32_e32 v6, v2432; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 v[2:5], v[30:31], v[26:27], v[16:19]433; CHECK-NEXT:    global_store_short v[0:1], v6, off434; CHECK-NEXT:    buffer_wbl2 sc0 sc1435; CHECK-NEXT:    s_waitcnt vmcnt(0)436; CHECK-NEXT:    buffer_inv sc0 sc1437; CHECK-NEXT:    s_nop 2438; CHECK-NEXT:    v_cvt_f16_f32_e32 v2, v2439; CHECK-NEXT:    global_store_short v[0:1], v2, off440; CHECK-NEXT:    s_endpgm441entry:442  %k0 = call <4 x float> asm sideeffect "; def $0", "=s"()443  %i2 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %k0, i32 0, i32 0, i32 0)444  %i4 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> splat (half 0xH3C00), <4 x float> %k0, i32 0, i32 0, i32 0)445  %i6 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> splat (half 0xH7E00), <4 x float> %k0, i32 0, i32 0, i32 0)446  %i5 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> splat (float 0x7FF8000000000000), i32 0, i32 0, i32 0)447  %k = call <4 x float> asm sideeffect "; def $0", "=v"()448  %i1 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %k, i32 0, i32 0, i32 0)449  %i7 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> splat (half 0xH3C00), <4 x float> %k, i32 0, i32 0, i32 0)450  %i17 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i1, i32 0, i32 0, i32 0)451  %i19 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i4, i32 0, i32 0, i32 0)452  %c_thread_buf.0 = extractelement <4 x float> %i19, i64 0453  %conv.0 = fptrunc float %c_thread_buf.0 to half454  store half %conv.0, ptr addrspace(1) null, align 2455  fence seq_cst456  %i22 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i5, i32 0, i32 0, i32 0)457  %c_thread_buf.1 = extractelement <4 x float> %i22, i64 0458  %conv1 = fptrunc float %c_thread_buf.1 to half459  store half %conv1, ptr addrspace(1) null, align 2460  fence seq_cst461  %i23 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i6, i32 0, i32 0, i32 0)462  %c_thread_buf.2 = extractelement <4 x float> %i23, i64 0463  %conv2 = fptrunc float %c_thread_buf.2 to half464  store half %conv2, ptr addrspace(1) null, align 2465  fence seq_cst466  %i25 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i2, i32 0, i32 0, i32 0)467  %c_thread_buf.3 = extractelement <4 x float> %i25, i64 0468  %conv3 = fptrunc float %c_thread_buf.3 to half469  store half %conv3, ptr addrspace(1) null, align 2470  fence seq_cst471  %i26 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i7, i32 0, i32 0, i32 0)472  %i27 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> splat (half 0xH3C00), <4 x half> zeroinitializer, <4 x float> %i26, i32 0, i32 0, i32 0)473  %c_thread_buf.4 = extractelement <4 x float> %i27, i64 0474  %conv4 = fptrunc float %c_thread_buf.4 to half475  store half %conv4, ptr addrspace(1) null, align 2476  fence seq_cst477  %i31 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> splat (half 0xH7E00), <4 x half> zeroinitializer, <4 x float> %i17, i32 0, i32 0, i32 0)478  %c_thread_buf.5 = extractelement <4 x float> %i31, i64 0479  %conv5 = fptrunc float %c_thread_buf.5 to half480  store half %conv5, ptr addrspace(1) null, align 2481  ret void482}483 484define void @test_rewrite_mfma_subreg_insert0(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {485; CHECK-LABEL: test_rewrite_mfma_subreg_insert0:486; CHECK:       ; %bb.0:487; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)488; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off489; CHECK-NEXT:    s_waitcnt vmcnt(0)490; CHECK-NEXT:    v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[0:3]491; CHECK-NEXT:    ;;#ASMSTART492; CHECK-NEXT:    ; use a[0:7]493; CHECK-NEXT:    ;;#ASMEND494; CHECK-NEXT:    s_setpc_b64 s[30:31]495  %src2 = load <4 x float>, ptr addrspace(1) %ptr496  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float %arg0, float %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)497  %insert.sub0 = shufflevector <4 x float> %mai, <4 x float> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 poison, i32 poison, i32 poison, i32 poison>498  call void asm sideeffect "; use $0", "a"(<8 x float> %insert.sub0)499  ret void500}501 502; odd offset503define void @test_rewrite_mfma_subreg_insert1(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {504; CHECK-LABEL: test_rewrite_mfma_subreg_insert1:505; CHECK:       ; %bb.0:506; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)507; CHECK-NEXT:    global_load_dwordx4 v[2:5], v[2:3], off508; CHECK-NEXT:    s_waitcnt vmcnt(0)509; CHECK-NEXT:    v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5]510; CHECK-NEXT:    s_nop 3511; CHECK-NEXT:    v_pk_mov_b32 v[0:1], v[0:1], v[2:3] op_sel:[1,0]512; CHECK-NEXT:    s_nop 0513; CHECK-NEXT:    v_accvgpr_write_b32 a0, v0514; CHECK-NEXT:    v_accvgpr_write_b32 a1, v1515; CHECK-NEXT:    v_accvgpr_write_b32 a2, v3516; CHECK-NEXT:    ;;#ASMSTART517; CHECK-NEXT:    ; use a[0:7]518; CHECK-NEXT:    ;;#ASMEND519; CHECK-NEXT:    s_setpc_b64 s[30:31]520  %src2 = load <4 x float>, ptr addrspace(1) %ptr521  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float %arg0, float %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)522  %insert.sub0 = shufflevector <4 x float> %mai, <4 x float> poison, <8 x i32> <i32 1, i32 2, i32 3, i32 4, i32 poison, i32 poison, i32 poison, i32 poison>523  call void asm sideeffect "; use $0", "a"(<8 x float> %insert.sub0)524  ret void525}526 527define void @test_rewrite_mfma_subreg_insert2(double %arg0, double %arg1, ptr addrspace(1) %ptr) #0 {528; CHECK-LABEL: test_rewrite_mfma_subreg_insert2:529; CHECK:       ; %bb.0:530; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)531; CHECK-NEXT:    global_load_dwordx2 a[0:1], v[4:5], off532; CHECK-NEXT:    s_waitcnt vmcnt(0)533; CHECK-NEXT:    v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], v[2:3], a[0:1]534; CHECK-NEXT:    ;;#ASMSTART535; CHECK-NEXT:    ; use a[0:3]536; CHECK-NEXT:    ;;#ASMEND537; CHECK-NEXT:    s_setpc_b64 s[30:31]538  %src2 = load double, ptr addrspace(1) %ptr539  %mai = call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %arg0, double %arg1, double %src2, i32 0, i32 0, i32 0)540  %insert.sub0 = insertelement <2 x double> poison, double %mai, i32 0541  call void asm sideeffect "; use $0", "a"(<2 x double> %insert.sub0)542  ret void543}544 545define amdgpu_kernel void @test_rewrite_mfma_direct_copy_from_agpr_class(ptr addrspace(1) %arg0, ptr addrspace(1) %arg1) #0 {546; CHECK-LABEL: test_rewrite_mfma_direct_copy_from_agpr_class:547; CHECK:       ; %bb.0:548; CHECK-NEXT:    v_accvgpr_write_b32 a34, 2.0549; CHECK-NEXT:    v_and_b32_e32 v0, 0x3ff, v0550; CHECK-NEXT:    v_lshlrev_b32_e32 v0, 7, v0551; CHECK-NEXT:    ;;#ASMSTART552; CHECK-NEXT:    ; def a[0:31]553; CHECK-NEXT:    ;;#ASMEND554; CHECK-NEXT:    v_accvgpr_write_b32 a32, v0555; CHECK-NEXT:    v_accvgpr_read_b32 v0, a0556; CHECK-NEXT:    v_accvgpr_read_b32 v1, a1557; CHECK-NEXT:    v_accvgpr_read_b32 v2, a2558; CHECK-NEXT:    v_accvgpr_read_b32 v3, a3559; CHECK-NEXT:    v_accvgpr_read_b32 v4, a4560; CHECK-NEXT:    v_accvgpr_read_b32 v5, a5561; CHECK-NEXT:    v_accvgpr_read_b32 v6, a6562; CHECK-NEXT:    v_accvgpr_read_b32 v7, a7563; CHECK-NEXT:    v_accvgpr_read_b32 v8, a8564; CHECK-NEXT:    v_accvgpr_read_b32 v9, a9565; CHECK-NEXT:    v_accvgpr_read_b32 v10, a10566; CHECK-NEXT:    v_accvgpr_read_b32 v11, a11567; CHECK-NEXT:    v_accvgpr_read_b32 v12, a12568; CHECK-NEXT:    v_accvgpr_read_b32 v13, a13569; CHECK-NEXT:    v_accvgpr_read_b32 v14, a14570; CHECK-NEXT:    v_accvgpr_read_b32 v15, a15571; CHECK-NEXT:    v_accvgpr_read_b32 v16, a16572; CHECK-NEXT:    v_accvgpr_read_b32 v17, a17573; CHECK-NEXT:    v_accvgpr_read_b32 v18, a18574; CHECK-NEXT:    v_accvgpr_read_b32 v19, a19575; CHECK-NEXT:    v_accvgpr_read_b32 v20, a20576; CHECK-NEXT:    v_accvgpr_read_b32 v21, a21577; CHECK-NEXT:    v_accvgpr_read_b32 v22, a22578; CHECK-NEXT:    v_accvgpr_read_b32 v23, a23579; CHECK-NEXT:    v_accvgpr_read_b32 v24, a24580; CHECK-NEXT:    v_accvgpr_read_b32 v25, a25581; CHECK-NEXT:    v_accvgpr_read_b32 v26, a26582; CHECK-NEXT:    v_accvgpr_read_b32 v27, a27583; CHECK-NEXT:    v_accvgpr_read_b32 v28, a28584; CHECK-NEXT:    v_accvgpr_read_b32 v29, a29585; CHECK-NEXT:    v_accvgpr_read_b32 v30, a30586; CHECK-NEXT:    v_accvgpr_read_b32 v31, a31587; CHECK-NEXT:    v_accvgpr_write_b32 a33, 4.0588; CHECK-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x0589; CHECK-NEXT:    s_nop 0590; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 v[32:63], a34, a33, v[0:31]591; CHECK-NEXT:    v_mov_b32_e32 v1, 0x41000000592; CHECK-NEXT:    v_accvgpr_read_b32 v0, a32593; CHECK-NEXT:    s_nop 15594; CHECK-NEXT:    v_mov_b64_e32 v[2:3], v[32:33]595; CHECK-NEXT:    v_mov_b64_e32 v[4:5], v[34:35]596; CHECK-NEXT:    v_mov_b64_e32 v[6:7], v[36:37]597; CHECK-NEXT:    v_mov_b64_e32 v[8:9], v[38:39]598; CHECK-NEXT:    v_mov_b64_e32 v[10:11], v[40:41]599; CHECK-NEXT:    v_mov_b64_e32 v[12:13], v[42:43]600; CHECK-NEXT:    v_mov_b64_e32 v[14:15], v[44:45]601; CHECK-NEXT:    v_mov_b64_e32 v[16:17], v[46:47]602; CHECK-NEXT:    v_mov_b64_e32 v[18:19], v[48:49]603; CHECK-NEXT:    v_mov_b64_e32 v[20:21], v[50:51]604; CHECK-NEXT:    v_mov_b64_e32 v[22:23], v[52:53]605; CHECK-NEXT:    v_mov_b64_e32 v[24:25], v[54:55]606; CHECK-NEXT:    v_mov_b64_e32 v[26:27], v[56:57]607; CHECK-NEXT:    v_mov_b64_e32 v[28:29], v[58:59]608; CHECK-NEXT:    v_mov_b64_e32 v[30:31], v[60:61]609; CHECK-NEXT:    v_mov_b64_e32 v[32:33], v[62:63]610; CHECK-NEXT:    s_waitcnt lgkmcnt(0)611; CHECK-NEXT:    global_store_dwordx4 v0, v[30:33], s[0:1] offset:112612; CHECK-NEXT:    global_store_dwordx4 v0, v[26:29], s[0:1] offset:96613; CHECK-NEXT:    global_store_dwordx4 v0, v[22:25], s[0:1] offset:80614; CHECK-NEXT:    global_store_dwordx4 v0, v[18:21], s[0:1] offset:64615; CHECK-NEXT:    global_store_dwordx4 v0, v[14:17], s[0:1] offset:48616; CHECK-NEXT:    global_store_dwordx4 v0, v[10:13], s[0:1] offset:32617; CHECK-NEXT:    global_store_dwordx4 v0, v[6:9], s[0:1] offset:16618; CHECK-NEXT:    global_store_dwordx4 v0, v[2:5], s[0:1]619; CHECK-NEXT:    s_nop 1620; CHECK-NEXT:    v_mov_b32_e32 v2, 0x41800000621; CHECK-NEXT:    s_nop 1622; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v1, v2, a[0:31]623; CHECK-NEXT:    s_nop 15624; CHECK-NEXT:    s_nop 1625; CHECK-NEXT:    global_store_dwordx4 v0, a[24:27], s[2:3] offset:96626; CHECK-NEXT:    global_store_dwordx4 v0, a[28:31], s[2:3] offset:112627; CHECK-NEXT:    global_store_dwordx4 v0, a[16:19], s[2:3] offset:64628; CHECK-NEXT:    global_store_dwordx4 v0, a[20:23], s[2:3] offset:80629; CHECK-NEXT:    global_store_dwordx4 v0, a[8:11], s[2:3] offset:32630; CHECK-NEXT:    global_store_dwordx4 v0, a[12:15], s[2:3] offset:48631; CHECK-NEXT:    global_store_dwordx4 v0, a[0:3], s[2:3]632; CHECK-NEXT:    global_store_dwordx4 v0, a[4:7], s[2:3] offset:16633; CHECK-NEXT:    s_endpgm634  %src2 = call <32 x float> asm sideeffect "; def $0", "=a"()635  %mai0 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 2.0, float 4.0, <32 x float> %src2, i32 0, i32 0, i32 0)636  %mai1 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 8.0, float 16.0, <32 x float> %src2, i32 0, i32 0, i32 0)637  %id = call i32 @llvm.amdgcn.workitem.id.x()638  %gep0 = getelementptr <32 x float>, ptr addrspace(1) %arg0, i32 %id639  store <32 x float> %mai0, ptr addrspace(1) %gep0, align 128640  %gep1 = getelementptr <32 x float>, ptr addrspace(1) %arg1, i32 %id641  store <32 x float> %mai1, ptr addrspace(1) %gep1, align 128642  ret void643}644 645define amdgpu_kernel void @test_rewrite_mfma_direct_copy_from_agpr_class_chain(ptr addrspace(1) %arg0) #0 {646; CHECK-LABEL: test_rewrite_mfma_direct_copy_from_agpr_class_chain:647; CHECK:       ; %bb.0:648; CHECK-NEXT:    v_mov_b32_e32 v1, 2.0649; CHECK-NEXT:    ;;#ASMSTART650; CHECK-NEXT:    ; def a[0:31]651; CHECK-NEXT:    ;;#ASMEND652; CHECK-NEXT:    v_mov_b32_e32 v34, 4.0653; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0654; CHECK-NEXT:    v_and_b32_e32 v0, 0x3ff, v0655; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v1, v34, a[0:31]656; CHECK-NEXT:    v_mov_b32_e32 v1, 0x41000000657; CHECK-NEXT:    v_mov_b32_e32 v34, 0x41800000658; CHECK-NEXT:    v_lshlrev_b32_e32 v0, 7, v0659; CHECK-NEXT:    s_nop 0660; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v1, v34, a[0:31]661; CHECK-NEXT:    s_waitcnt lgkmcnt(0)662; CHECK-NEXT:    s_nop 15663; CHECK-NEXT:    s_nop 0664; CHECK-NEXT:    global_store_dwordx4 v0, a[28:31], s[0:1] offset:112665; CHECK-NEXT:    global_store_dwordx4 v0, a[24:27], s[0:1] offset:96666; CHECK-NEXT:    global_store_dwordx4 v0, a[20:23], s[0:1] offset:80667; CHECK-NEXT:    global_store_dwordx4 v0, a[16:19], s[0:1] offset:64668; CHECK-NEXT:    global_store_dwordx4 v0, a[12:15], s[0:1] offset:48669; CHECK-NEXT:    global_store_dwordx4 v0, a[8:11], s[0:1] offset:32670; CHECK-NEXT:    global_store_dwordx4 v0, a[4:7], s[0:1] offset:16671; CHECK-NEXT:    global_store_dwordx4 v0, a[0:3], s[0:1]672; CHECK-NEXT:    s_endpgm673  %src2 = call <32 x float> asm sideeffect "; def $0", "=a"()674  %mai0 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 2.0, float 4.0, <32 x float> %src2, i32 0, i32 0, i32 0)675  %mai1 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 8.0, float 16.0, <32 x float> %mai0, i32 0, i32 0, i32 0)676  %id = call i32 @llvm.amdgcn.workitem.id.x()677  %gep0 = getelementptr <32 x float>, ptr addrspace(1) %arg0, i32 %id678  store <32 x float> %mai1, ptr addrspace(1) %gep0, align 128679  ret void680}681 682; Untied case683define void @test_rewrite_mfma_copy_from_agpr_class_f64_4x4x4f64(double %arg0, double %arg1, ptr addrspace(1) %ptr) #0 {684; CHECK-LABEL: test_rewrite_mfma_copy_from_agpr_class_f64_4x4x4f64:685; CHECK:       ; %bb.0:686; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)687; CHECK-NEXT:    ;;#ASMSTART688; CHECK-NEXT:    ; def a[0:1]689; CHECK-NEXT:    ;;#ASMEND690; CHECK-NEXT:    v_and_b32_e32 v8, 0x3ff, v31691; CHECK-NEXT:    v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], v[2:3], a[0:1]692; CHECK-NEXT:    v_lshlrev_b32_e32 v2, 3, v8693; CHECK-NEXT:    v_mov_b32_e32 v3, 0694; CHECK-NEXT:    v_lshl_add_u64 v[2:3], v[4:5], 0, v[2:3]695; CHECK-NEXT:    s_nop 5696; CHECK-NEXT:    global_store_dwordx2 v[2:3], a[0:1], off697; CHECK-NEXT:    s_waitcnt vmcnt(0)698; CHECK-NEXT:    s_setpc_b64 s[30:31]699  %src2 = call double asm sideeffect "; def $0", "=a"()700  %mai = call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %arg0, double %arg1, double %src2, i32 0, i32 0, i32 0)701  %id = call i32 @llvm.amdgcn.workitem.id.x()702  %gep0 = getelementptr double, ptr addrspace(1) %ptr, i32 %id703  store double %mai, ptr addrspace(1) %gep0, align 8704  ret void705}706 707define void @test_rewrite_mfma_copy_from_agpr_class_f64_4x4x4f64_chain(double %arg0, double %arg1, double %arg2, double %arg3, ptr addrspace(1) %ptr) #0 {708; CHECK-LABEL: test_rewrite_mfma_copy_from_agpr_class_f64_4x4x4f64_chain:709; CHECK:       ; %bb.0:710; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)711; CHECK-NEXT:    ;;#ASMSTART712; CHECK-NEXT:    ; def a[0:1]713; CHECK-NEXT:    ;;#ASMEND714; CHECK-NEXT:    s_nop 0715; CHECK-NEXT:    v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], v[2:3], a[0:1]716; CHECK-NEXT:    v_and_b32_e32 v2, 0x3ff, v31717; CHECK-NEXT:    v_lshlrev_b32_e32 v2, 3, v2718; CHECK-NEXT:    v_mov_b32_e32 v3, 0719; CHECK-NEXT:    v_lshl_add_u64 v[2:3], v[8:9], 0, v[2:3]720; CHECK-NEXT:    v_mfma_f64_4x4x4_4b_f64 a[0:1], v[4:5], v[6:7], a[0:1]721; CHECK-NEXT:    s_nop 8722; CHECK-NEXT:    global_store_dwordx2 v[2:3], a[0:1], off723; CHECK-NEXT:    s_waitcnt vmcnt(0)724; CHECK-NEXT:    s_setpc_b64 s[30:31]725  %src2 = call double asm sideeffect "; def $0", "=a"()726  %mai0 = call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %arg0, double %arg1, double %src2, i32 0, i32 0, i32 0)727  %mai1 = call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %arg2, double %arg3, double %mai0, i32 0, i32 0, i32 0)728  %id = call i32 @llvm.amdgcn.workitem.id.x()729  %gep0 = getelementptr double, ptr addrspace(1) %ptr, i32 %id730  store double %mai1, ptr addrspace(1) %gep0, align 8731  ret void732}733 734define amdgpu_kernel void @test_rewrite_mfma_direct_copy_from_agpr_class_subreg(ptr addrspace(1) %arg) #0 {735; CHECK-LABEL: test_rewrite_mfma_direct_copy_from_agpr_class_subreg:736; CHECK:       ; %bb.0:737; CHECK-NEXT:    v_mov_b32_e32 v1, 2.0738; CHECK-NEXT:    ;;#ASMSTART739; CHECK-NEXT:    ; def a[0:31]740; CHECK-NEXT:    ;;#ASMEND741; CHECK-NEXT:    v_mov_b32_e32 v18, 4.0742; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0743; CHECK-NEXT:    v_and_b32_e32 v0, 0x3ff, v0744; CHECK-NEXT:    v_mfma_f32_16x16x1_4b_f32 a[0:15], v1, v18, a[0:15]745; CHECK-NEXT:    v_lshlrev_b32_e32 v0, 6, v0746; CHECK-NEXT:    s_waitcnt lgkmcnt(0)747; CHECK-NEXT:    s_nop 7748; CHECK-NEXT:    global_store_dwordx4 v0, a[12:15], s[0:1] offset:48749; CHECK-NEXT:    global_store_dwordx4 v0, a[8:11], s[0:1] offset:32750; CHECK-NEXT:    global_store_dwordx4 v0, a[4:7], s[0:1] offset:16751; CHECK-NEXT:    global_store_dwordx4 v0, a[0:3], s[0:1]752; CHECK-NEXT:    s_endpgm753  %def = call <32 x float> asm sideeffect "; def $0", "=a"()754  %src2 = shufflevector <32 x float> %def, <32 x float> poison, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>755  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 2.0, float 4.0, <16 x float> %src2, i32 0, i32 0, i32 0)756  %id = call i32 @llvm.amdgcn.workitem.id.x()757  %gep = getelementptr <16 x float>, ptr addrspace(1) %arg, i32 %id758  store <16 x float> %mai, ptr addrspace(1) %gep, align 64759  ret void760}761 762define amdgpu_kernel void @test_rewrite_mfma_direct_copy_from_agpr_class_subreg_odd(ptr addrspace(1) %arg) #0 {763; CHECK-LABEL: test_rewrite_mfma_direct_copy_from_agpr_class_subreg_odd:764; CHECK:       ; %bb.0:765; CHECK-NEXT:    v_mov_b32_e32 v1, 2.0766; CHECK-NEXT:    ;;#ASMSTART767; CHECK-NEXT:    ; def a[0:31]768; CHECK-NEXT:    ;;#ASMEND769; CHECK-NEXT:    v_mov_b32_e32 v18, 4.0770; CHECK-NEXT:    v_accvgpr_mov_b32 a0, a1771; CHECK-NEXT:    v_accvgpr_mov_b32 a1, a2772; CHECK-NEXT:    v_accvgpr_mov_b32 a2, a3773; CHECK-NEXT:    v_accvgpr_mov_b32 a3, a4774; CHECK-NEXT:    v_accvgpr_mov_b32 a4, a5775; CHECK-NEXT:    v_accvgpr_mov_b32 a5, a6776; CHECK-NEXT:    v_accvgpr_mov_b32 a6, a7777; CHECK-NEXT:    v_accvgpr_mov_b32 a7, a8778; CHECK-NEXT:    v_accvgpr_mov_b32 a8, a9779; CHECK-NEXT:    v_accvgpr_mov_b32 a9, a10780; CHECK-NEXT:    v_accvgpr_mov_b32 a10, a11781; CHECK-NEXT:    v_accvgpr_mov_b32 a11, a12782; CHECK-NEXT:    v_accvgpr_mov_b32 a12, a13783; CHECK-NEXT:    v_accvgpr_mov_b32 a13, a14784; CHECK-NEXT:    v_accvgpr_mov_b32 a14, a15785; CHECK-NEXT:    v_accvgpr_mov_b32 a15, a16786; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0787; CHECK-NEXT:    v_and_b32_e32 v0, 0x3ff, v0788; CHECK-NEXT:    v_mfma_f32_16x16x1_4b_f32 a[0:15], v1, v18, a[0:15]789; CHECK-NEXT:    v_lshlrev_b32_e32 v0, 6, v0790; CHECK-NEXT:    s_waitcnt lgkmcnt(0)791; CHECK-NEXT:    s_nop 7792; CHECK-NEXT:    global_store_dwordx4 v0, a[12:15], s[0:1] offset:48793; CHECK-NEXT:    global_store_dwordx4 v0, a[8:11], s[0:1] offset:32794; CHECK-NEXT:    global_store_dwordx4 v0, a[4:7], s[0:1] offset:16795; CHECK-NEXT:    global_store_dwordx4 v0, a[0:3], s[0:1]796; CHECK-NEXT:    s_endpgm797  %def = call <32 x float> asm sideeffect "; def $0", "=a"()798  %src2 = shufflevector <32 x float> %def, <32 x float> poison, <16 x i32> <i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15, i32 16>799  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 2.0, float 4.0, <16 x float> %src2, i32 0, i32 0, i32 0)800  %id = call i32 @llvm.amdgcn.workitem.id.x()801  %gep = getelementptr <16 x float>, ptr addrspace(1) %arg, i32 %id802  store <16 x float> %mai, ptr addrspace(1) %gep, align 64803  ret void804}805 806; a->v->mfma->a807define amdgpu_kernel void @test_rewrite_mfma_direct_copy_from_agpr_class_copy_back() #0 {808; CHECK-LABEL: test_rewrite_mfma_direct_copy_from_agpr_class_copy_back:809; CHECK:       ; %bb.0:810; CHECK-NEXT:    v_mov_b32_e32 v32, 2.0811; CHECK-NEXT:    ;;#ASMSTART812; CHECK-NEXT:    ; def a[0:31]813; CHECK-NEXT:    ;;#ASMEND814; CHECK-NEXT:    v_mov_b32_e32 v33, 4.0815; CHECK-NEXT:    s_nop 1816; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v32, v33, a[0:31]817; CHECK-NEXT:    ;;#ASMSTART818; CHECK-NEXT:    ; use a[0:31]819; CHECK-NEXT:    ;;#ASMEND820; CHECK-NEXT:    s_endpgm821  %src2 = call <32 x float> asm sideeffect "; def $0", "=a"()822  %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 2.0, float 4.0, <32 x float> %src2, i32 0, i32 0, i32 0)823  call void asm sideeffect "; use $0", "a"(<32 x float> %mai)824  ret void825}826 827;---------------------------------------------------------------------828; Comprehensively test all MFMA intrinsics are in the rewrite table829;---------------------------------------------------------------------830 831define void @test_rewrite_mfma_f32_32x32x1f32(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {832; CHECK-LABEL: test_rewrite_mfma_f32_32x32x1f32:833; CHECK:       ; %bb.0:834; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)835; CHECK-NEXT:    global_load_dwordx4 a[28:31], v[2:3], off offset:112836; CHECK-NEXT:    global_load_dwordx4 a[24:27], v[2:3], off offset:96837; CHECK-NEXT:    global_load_dwordx4 a[20:23], v[2:3], off offset:80838; CHECK-NEXT:    global_load_dwordx4 a[16:19], v[2:3], off offset:64839; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[2:3], off offset:48840; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[2:3], off offset:32841; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[2:3], off offset:16842; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off843; CHECK-NEXT:    s_waitcnt vmcnt(0)844; CHECK-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[0:31]845; CHECK-NEXT:    ;;#ASMSTART846; CHECK-NEXT:    ; use a[0:31]847; CHECK-NEXT:    ;;#ASMEND848; CHECK-NEXT:    s_setpc_b64 s[30:31]849  %src2 = load <32 x float>, ptr addrspace(1) %ptr850  %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float %arg0, float %arg1, <32 x float> %src2, i32 0, i32 0, i32 0)851  call void asm sideeffect "; use $0", "a"(<32 x float> %mai)852  ret void853}854 855define void @test_rewrite_mfma_f32_16x16x1f32(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {856; CHECK-LABEL: test_rewrite_mfma_f32_16x16x1f32:857; CHECK:       ; %bb.0:858; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)859; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[2:3], off offset:48860; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[2:3], off offset:32861; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[2:3], off offset:16862; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off863; CHECK-NEXT:    s_waitcnt vmcnt(0)864; CHECK-NEXT:    v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[0:15]865; CHECK-NEXT:    ;;#ASMSTART866; CHECK-NEXT:    ; use a[0:15]867; CHECK-NEXT:    ;;#ASMEND868; CHECK-NEXT:    s_setpc_b64 s[30:31]869  %src2 = load <16 x float>, ptr addrspace(1) %ptr870  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float %arg0, float %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)871  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)872  ret void873}874 875define void @test_rewrite_mfma_f32_4x4x1f32(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {876; CHECK-LABEL: test_rewrite_mfma_f32_4x4x1f32:877; CHECK:       ; %bb.0:878; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)879; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off880; CHECK-NEXT:    s_waitcnt vmcnt(0)881; CHECK-NEXT:    v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[0:3]882; CHECK-NEXT:    ;;#ASMSTART883; CHECK-NEXT:    ; use a[0:3]884; CHECK-NEXT:    ;;#ASMEND885; CHECK-NEXT:    s_setpc_b64 s[30:31]886  %src2 = load <4 x float>, ptr addrspace(1) %ptr887  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float %arg0, float %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)888  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)889  ret void890}891 892define void @test_rewrite_mfma_f32_32x32x2f32(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {893; CHECK-LABEL: test_rewrite_mfma_f32_32x32x2f32:894; CHECK:       ; %bb.0:895; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)896; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[2:3], off offset:48897; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[2:3], off offset:32898; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[2:3], off offset:16899; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off900; CHECK-NEXT:    s_waitcnt vmcnt(0)901; CHECK-NEXT:    v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[0:15]902; CHECK-NEXT:    ;;#ASMSTART903; CHECK-NEXT:    ; use a[0:15]904; CHECK-NEXT:    ;;#ASMEND905; CHECK-NEXT:    s_setpc_b64 s[30:31]906  %src2 = load <16 x float>, ptr addrspace(1) %ptr907  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x2f32(float %arg0, float %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)908  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)909  ret void910}911 912define void @test_rewrite_mfma_f32_16x16x4f32(float %arg0, float %arg1, ptr addrspace(1) %ptr) #0 {913; CHECK-LABEL: test_rewrite_mfma_f32_16x16x4f32:914; CHECK:       ; %bb.0:915; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)916; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off917; CHECK-NEXT:    s_waitcnt vmcnt(0)918; CHECK-NEXT:    v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[0:3]919; CHECK-NEXT:    ;;#ASMSTART920; CHECK-NEXT:    ; use a[0:3]921; CHECK-NEXT:    ;;#ASMEND922; CHECK-NEXT:    s_setpc_b64 s[30:31]923  %src2 = load <4 x float>, ptr addrspace(1) %ptr924  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x4f32(float %arg0, float %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)925  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)926  ret void927}928 929define void @test_rewrite_mfma_f32_32x32x4f16(<4 x half> %arg0, <4 x half> %arg1, ptr addrspace(1) %ptr) #0 {930; CHECK-LABEL: test_rewrite_mfma_f32_32x32x4f16:931; CHECK:       ; %bb.0:932; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)933; CHECK-NEXT:    global_load_dwordx4 a[28:31], v[4:5], off offset:112934; CHECK-NEXT:    global_load_dwordx4 a[24:27], v[4:5], off offset:96935; CHECK-NEXT:    global_load_dwordx4 a[20:23], v[4:5], off offset:80936; CHECK-NEXT:    global_load_dwordx4 a[16:19], v[4:5], off offset:64937; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:48938; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:32939; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:16940; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off941; CHECK-NEXT:    s_waitcnt vmcnt(0)942; CHECK-NEXT:    v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[0:31]943; CHECK-NEXT:    ;;#ASMSTART944; CHECK-NEXT:    ; use a[0:31]945; CHECK-NEXT:    ;;#ASMEND946; CHECK-NEXT:    s_setpc_b64 s[30:31]947  %src2 = load <32 x float>, ptr addrspace(1) %ptr948  %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half> %arg0, <4 x half> %arg1, <32 x float> %src2, i32 0, i32 0, i32 0)949  call void asm sideeffect "; use $0", "a"(<32 x float> %mai)950  ret void951}952 953define void @test_rewrite_mfma_f32_16x16x4f16(<4 x half> %arg0, <4 x half> %arg1, ptr addrspace(1) %ptr) #0 {954; CHECK-LABEL: test_rewrite_mfma_f32_16x16x4f16:955; CHECK:       ; %bb.0:956; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)957; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:48958; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:32959; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:16960; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off961; CHECK-NEXT:    s_waitcnt vmcnt(0)962; CHECK-NEXT:    v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[0:15]963; CHECK-NEXT:    ;;#ASMSTART964; CHECK-NEXT:    ; use a[0:15]965; CHECK-NEXT:    ;;#ASMEND966; CHECK-NEXT:    s_setpc_b64 s[30:31]967  %src2 = load <16 x float>, ptr addrspace(1) %ptr968  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half> %arg0, <4 x half> %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)969  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)970  ret void971}972 973define void @test_rewrite_mfma_f32_4x4x4f16(<4 x half> %arg0, <4 x half> %arg1, ptr addrspace(1) %ptr) #0 {974; CHECK-LABEL: test_rewrite_mfma_f32_4x4x4f16:975; CHECK:       ; %bb.0:976; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)977; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off978; CHECK-NEXT:    s_waitcnt vmcnt(0)979; CHECK-NEXT:    v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[0:3]980; CHECK-NEXT:    ;;#ASMSTART981; CHECK-NEXT:    ; use a[0:3]982; CHECK-NEXT:    ;;#ASMEND983; CHECK-NEXT:    s_setpc_b64 s[30:31]984  %src2 = load <4 x float>, ptr addrspace(1) %ptr985  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half> %arg0, <4 x half> %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)986  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)987  ret void988}989 990define void @test_rewrite_mfma_f32_32x32x8f16(<4 x half> %arg0, <4 x half> %arg1, ptr addrspace(1) %ptr) #0 {991; CHECK-LABEL: test_rewrite_mfma_f32_32x32x8f16:992; CHECK:       ; %bb.0:993; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)994; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:48995; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:32996; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:16997; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off998; CHECK-NEXT:    s_waitcnt vmcnt(0)999; CHECK-NEXT:    v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[0:15]1000; CHECK-NEXT:    ;;#ASMSTART1001; CHECK-NEXT:    ; use a[0:15]1002; CHECK-NEXT:    ;;#ASMEND1003; CHECK-NEXT:    s_setpc_b64 s[30:31]1004  %src2 = load <16 x float>, ptr addrspace(1) %ptr1005  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> %arg0, <4 x half> %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)1006  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1007  ret void1008}1009 1010define void @test_rewrite_mfma_f32_16x16x16f16(<4 x half> %arg0, <4 x half> %arg1, ptr addrspace(1) %ptr) #0 {1011; CHECK-LABEL: test_rewrite_mfma_f32_16x16x16f16:1012; CHECK:       ; %bb.0:1013; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1014; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1015; CHECK-NEXT:    s_waitcnt vmcnt(0)1016; CHECK-NEXT:    v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[0:3]1017; CHECK-NEXT:    ;;#ASMSTART1018; CHECK-NEXT:    ; use a[0:3]1019; CHECK-NEXT:    ;;#ASMEND1020; CHECK-NEXT:    s_setpc_b64 s[30:31]1021  %src2 = load <4 x float>, ptr addrspace(1) %ptr1022  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> %arg0, <4 x half> %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)1023  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1024  ret void1025}1026 1027define void @test_rewrite_mfma_i32_32x32x4i8(i32 %arg0, i32 %arg1, ptr addrspace(1) %ptr) #0 {1028; CHECK-LABEL: test_rewrite_mfma_i32_32x32x4i8:1029; CHECK:       ; %bb.0:1030; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1031; CHECK-NEXT:    global_load_dwordx4 a[28:31], v[2:3], off offset:1121032; CHECK-NEXT:    global_load_dwordx4 a[24:27], v[2:3], off offset:961033; CHECK-NEXT:    global_load_dwordx4 a[20:23], v[2:3], off offset:801034; CHECK-NEXT:    global_load_dwordx4 a[16:19], v[2:3], off offset:641035; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[2:3], off offset:481036; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[2:3], off offset:321037; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[2:3], off offset:161038; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off1039; CHECK-NEXT:    s_waitcnt vmcnt(0)1040; CHECK-NEXT:    v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[0:31]1041; CHECK-NEXT:    ;;#ASMSTART1042; CHECK-NEXT:    ; use a[0:31]1043; CHECK-NEXT:    ;;#ASMEND1044; CHECK-NEXT:    s_setpc_b64 s[30:31]1045  %src2 = load <32 x i32>, ptr addrspace(1) %ptr1046  %mai = call <32 x i32> @llvm.amdgcn.mfma.i32.32x32x4i8(i32 %arg0, i32 %arg1, <32 x i32> %src2, i32 0, i32 0, i32 0)1047  call void asm sideeffect "; use $0", "a"(<32 x i32> %mai)1048  ret void1049}1050 1051define void @test_rewrite_mfma_i32_16x16x4i8(i32 %arg0, i32 %arg1, ptr addrspace(1) %ptr) #0 {1052; CHECK-LABEL: test_rewrite_mfma_i32_16x16x4i8:1053; CHECK:       ; %bb.0:1054; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1055; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[2:3], off offset:481056; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[2:3], off offset:321057; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[2:3], off offset:161058; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off1059; CHECK-NEXT:    s_waitcnt vmcnt(0)1060; CHECK-NEXT:    v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[0:15]1061; CHECK-NEXT:    ;;#ASMSTART1062; CHECK-NEXT:    ; use a[0:15]1063; CHECK-NEXT:    ;;#ASMEND1064; CHECK-NEXT:    s_setpc_b64 s[30:31]1065  %src2 = load <16 x i32>, ptr addrspace(1) %ptr1066  %mai = call <16 x i32> @llvm.amdgcn.mfma.i32.16x16x4i8(i32 %arg0, i32 %arg1, <16 x i32> %src2, i32 0, i32 0, i32 0)1067  call void asm sideeffect "; use $0", "a"(<16 x i32> %mai)1068  ret void1069}1070 1071define void @test_rewrite_mfma_i32_4x4x4i8(i32 %arg0, i32 %arg1, ptr addrspace(1) %ptr) #0 {1072; CHECK-LABEL: test_rewrite_mfma_i32_4x4x4i8:1073; CHECK:       ; %bb.0:1074; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1075; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[2:3], off1076; CHECK-NEXT:    s_waitcnt vmcnt(0)1077; CHECK-NEXT:    v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[0:3]1078; CHECK-NEXT:    ;;#ASMSTART1079; CHECK-NEXT:    ; use a[0:3]1080; CHECK-NEXT:    ;;#ASMEND1081; CHECK-NEXT:    s_setpc_b64 s[30:31]1082  %src2 = load <4 x i32>, ptr addrspace(1) %ptr1083  %mai = call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 %arg0, i32 %arg1, <4 x i32> %src2, i32 0, i32 0, i32 0)1084  call void asm sideeffect "; use $0", "a"(<4 x i32> %mai)1085  ret void1086}1087 1088;---------------------------------------------------------------------1089; gfx90a intrinsics1090;---------------------------------------------------------------------1091 1092define void @test_rewrite_mfma_f32_32x32x4bf16_1k(<4 x i16> %arg0, <4 x i16> %arg1, ptr addrspace(1) %ptr) #0 {1093; CHECK-LABEL: test_rewrite_mfma_f32_32x32x4bf16_1k:1094; CHECK:       ; %bb.0:1095; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1096; CHECK-NEXT:    global_load_dwordx4 a[28:31], v[4:5], off offset:1121097; CHECK-NEXT:    global_load_dwordx4 a[24:27], v[4:5], off offset:961098; CHECK-NEXT:    global_load_dwordx4 a[20:23], v[4:5], off offset:801099; CHECK-NEXT:    global_load_dwordx4 a[16:19], v[4:5], off offset:641100; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:481101; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:321102; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161103; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1104; CHECK-NEXT:    s_waitcnt vmcnt(0)1105; CHECK-NEXT:    v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[0:1], v[2:3], a[0:31]1106; CHECK-NEXT:    ;;#ASMSTART1107; CHECK-NEXT:    ; use a[0:31]1108; CHECK-NEXT:    ;;#ASMEND1109; CHECK-NEXT:    s_setpc_b64 s[30:31]1110  %src2 = load <32 x float>, ptr addrspace(1) %ptr1111  %mai = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16.1k(<4 x i16> %arg0, <4 x i16> %arg1, <32 x float> %src2, i32 0, i32 0, i32 0)1112  call void asm sideeffect "; use $0", "a"(<32 x float> %mai)1113  ret void1114}1115 1116define void @test_rewrite_mfma_f32_16x16x4bf16_1k(<4 x i16> %arg0, <4 x i16> %arg1, ptr addrspace(1) %ptr) #0 {1117; CHECK-LABEL: test_rewrite_mfma_f32_16x16x4bf16_1k:1118; CHECK:       ; %bb.0:1119; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1120; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:481121; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:321122; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161123; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1124; CHECK-NEXT:    s_waitcnt vmcnt(0)1125; CHECK-NEXT:    v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[0:1], v[2:3], a[0:15]1126; CHECK-NEXT:    ;;#ASMSTART1127; CHECK-NEXT:    ; use a[0:15]1128; CHECK-NEXT:    ;;#ASMEND1129; CHECK-NEXT:    s_setpc_b64 s[30:31]1130  %src2 = load <16 x float>, ptr addrspace(1) %ptr1131  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4bf16.1k(<4 x i16> %arg0, <4 x i16> %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)1132  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1133  ret void1134}1135 1136define void @test_rewrite_mfma_f32_4x4x4bf16_1k(<4 x i16> %arg0, <4 x i16> %arg1, ptr addrspace(1) %ptr) #0 {1137; CHECK-LABEL: test_rewrite_mfma_f32_4x4x4bf16_1k:1138; CHECK:       ; %bb.0:1139; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1140; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1141; CHECK-NEXT:    s_waitcnt vmcnt(0)1142; CHECK-NEXT:    v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[0:1], v[2:3], a[0:3]1143; CHECK-NEXT:    ;;#ASMSTART1144; CHECK-NEXT:    ; use a[0:3]1145; CHECK-NEXT:    ;;#ASMEND1146; CHECK-NEXT:    s_setpc_b64 s[30:31]1147  %src2 = load <4 x float>, ptr addrspace(1) %ptr1148  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4bf16.1k(<4 x i16> %arg0, <4 x i16> %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)1149  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1150  ret void1151}1152 1153define void @test_rewrite_mfma_f32_32x32x8bf16_1k(<4 x i16> %arg0, <4 x i16> %arg1, ptr addrspace(1) %ptr) #0 {1154; CHECK-LABEL: test_rewrite_mfma_f32_32x32x8bf16_1k:1155; CHECK:       ; %bb.0:1156; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1157; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:481158; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:321159; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161160; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1161; CHECK-NEXT:    s_waitcnt vmcnt(0)1162; CHECK-NEXT:    v_mfma_f32_32x32x8_bf16 a[0:15], v[0:1], v[2:3], a[0:15]1163; CHECK-NEXT:    ;;#ASMSTART1164; CHECK-NEXT:    ; use a[0:15]1165; CHECK-NEXT:    ;;#ASMEND1166; CHECK-NEXT:    s_setpc_b64 s[30:31]1167  %src2 = load <16 x float>, ptr addrspace(1) %ptr1168  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8bf16.1k(<4 x i16> %arg0, <4 x i16> %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)1169  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1170  ret void1171}1172 1173define void @test_rewrite_mfma_f32_16x16x16bf16_1k(<4 x i16> %arg0, <4 x i16> %arg1, ptr addrspace(1) %ptr) #0 {1174; CHECK-LABEL: test_rewrite_mfma_f32_16x16x16bf16_1k:1175; CHECK:       ; %bb.0:1176; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1177; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1178; CHECK-NEXT:    s_waitcnt vmcnt(0)1179; CHECK-NEXT:    v_mfma_f32_16x16x16_bf16 a[0:3], v[0:1], v[2:3], a[0:3]1180; CHECK-NEXT:    ;;#ASMSTART1181; CHECK-NEXT:    ; use a[0:3]1182; CHECK-NEXT:    ;;#ASMEND1183; CHECK-NEXT:    s_setpc_b64 s[30:31]1184  %src2 = load <4 x float>, ptr addrspace(1) %ptr1185  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> %arg0, <4 x i16> %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)1186  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1187  ret void1188}1189 1190define void @test_rewrite_mfma_f64_16x16x4f64(double %arg0, double %arg1, ptr addrspace(1) %ptr) #0 {1191; CHECK-LABEL: test_rewrite_mfma_f64_16x16x4f64:1192; CHECK:       ; %bb.0:1193; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1194; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161195; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1196; CHECK-NEXT:    s_waitcnt vmcnt(0)1197; CHECK-NEXT:    v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]1198; CHECK-NEXT:    ;;#ASMSTART1199; CHECK-NEXT:    ; use a[0:7]1200; CHECK-NEXT:    ;;#ASMEND1201; CHECK-NEXT:    s_setpc_b64 s[30:31]1202  %src2 = load <4 x double>, ptr addrspace(1) %ptr1203  %mai = call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double %arg0, double %arg1, <4 x double> %src2, i32 0, i32 0, i32 0)1204  call void asm sideeffect "; use $0", "a"(<4 x double> %mai)1205  ret void1206}1207 1208define void @test_rewrite_mfma_f64_4x4xf64(double %arg0, double %arg1, ptr addrspace(1) %ptr) #0 {1209; CHECK-LABEL: test_rewrite_mfma_f64_4x4xf64:1210; CHECK:       ; %bb.0:1211; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1212; CHECK-NEXT:    global_load_dwordx2 a[0:1], v[4:5], off1213; CHECK-NEXT:    s_waitcnt vmcnt(0)1214; CHECK-NEXT:    v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], v[2:3], a[0:1]1215; CHECK-NEXT:    ;;#ASMSTART1216; CHECK-NEXT:    ; use a[0:1]1217; CHECK-NEXT:    ;;#ASMEND1218; CHECK-NEXT:    s_setpc_b64 s[30:31]1219  %src2 = load double, ptr addrspace(1) %ptr1220  %mai = call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %arg0, double %arg1, double %src2, i32 0, i32 0, i32 0)1221  call void asm sideeffect "; use $0", "a"(double %mai)1222  ret void1223}1224 1225;---------------------------------------------------------------------1226; gfx942 intrinsics1227;---------------------------------------------------------------------1228 1229define void @test_rewrite_mfma_i32_16x16x32_i8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1230; CHECK-LABEL: test_rewrite_mfma_i32_16x16x32_i8:1231; CHECK:       ; %bb.0:1232; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1233; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1234; CHECK-NEXT:    s_waitcnt vmcnt(0)1235; CHECK-NEXT:    v_mfma_i32_16x16x32_i8 a[0:3], v[0:1], v[2:3], a[0:3]1236; CHECK-NEXT:    ;;#ASMSTART1237; CHECK-NEXT:    ; use a[0:3]1238; CHECK-NEXT:    ;;#ASMEND1239; CHECK-NEXT:    s_setpc_b64 s[30:31]1240  %src2 = load <4 x i32>, ptr addrspace(1) %ptr1241  %mai = call <4 x i32> @llvm.amdgcn.mfma.i32.16x16x32.i8(i64 %arg0, i64 %arg1, <4 x i32> %src2, i32 0, i32 0, i32 0)1242  call void asm sideeffect "; use $0", "a"(<4 x i32> %mai)1243  ret void1244}1245 1246define void @test_rewrite_mfma_i32_32x32x16_i8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1247; CHECK-LABEL: test_rewrite_mfma_i32_32x32x16_i8:1248; CHECK:       ; %bb.0:1249; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1250; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:481251; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:321252; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161253; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1254; CHECK-NEXT:    s_waitcnt vmcnt(0)1255; CHECK-NEXT:    v_mfma_i32_32x32x16_i8 a[0:15], v[0:1], v[2:3], a[0:15]1256; CHECK-NEXT:    ;;#ASMSTART1257; CHECK-NEXT:    ; use a[0:15]1258; CHECK-NEXT:    ;;#ASMEND1259; CHECK-NEXT:    s_setpc_b64 s[30:31]1260  %src2 = load <16 x i32>, ptr addrspace(1) %ptr1261  %mai = call <16 x i32> @llvm.amdgcn.mfma.i32.32x32x16.i8(i64 %arg0, i64 %arg1, <16 x i32> %src2, i32 0, i32 0, i32 0)1262  call void asm sideeffect "; use $0", "a"(<16 x i32> %mai)1263  ret void1264}1265 1266define void @test_rewrite_mfma_f32_16x16x8_xf32(<2 x float> %arg0, <2 x float> %arg1, ptr addrspace(1) %ptr) #0 {1267; CHECK-LABEL: test_rewrite_mfma_f32_16x16x8_xf32:1268; CHECK:       ; %bb.0:1269; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1270; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1271; CHECK-NEXT:    s_waitcnt vmcnt(0)1272; CHECK-NEXT:    v_mfma_f32_16x16x8_xf32 a[0:3], v[0:1], v[2:3], a[0:3]1273; CHECK-NEXT:    ;;#ASMSTART1274; CHECK-NEXT:    ; use a[0:3]1275; CHECK-NEXT:    ;;#ASMEND1276; CHECK-NEXT:    s_setpc_b64 s[30:31]1277  %src2 = load <4 x float>, ptr addrspace(1) %ptr1278  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8.xf32(<2 x float> %arg0, <2 x float> %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)1279  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1280  ret void1281}1282 1283define void @test_rewrite_mfma_f32_32x32x4_xf32(<2 x float> %arg0, <2 x float> %arg1, ptr addrspace(1) %ptr) #0 {1284; CHECK-LABEL: test_rewrite_mfma_f32_32x32x4_xf32:1285; CHECK:       ; %bb.0:1286; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1287; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:481288; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:321289; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161290; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1291; CHECK-NEXT:    s_waitcnt vmcnt(0)1292; CHECK-NEXT:    v_mfma_f32_32x32x4_xf32 a[0:15], v[0:1], v[2:3], a[0:15]1293; CHECK-NEXT:    ;;#ASMSTART1294; CHECK-NEXT:    ; use a[0:15]1295; CHECK-NEXT:    ;;#ASMEND1296; CHECK-NEXT:    s_setpc_b64 s[30:31]1297  %src2 = load <16 x float>, ptr addrspace(1) %ptr1298  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4.xf32(<2 x float> %arg0, <2 x float> %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)1299  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1300  ret void1301}1302 1303define void @test_rewrite_mfma_f32_16x16x32_bf8_bf8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1304; CHECK-LABEL: test_rewrite_mfma_f32_16x16x32_bf8_bf8:1305; CHECK:       ; %bb.0:1306; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1307; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1308; CHECK-NEXT:    s_waitcnt vmcnt(0)1309; CHECK-NEXT:    v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[0:1], v[2:3], a[0:3]1310; CHECK-NEXT:    ;;#ASMSTART1311; CHECK-NEXT:    ; use a[0:3]1312; CHECK-NEXT:    ;;#ASMEND1313; CHECK-NEXT:    s_setpc_b64 s[30:31]1314  %src2 = load <4 x float>, ptr addrspace(1) %ptr1315  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.bf8.bf8(i64 %arg0, i64 %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)1316  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1317  ret void1318}1319 1320define void @test_rewrite_mfma_f32_16x16x32_bf8_fp8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1321; CHECK-LABEL: test_rewrite_mfma_f32_16x16x32_bf8_fp8:1322; CHECK:       ; %bb.0:1323; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1324; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1325; CHECK-NEXT:    s_waitcnt vmcnt(0)1326; CHECK-NEXT:    v_mfma_f32_16x16x32_bf8_fp8 a[0:3], v[0:1], v[2:3], a[0:3]1327; CHECK-NEXT:    ;;#ASMSTART1328; CHECK-NEXT:    ; use a[0:3]1329; CHECK-NEXT:    ;;#ASMEND1330; CHECK-NEXT:    s_setpc_b64 s[30:31]1331  %src2 = load <4 x float>, ptr addrspace(1) %ptr1332  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.bf8.fp8(i64 %arg0, i64 %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)1333  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1334  ret void1335}1336 1337define void @test_rewrite_mfma_f32_16x16x32_fp8_bf8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1338; CHECK-LABEL: test_rewrite_mfma_f32_16x16x32_fp8_bf8:1339; CHECK:       ; %bb.0:1340; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1341; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1342; CHECK-NEXT:    s_waitcnt vmcnt(0)1343; CHECK-NEXT:    v_mfma_f32_16x16x32_fp8_bf8 a[0:3], v[0:1], v[2:3], a[0:3]1344; CHECK-NEXT:    ;;#ASMSTART1345; CHECK-NEXT:    ; use a[0:3]1346; CHECK-NEXT:    ;;#ASMEND1347; CHECK-NEXT:    s_setpc_b64 s[30:31]1348  %src2 = load <4 x float>, ptr addrspace(1) %ptr1349  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.fp8.bf8(i64 %arg0, i64 %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)1350  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1351  ret void1352}1353 1354define void @test_rewrite_mfma_f32_16x16x32_fp8_fp8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1355; CHECK-LABEL: test_rewrite_mfma_f32_16x16x32_fp8_fp8:1356; CHECK:       ; %bb.0:1357; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1358; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1359; CHECK-NEXT:    s_waitcnt vmcnt(0)1360; CHECK-NEXT:    v_mfma_f32_16x16x32_fp8_fp8 a[0:3], v[0:1], v[2:3], a[0:3]1361; CHECK-NEXT:    ;;#ASMSTART1362; CHECK-NEXT:    ; use a[0:3]1363; CHECK-NEXT:    ;;#ASMEND1364; CHECK-NEXT:    s_setpc_b64 s[30:31]1365  %src2 = load <4 x float>, ptr addrspace(1) %ptr1366  %mai = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.fp8.fp8(i64 %arg0, i64 %arg1, <4 x float> %src2, i32 0, i32 0, i32 0)1367  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1368  ret void1369}1370 1371define void @test_rewrite_mfma_f32_32x32x16_bf8_bf8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1372; CHECK-LABEL: test_rewrite_mfma_f32_32x32x16_bf8_bf8:1373; CHECK:       ; %bb.0:1374; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1375; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:481376; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:321377; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161378; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1379; CHECK-NEXT:    s_waitcnt vmcnt(0)1380; CHECK-NEXT:    v_mfma_f32_32x32x16_bf8_bf8 a[0:15], v[0:1], v[2:3], a[0:15]1381; CHECK-NEXT:    ;;#ASMSTART1382; CHECK-NEXT:    ; use a[0:15]1383; CHECK-NEXT:    ;;#ASMEND1384; CHECK-NEXT:    s_setpc_b64 s[30:31]1385  %src2 = load <16 x float>, ptr addrspace(1) %ptr1386  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf8.bf8(i64 %arg0, i64 %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)1387  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1388  ret void1389}1390 1391define void @test_rewrite_mfma_f32_32x32x16_bf8_fp8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1392; CHECK-LABEL: test_rewrite_mfma_f32_32x32x16_bf8_fp8:1393; CHECK:       ; %bb.0:1394; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1395; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:481396; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:321397; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161398; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1399; CHECK-NEXT:    s_waitcnt vmcnt(0)1400; CHECK-NEXT:    v_mfma_f32_32x32x16_bf8_fp8 a[0:15], v[0:1], v[2:3], a[0:15]1401; CHECK-NEXT:    ;;#ASMSTART1402; CHECK-NEXT:    ; use a[0:15]1403; CHECK-NEXT:    ;;#ASMEND1404; CHECK-NEXT:    s_setpc_b64 s[30:31]1405  %src2 = load <16 x float>, ptr addrspace(1) %ptr1406  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf8.fp8(i64 %arg0, i64 %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)1407  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1408  ret void1409}1410 1411define void @test_rewrite_mfma_f32_32x32x16_fp8_bf8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1412; CHECK-LABEL: test_rewrite_mfma_f32_32x32x16_fp8_bf8:1413; CHECK:       ; %bb.0:1414; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1415; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:481416; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:321417; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161418; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1419; CHECK-NEXT:    s_waitcnt vmcnt(0)1420; CHECK-NEXT:    v_mfma_f32_32x32x16_fp8_bf8 a[0:15], v[0:1], v[2:3], a[0:15]1421; CHECK-NEXT:    ;;#ASMSTART1422; CHECK-NEXT:    ; use a[0:15]1423; CHECK-NEXT:    ;;#ASMEND1424; CHECK-NEXT:    s_setpc_b64 s[30:31]1425  %src2 = load <16 x float>, ptr addrspace(1) %ptr1426  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.fp8.bf8(i64 %arg0, i64 %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)1427  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1428  ret void1429}1430 1431define void @test_rewrite_mfma_f32_32x32x16_fp8_fp8(i64 %arg0, i64 %arg1, ptr addrspace(1) %ptr) #0 {1432; CHECK-LABEL: test_rewrite_mfma_f32_32x32x16_fp8_fp8:1433; CHECK:       ; %bb.0:1434; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1435; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[4:5], off offset:481436; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[4:5], off offset:321437; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[4:5], off offset:161438; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[4:5], off1439; CHECK-NEXT:    s_waitcnt vmcnt(0)1440; CHECK-NEXT:    v_mfma_f32_32x32x16_fp8_fp8 a[0:15], v[0:1], v[2:3], a[0:15]1441; CHECK-NEXT:    ;;#ASMSTART1442; CHECK-NEXT:    ; use a[0:15]1443; CHECK-NEXT:    ;;#ASMEND1444; CHECK-NEXT:    s_setpc_b64 s[30:31]1445  %src2 = load <16 x float>, ptr addrspace(1) %ptr1446  %mai = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.fp8.fp8(i64 %arg0, i64 %arg1, <16 x float> %src2, i32 0, i32 0, i32 0)1447  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1448  ret void1449}1450 1451define void @test_rewrite_smfmac_f32_16x16x32_f16(<4 x half> %arg0, <8 x half> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1452; CHECK-LABEL: test_rewrite_smfmac_f32_16x16x32_f16:1453; CHECK:       ; %bb.0:1454; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1455; CHECK-NEXT:    v_mov_b32_e32 v9, v81456; CHECK-NEXT:    v_mov_b32_e32 v8, v71457; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1458; CHECK-NEXT:    s_waitcnt vmcnt(0)1459; CHECK-NEXT:    v_smfmac_f32_16x16x32_f16 a[0:3], v[0:1], v[2:5], v61460; CHECK-NEXT:    ;;#ASMSTART1461; CHECK-NEXT:    ; use a[0:3]1462; CHECK-NEXT:    ;;#ASMEND1463; CHECK-NEXT:    s_setpc_b64 s[30:31]1464  %src2 = load <4 x float>, ptr addrspace(1) %ptr1465  %mai = call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x32.f16(<4 x half> %arg0, <8 x half> %arg1, <4 x float> %src2, i32 %arg2, i32 0, i32 0)1466  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1467  ret void1468}1469 1470define void @test_rewrite_smfmac_f32_32x32x16_f16(<4 x half> %arg0, <8 x half> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1471; CHECK-LABEL: test_rewrite_smfmac_f32_32x32x16_f16:1472; CHECK:       ; %bb.0:1473; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1474; CHECK-NEXT:    v_mov_b32_e32 v9, v81475; CHECK-NEXT:    v_mov_b32_e32 v8, v71476; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[8:9], off offset:481477; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[8:9], off offset:321478; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[8:9], off offset:161479; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1480; CHECK-NEXT:    s_waitcnt vmcnt(0)1481; CHECK-NEXT:    v_smfmac_f32_32x32x16_f16 a[0:15], v[0:1], v[2:5], v61482; CHECK-NEXT:    ;;#ASMSTART1483; CHECK-NEXT:    ; use a[0:15]1484; CHECK-NEXT:    ;;#ASMEND1485; CHECK-NEXT:    s_setpc_b64 s[30:31]1486  %src2 = load <16 x float>, ptr addrspace(1) %ptr1487  %mai = call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x16.f16(<4 x half> %arg0, <8 x half> %arg1, <16 x float> %src2, i32 %arg2, i32 0, i32 0)1488  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1489  ret void1490}1491 1492define void @test_rewrite_smfmac_f32_16x16x32_bf16(<4 x i16> %arg0, <8 x i16> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1493; CHECK-LABEL: test_rewrite_smfmac_f32_16x16x32_bf16:1494; CHECK:       ; %bb.0:1495; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1496; CHECK-NEXT:    v_mov_b32_e32 v9, v81497; CHECK-NEXT:    v_mov_b32_e32 v8, v71498; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1499; CHECK-NEXT:    s_waitcnt vmcnt(0)1500; CHECK-NEXT:    v_smfmac_f32_16x16x32_bf16 a[0:3], v[0:1], v[2:5], v61501; CHECK-NEXT:    ;;#ASMSTART1502; CHECK-NEXT:    ; use a[0:3]1503; CHECK-NEXT:    ;;#ASMEND1504; CHECK-NEXT:    s_setpc_b64 s[30:31]1505  %src2 = load <4 x float>, ptr addrspace(1) %ptr1506  %mai = call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x32.bf16(<4 x i16> %arg0, <8 x i16> %arg1, <4 x float> %src2, i32 %arg2, i32 0, i32 0)1507  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1508  ret void1509}1510 1511define void @test_rewrite_smfmac_f32_32x32x16_bf16(<4 x i16> %arg0, <8 x i16> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1512; CHECK-LABEL: test_rewrite_smfmac_f32_32x32x16_bf16:1513; CHECK:       ; %bb.0:1514; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1515; CHECK-NEXT:    v_mov_b32_e32 v9, v81516; CHECK-NEXT:    v_mov_b32_e32 v8, v71517; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[8:9], off offset:481518; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[8:9], off offset:321519; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[8:9], off offset:161520; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1521; CHECK-NEXT:    s_waitcnt vmcnt(0)1522; CHECK-NEXT:    v_smfmac_f32_32x32x16_bf16 a[0:15], v[0:1], v[2:5], v61523; CHECK-NEXT:    ;;#ASMSTART1524; CHECK-NEXT:    ; use a[0:15]1525; CHECK-NEXT:    ;;#ASMEND1526; CHECK-NEXT:    s_setpc_b64 s[30:31]1527  %src2 = load <16 x float>, ptr addrspace(1) %ptr1528  %mai = call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x16.bf16(<4 x i16> %arg0, <8 x i16> %arg1, <16 x float> %src2, i32 %arg2, i32 0, i32 0)1529  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1530  ret void1531}1532 1533define void @test_rewrite_smfmac_i32_16x16x64_i8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1534; CHECK-LABEL: test_rewrite_smfmac_i32_16x16x64_i8:1535; CHECK:       ; %bb.0:1536; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1537; CHECK-NEXT:    v_mov_b32_e32 v9, v81538; CHECK-NEXT:    v_mov_b32_e32 v8, v71539; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1540; CHECK-NEXT:    s_waitcnt vmcnt(0)1541; CHECK-NEXT:    v_smfmac_i32_16x16x64_i8 a[0:3], v[0:1], v[2:5], v61542; CHECK-NEXT:    ;;#ASMSTART1543; CHECK-NEXT:    ; use a[0:3]1544; CHECK-NEXT:    ;;#ASMEND1545; CHECK-NEXT:    s_setpc_b64 s[30:31]1546  %src2 = load <4 x i32>, ptr addrspace(1) %ptr1547  %mai = call <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x64.i8(<2 x i32> %arg0, <4 x i32> %arg1, <4 x i32> %src2, i32 %arg2, i32 0, i32 0)1548  call void asm sideeffect "; use $0", "a"(<4 x i32> %mai)1549  ret void1550}1551 1552define void @test_rewrite_smfmac_i32_32x32x32_i8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1553; CHECK-LABEL: test_rewrite_smfmac_i32_32x32x32_i8:1554; CHECK:       ; %bb.0:1555; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1556; CHECK-NEXT:    v_mov_b32_e32 v9, v81557; CHECK-NEXT:    v_mov_b32_e32 v8, v71558; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[8:9], off offset:481559; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[8:9], off offset:321560; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[8:9], off offset:161561; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1562; CHECK-NEXT:    s_waitcnt vmcnt(0)1563; CHECK-NEXT:    v_smfmac_i32_32x32x32_i8 a[0:15], v[0:1], v[2:5], v61564; CHECK-NEXT:    ;;#ASMSTART1565; CHECK-NEXT:    ; use a[0:15]1566; CHECK-NEXT:    ;;#ASMEND1567; CHECK-NEXT:    s_setpc_b64 s[30:31]1568  %src2 = load <16 x i32>, ptr addrspace(1) %ptr1569  %mai = call <16 x i32> @llvm.amdgcn.smfmac.i32.32x32x32.i8(<2 x i32> %arg0, <4 x i32> %arg1, <16 x i32> %src2, i32 %arg2, i32 0, i32 0)1570  call void asm sideeffect "; use $0", "a"(<16 x i32> %mai)1571  ret void1572}1573 1574define void @test_rewrite_smfmac_16x16x64_bf8_bf8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1575; CHECK-LABEL: test_rewrite_smfmac_16x16x64_bf8_bf8:1576; CHECK:       ; %bb.0:1577; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1578; CHECK-NEXT:    v_mov_b32_e32 v9, v81579; CHECK-NEXT:    v_mov_b32_e32 v8, v71580; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1581; CHECK-NEXT:    s_waitcnt vmcnt(0)1582; CHECK-NEXT:    v_smfmac_f32_16x16x64_bf8_bf8 a[0:3], v[0:1], v[2:5], v61583; CHECK-NEXT:    ;;#ASMSTART1584; CHECK-NEXT:    ; use a[0:3]1585; CHECK-NEXT:    ;;#ASMEND1586; CHECK-NEXT:    s_setpc_b64 s[30:31]1587  %src2 = load <4 x float>, ptr addrspace(1) %ptr1588  %mai = tail call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x64.bf8.bf8(<2 x i32> %arg0, <4 x i32> %arg1, <4 x float> %src2, i32 %arg2, i32 0, i32 0)1589  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1590  ret void1591}1592 1593define void @test_rewrite_smfmac_16x16x64_bf8_fp8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1594; CHECK-LABEL: test_rewrite_smfmac_16x16x64_bf8_fp8:1595; CHECK:       ; %bb.0:1596; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1597; CHECK-NEXT:    v_mov_b32_e32 v9, v81598; CHECK-NEXT:    v_mov_b32_e32 v8, v71599; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1600; CHECK-NEXT:    s_waitcnt vmcnt(0)1601; CHECK-NEXT:    v_smfmac_f32_16x16x64_bf8_fp8 a[0:3], v[0:1], v[2:5], v61602; CHECK-NEXT:    ;;#ASMSTART1603; CHECK-NEXT:    ; use a[0:3]1604; CHECK-NEXT:    ;;#ASMEND1605; CHECK-NEXT:    s_setpc_b64 s[30:31]1606  %src2 = load <4 x float>, ptr addrspace(1) %ptr1607  %mai = tail call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x64.bf8.fp8(<2 x i32> %arg0, <4 x i32> %arg1, <4 x float> %src2, i32 %arg2, i32 0, i32 0)1608  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1609  ret void1610}1611 1612define void @test_rewrite_smfmac_16x16x64_fp8_bf8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1613; CHECK-LABEL: test_rewrite_smfmac_16x16x64_fp8_bf8:1614; CHECK:       ; %bb.0:1615; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1616; CHECK-NEXT:    v_mov_b32_e32 v9, v81617; CHECK-NEXT:    v_mov_b32_e32 v8, v71618; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1619; CHECK-NEXT:    s_waitcnt vmcnt(0)1620; CHECK-NEXT:    v_smfmac_f32_16x16x64_fp8_bf8 a[0:3], v[0:1], v[2:5], v61621; CHECK-NEXT:    ;;#ASMSTART1622; CHECK-NEXT:    ; use a[0:3]1623; CHECK-NEXT:    ;;#ASMEND1624; CHECK-NEXT:    s_setpc_b64 s[30:31]1625  %src2 = load <4 x float>, ptr addrspace(1) %ptr1626  %mai = tail call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x64.fp8.bf8(<2 x i32> %arg0, <4 x i32> %arg1, <4 x float> %src2, i32 %arg2, i32 0, i32 0)1627  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1628  ret void1629}1630 1631define void @test_rewrite_smfmac_16x16x64_fp8_fp8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1632; CHECK-LABEL: test_rewrite_smfmac_16x16x64_fp8_fp8:1633; CHECK:       ; %bb.0:1634; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1635; CHECK-NEXT:    v_mov_b32_e32 v9, v81636; CHECK-NEXT:    v_mov_b32_e32 v8, v71637; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1638; CHECK-NEXT:    s_waitcnt vmcnt(0)1639; CHECK-NEXT:    v_smfmac_f32_16x16x64_fp8_fp8 a[0:3], v[0:1], v[2:5], v61640; CHECK-NEXT:    ;;#ASMSTART1641; CHECK-NEXT:    ; use a[0:3]1642; CHECK-NEXT:    ;;#ASMEND1643; CHECK-NEXT:    s_setpc_b64 s[30:31]1644  %src2 = load <4 x float>, ptr addrspace(1) %ptr1645  %mai = tail call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x64.fp8.fp8(<2 x i32> %arg0, <4 x i32> %arg1, <4 x float> %src2, i32 %arg2, i32 0, i32 0)1646  call void asm sideeffect "; use $0", "a"(<4 x float> %mai)1647  ret void1648}1649 1650define void @test_rewrite_smfmac_32x32x32_bf8_bf8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1651; CHECK-LABEL: test_rewrite_smfmac_32x32x32_bf8_bf8:1652; CHECK:       ; %bb.0:1653; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1654; CHECK-NEXT:    v_mov_b32_e32 v9, v81655; CHECK-NEXT:    v_mov_b32_e32 v8, v71656; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[8:9], off offset:481657; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[8:9], off offset:321658; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[8:9], off offset:161659; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1660; CHECK-NEXT:    s_waitcnt vmcnt(0)1661; CHECK-NEXT:    v_smfmac_f32_32x32x32_bf8_bf8 a[0:15], v[0:1], v[2:5], v61662; CHECK-NEXT:    ;;#ASMSTART1663; CHECK-NEXT:    ; use a[0:15]1664; CHECK-NEXT:    ;;#ASMEND1665; CHECK-NEXT:    s_setpc_b64 s[30:31]1666  %src2 = load <16 x float>, ptr addrspace(1) %ptr1667  %mai = tail call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.bf8.bf8(<2 x i32> %arg0, <4 x i32> %arg1, <16 x float> %src2, i32 %arg2, i32 0, i32 0)1668  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1669  ret void1670}1671 1672define void @test_rewrite_smfmac_32x32x32_bf8_fp8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1673; CHECK-LABEL: test_rewrite_smfmac_32x32x32_bf8_fp8:1674; CHECK:       ; %bb.0:1675; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1676; CHECK-NEXT:    v_mov_b32_e32 v9, v81677; CHECK-NEXT:    v_mov_b32_e32 v8, v71678; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[8:9], off offset:481679; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[8:9], off offset:321680; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[8:9], off offset:161681; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1682; CHECK-NEXT:    s_waitcnt vmcnt(0)1683; CHECK-NEXT:    v_smfmac_f32_32x32x32_bf8_fp8 a[0:15], v[0:1], v[2:5], v61684; CHECK-NEXT:    ;;#ASMSTART1685; CHECK-NEXT:    ; use a[0:15]1686; CHECK-NEXT:    ;;#ASMEND1687; CHECK-NEXT:    s_setpc_b64 s[30:31]1688  %src2 = load <16 x float>, ptr addrspace(1) %ptr1689  %mai = tail call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.bf8.fp8(<2 x i32> %arg0, <4 x i32> %arg1, <16 x float> %src2, i32 %arg2, i32 0, i32 0)1690  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1691  ret void1692}1693 1694define void @test_rewrite_smfmac_32x32x32_fp8_bf8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1695; CHECK-LABEL: test_rewrite_smfmac_32x32x32_fp8_bf8:1696; CHECK:       ; %bb.0:1697; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1698; CHECK-NEXT:    v_mov_b32_e32 v9, v81699; CHECK-NEXT:    v_mov_b32_e32 v8, v71700; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[8:9], off offset:481701; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[8:9], off offset:321702; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[8:9], off offset:161703; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1704; CHECK-NEXT:    s_waitcnt vmcnt(0)1705; CHECK-NEXT:    v_smfmac_f32_32x32x32_fp8_bf8 a[0:15], v[0:1], v[2:5], v61706; CHECK-NEXT:    ;;#ASMSTART1707; CHECK-NEXT:    ; use a[0:15]1708; CHECK-NEXT:    ;;#ASMEND1709; CHECK-NEXT:    s_setpc_b64 s[30:31]1710  %src2 = load <16 x float>, ptr addrspace(1) %ptr1711  %mai = tail call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.fp8.bf8(<2 x i32> %arg0, <4 x i32> %arg1, <16 x float> %src2, i32 %arg2, i32 0, i32 0)1712  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1713  ret void1714}1715 1716define void @test_rewrite_smfmac_32x32x32_fp8_fp8(<2 x i32> %arg0, <4 x i32> %arg1, i32 %arg2, ptr addrspace(1) %ptr) #0 {1717; CHECK-LABEL: test_rewrite_smfmac_32x32x32_fp8_fp8:1718; CHECK:       ; %bb.0:1719; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1720; CHECK-NEXT:    v_mov_b32_e32 v9, v81721; CHECK-NEXT:    v_mov_b32_e32 v8, v71722; CHECK-NEXT:    global_load_dwordx4 a[12:15], v[8:9], off offset:481723; CHECK-NEXT:    global_load_dwordx4 a[8:11], v[8:9], off offset:321724; CHECK-NEXT:    global_load_dwordx4 a[4:7], v[8:9], off offset:161725; CHECK-NEXT:    global_load_dwordx4 a[0:3], v[8:9], off1726; CHECK-NEXT:    s_waitcnt vmcnt(0)1727; CHECK-NEXT:    v_smfmac_f32_32x32x32_fp8_fp8 a[0:15], v[0:1], v[2:5], v61728; CHECK-NEXT:    ;;#ASMSTART1729; CHECK-NEXT:    ; use a[0:15]1730; CHECK-NEXT:    ;;#ASMEND1731; CHECK-NEXT:    s_setpc_b64 s[30:31]1732  %src2 = load <16 x float>, ptr addrspace(1) %ptr1733  %mai = tail call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.fp8.fp8(<2 x i32> %arg0, <4 x i32> %arg1, <16 x float> %src2, i32 %arg2, i32 0, i32 0)1734  call void asm sideeffect "; use $0", "a"(<16 x float> %mai)1735  ret void1736}1737 1738declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg) #21739declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32 immarg, i32 immarg, i32 immarg) #21740declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32 immarg, i32 immarg, i32 immarg) #21741declare noundef range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.x() #31742 1743attributes #0 = { nounwind "amdgpu-flat-work-group-size"="1,256" "amdgpu-waves-per-eu"="4,4" }1744attributes #1 = { mustprogress nofree norecurse nounwind willreturn "amdgpu-waves-per-eu"="8,8" }1745attributes #2 = { convergent nocallback nofree nosync nounwind willreturn memory(none) }1746attributes #3 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }1747attributes #4 = { nounwind noinline optnone }1748