321 lines · plain
1; RUN: llc -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck -enable-var-scope --check-prefix=GCN %s2; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -early-live-intervals < %s | FileCheck -enable-var-scope --check-prefix=GCN %s3 4declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)5declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32, i32, i32)6declare i32 @llvm.amdgcn.workitem.id.x()7 8; GCN-LABEL: {{^}}test_load_mfma_store16:9; GCN-COUNT-8: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]10; GCN-NOT: v_accvgpr_write11; GCN: v_mfma_f32_32x32x1f3212; GCN-NEXT: s_nop 1513; GCN-NEXT: s_nop 214; GCN-NOT: v_accvgpr_read15; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, a[{{[0-9:]+}}], s[{{[0-9:]+}}]16define amdgpu_kernel void @test_load_mfma_store16(ptr addrspace(1) %arg) #0 {17bb:18 %tid = call i32 @llvm.amdgcn.workitem.id.x()19 %gep = getelementptr inbounds <32 x float>, ptr addrspace(1) %arg, i32 %tid20 %in.1 = load <32 x float>, ptr addrspace(1) %gep21 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)22 store <32 x float> %mai.1, ptr addrspace(1) %gep23 ret void24}25 26; GCN-LABEL: {{^}}test_load1_mfma_store1:27; GCN: global_load_dword a{{[0-9]+}}, v{{[0-9:]+}}, s[{{[0-9:]+}}]28; GCN-NOT: v_accvgpr_read29; GCN: v_mfma_f32_32x32x1f32 a[[[N:[0-9]+]]:30; GCN-NEXT: s_nop 1531; GCN-NEXT: s_nop 232; GCN-NOT: v_accvgpr_read33; GCN-NEXT: global_store_dword v{{[0-9:]+}}, a[[N]], s[{{[0-9:]+}}]34define amdgpu_kernel void @test_load1_mfma_store1(ptr addrspace(1) %arg) #0 {35bb:36 %tid = call i32 @llvm.amdgcn.workitem.id.x()37 %gep = getelementptr inbounds float, ptr addrspace(1) %arg, i32 %tid38 %in.1 = load float, ptr addrspace(1) %gep39 %init = insertelement <32 x float> zeroinitializer, float %in.1, i32 040 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %init, i32 1, i32 2, i32 3)41 %elt = extractelement <32 x float> %mai.1, i32 042 store float %elt, ptr addrspace(1) %gep43 ret void44}45 46; GCN-LABEL: {{^}}test_load4_mfma_store4:47; GCN: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]48; GCN-NOT: v_accvgpr_write49; GCN: v_mfma_i32_4x4x4i8 [[A:a\[[0-9:]+\]]]50; GCN-NEXT: s_nop 451; GCN-NOT: v_accvgpr_read52; GCN-NEXT: global_store_dwordx4 v{{[0-9:]+}}, [[A]], s[{{[0-9:]+}}]53define amdgpu_kernel void @test_load4_mfma_store4(ptr addrspace(1) %arg) #0 {54bb:55 %tid = call i32 @llvm.amdgcn.workitem.id.x()56 %gep = getelementptr inbounds <4 x i32>, ptr addrspace(1) %arg, i32 %tid57 %in.1 = load <4 x i32>, ptr addrspace(1) %gep58 %mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %in.1, i32 0, i32 0, i32 0)59 store <4 x i32> %mai.1, ptr addrspace(1) %gep60 ret void61}62 63; GCN-LABEL: {{^}}test_load_store:64; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]65; GCN-NOT: v_accvgpr66; GCN-COUNT-8: global_store_dwordx4 v[{{[0-9:]+}}], v[{{[0-9:]+}}]67define amdgpu_kernel void @test_load_store(ptr addrspace(1) %arg) #0 {68bb:69 %tid = call i32 @llvm.amdgcn.workitem.id.x()70 %gep.1 = getelementptr inbounds <32 x float>, ptr addrspace(1) %arg, i32 %tid71 %gep.2 = getelementptr inbounds <32 x float>, ptr addrspace(1) %gep.1, i32 3272 %in.1 = load <32 x float>, ptr addrspace(1) %gep.173 store <32 x float> %in.1, ptr addrspace(1) %gep.274 ret void75}76 77; GCN-LABEL: {{^}}test_load_add_mfma_store:78; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]79; GCN-COUNT-32: v_accvgpr_write80; GCN: v_mfma_f32_32x32x1f3281; GCN-NEXT: s_nop 1582; GCN-NEXT: s_nop 283; GCN-NOT: v_accvgpr_read84; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, a[{{[0-9:]+}}]85define amdgpu_kernel void @test_load_add_mfma_store(ptr addrspace(1) %arg) #0 {86bb:87 %tid = call i32 @llvm.amdgcn.workitem.id.x()88 %gep = getelementptr inbounds <32 x float>, ptr addrspace(1) %arg, i32 %tid89 %in.1 = load <32 x float>, ptr addrspace(1) %gep90 %add.1 = fadd <32 x float> %in.1, %in.191 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %add.1, i32 1, i32 2, i32 3)92 store <32 x float> %mai.1, ptr addrspace(1) %gep93 ret void94}95 96; GCN-LABEL: {{^}}test_load_add_store:97; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]98; GCN-NOT: v_accvgpr99; GCN-COUNT-16: v_pk_add_f32100; GCN-NOT: v_accvgpr101; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}]102define amdgpu_kernel void @test_load_add_store(ptr addrspace(1) %arg) #0 {103bb:104 %tid = call i32 @llvm.amdgcn.workitem.id.x()105 %gep = getelementptr inbounds <32 x float>, ptr addrspace(1) %arg, i32 %tid106 %in.1 = load <32 x float>, ptr addrspace(1) %gep107 %add.1 = fadd <32 x float> %in.1, %in.1108 store <32 x float> %add.1, ptr addrspace(1) %gep109 ret void110}111 112; GCN-LABEL: {{^}}test_load_mfma_add_store:113; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]114; GCN-COUNT-32: v_accvgpr_write115; GCN: v_mfma_f32_32x32x1f32116; GCN-COUNT-32: v_accvgpr_read117; GCN: v_pk_add_f32118; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}]119define amdgpu_kernel void @test_load_mfma_add_store(ptr addrspace(1) %arg) #0 {120bb:121 %tid = call i32 @llvm.amdgcn.workitem.id.x()122 %gep = getelementptr inbounds <32 x float>, ptr addrspace(1) %arg, i32 %tid123 %in.1 = load <32 x float>, ptr addrspace(1) %gep124 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)125 %add.1 = fadd <32 x float> %mai.1, %in.1126 store <32 x float> %add.1, ptr addrspace(1) %gep127 ret void128}129 130; GCN-LABEL: {{^}}test_load_add_mfma_mul_store:131; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]132; GCN: v_pk_add_f32133; GCN-COUNT-32: v_accvgpr_write134; GCN: v_mfma_f32_32x32x1f32135; GCN-COUNT-32: v_accvgpr_read136; GCN: v_pk_mul_f32137; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}]138define amdgpu_kernel void @test_load_add_mfma_mul_store(ptr addrspace(1) %arg) #0 {139bb:140 %tid = call i32 @llvm.amdgcn.workitem.id.x()141 %gep = getelementptr inbounds <32 x float>, ptr addrspace(1) %arg, i32 %tid142 %in.1 = load <32 x float>, ptr addrspace(1) %gep143 %add.1 = fadd <32 x float> %in.1, %in.1144 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %add.1, i32 1, i32 2, i32 3)145 %mul.1 = fmul <32 x float> %mai.1, %mai.1146 store <32 x float> %mul.1, ptr addrspace(1) %gep147 ret void148}149 150; GCN-LABEL: {{^}}test_mixeduse_load_add_mfma_mul_store:151; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]152; GCN-COUNT-32: v_accvgpr_write153; GCN: v_mfma_f32_32x32x1f32154; GCN-COUNT-32: v_accvgpr_read155; GCN: v_pk_mul_f32156; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}]157define amdgpu_kernel void @test_mixeduse_load_add_mfma_mul_store(ptr addrspace(1) %arg) #0 {158bb:159 %tid = call i32 @llvm.amdgcn.workitem.id.x()160 %gep = getelementptr inbounds <32 x float>, ptr addrspace(1) %arg, i32 %tid161 %in.1 = load <32 x float>, ptr addrspace(1) %gep162 %add.1 = fadd <32 x float> %in.1, %in.1163 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %add.1, i32 1, i32 2, i32 3)164 %mul.1 = fmul <32 x float> %mai.1, %in.1165 store <32 x float> %mul.1, ptr addrspace(1) %gep166 ret void167}168 169; GCN-LABEL: {{^}}test_multiuse_load_mfma_mfma_store:170; GCN-COUNT-8: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]171; GCN-NOT: v_accvgpr_write172; GCN: v_mfma_f32_32x32x1f32173; GCN-NOT: v_accvgpr_read174; GCN-COUNT-8: global_store_dwordx4 v[{{[0-9:]+}}], a[{{[0-9:]+}}]175define amdgpu_kernel void @test_multiuse_load_mfma_mfma_store(ptr addrspace(1) %arg) #0 {176bb:177 %tid = call i32 @llvm.amdgcn.workitem.id.x()178 %gep.1 = getelementptr inbounds <32 x float>, ptr addrspace(1) %arg, i32 %tid179 %gep.2 = getelementptr inbounds <32 x float>, ptr addrspace(1) %gep.1, i32 32180 %in.1 = load <32 x float>, ptr addrspace(1) %gep.1181 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)182 %mai.2 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 0, i32 0, i32 0)183 store <32 x float> %mai.1, ptr addrspace(1) %gep.1184 store <32 x float> %mai.2, ptr addrspace(1) %gep.2185 ret void186}187 188; NB: for atomics both vdata and vdst shall be either VGPR or AGPR189; GCN-LABEL: {{^}}test_atomic_mfma_4xi32_atomic_store:190; GCN: v_accvgpr_write_b32 [[A_ZERO:a[0-9]+]], 0191; GCN: global_atomic_sub [[IN:v[0-9]+]], v{{[0-9:]+}}, v{{[0-9]+}}, s[{{[0-9:]+}}] glc192; GCN-DAG: v_accvgpr_mov_b32 a{{[0-9]+}}, [[A_ZERO]]193; GCN-DAG: v_accvgpr_mov_b32 a{{[0-9]+}}, [[A_ZERO]]194; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, [[IN]]195; GCN: v_mfma_i32_4x4x4i8 a[[[N:[0-9]+]]:196; GCN: v_accvgpr_read_b32 [[V:v[0-9]+]], a[[N]]{{$}}197; GCN: global_atomic_add v{{[0-9]+}}, v{{[0-9:]+}}, [[V]], s[{{[0-9:]+}}] glc198; GCN: global_store_dword v{{[0-9]+}}, v{{[0-9]+}},199define amdgpu_kernel void @test_atomic_mfma_4xi32_atomic_store(ptr addrspace(1) %arg) #0 {200bb:201 %tid = call i32 @llvm.amdgcn.workitem.id.x()202 %gep = getelementptr inbounds i32, ptr addrspace(1) %arg, i32 %tid203 %in.1 = atomicrmw volatile sub ptr addrspace(1) %gep, i32 1 syncscope("agent") seq_cst, !amdgpu.no.remote.memory !0204 %tmp0 = insertelement <4 x i32> poison, i32 %in.1, i32 0205 %tmp1 = insertelement <4 x i32> %tmp0, i32 0, i32 1206 %tmp2 = insertelement <4 x i32> %tmp1, i32 0, i32 2207 %tmp3 = insertelement <4 x i32> %tmp2, i32 0, i32 3208 %mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %tmp3, i32 0, i32 0, i32 0)209 %elt = extractelement <4 x i32> %mai.1, i32 0210 %val = atomicrmw volatile add ptr addrspace(1) %gep, i32 %elt syncscope("agent") seq_cst211 store i32 %val, ptr addrspace(1) %arg212 ret void213}214 215; GCN-LABEL: {{^}}test_atomic_mfma_4xi32_atomic64_store:216; GCN: global_atomic_sub_x2 v[{{[0-9:]+}}], v{{[0-9:]+}}, v[{{[0-9:]+}}], s[{{[0-9:]+}}] glc217; GCN: v_accvgpr_write_b32 [[A_ZERO:a[0-9]+]], 0218; GCN: v_accvgpr_mov_b32 a{{[0-9]+}}, [[A_ZERO]]219; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}220; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}221; GCN: v_mfma_i32_4x4x4i8 a[[[N:[0-9]+]]:222; GCN: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}223; GCN: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}224; GCN: global_atomic_add_x2 v[{{[0-9:]+}}], v{{[0-9:]+}}, v[{{[0-9:]+}}], s[{{[0-9:]+}}] glc225define amdgpu_kernel void @test_atomic_mfma_4xi32_atomic64_store(ptr addrspace(1) %arg) #0 {226bb:227 %tid = call i32 @llvm.amdgcn.workitem.id.x()228 %gep = getelementptr inbounds i64, ptr addrspace(1) %arg, i32 %tid229 %in.1 = atomicrmw volatile sub ptr addrspace(1) %gep, i64 1 syncscope("agent") seq_cst, !amdgpu.no.remote.memory !0230 %tmp0 = insertelement <2 x i64> poison, i64 %in.1, i32 0231 %tmp1 = insertelement <2 x i64> %tmp0, i64 0, i32 1232 %tmp2 = bitcast <2 x i64> %tmp1 to <4 x i32>233 %mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %tmp2, i32 0, i32 0, i32 0)234 %elt.1 = extractelement <4 x i32> %mai.1, i32 0235 %elt.2 = extractelement <4 x i32> %mai.1, i32 1236 %v2.1 = insertelement <2 x i32> poison, i32 %elt.1, i32 0237 %v2.2 = insertelement <2 x i32> %v2.1, i32 %elt.2, i32 1238 %v2 = bitcast <2 x i32> %v2.2 to i64239 %val = atomicrmw volatile add ptr addrspace(1) %gep, i64 %v2 syncscope("agent") seq_cst240 store i64 %val, ptr addrspace(1) %arg241 ret void242}243 244; NB: both data operands should be VGPR or AGPR245; GCN-LABEL: {{^}}test_load_mfma_ds2_store:246; GCN-DAG: ds_read_b128 [[IN:a\[[0-9:]+\]]], v{{[0-9:]+}}247; GCN-NOT: v_accvgpr_write248; GCN-DAG: v_mfma_i32_4x4x4i8 a[[[N:[0-9]+]]:{{[0-9]+}}], v{{[0-9:]+}}, v{{[0-9:]+}}, [[IN]]249; GCN-DAG: ds_write_b32 v{{[0-9]+}}, v{{[0-9]+}}250; GCN-NOT: v_accvgpr_read251; GCN: ds_write_b32 v{{[0-9]+}}, a[[N]] offset:128252define amdgpu_kernel void @test_load_mfma_ds2_store(ptr addrspace(3) %arg) #0 {253bb:254 %tid = call i32 @llvm.amdgcn.workitem.id.x()255 %gep.1 = getelementptr inbounds <4 x i32>, ptr addrspace(3) %arg, i32 %tid256 %in.1 = load <4 x i32>, ptr addrspace(3) %gep.1257 %mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %in.1, i32 0, i32 0, i32 0)258 %elt = extractelement <4 x i32> %mai.1, i32 0259 %gep.2 = getelementptr inbounds i32, ptr addrspace(3) %arg, i32 32260 store i32 1, ptr addrspace(3) %arg261 store i32 %elt, ptr addrspace(3) %gep.2262 ret void263}264 265; GCN-LABEL: {{^}}test_mfma_loop_4xi32:266; GCN: global_load_dwordx4 [[IN:a\[[0-9:]+\]]], v{{[0-9:]+}}, s[{{[0-9:]+}}]267; GCN-NOT: v_accvgpr_write268; GCN: v_mfma_i32_4x4x4i8 [[RES:a\[[0-9:]+\]]], v{{[0-9:]+}}, v{{[0-9:]+}}, [[IN]]269; GCN-NOT: v_accvgpr_read270; GCN: global_store_dwordx4 v[{{[0-9:]+}}], [[RES]],271define amdgpu_kernel void @test_mfma_loop_4xi32(ptr addrspace(1) %arg) #0 {272entry:273 %tid = call i32 @llvm.amdgcn.workitem.id.x()274 %gep = getelementptr inbounds <4 x i32>, ptr addrspace(1) %arg, i32 %tid275 %in = load <4 x i32>, ptr addrspace(1) %gep276 br label %for.cond.preheader277 278for.cond.preheader:279 %phi = phi <4 x i32> [ %in, %entry ], [ %mai.1, %for.cond.preheader ]280 %c = phi i32 [ 0, %entry ], [ %inc, %for.cond.preheader ]281 %mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %phi, i32 0, i32 0, i32 0)282 %inc = add nuw nsw i32 %c, 1283 %cc = icmp eq i32 %inc, 16284 br i1 %cc, label %exit, label %for.cond.preheader285 286exit:287 store <4 x i32> %mai.1, ptr addrspace(1) %gep288 ret void289}290 291; GCN-LABEL: {{^}}test_mfma_loop_32xfloat:292; GCN-COUNT-8: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}]293; GCN-NOT: v_accvgpr_write294; GCN: v_mfma_f32_32x32x1f32295; GCN-NOT: v_accvgpr_read296; GCN-COUNT-8: global_store_dwordx4 v[{{[0-9:]+}}], a[{{[0-9:]+}}],297; GCN: s_endpgm298define amdgpu_kernel void @test_mfma_loop_32xfloat(ptr addrspace(1) %arg) #0 {299entry:300 %tid = call i32 @llvm.amdgcn.workitem.id.x()301 %gep = getelementptr inbounds <32 x float>, ptr addrspace(1) %arg, i32 %tid302 %in = load <32 x float>, ptr addrspace(1) %gep303 br label %for.cond.preheader304 305for.cond.preheader:306 %phi = phi <32 x float> [ %in, %entry ], [ %mai.1, %for.cond.preheader ]307 %c = phi i32 [ 0, %entry ], [ %inc, %for.cond.preheader ]308 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %phi, i32 0, i32 0, i32 0)309 %inc = add nuw nsw i32 %c, 1310 %cc = icmp eq i32 %inc, 16311 br i1 %cc, label %exit, label %for.cond.preheader312 313exit:314 store <32 x float> %mai.1, ptr addrspace(1) %gep315 ret void316}317 318attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }319 320!0 = !{}321