549 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck --check-prefixes=GCN,GFX908 %s3; RUN: llc -mtriple=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug < %s | FileCheck --check-prefixes=GCN,GFX908 %s4; RUN: llc -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck --check-prefixes=GCN,GFX90A %s5 6declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16>, <2 x i16>, <32 x float>, i32, i32, i32)7declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16>, <2 x i16>, <16 x float>, i32, i32, i32)8declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16>, <2 x i16>, <4 x float>, i32, i32, i32)9declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16>, <2 x i16>, <16 x float>, i32, i32, i32)10declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16>, <2 x i16>, <4 x float>, i32, i32, i32)11declare i32 @llvm.amdgcn.workitem.id.x()12 13define amdgpu_kernel void @test_mfma_f32_32x32x2bf16(ptr addrspace(1) %arg) #0 {14; GFX908-LABEL: test_mfma_f32_32x32x2bf16:15; GFX908: ; %bb.0: ; %bb16; GFX908-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x2417; GFX908-NEXT: v_mov_b32_e32 v4, 018; GFX908-NEXT: s_waitcnt lgkmcnt(0)19; GFX908-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x020; GFX908-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x4021; GFX908-NEXT: s_waitcnt lgkmcnt(0)22; GFX908-NEXT: v_mov_b32_e32 v0, s1623; GFX908-NEXT: v_mov_b32_e32 v1, s1724; GFX908-NEXT: v_mov_b32_e32 v2, s1825; GFX908-NEXT: v_accvgpr_write_b32 a0, v026; GFX908-NEXT: v_accvgpr_write_b32 a1, v127; GFX908-NEXT: v_accvgpr_write_b32 a2, v228; GFX908-NEXT: v_mov_b32_e32 v0, s2129; GFX908-NEXT: v_mov_b32_e32 v1, s2230; GFX908-NEXT: v_mov_b32_e32 v2, s2331; GFX908-NEXT: v_accvgpr_write_b32 a5, v032; GFX908-NEXT: v_accvgpr_write_b32 a6, v133; GFX908-NEXT: v_accvgpr_write_b32 a7, v234; GFX908-NEXT: v_mov_b32_e32 v0, s2435; GFX908-NEXT: v_mov_b32_e32 v1, s2536; GFX908-NEXT: v_mov_b32_e32 v2, s2637; GFX908-NEXT: v_accvgpr_write_b32 a8, v038; GFX908-NEXT: v_accvgpr_write_b32 a9, v139; GFX908-NEXT: v_accvgpr_write_b32 a10, v240; GFX908-NEXT: v_mov_b32_e32 v0, s2741; GFX908-NEXT: v_mov_b32_e32 v1, s2842; GFX908-NEXT: v_mov_b32_e32 v2, s2943; GFX908-NEXT: v_accvgpr_write_b32 a11, v044; GFX908-NEXT: v_accvgpr_write_b32 a12, v145; GFX908-NEXT: v_accvgpr_write_b32 a13, v246; GFX908-NEXT: v_mov_b32_e32 v0, s3047; GFX908-NEXT: v_mov_b32_e32 v1, s3148; GFX908-NEXT: v_mov_b32_e32 v2, s049; GFX908-NEXT: v_accvgpr_write_b32 a14, v050; GFX908-NEXT: v_accvgpr_write_b32 a15, v151; GFX908-NEXT: v_accvgpr_write_b32 a16, v252; GFX908-NEXT: v_mov_b32_e32 v0, s153; GFX908-NEXT: v_mov_b32_e32 v1, s254; GFX908-NEXT: v_mov_b32_e32 v2, s355; GFX908-NEXT: v_accvgpr_write_b32 a17, v056; GFX908-NEXT: v_accvgpr_write_b32 a18, v157; GFX908-NEXT: v_accvgpr_write_b32 a19, v258; GFX908-NEXT: v_mov_b32_e32 v0, s459; GFX908-NEXT: v_mov_b32_e32 v1, s560; GFX908-NEXT: v_mov_b32_e32 v2, s661; GFX908-NEXT: v_accvgpr_write_b32 a20, v062; GFX908-NEXT: v_accvgpr_write_b32 a21, v163; GFX908-NEXT: v_accvgpr_write_b32 a22, v264; GFX908-NEXT: v_mov_b32_e32 v0, s765; GFX908-NEXT: v_mov_b32_e32 v1, s866; GFX908-NEXT: v_mov_b32_e32 v2, s967; GFX908-NEXT: v_mov_b32_e32 v3, s1968; GFX908-NEXT: v_accvgpr_write_b32 a23, v069; GFX908-NEXT: v_accvgpr_write_b32 a24, v170; GFX908-NEXT: v_accvgpr_write_b32 a25, v271; GFX908-NEXT: v_mov_b32_e32 v0, s1072; GFX908-NEXT: v_mov_b32_e32 v1, s1173; GFX908-NEXT: v_mov_b32_e32 v2, s1274; GFX908-NEXT: v_mov_b32_e32 v5, s2075; GFX908-NEXT: v_accvgpr_write_b32 a3, v376; GFX908-NEXT: v_accvgpr_write_b32 a26, v077; GFX908-NEXT: v_accvgpr_write_b32 a27, v178; GFX908-NEXT: v_accvgpr_write_b32 a28, v279; GFX908-NEXT: v_mov_b32_e32 v0, s1380; GFX908-NEXT: v_mov_b32_e32 v1, s1481; GFX908-NEXT: v_mov_b32_e32 v2, s1582; GFX908-NEXT: v_mov_b32_e32 v3, 183; GFX908-NEXT: v_accvgpr_write_b32 a4, v584; GFX908-NEXT: v_accvgpr_write_b32 a29, v085; GFX908-NEXT: v_accvgpr_write_b32 a30, v186; GFX908-NEXT: v_accvgpr_write_b32 a31, v287; GFX908-NEXT: v_mov_b32_e32 v0, 288; GFX908-NEXT: s_nop 189; GFX908-NEXT: v_mfma_f32_32x32x2bf16 a[0:31], v3, v0, a[0:31] cbsz:1 abid:2 blgp:390; GFX908-NEXT: s_nop 1591; GFX908-NEXT: s_nop 192; GFX908-NEXT: v_accvgpr_read_b32 v3, a2793; GFX908-NEXT: v_accvgpr_read_b32 v2, a2694; GFX908-NEXT: v_accvgpr_read_b32 v1, a2595; GFX908-NEXT: v_accvgpr_read_b32 v0, a2496; GFX908-NEXT: s_nop 197; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:9698; GFX908-NEXT: s_nop 099; GFX908-NEXT: v_accvgpr_read_b32 v3, a31100; GFX908-NEXT: v_accvgpr_read_b32 v2, a30101; GFX908-NEXT: v_accvgpr_read_b32 v1, a29102; GFX908-NEXT: v_accvgpr_read_b32 v0, a28103; GFX908-NEXT: s_nop 1104; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:112105; GFX908-NEXT: s_nop 0106; GFX908-NEXT: v_accvgpr_read_b32 v3, a19107; GFX908-NEXT: v_accvgpr_read_b32 v2, a18108; GFX908-NEXT: v_accvgpr_read_b32 v1, a17109; GFX908-NEXT: v_accvgpr_read_b32 v0, a16110; GFX908-NEXT: s_nop 1111; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:64112; GFX908-NEXT: s_nop 0113; GFX908-NEXT: v_accvgpr_read_b32 v3, a23114; GFX908-NEXT: v_accvgpr_read_b32 v2, a22115; GFX908-NEXT: v_accvgpr_read_b32 v1, a21116; GFX908-NEXT: v_accvgpr_read_b32 v0, a20117; GFX908-NEXT: s_nop 1118; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:80119; GFX908-NEXT: s_nop 0120; GFX908-NEXT: v_accvgpr_read_b32 v3, a11121; GFX908-NEXT: v_accvgpr_read_b32 v2, a10122; GFX908-NEXT: v_accvgpr_read_b32 v1, a9123; GFX908-NEXT: v_accvgpr_read_b32 v0, a8124; GFX908-NEXT: s_nop 1125; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:32126; GFX908-NEXT: s_nop 0127; GFX908-NEXT: v_accvgpr_read_b32 v3, a15128; GFX908-NEXT: v_accvgpr_read_b32 v2, a14129; GFX908-NEXT: v_accvgpr_read_b32 v1, a13130; GFX908-NEXT: v_accvgpr_read_b32 v0, a12131; GFX908-NEXT: s_nop 1132; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:48133; GFX908-NEXT: s_nop 0134; GFX908-NEXT: v_accvgpr_read_b32 v3, a3135; GFX908-NEXT: v_accvgpr_read_b32 v2, a2136; GFX908-NEXT: v_accvgpr_read_b32 v1, a1137; GFX908-NEXT: v_accvgpr_read_b32 v0, a0138; GFX908-NEXT: s_nop 1139; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35]140; GFX908-NEXT: s_nop 0141; GFX908-NEXT: v_accvgpr_read_b32 v3, a7142; GFX908-NEXT: v_accvgpr_read_b32 v2, a6143; GFX908-NEXT: v_accvgpr_read_b32 v1, a5144; GFX908-NEXT: v_accvgpr_read_b32 v0, a4145; GFX908-NEXT: s_nop 1146; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:16147; GFX908-NEXT: s_endpgm148;149; GFX90A-LABEL: test_mfma_f32_32x32x2bf16:150; GFX90A: ; %bb.0: ; %bb151; GFX90A-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24152; GFX90A-NEXT: v_mov_b32_e32 v1, 1153; GFX90A-NEXT: v_mov_b32_e32 v2, 2154; GFX90A-NEXT: v_mov_b32_e32 v0, 0155; GFX90A-NEXT: s_waitcnt lgkmcnt(0)156; GFX90A-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0157; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40158; GFX90A-NEXT: s_waitcnt lgkmcnt(0)159; GFX90A-NEXT: v_accvgpr_write_b32 a0, s16160; GFX90A-NEXT: v_accvgpr_write_b32 a1, s17161; GFX90A-NEXT: v_accvgpr_write_b32 a2, s18162; GFX90A-NEXT: v_accvgpr_write_b32 a3, s19163; GFX90A-NEXT: v_accvgpr_write_b32 a4, s20164; GFX90A-NEXT: v_accvgpr_write_b32 a5, s21165; GFX90A-NEXT: v_accvgpr_write_b32 a6, s22166; GFX90A-NEXT: v_accvgpr_write_b32 a7, s23167; GFX90A-NEXT: v_accvgpr_write_b32 a8, s24168; GFX90A-NEXT: v_accvgpr_write_b32 a9, s25169; GFX90A-NEXT: v_accvgpr_write_b32 a10, s26170; GFX90A-NEXT: v_accvgpr_write_b32 a11, s27171; GFX90A-NEXT: v_accvgpr_write_b32 a12, s28172; GFX90A-NEXT: v_accvgpr_write_b32 a13, s29173; GFX90A-NEXT: v_accvgpr_write_b32 a14, s30174; GFX90A-NEXT: v_accvgpr_write_b32 a15, s31175; GFX90A-NEXT: v_accvgpr_write_b32 a16, s0176; GFX90A-NEXT: v_accvgpr_write_b32 a17, s1177; GFX90A-NEXT: v_accvgpr_write_b32 a18, s2178; GFX90A-NEXT: v_accvgpr_write_b32 a19, s3179; GFX90A-NEXT: v_accvgpr_write_b32 a20, s4180; GFX90A-NEXT: v_accvgpr_write_b32 a21, s5181; GFX90A-NEXT: v_accvgpr_write_b32 a22, s6182; GFX90A-NEXT: v_accvgpr_write_b32 a23, s7183; GFX90A-NEXT: v_accvgpr_write_b32 a24, s8184; GFX90A-NEXT: v_accvgpr_write_b32 a25, s9185; GFX90A-NEXT: v_accvgpr_write_b32 a26, s10186; GFX90A-NEXT: v_accvgpr_write_b32 a27, s11187; GFX90A-NEXT: v_accvgpr_write_b32 a28, s12188; GFX90A-NEXT: v_accvgpr_write_b32 a29, s13189; GFX90A-NEXT: v_accvgpr_write_b32 a30, s14190; GFX90A-NEXT: v_accvgpr_write_b32 a31, s15191; GFX90A-NEXT: s_nop 1192; GFX90A-NEXT: v_mfma_f32_32x32x2bf16 a[0:31], v1, v2, a[0:31] cbsz:1 abid:2 blgp:3193; GFX90A-NEXT: s_nop 15194; GFX90A-NEXT: s_nop 2195; GFX90A-NEXT: global_store_dwordx4 v0, a[24:27], s[34:35] offset:96196; GFX90A-NEXT: global_store_dwordx4 v0, a[28:31], s[34:35] offset:112197; GFX90A-NEXT: global_store_dwordx4 v0, a[16:19], s[34:35] offset:64198; GFX90A-NEXT: global_store_dwordx4 v0, a[20:23], s[34:35] offset:80199; GFX90A-NEXT: global_store_dwordx4 v0, a[8:11], s[34:35] offset:32200; GFX90A-NEXT: global_store_dwordx4 v0, a[12:15], s[34:35] offset:48201; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[34:35]202; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[34:35] offset:16203; GFX90A-NEXT: s_endpgm204bb:205 %in.1 = load <32 x float>, ptr addrspace(1) %arg206 %a = bitcast i32 1 to <2 x i16>207 %b = bitcast i32 2 to <2 x i16>208 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16> %a, <2 x i16> %b, <32 x float> %in.1, i32 1, i32 2, i32 3)209 store <32 x float> %mai.1, ptr addrspace(1) %arg210 ret void211}212 213define amdgpu_kernel void @test_mfma_f32_16x16x2bf16(ptr addrspace(1) %arg) #0 {214; GFX908-LABEL: test_mfma_f32_16x16x2bf16:215; GFX908: ; %bb.0: ; %bb216; GFX908-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24217; GFX908-NEXT: v_mov_b32_e32 v0, 1218; GFX908-NEXT: v_mov_b32_e32 v12, 0219; GFX908-NEXT: s_waitcnt lgkmcnt(0)220; GFX908-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0221; GFX908-NEXT: s_waitcnt lgkmcnt(0)222; GFX908-NEXT: v_mov_b32_e32 v13, s0223; GFX908-NEXT: v_mov_b32_e32 v1, s1224; GFX908-NEXT: v_mov_b32_e32 v2, s2225; GFX908-NEXT: v_accvgpr_write_b32 a0, v13226; GFX908-NEXT: v_mov_b32_e32 v13, s3227; GFX908-NEXT: v_accvgpr_write_b32 a1, v1228; GFX908-NEXT: v_accvgpr_write_b32 a2, v2229; GFX908-NEXT: v_accvgpr_write_b32 a3, v13230; GFX908-NEXT: v_mov_b32_e32 v1, s4231; GFX908-NEXT: v_mov_b32_e32 v2, s5232; GFX908-NEXT: v_mov_b32_e32 v13, s6233; GFX908-NEXT: v_accvgpr_write_b32 a4, v1234; GFX908-NEXT: v_accvgpr_write_b32 a5, v2235; GFX908-NEXT: v_accvgpr_write_b32 a6, v13236; GFX908-NEXT: v_mov_b32_e32 v1, s7237; GFX908-NEXT: v_mov_b32_e32 v2, s8238; GFX908-NEXT: v_mov_b32_e32 v13, s9239; GFX908-NEXT: v_accvgpr_write_b32 a7, v1240; GFX908-NEXT: v_accvgpr_write_b32 a8, v2241; GFX908-NEXT: v_accvgpr_write_b32 a9, v13242; GFX908-NEXT: v_mov_b32_e32 v1, s10243; GFX908-NEXT: v_mov_b32_e32 v2, s11244; GFX908-NEXT: v_mov_b32_e32 v13, s12245; GFX908-NEXT: v_accvgpr_write_b32 a10, v1246; GFX908-NEXT: v_accvgpr_write_b32 a11, v2247; GFX908-NEXT: v_accvgpr_write_b32 a12, v13248; GFX908-NEXT: v_mov_b32_e32 v1, s13249; GFX908-NEXT: v_mov_b32_e32 v2, s14250; GFX908-NEXT: v_mov_b32_e32 v13, s15251; GFX908-NEXT: v_accvgpr_write_b32 a13, v1252; GFX908-NEXT: v_accvgpr_write_b32 a14, v2253; GFX908-NEXT: v_accvgpr_write_b32 a15, v13254; GFX908-NEXT: v_mov_b32_e32 v1, 2255; GFX908-NEXT: s_nop 1256; GFX908-NEXT: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3257; GFX908-NEXT: s_nop 9258; GFX908-NEXT: v_accvgpr_read_b32 v3, a15259; GFX908-NEXT: v_accvgpr_read_b32 v2, a14260; GFX908-NEXT: v_accvgpr_read_b32 v1, a13261; GFX908-NEXT: v_accvgpr_read_b32 v0, a12262; GFX908-NEXT: v_accvgpr_read_b32 v7, a11263; GFX908-NEXT: v_accvgpr_read_b32 v6, a10264; GFX908-NEXT: v_accvgpr_read_b32 v5, a9265; GFX908-NEXT: v_accvgpr_read_b32 v4, a8266; GFX908-NEXT: v_accvgpr_read_b32 v11, a7267; GFX908-NEXT: v_accvgpr_read_b32 v10, a6268; GFX908-NEXT: v_accvgpr_read_b32 v9, a5269; GFX908-NEXT: v_accvgpr_read_b32 v8, a4270; GFX908-NEXT: global_store_dwordx4 v12, v[0:3], s[16:17] offset:48271; GFX908-NEXT: global_store_dwordx4 v12, v[4:7], s[16:17] offset:32272; GFX908-NEXT: global_store_dwordx4 v12, v[8:11], s[16:17] offset:16273; GFX908-NEXT: v_accvgpr_read_b32 v3, a3274; GFX908-NEXT: v_accvgpr_read_b32 v2, a2275; GFX908-NEXT: v_accvgpr_read_b32 v1, a1276; GFX908-NEXT: v_accvgpr_read_b32 v0, a0277; GFX908-NEXT: s_nop 1278; GFX908-NEXT: global_store_dwordx4 v12, v[0:3], s[16:17]279; GFX908-NEXT: s_endpgm280;281; GFX90A-LABEL: test_mfma_f32_16x16x2bf16:282; GFX90A: ; %bb.0: ; %bb283; GFX90A-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24284; GFX90A-NEXT: v_mov_b32_e32 v0, 1285; GFX90A-NEXT: v_mov_b32_e32 v1, 2286; GFX90A-NEXT: s_waitcnt lgkmcnt(0)287; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0288; GFX90A-NEXT: s_waitcnt lgkmcnt(0)289; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0290; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1291; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2292; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3293; GFX90A-NEXT: v_accvgpr_write_b32 a4, s4294; GFX90A-NEXT: v_accvgpr_write_b32 a5, s5295; GFX90A-NEXT: v_accvgpr_write_b32 a6, s6296; GFX90A-NEXT: v_accvgpr_write_b32 a7, s7297; GFX90A-NEXT: v_accvgpr_write_b32 a8, s8298; GFX90A-NEXT: v_accvgpr_write_b32 a9, s9299; GFX90A-NEXT: v_accvgpr_write_b32 a10, s10300; GFX90A-NEXT: v_accvgpr_write_b32 a11, s11301; GFX90A-NEXT: v_accvgpr_write_b32 a12, s12302; GFX90A-NEXT: v_accvgpr_write_b32 a13, s13303; GFX90A-NEXT: v_accvgpr_write_b32 a14, s14304; GFX90A-NEXT: v_accvgpr_write_b32 a15, s15305; GFX90A-NEXT: s_nop 1306; GFX90A-NEXT: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3307; GFX90A-NEXT: v_mov_b32_e32 v0, 0308; GFX90A-NEXT: s_nop 9309; GFX90A-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48310; GFX90A-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32311; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16312; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]313; GFX90A-NEXT: s_endpgm314bb:315 %in.1 = load <16 x float>, ptr addrspace(1) %arg316 %a = bitcast i32 1 to <2 x i16>317 %b = bitcast i32 2 to <2 x i16>318 %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %in.1, i32 1, i32 2, i32 3)319 store <16 x float> %mai.1, ptr addrspace(1) %arg320 ret void321}322 323define amdgpu_kernel void @test_mfma_f32_4x4x2bf16(ptr addrspace(1) %arg) #0 {324; GFX908-LABEL: test_mfma_f32_4x4x2bf16:325; GFX908: ; %bb.0: ; %bb326; GFX908-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24327; GFX908-NEXT: v_mov_b32_e32 v0, 1328; GFX908-NEXT: v_mov_b32_e32 v1, 2329; GFX908-NEXT: v_mov_b32_e32 v4, 0330; GFX908-NEXT: s_waitcnt lgkmcnt(0)331; GFX908-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0332; GFX908-NEXT: s_waitcnt lgkmcnt(0)333; GFX908-NEXT: v_mov_b32_e32 v5, s0334; GFX908-NEXT: v_mov_b32_e32 v2, s1335; GFX908-NEXT: v_mov_b32_e32 v3, s2336; GFX908-NEXT: v_accvgpr_write_b32 a0, v5337; GFX908-NEXT: v_mov_b32_e32 v5, s3338; GFX908-NEXT: v_accvgpr_write_b32 a1, v2339; GFX908-NEXT: v_accvgpr_write_b32 a2, v3340; GFX908-NEXT: v_accvgpr_write_b32 a3, v5341; GFX908-NEXT: s_nop 0342; GFX908-NEXT: v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, a[0:3] cbsz:1 abid:2 blgp:3343; GFX908-NEXT: s_nop 3344; GFX908-NEXT: v_accvgpr_read_b32 v0, a0345; GFX908-NEXT: v_accvgpr_read_b32 v1, a1346; GFX908-NEXT: v_accvgpr_read_b32 v2, a2347; GFX908-NEXT: v_accvgpr_read_b32 v3, a3348; GFX908-NEXT: s_nop 1349; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]350; GFX908-NEXT: s_endpgm351;352; GFX90A-LABEL: test_mfma_f32_4x4x2bf16:353; GFX90A: ; %bb.0: ; %bb354; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24355; GFX90A-NEXT: v_mov_b32_e32 v0, 1356; GFX90A-NEXT: v_mov_b32_e32 v2, 2357; GFX90A-NEXT: v_mov_b32_e32 v1, 0358; GFX90A-NEXT: s_waitcnt lgkmcnt(0)359; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0360; GFX90A-NEXT: s_waitcnt lgkmcnt(0)361; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0362; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1363; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2364; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3365; GFX90A-NEXT: s_nop 1366; GFX90A-NEXT: v_mfma_f32_4x4x2bf16 a[0:3], v0, v2, a[0:3] cbsz:1 abid:2 blgp:3367; GFX90A-NEXT: s_nop 4368; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[6:7]369; GFX90A-NEXT: s_endpgm370bb:371 %in.1 = load <4 x float>, ptr addrspace(1) %arg372 %a = bitcast i32 1 to <2 x i16>373 %b = bitcast i32 2 to <2 x i16>374 %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %in.1, i32 1, i32 2, i32 3)375 store <4 x float> %mai.1, ptr addrspace(1) %arg376 ret void377}378 379define amdgpu_kernel void @test_mfma_f32_32x32x4bf16(ptr addrspace(1) %arg) #0 {380; GFX908-LABEL: test_mfma_f32_32x32x4bf16:381; GFX908: ; %bb.0: ; %bb382; GFX908-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24383; GFX908-NEXT: v_mov_b32_e32 v0, 1384; GFX908-NEXT: v_mov_b32_e32 v16, 0385; GFX908-NEXT: s_waitcnt lgkmcnt(0)386; GFX908-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0387; GFX908-NEXT: s_waitcnt lgkmcnt(0)388; GFX908-NEXT: v_mov_b32_e32 v17, s0389; GFX908-NEXT: v_mov_b32_e32 v1, s1390; GFX908-NEXT: v_mov_b32_e32 v2, s2391; GFX908-NEXT: v_accvgpr_write_b32 a0, v17392; GFX908-NEXT: v_mov_b32_e32 v17, s3393; GFX908-NEXT: v_accvgpr_write_b32 a1, v1394; GFX908-NEXT: v_accvgpr_write_b32 a2, v2395; GFX908-NEXT: v_accvgpr_write_b32 a3, v17396; GFX908-NEXT: v_mov_b32_e32 v1, s4397; GFX908-NEXT: v_mov_b32_e32 v2, s5398; GFX908-NEXT: v_mov_b32_e32 v17, s6399; GFX908-NEXT: v_accvgpr_write_b32 a4, v1400; GFX908-NEXT: v_accvgpr_write_b32 a5, v2401; GFX908-NEXT: v_accvgpr_write_b32 a6, v17402; GFX908-NEXT: v_mov_b32_e32 v1, s7403; GFX908-NEXT: v_mov_b32_e32 v2, s8404; GFX908-NEXT: v_mov_b32_e32 v17, s9405; GFX908-NEXT: v_accvgpr_write_b32 a7, v1406; GFX908-NEXT: v_accvgpr_write_b32 a8, v2407; GFX908-NEXT: v_accvgpr_write_b32 a9, v17408; GFX908-NEXT: v_mov_b32_e32 v1, s10409; GFX908-NEXT: v_mov_b32_e32 v2, s11410; GFX908-NEXT: v_mov_b32_e32 v17, s12411; GFX908-NEXT: v_accvgpr_write_b32 a10, v1412; GFX908-NEXT: v_accvgpr_write_b32 a11, v2413; GFX908-NEXT: v_accvgpr_write_b32 a12, v17414; GFX908-NEXT: v_mov_b32_e32 v1, s13415; GFX908-NEXT: v_mov_b32_e32 v2, s14416; GFX908-NEXT: v_mov_b32_e32 v17, s15417; GFX908-NEXT: v_accvgpr_write_b32 a13, v1418; GFX908-NEXT: v_accvgpr_write_b32 a14, v2419; GFX908-NEXT: v_accvgpr_write_b32 a15, v17420; GFX908-NEXT: v_mov_b32_e32 v1, 2421; GFX908-NEXT: s_nop 1422; GFX908-NEXT: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3423; GFX908-NEXT: s_nop 15424; GFX908-NEXT: s_nop 1425; GFX908-NEXT: v_accvgpr_read_b32 v3, a15426; GFX908-NEXT: v_accvgpr_read_b32 v2, a14427; GFX908-NEXT: v_accvgpr_read_b32 v1, a13428; GFX908-NEXT: v_accvgpr_read_b32 v0, a12429; GFX908-NEXT: v_accvgpr_read_b32 v7, a11430; GFX908-NEXT: v_accvgpr_read_b32 v6, a10431; GFX908-NEXT: v_accvgpr_read_b32 v5, a9432; GFX908-NEXT: v_accvgpr_read_b32 v4, a8433; GFX908-NEXT: v_accvgpr_read_b32 v11, a7434; GFX908-NEXT: v_accvgpr_read_b32 v10, a6435; GFX908-NEXT: v_accvgpr_read_b32 v9, a5436; GFX908-NEXT: v_accvgpr_read_b32 v8, a4437; GFX908-NEXT: v_accvgpr_read_b32 v15, a3438; GFX908-NEXT: v_accvgpr_read_b32 v14, a2439; GFX908-NEXT: v_accvgpr_read_b32 v13, a1440; GFX908-NEXT: v_accvgpr_read_b32 v12, a0441; GFX908-NEXT: global_store_dwordx4 v16, v[0:3], s[16:17] offset:48442; GFX908-NEXT: global_store_dwordx4 v16, v[4:7], s[16:17] offset:32443; GFX908-NEXT: global_store_dwordx4 v16, v[8:11], s[16:17] offset:16444; GFX908-NEXT: global_store_dwordx4 v16, v[12:15], s[16:17]445; GFX908-NEXT: s_endpgm446;447; GFX90A-LABEL: test_mfma_f32_32x32x4bf16:448; GFX90A: ; %bb.0: ; %bb449; GFX90A-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24450; GFX90A-NEXT: v_mov_b32_e32 v0, 1451; GFX90A-NEXT: v_mov_b32_e32 v1, 2452; GFX90A-NEXT: s_waitcnt lgkmcnt(0)453; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0454; GFX90A-NEXT: s_waitcnt lgkmcnt(0)455; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0456; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1457; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2458; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3459; GFX90A-NEXT: v_accvgpr_write_b32 a4, s4460; GFX90A-NEXT: v_accvgpr_write_b32 a5, s5461; GFX90A-NEXT: v_accvgpr_write_b32 a6, s6462; GFX90A-NEXT: v_accvgpr_write_b32 a7, s7463; GFX90A-NEXT: v_accvgpr_write_b32 a8, s8464; GFX90A-NEXT: v_accvgpr_write_b32 a9, s9465; GFX90A-NEXT: v_accvgpr_write_b32 a10, s10466; GFX90A-NEXT: v_accvgpr_write_b32 a11, s11467; GFX90A-NEXT: v_accvgpr_write_b32 a12, s12468; GFX90A-NEXT: v_accvgpr_write_b32 a13, s13469; GFX90A-NEXT: v_accvgpr_write_b32 a14, s14470; GFX90A-NEXT: v_accvgpr_write_b32 a15, s15471; GFX90A-NEXT: s_nop 1472; GFX90A-NEXT: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3473; GFX90A-NEXT: v_mov_b32_e32 v0, 0474; GFX90A-NEXT: s_nop 15475; GFX90A-NEXT: s_nop 1476; GFX90A-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48477; GFX90A-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32478; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16479; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]480; GFX90A-NEXT: s_endpgm481bb:482 %in.1 = load <16 x float>, ptr addrspace(1) %arg483 %a = bitcast i32 1 to <2 x i16>484 %b = bitcast i32 2 to <2 x i16>485 %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %in.1, i32 1, i32 2, i32 3)486 store <16 x float> %mai.1, ptr addrspace(1) %arg487 ret void488}489 490define amdgpu_kernel void @test_mfma_f32_16x16x8bf16(ptr addrspace(1) %arg) #0 {491; GFX908-LABEL: test_mfma_f32_16x16x8bf16:492; GFX908: ; %bb.0: ; %bb493; GFX908-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24494; GFX908-NEXT: v_mov_b32_e32 v0, 1495; GFX908-NEXT: v_mov_b32_e32 v1, 2496; GFX908-NEXT: v_mov_b32_e32 v4, 0497; GFX908-NEXT: s_waitcnt lgkmcnt(0)498; GFX908-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0499; GFX908-NEXT: s_waitcnt lgkmcnt(0)500; GFX908-NEXT: v_mov_b32_e32 v5, s0501; GFX908-NEXT: v_mov_b32_e32 v2, s1502; GFX908-NEXT: v_mov_b32_e32 v3, s2503; GFX908-NEXT: v_accvgpr_write_b32 a0, v5504; GFX908-NEXT: v_mov_b32_e32 v5, s3505; GFX908-NEXT: v_accvgpr_write_b32 a1, v2506; GFX908-NEXT: v_accvgpr_write_b32 a2, v3507; GFX908-NEXT: v_accvgpr_write_b32 a3, v5508; GFX908-NEXT: s_nop 0509; GFX908-NEXT: v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, a[0:3] cbsz:1 abid:2 blgp:3510; GFX908-NEXT: s_nop 9511; GFX908-NEXT: v_accvgpr_read_b32 v0, a0512; GFX908-NEXT: v_accvgpr_read_b32 v1, a1513; GFX908-NEXT: v_accvgpr_read_b32 v2, a2514; GFX908-NEXT: v_accvgpr_read_b32 v3, a3515; GFX908-NEXT: s_nop 1516; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]517; GFX908-NEXT: s_endpgm518;519; GFX90A-LABEL: test_mfma_f32_16x16x8bf16:520; GFX90A: ; %bb.0: ; %bb521; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24522; GFX90A-NEXT: v_mov_b32_e32 v0, 1523; GFX90A-NEXT: v_mov_b32_e32 v2, 2524; GFX90A-NEXT: v_mov_b32_e32 v1, 0525; GFX90A-NEXT: s_waitcnt lgkmcnt(0)526; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0527; GFX90A-NEXT: s_waitcnt lgkmcnt(0)528; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0529; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1530; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2531; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3532; GFX90A-NEXT: s_nop 1533; GFX90A-NEXT: v_mfma_f32_16x16x8bf16 a[0:3], v0, v2, a[0:3] cbsz:1 abid:2 blgp:3534; GFX90A-NEXT: s_nop 10535; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[6:7]536; GFX90A-NEXT: s_endpgm537bb:538 %in.1 = load <4 x float>, ptr addrspace(1) %arg539 %a = bitcast i32 1 to <2 x i16>540 %b = bitcast i32 2 to <2 x i16>541 %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %in.1, i32 1, i32 2, i32 3)542 store <4 x float> %mai.1, ptr addrspace(1) %arg543 ret void544}545 546attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }547;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:548; GCN: {{.*}}549