377 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py2; RUN: llc -verify-machineinstrs -global-isel -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck --check-prefixes=GCN %s3 4declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16.1k(<4 x i16>, <4 x i16>, <32 x float>, i32, i32, i32)5declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x4bf16.1k(<4 x i16>, <4 x i16>, <16 x float>, i32, i32, i32)6declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x4bf16.1k(<4 x i16>, <4 x i16>, <4 x float>, i32, i32, i32)7declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x8bf16.1k(<4 x i16>, <4 x i16>, <16 x float>, i32, i32, i32)8declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16>, <4 x i16>, <4 x float>, i32, i32, i32)9declare <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double, double, <4 x double>, i32, i32, i32)10declare double @llvm.amdgcn.mfma.f64.4x4x4f64(double, double, double, i32, i32, i32)11declare i32 @llvm.amdgcn.workitem.id.x()12 13define amdgpu_kernel void @test_mfma_f32_32x32x4bf16_1k(ptr addrspace(1) %arg) #0 {14; GCN-LABEL: test_mfma_f32_32x32x4bf16_1k:15; GCN: ; %bb.0: ; %bb16; GCN-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x2417; GCN-NEXT: s_mov_b64 s[36:37], 118; GCN-NEXT: v_pk_mov_b32 v[0:1], s[36:37], s[36:37] op_sel:[0,1]19; GCN-NEXT: s_mov_b32 s38, 220; GCN-NEXT: s_mov_b32 s39, s3721; GCN-NEXT: s_waitcnt lgkmcnt(0)22; GCN-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x023; GCN-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x4024; GCN-NEXT: v_pk_mov_b32 v[2:3], s[38:39], s[38:39] op_sel:[0,1]25; GCN-NEXT: s_waitcnt lgkmcnt(0)26; GCN-NEXT: v_accvgpr_write_b32 a0, s027; GCN-NEXT: v_accvgpr_write_b32 a16, s1628; GCN-NEXT: v_accvgpr_write_b32 a1, s129; GCN-NEXT: v_accvgpr_write_b32 a2, s230; GCN-NEXT: v_accvgpr_write_b32 a3, s331; GCN-NEXT: v_accvgpr_write_b32 a4, s432; GCN-NEXT: v_accvgpr_write_b32 a5, s533; GCN-NEXT: v_accvgpr_write_b32 a6, s634; GCN-NEXT: v_accvgpr_write_b32 a7, s735; GCN-NEXT: v_accvgpr_write_b32 a8, s836; GCN-NEXT: v_accvgpr_write_b32 a9, s937; GCN-NEXT: v_accvgpr_write_b32 a10, s1038; GCN-NEXT: v_accvgpr_write_b32 a11, s1139; GCN-NEXT: v_accvgpr_write_b32 a12, s1240; GCN-NEXT: v_accvgpr_write_b32 a13, s1341; GCN-NEXT: v_accvgpr_write_b32 a14, s1442; GCN-NEXT: v_accvgpr_write_b32 a15, s1543; GCN-NEXT: v_accvgpr_write_b32 a17, s1744; GCN-NEXT: v_accvgpr_write_b32 a18, s1845; GCN-NEXT: v_accvgpr_write_b32 a19, s1946; GCN-NEXT: v_accvgpr_write_b32 a20, s2047; GCN-NEXT: v_accvgpr_write_b32 a21, s2148; GCN-NEXT: v_accvgpr_write_b32 a22, s2249; GCN-NEXT: v_accvgpr_write_b32 a23, s2350; GCN-NEXT: v_accvgpr_write_b32 a24, s2451; GCN-NEXT: v_accvgpr_write_b32 a25, s2552; GCN-NEXT: v_accvgpr_write_b32 a26, s2653; GCN-NEXT: v_accvgpr_write_b32 a27, s2754; GCN-NEXT: v_accvgpr_write_b32 a28, s2855; GCN-NEXT: v_accvgpr_write_b32 a29, s2956; GCN-NEXT: v_accvgpr_write_b32 a30, s3057; GCN-NEXT: v_accvgpr_write_b32 a31, s3158; GCN-NEXT: s_nop 159; GCN-NEXT: v_mfma_f32_32x32x4bf16_1k a[0:31], v[0:1], v[2:3], a[0:31] cbsz:1 abid:2 blgp:360; GCN-NEXT: v_mov_b32_e32 v0, 061; GCN-NEXT: s_nop 1562; GCN-NEXT: s_nop 163; GCN-NEXT: global_store_dwordx4 v0, a[0:3], s[34:35]64; GCN-NEXT: global_store_dwordx4 v0, a[4:7], s[34:35] offset:1665; GCN-NEXT: global_store_dwordx4 v0, a[8:11], s[34:35] offset:3266; GCN-NEXT: global_store_dwordx4 v0, a[12:15], s[34:35] offset:4867; GCN-NEXT: global_store_dwordx4 v0, a[16:19], s[34:35] offset:6468; GCN-NEXT: global_store_dwordx4 v0, a[20:23], s[34:35] offset:8069; GCN-NEXT: global_store_dwordx4 v0, a[24:27], s[34:35] offset:9670; GCN-NEXT: global_store_dwordx4 v0, a[28:31], s[34:35] offset:11271; GCN-NEXT: s_endpgm72bb:73 %in.1 = load <32 x float>, ptr addrspace(1) %arg74 %a = bitcast i64 1 to <4 x i16>75 %b = bitcast i64 2 to <4 x i16>76 %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)77 store <32 x float> %mai.1, ptr addrspace(1) %arg78 ret void79}80 81define amdgpu_kernel void @test_mfma_f32_16x16x4bf16_1k(ptr addrspace(1) %arg) #0 {82; GCN-LABEL: test_mfma_f32_16x16x4bf16_1k:83; GCN: ; %bb.0: ; %bb84; GCN-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x2485; GCN-NEXT: s_mov_b64 s[18:19], 186; GCN-NEXT: v_pk_mov_b32 v[0:1], s[18:19], s[18:19] op_sel:[0,1]87; GCN-NEXT: s_mov_b32 s18, 288; GCN-NEXT: v_pk_mov_b32 v[2:3], s[18:19], s[18:19] op_sel:[0,1]89; GCN-NEXT: s_waitcnt lgkmcnt(0)90; GCN-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x091; GCN-NEXT: s_waitcnt lgkmcnt(0)92; GCN-NEXT: v_accvgpr_write_b32 a0, s093; GCN-NEXT: v_accvgpr_write_b32 a1, s194; GCN-NEXT: v_accvgpr_write_b32 a2, s295; GCN-NEXT: v_accvgpr_write_b32 a3, s396; GCN-NEXT: v_accvgpr_write_b32 a4, s497; GCN-NEXT: v_accvgpr_write_b32 a5, s598; GCN-NEXT: v_accvgpr_write_b32 a6, s699; GCN-NEXT: v_accvgpr_write_b32 a7, s7100; GCN-NEXT: v_accvgpr_write_b32 a8, s8101; GCN-NEXT: v_accvgpr_write_b32 a9, s9102; GCN-NEXT: v_accvgpr_write_b32 a10, s10103; GCN-NEXT: v_accvgpr_write_b32 a11, s11104; GCN-NEXT: v_accvgpr_write_b32 a12, s12105; GCN-NEXT: v_accvgpr_write_b32 a13, s13106; GCN-NEXT: v_accvgpr_write_b32 a14, s14107; GCN-NEXT: v_accvgpr_write_b32 a15, s15108; GCN-NEXT: s_nop 1109; GCN-NEXT: v_mfma_f32_16x16x4bf16_1k a[0:15], v[0:1], v[2:3], a[0:15] cbsz:1 abid:2 blgp:3110; GCN-NEXT: v_mov_b32_e32 v0, 0111; GCN-NEXT: s_nop 9112; GCN-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]113; GCN-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16114; GCN-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32115; GCN-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48116; GCN-NEXT: s_endpgm117bb:118 %in.1 = load <16 x float>, ptr addrspace(1) %arg119 %a = bitcast i64 1 to <4 x i16>120 %b = bitcast i64 2 to <4 x i16>121 %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)122 store <16 x float> %mai.1, ptr addrspace(1) %arg123 ret void124}125 126define amdgpu_kernel void @test_mfma_f32_4x4x4bf16_1k(ptr addrspace(1) %arg) #0 {127; GCN-LABEL: test_mfma_f32_4x4x4bf16_1k:128; GCN: ; %bb.0: ; %bb129; GCN-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24130; GCN-NEXT: s_mov_b64 s[4:5], 1131; GCN-NEXT: v_pk_mov_b32 v[0:1], s[4:5], s[4:5] op_sel:[0,1]132; GCN-NEXT: s_mov_b32 s4, 2133; GCN-NEXT: v_pk_mov_b32 v[2:3], s[4:5], s[4:5] op_sel:[0,1]134; GCN-NEXT: s_waitcnt lgkmcnt(0)135; GCN-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0136; GCN-NEXT: s_waitcnt lgkmcnt(0)137; GCN-NEXT: v_accvgpr_write_b32 a0, s0138; GCN-NEXT: v_accvgpr_write_b32 a1, s1139; GCN-NEXT: v_accvgpr_write_b32 a2, s2140; GCN-NEXT: v_accvgpr_write_b32 a3, s3141; GCN-NEXT: s_nop 1142; GCN-NEXT: v_mfma_f32_4x4x4bf16_1k a[0:3], v[0:1], v[2:3], a[0:3] cbsz:1 abid:2 blgp:3143; GCN-NEXT: v_mov_b32_e32 v0, 0144; GCN-NEXT: s_nop 3145; GCN-NEXT: global_store_dwordx4 v0, a[0:3], s[6:7]146; GCN-NEXT: s_endpgm147bb:148 %in.1 = load <4 x float>, ptr addrspace(1) %arg149 %a = bitcast i64 1 to <4 x i16>150 %b = bitcast i64 2 to <4 x i16>151 %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)152 store <4 x float> %mai.1, ptr addrspace(1) %arg153 ret void154}155 156define amdgpu_kernel void @test_mfma_f32_32x32x8bf16_1k(ptr addrspace(1) %arg) #0 {157; GCN-LABEL: test_mfma_f32_32x32x8bf16_1k:158; GCN: ; %bb.0: ; %bb159; GCN-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24160; GCN-NEXT: s_mov_b64 s[18:19], 1161; GCN-NEXT: v_pk_mov_b32 v[0:1], s[18:19], s[18:19] op_sel:[0,1]162; GCN-NEXT: s_mov_b32 s18, 2163; GCN-NEXT: v_pk_mov_b32 v[2:3], s[18:19], s[18:19] op_sel:[0,1]164; GCN-NEXT: s_waitcnt lgkmcnt(0)165; GCN-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0166; GCN-NEXT: s_waitcnt lgkmcnt(0)167; GCN-NEXT: v_accvgpr_write_b32 a0, s0168; GCN-NEXT: v_accvgpr_write_b32 a1, s1169; GCN-NEXT: v_accvgpr_write_b32 a2, s2170; GCN-NEXT: v_accvgpr_write_b32 a3, s3171; GCN-NEXT: v_accvgpr_write_b32 a4, s4172; GCN-NEXT: v_accvgpr_write_b32 a5, s5173; GCN-NEXT: v_accvgpr_write_b32 a6, s6174; GCN-NEXT: v_accvgpr_write_b32 a7, s7175; GCN-NEXT: v_accvgpr_write_b32 a8, s8176; GCN-NEXT: v_accvgpr_write_b32 a9, s9177; GCN-NEXT: v_accvgpr_write_b32 a10, s10178; GCN-NEXT: v_accvgpr_write_b32 a11, s11179; GCN-NEXT: v_accvgpr_write_b32 a12, s12180; GCN-NEXT: v_accvgpr_write_b32 a13, s13181; GCN-NEXT: v_accvgpr_write_b32 a14, s14182; GCN-NEXT: v_accvgpr_write_b32 a15, s15183; GCN-NEXT: s_nop 1184; GCN-NEXT: v_mfma_f32_32x32x8bf16_1k a[0:15], v[0:1], v[2:3], a[0:15] cbsz:1 abid:2 blgp:3185; GCN-NEXT: v_mov_b32_e32 v0, 0186; GCN-NEXT: s_nop 15187; GCN-NEXT: s_nop 1188; GCN-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]189; GCN-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16190; GCN-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32191; GCN-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48192; GCN-NEXT: s_endpgm193bb:194 %in.1 = load <16 x float>, ptr addrspace(1) %arg195 %a = bitcast i64 1 to <4 x i16>196 %b = bitcast i64 2 to <4 x i16>197 %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)198 store <16 x float> %mai.1, ptr addrspace(1) %arg199 ret void200}201 202define amdgpu_kernel void @test_mfma_f32_16x16x16bf16_1k(ptr addrspace(1) %arg) #0 {203; GCN-LABEL: test_mfma_f32_16x16x16bf16_1k:204; GCN: ; %bb.0: ; %bb205; GCN-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24206; GCN-NEXT: s_mov_b64 s[4:5], 1207; GCN-NEXT: v_pk_mov_b32 v[0:1], s[4:5], s[4:5] op_sel:[0,1]208; GCN-NEXT: s_mov_b32 s4, 2209; GCN-NEXT: v_pk_mov_b32 v[2:3], s[4:5], s[4:5] op_sel:[0,1]210; GCN-NEXT: s_waitcnt lgkmcnt(0)211; GCN-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0212; GCN-NEXT: s_waitcnt lgkmcnt(0)213; GCN-NEXT: v_accvgpr_write_b32 a0, s0214; GCN-NEXT: v_accvgpr_write_b32 a1, s1215; GCN-NEXT: v_accvgpr_write_b32 a2, s2216; GCN-NEXT: v_accvgpr_write_b32 a3, s3217; GCN-NEXT: s_nop 1218; GCN-NEXT: v_mfma_f32_16x16x16bf16_1k a[0:3], v[0:1], v[2:3], a[0:3] cbsz:1 abid:2 blgp:3219; GCN-NEXT: v_mov_b32_e32 v0, 0220; GCN-NEXT: s_nop 9221; GCN-NEXT: global_store_dwordx4 v0, a[0:3], s[6:7]222; GCN-NEXT: s_endpgm223bb:224 %in.1 = load <4 x float>, ptr addrspace(1) %arg225 %a = bitcast i64 1 to <4 x i16>226 %b = bitcast i64 2 to <4 x i16>227 %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)228 store <4 x float> %mai.1, ptr addrspace(1) %arg229 ret void230}231 232define amdgpu_kernel void @test_mfma_f64_4x4x4f64(ptr addrspace(1) %arg, double %a, double %b) #0 {233; GCN-LABEL: test_mfma_f64_4x4x4f64:234; GCN: ; %bb.0: ; %bb235; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24236; GCN-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34237; GCN-NEXT: s_waitcnt lgkmcnt(0)238; GCN-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]239; GCN-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]240; GCN-NEXT: s_nop 1241; GCN-NEXT: v_mfma_f64_4x4x4f64 a[0:1], v[0:1], v[2:3], 0242; GCN-NEXT: s_nop 3243; GCN-NEXT: v_mfma_f64_4x4x4f64 a[0:1], v[0:1], v[2:3], a[0:1] cbsz:1 abid:2 blgp:3244; GCN-NEXT: v_mov_b32_e32 v0, 0245; GCN-NEXT: s_nop 7246; GCN-NEXT: global_store_dwordx2 v0, a[0:1], s[0:1]247; GCN-NEXT: s_endpgm248bb:249 %mai.1 = tail call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %a, double %b, double 0.0, i32 0, i32 0, i32 0)250 %mai.2 = tail call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %a, double %b, double %mai.1, i32 1, i32 2, i32 3)251 store double %mai.2, ptr addrspace(1) %arg252 ret void253}254 255define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, double %a, double %b) #0 {256; GCN-LABEL: test_mfma_f64_16x16x4f64:257; GCN: ; %bb.0: ; %bb258; GCN-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24259; GCN-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34260; GCN-NEXT: s_waitcnt lgkmcnt(0)261; GCN-NEXT: v_pk_mov_b32 v[0:1], s[10:11], s[10:11] op_sel:[0,1]262; GCN-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0263; GCN-NEXT: v_pk_mov_b32 v[2:3], s[12:13], s[12:13] op_sel:[0,1]264; GCN-NEXT: s_waitcnt lgkmcnt(0)265; GCN-NEXT: v_accvgpr_write_b32 a0, s0266; GCN-NEXT: v_accvgpr_write_b32 a1, s1267; GCN-NEXT: v_accvgpr_write_b32 a2, s2268; GCN-NEXT: v_accvgpr_write_b32 a3, s3269; GCN-NEXT: v_accvgpr_write_b32 a4, s4270; GCN-NEXT: v_accvgpr_write_b32 a5, s5271; GCN-NEXT: v_accvgpr_write_b32 a6, s6272; GCN-NEXT: v_accvgpr_write_b32 a7, s7273; GCN-NEXT: s_nop 1274; GCN-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:3275; GCN-NEXT: v_mov_b32_e32 v0, 0276; GCN-NEXT: s_nop 15277; GCN-NEXT: s_nop 0278; GCN-NEXT: global_store_dwordx4 v0, a[0:3], s[8:9]279; GCN-NEXT: global_store_dwordx4 v0, a[4:7], s[8:9] offset:16280; GCN-NEXT: s_endpgm281bb:282 %in.1 = load <4 x double>, ptr addrspace(1) %arg283 %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)284 store <4 x double> %mai.1, ptr addrspace(1) %arg285 ret void286}287 288define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_imm(ptr addrspace(1) %arg, double %a, double %b) #0 {289; GCN-LABEL: test_mfma_f64_16x16x4f64_splat_imm:290; GCN: ; %bb.0: ; %bb291; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24292; GCN-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34293; GCN-NEXT: s_waitcnt lgkmcnt(0)294; GCN-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]295; GCN-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]296; GCN-NEXT: s_nop 1297; GCN-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], 0298; GCN-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:3299; GCN-NEXT: v_mov_b32_e32 v0, 0300; GCN-NEXT: s_nop 15301; GCN-NEXT: s_nop 0302; GCN-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]303; GCN-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:16304; GCN-NEXT: s_endpgm305bb:306 %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 0.0>, i32 0, i32 0, i32 0)307 %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)308 store <4 x double> %mai.2, ptr addrspace(1) %arg309 ret void310}311 312define amdgpu_kernel void @test_mfma_f64_16x16x4f64_imm(ptr addrspace(1) %arg, double %a, double %b) #0 {313; GCN-LABEL: test_mfma_f64_16x16x4f64_imm:314; GCN: ; %bb.0: ; %bb315; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24316; GCN-NEXT: s_load_dwordx2 s[10:11], s[4:5], 0x34317; GCN-NEXT: s_mov_b64 s[6:7], 1.0318; GCN-NEXT: s_mov_b64 s[8:9], 0319; GCN-NEXT: v_accvgpr_write_b32 a0, s8320; GCN-NEXT: s_waitcnt lgkmcnt(0)321; GCN-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]322; GCN-NEXT: v_accvgpr_write_b32 a2, s8323; GCN-NEXT: v_accvgpr_write_b32 a4, s8324; GCN-NEXT: v_accvgpr_write_b32 a6, s6325; GCN-NEXT: v_accvgpr_write_b32 a1, s9326; GCN-NEXT: v_accvgpr_write_b32 a3, s9327; GCN-NEXT: v_accvgpr_write_b32 a5, s9328; GCN-NEXT: v_accvgpr_write_b32 a7, s7329; GCN-NEXT: v_pk_mov_b32 v[2:3], s[10:11], s[10:11] op_sel:[0,1]330; GCN-NEXT: s_nop 1331; GCN-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]332; GCN-NEXT: v_mov_b32_e32 v0, 0333; GCN-NEXT: s_nop 15334; GCN-NEXT: s_nop 0335; GCN-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]336; GCN-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:16337; GCN-NEXT: s_endpgm338bb:339 %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)340 store <4 x double> %mai.1, ptr addrspace(1) %arg341 ret void342}343 344define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_lit(ptr addrspace(1) %arg, double %a, double %b) #0 {345; GCN-LABEL: test_mfma_f64_16x16x4f64_splat_lit:346; GCN: ; %bb.0: ; %bb347; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24348; GCN-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x34349; GCN-NEXT: s_mov_b32 s6, 0350; GCN-NEXT: s_mov_b32 s7, 0x405ec000351; GCN-NEXT: v_accvgpr_write_b32 a0, s6352; GCN-NEXT: s_waitcnt lgkmcnt(0)353; GCN-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]354; GCN-NEXT: v_accvgpr_write_b32 a2, s6355; GCN-NEXT: v_accvgpr_write_b32 a4, s6356; GCN-NEXT: v_accvgpr_write_b32 a6, s6357; GCN-NEXT: v_accvgpr_write_b32 a1, s7358; GCN-NEXT: v_accvgpr_write_b32 a3, s7359; GCN-NEXT: v_accvgpr_write_b32 a5, s7360; GCN-NEXT: v_accvgpr_write_b32 a7, s7361; GCN-NEXT: v_pk_mov_b32 v[2:3], s[8:9], s[8:9] op_sel:[0,1]362; GCN-NEXT: s_nop 1363; GCN-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]364; GCN-NEXT: v_mov_b32_e32 v0, 0365; GCN-NEXT: s_nop 15366; GCN-NEXT: s_nop 0367; GCN-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]368; GCN-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:16369; GCN-NEXT: s_endpgm370bb:371 %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)372 store <4 x double> %mai.1, ptr addrspace(1) %arg373 ret void374}375 376attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }377