176 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 <16 x i32> @llvm.amdgcn.mfma.i32.32x32x8i8(i32, i32, <16 x i32>, i32, i32, i32)7declare <4 x i32> @llvm.amdgcn.mfma.i32.16x16x16i8(i32, i32, <4 x i32>, i32, i32, i32)8 9define amdgpu_kernel void @test_mfma_i32_32x32x8i8(ptr addrspace(1) %arg) #0 {10; GFX908-LABEL: test_mfma_i32_32x32x8i8:11; GFX908: ; %bb.0: ; %bb12; GFX908-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x2413; GFX908-NEXT: v_mov_b32_e32 v0, 114; GFX908-NEXT: v_mov_b32_e32 v16, 015; GFX908-NEXT: s_waitcnt lgkmcnt(0)16; GFX908-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x017; GFX908-NEXT: s_waitcnt lgkmcnt(0)18; GFX908-NEXT: v_mov_b32_e32 v17, s019; GFX908-NEXT: v_mov_b32_e32 v1, s120; GFX908-NEXT: v_mov_b32_e32 v2, s221; GFX908-NEXT: v_accvgpr_write_b32 a0, v1722; GFX908-NEXT: v_mov_b32_e32 v17, s323; GFX908-NEXT: v_accvgpr_write_b32 a1, v124; GFX908-NEXT: v_accvgpr_write_b32 a2, v225; GFX908-NEXT: v_accvgpr_write_b32 a3, v1726; GFX908-NEXT: v_mov_b32_e32 v1, s427; GFX908-NEXT: v_mov_b32_e32 v2, s528; GFX908-NEXT: v_mov_b32_e32 v17, s629; GFX908-NEXT: v_accvgpr_write_b32 a4, v130; GFX908-NEXT: v_accvgpr_write_b32 a5, v231; GFX908-NEXT: v_accvgpr_write_b32 a6, v1732; GFX908-NEXT: v_mov_b32_e32 v1, s733; GFX908-NEXT: v_mov_b32_e32 v2, s834; GFX908-NEXT: v_mov_b32_e32 v17, s935; GFX908-NEXT: v_accvgpr_write_b32 a7, v136; GFX908-NEXT: v_accvgpr_write_b32 a8, v237; GFX908-NEXT: v_accvgpr_write_b32 a9, v1738; GFX908-NEXT: v_mov_b32_e32 v1, s1039; GFX908-NEXT: v_mov_b32_e32 v2, s1140; GFX908-NEXT: v_mov_b32_e32 v17, s1241; GFX908-NEXT: v_accvgpr_write_b32 a10, v142; GFX908-NEXT: v_accvgpr_write_b32 a11, v243; GFX908-NEXT: v_accvgpr_write_b32 a12, v1744; GFX908-NEXT: v_mov_b32_e32 v1, s1345; GFX908-NEXT: v_mov_b32_e32 v2, s1446; GFX908-NEXT: v_mov_b32_e32 v17, s1547; GFX908-NEXT: v_accvgpr_write_b32 a13, v148; GFX908-NEXT: v_accvgpr_write_b32 a14, v249; GFX908-NEXT: v_accvgpr_write_b32 a15, v1750; GFX908-NEXT: v_mov_b32_e32 v1, 251; GFX908-NEXT: s_nop 152; GFX908-NEXT: v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:353; GFX908-NEXT: s_nop 1554; GFX908-NEXT: s_nop 155; GFX908-NEXT: v_accvgpr_read_b32 v15, a1556; GFX908-NEXT: v_accvgpr_read_b32 v14, a1457; GFX908-NEXT: v_accvgpr_read_b32 v13, a1358; GFX908-NEXT: v_accvgpr_read_b32 v12, a1259; GFX908-NEXT: v_accvgpr_read_b32 v3, a360; GFX908-NEXT: v_accvgpr_read_b32 v2, a261; GFX908-NEXT: v_accvgpr_read_b32 v1, a162; GFX908-NEXT: v_accvgpr_read_b32 v0, a063; GFX908-NEXT: v_accvgpr_read_b32 v7, a764; GFX908-NEXT: v_accvgpr_read_b32 v6, a665; GFX908-NEXT: v_accvgpr_read_b32 v5, a566; GFX908-NEXT: v_accvgpr_read_b32 v4, a467; GFX908-NEXT: v_accvgpr_read_b32 v11, a1168; GFX908-NEXT: v_accvgpr_read_b32 v10, a1069; GFX908-NEXT: v_accvgpr_read_b32 v9, a970; GFX908-NEXT: v_accvgpr_read_b32 v8, a871; GFX908-NEXT: global_store_dwordx4 v16, v[12:15], s[16:17] offset:4872; GFX908-NEXT: s_nop 073; GFX908-NEXT: global_store_dwordx4 v16, v[8:11], s[16:17] offset:3274; GFX908-NEXT: global_store_dwordx4 v16, v[4:7], s[16:17] offset:1675; GFX908-NEXT: global_store_dwordx4 v16, v[0:3], s[16:17]76; GFX908-NEXT: s_endpgm77;78; GFX90A-LABEL: test_mfma_i32_32x32x8i8:79; GFX90A: ; %bb.0: ; %bb80; GFX90A-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x2481; GFX90A-NEXT: v_mov_b32_e32 v0, 182; GFX90A-NEXT: v_mov_b32_e32 v1, 283; GFX90A-NEXT: s_waitcnt lgkmcnt(0)84; GFX90A-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x085; GFX90A-NEXT: s_waitcnt lgkmcnt(0)86; GFX90A-NEXT: v_accvgpr_write_b32 a0, s087; GFX90A-NEXT: v_accvgpr_write_b32 a1, s188; GFX90A-NEXT: v_accvgpr_write_b32 a2, s289; GFX90A-NEXT: v_accvgpr_write_b32 a3, s390; GFX90A-NEXT: v_accvgpr_write_b32 a4, s491; GFX90A-NEXT: v_accvgpr_write_b32 a5, s592; GFX90A-NEXT: v_accvgpr_write_b32 a6, s693; GFX90A-NEXT: v_accvgpr_write_b32 a7, s794; GFX90A-NEXT: v_accvgpr_write_b32 a8, s895; GFX90A-NEXT: v_accvgpr_write_b32 a9, s996; GFX90A-NEXT: v_accvgpr_write_b32 a10, s1097; GFX90A-NEXT: v_accvgpr_write_b32 a11, s1198; GFX90A-NEXT: v_accvgpr_write_b32 a12, s1299; GFX90A-NEXT: v_accvgpr_write_b32 a13, s13100; GFX90A-NEXT: v_accvgpr_write_b32 a14, s14101; GFX90A-NEXT: v_accvgpr_write_b32 a15, s15102; GFX90A-NEXT: s_nop 1103; GFX90A-NEXT: v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3104; GFX90A-NEXT: v_mov_b32_e32 v0, 0105; GFX90A-NEXT: s_nop 15106; GFX90A-NEXT: s_nop 1107; GFX90A-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48108; GFX90A-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32109; GFX90A-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16110; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]111; GFX90A-NEXT: s_endpgm112bb:113 %in.1 = load <16 x i32>, ptr addrspace(1) %arg114 %mai.1 = tail call <16 x i32> @llvm.amdgcn.mfma.i32.32x32x8i8(i32 1, i32 2, <16 x i32> %in.1, i32 1, i32 2, i32 3)115 store <16 x i32> %mai.1, ptr addrspace(1) %arg116 ret void117}118 119define amdgpu_kernel void @test_mfma_i32_16x16x16i8(ptr addrspace(1) %arg) #0 {120; GFX908-LABEL: test_mfma_i32_16x16x16i8:121; GFX908: ; %bb.0: ; %bb122; GFX908-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24123; GFX908-NEXT: v_mov_b32_e32 v0, 1124; GFX908-NEXT: v_mov_b32_e32 v1, 2125; GFX908-NEXT: v_mov_b32_e32 v4, 0126; GFX908-NEXT: s_waitcnt lgkmcnt(0)127; GFX908-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0128; GFX908-NEXT: s_waitcnt lgkmcnt(0)129; GFX908-NEXT: v_mov_b32_e32 v5, s0130; GFX908-NEXT: v_mov_b32_e32 v2, s1131; GFX908-NEXT: v_mov_b32_e32 v3, s2132; GFX908-NEXT: v_accvgpr_write_b32 a0, v5133; GFX908-NEXT: v_mov_b32_e32 v5, s3134; GFX908-NEXT: v_accvgpr_write_b32 a1, v2135; GFX908-NEXT: v_accvgpr_write_b32 a2, v3136; GFX908-NEXT: v_accvgpr_write_b32 a3, v5137; GFX908-NEXT: s_nop 0138; GFX908-NEXT: v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[0:3] cbsz:1 abid:2 blgp:3139; GFX908-NEXT: s_nop 9140; GFX908-NEXT: v_accvgpr_read_b32 v0, a0141; GFX908-NEXT: v_accvgpr_read_b32 v1, a1142; GFX908-NEXT: v_accvgpr_read_b32 v2, a2143; GFX908-NEXT: v_accvgpr_read_b32 v3, a3144; GFX908-NEXT: s_nop 1145; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]146; GFX908-NEXT: s_endpgm147;148; GFX90A-LABEL: test_mfma_i32_16x16x16i8:149; GFX90A: ; %bb.0: ; %bb150; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24151; GFX90A-NEXT: v_mov_b32_e32 v0, 1152; GFX90A-NEXT: v_mov_b32_e32 v2, 2153; GFX90A-NEXT: v_mov_b32_e32 v1, 0154; GFX90A-NEXT: s_waitcnt lgkmcnt(0)155; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0156; GFX90A-NEXT: s_waitcnt lgkmcnt(0)157; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0158; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1159; GFX90A-NEXT: v_accvgpr_write_b32 a2, s2160; GFX90A-NEXT: v_accvgpr_write_b32 a3, s3161; GFX90A-NEXT: s_nop 1162; GFX90A-NEXT: v_mfma_i32_16x16x16i8 a[0:3], v0, v2, a[0:3] cbsz:1 abid:2 blgp:3163; GFX90A-NEXT: s_nop 10164; GFX90A-NEXT: global_store_dwordx4 v1, a[0:3], s[6:7]165; GFX90A-NEXT: s_endpgm166bb:167 %in.1 = load <4 x i32>, ptr addrspace(1) %arg168 %mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.16x16x16i8(i32 1, i32 2, <4 x i32> %in.1, i32 1, i32 2, i32 3)169 store <4 x i32> %mai.1, ptr addrspace(1) %arg170 ret void171}172 173attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }174;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:175; GCN: {{.*}}176