| ; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefix=GCN %s |
| |
| declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32) |
| declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32, i32, i32) |
| declare i32 @llvm.amdgcn.workitem.id.x() |
| |
| ; GCN-LABEL: {{^}}test_load_mfma_store16: |
| ; GCN-COUNT-8: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-NOT: v_accvgpr_write |
| ; GCN: v_mfma_f32_32x32x1f32 |
| ; GCN-NEXT: s_nop 7 |
| ; GCN-NEXT: s_nop 7 |
| ; GCN-NEXT: s_nop 2 |
| ; GCN-NOT: v_accvgpr_read |
| ; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, a[{{[0-9:]+}}], s[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_load_mfma_store16(<32 x float> addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid |
| %in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep |
| %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) |
| store <32 x float> %mai.1, <32 x float> addrspace(1)* %gep |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_load1_mfma_store1: |
| ; GCN: global_load_dword a{{[0-9]+}}, v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-NOT: v_accvgpr_read |
| ; GCN: v_mfma_f32_32x32x1f32 a{{\[}}[[N:[0-9]+]]: |
| ; GCN-NEXT: s_nop 7 |
| ; GCN-NEXT: s_nop 7 |
| ; GCN-NEXT: s_nop 2 |
| ; GCN-NOT: v_accvgpr_read |
| ; GCN-NEXT: global_store_dword v{{[0-9:]+}}, a[[N]], s[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_load1_mfma_store1(float addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds float, float addrspace(1)* %arg, i32 %tid |
| %in.1 = load float, float addrspace(1)* %gep |
| %init = insertelement <32 x float> zeroinitializer, float %in.1, i32 0 |
| %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) |
| %elt = extractelement <32 x float> %mai.1, i32 0 |
| store float %elt, float addrspace(1)* %gep |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_load4_mfma_store4: |
| ; GCN: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-NOT: v_accvgpr_write |
| ; GCN: v_mfma_i32_4x4x4i8 [[A:a\[[0-9:]+\]]] |
| ; GCN-NEXT: s_nop 4 |
| ; GCN-NOT: v_accvgpr_read |
| ; GCN-NEXT: global_store_dwordx4 v{{[0-9:]+}}, [[A]], s[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_load4_mfma_store4(<4 x i32> addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds <4 x i32>, <4 x i32> addrspace(1)* %arg, i32 %tid |
| %in.1 = load <4 x i32>, <4 x i32> addrspace(1)* %gep |
| %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) |
| store <4 x i32> %mai.1, <4 x i32> addrspace(1)* %gep |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_load_store: |
| ; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-NOT: v_accvgpr |
| ; GCN-COUNT-8: global_store_dwordx4 v[{{[0-9:]+}}], v[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_load_store(<32 x float> addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep.1 = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid |
| %gep.2 = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %gep.1, i32 32 |
| %in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep.1 |
| store <32 x float> %in.1, <32 x float> addrspace(1)* %gep.2 |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_load_add_mfma_store: |
| ; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-COUNT-32: v_accvgpr_write |
| ; GCN: v_mfma_f32_32x32x1f32 |
| ; GCN-NEXT: s_nop 7 |
| ; GCN-NEXT: s_nop 7 |
| ; GCN-NEXT: s_nop 2 |
| ; GCN-NOT: v_accvgpr_read |
| ; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, a[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_load_add_mfma_store(<32 x float> addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid |
| %in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep |
| %add.1 = fadd <32 x float> %in.1, %in.1 |
| %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) |
| store <32 x float> %mai.1, <32 x float> addrspace(1)* %gep |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_load_add_store: |
| ; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-NOT: v_accvgpr |
| ; GCN-COUNT-16: v_pk_add_f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_load_add_store(<32 x float> addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid |
| %in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep |
| %add.1 = fadd <32 x float> %in.1, %in.1 |
| store <32 x float> %add.1, <32 x float> addrspace(1)* %gep |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_load_mfma_add_store: |
| ; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-COUNT-32: v_accvgpr_write |
| ; GCN: v_mfma_f32_32x32x1f32 |
| ; GCN-COUNT-32: v_accvgpr_read |
| ; GCN: v_pk_add_f32 |
| ; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_load_mfma_add_store(<32 x float> addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid |
| %in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep |
| %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) |
| %add.1 = fadd <32 x float> %mai.1, %in.1 |
| store <32 x float> %add.1, <32 x float> addrspace(1)* %gep |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_load_add_mfma_mul_store: |
| ; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN: v_pk_add_f32 |
| ; GCN-COUNT-32: v_accvgpr_write |
| ; GCN: v_mfma_f32_32x32x1f32 |
| ; GCN-COUNT-32: v_accvgpr_read |
| ; GCN: v_pk_mul_f32 |
| ; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_load_add_mfma_mul_store(<32 x float> addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid |
| %in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep |
| %add.1 = fadd <32 x float> %in.1, %in.1 |
| %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) |
| %mul.1 = fmul <32 x float> %mai.1, %mai.1 |
| store <32 x float> %mul.1, <32 x float> addrspace(1)* %gep |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mixeduse_load_add_mfma_mul_store: |
| ; GCN-COUNT-8: global_load_dwordx4 v[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-COUNT-32: v_accvgpr_write |
| ; GCN: v_mfma_f32_32x32x1f32 |
| ; GCN-COUNT-32: v_accvgpr_read |
| ; GCN: v_pk_mul_f32 |
| ; GCN-COUNT-8: global_store_dwordx4 v{{[0-9:]+}}, v[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_mixeduse_load_add_mfma_mul_store(<32 x float> addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid |
| %in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep |
| %add.1 = fadd <32 x float> %in.1, %in.1 |
| %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) |
| %mul.1 = fmul <32 x float> %mai.1, %in.1 |
| store <32 x float> %mul.1, <32 x float> addrspace(1)* %gep |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_multiuse_load_mfma_mfma_store: |
| ; GCN-COUNT-8: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-NOT: v_accvgpr_write |
| ; GCN: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr_read |
| ; GCN-COUNT-8: global_store_dwordx4 v[{{[0-9:]+}}], a[{{[0-9:]+}}] |
| define amdgpu_kernel void @test_multiuse_load_mfma_mfma_store(<32 x float> addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep.1 = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid |
| %gep.2 = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %gep.1, i32 32 |
| %in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep.1 |
| %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) |
| %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) |
| store <32 x float> %mai.1, <32 x float> addrspace(1)* %gep.1 |
| store <32 x float> %mai.2, <32 x float> addrspace(1)* %gep.2 |
| ret void |
| } |
| |
| ; NB: for atomics both vdata and vdst shall be either VGPR or AGPR |
| ; GCN-LABEL: {{^}}test_atomic_mfma_4xi32_atomic_store: |
| ; GCN: global_atomic_sub [[IN:v[0-9]+]], v{{[0-9:]+}}, v{{[0-9]+}}, s[{{[0-9:]+}}] glc |
| ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, [[IN]] |
| ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} |
| ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} |
| ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} |
| ; GCN: v_mfma_i32_4x4x4i8 a{{\[}}[[N:[0-9]+]]: |
| ; GCN: v_accvgpr_read_b32 [[V:v[0-9]+]], a[[N]]{{$}} |
| ; GCN: global_atomic_add v{{[0-9]+}}, v{{[0-9:]+}}, [[V]], s[{{[0-9:]+}}] glc |
| ; GCN: global_store_dword v{{[0-9]+}}, v{{[0-9]+}}, |
| define amdgpu_kernel void @test_atomic_mfma_4xi32_atomic_store(i32 addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds i32, i32 addrspace(1)* %arg, i32 %tid |
| %in.1 = atomicrmw volatile sub i32 addrspace(1)* %gep, i32 1 seq_cst |
| %tmp0 = insertelement <4 x i32> undef, i32 %in.1, i32 0 |
| %tmp1 = insertelement <4 x i32> %tmp0, i32 0, i32 1 |
| %tmp2 = insertelement <4 x i32> %tmp1, i32 0, i32 2 |
| %tmp3 = insertelement <4 x i32> %tmp2, i32 0, i32 3 |
| %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) |
| %elt = extractelement <4 x i32> %mai.1, i32 0 |
| %val = atomicrmw volatile add i32 addrspace(1)* %gep, i32 %elt seq_cst |
| store i32 %val, i32 addrspace(1)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_atomic_mfma_4xi32_atomic64_store: |
| ; GCN: global_atomic_sub_x2 v[{{[0-9:]+}}], v{{[0-9:]+}}, v[{{[0-9:]+}}], s[{{[0-9:]+}}] glc |
| ; GCN-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}} |
| ; GCN: v_mfma_i32_4x4x4i8 a{{\[}}[[N:[0-9]+]]: |
| ; GCN: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}} |
| ; GCN: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}} |
| ; GCN: global_atomic_add_x2 v[{{[0-9:]+}}], v{{[0-9:]+}}, v[{{[0-9:]+}}], s[{{[0-9:]+}}] glc |
| define amdgpu_kernel void @test_atomic_mfma_4xi32_atomic64_store(i64 addrspace(1)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds i64, i64 addrspace(1)* %arg, i32 %tid |
| %in.1 = atomicrmw volatile sub i64 addrspace(1)* %gep, i64 1 seq_cst |
| %tmp0 = insertelement <2 x i64> undef, i64 %in.1, i32 0 |
| %tmp1 = insertelement <2 x i64> %tmp0, i64 0, i32 1 |
| %tmp2 = bitcast <2 x i64> %tmp0 to <4 x i32> |
| %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) |
| %elt.1 = extractelement <4 x i32> %mai.1, i32 0 |
| %elt.2 = extractelement <4 x i32> %mai.1, i32 1 |
| %v2.1 = insertelement <2 x i32> undef, i32 %elt.1, i32 0 |
| %v2.2 = insertelement <2 x i32> %v2.1, i32 %elt.2, i32 1 |
| %v2 = bitcast <2 x i32> %v2.2 to i64 |
| %val = atomicrmw volatile add i64 addrspace(1)* %gep, i64 %v2 seq_cst |
| store i64 %val, i64 addrspace(1)* %arg |
| ret void |
| } |
| |
| ; NB: both data operands should be VGPR or AGPR |
| ; GCN-LABEL: {{^}}test_load_mfma_ds2_store: |
| ; GCN-DAG: ds_read_b128 [[IN:a\[[0-9:]+\]]], v{{[0-9:]+}} |
| ; GCN-NOT: v_accvgpr_write |
| ; GCN-DAG: v_mfma_i32_4x4x4i8 a{{\[}}[[N:[0-9]+]]:{{[0-9]+}}], v{{[0-9:]+}}, v{{[0-9:]+}}, [[IN]] |
| ; GCN-DAG: ds_write_b32 v{{[0-9]+}}, v{{[0-9]+}} |
| ; GCN-NOT: v_accvgpr_read |
| ; GCN: ds_write_b32 v{{[0-9]+}}, a[[N]] offset:128 |
| define amdgpu_kernel void @test_load_mfma_ds2_store(<4 x i32> addrspace(3)* %arg) { |
| bb: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep.1 = getelementptr inbounds <4 x i32>, <4 x i32> addrspace(3)* %arg, i32 %tid |
| %in.1 = load <4 x i32>, <4 x i32> addrspace(3)* %gep.1 |
| %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) |
| %elt = extractelement <4 x i32> %mai.1, i32 0 |
| %ptr = bitcast <4 x i32> addrspace(3)* %arg to i32 addrspace(3)* |
| %gep.2 = getelementptr inbounds i32, i32 addrspace(3)* %ptr, i32 32 |
| store i32 1, i32 addrspace(3)* %ptr |
| store i32 %elt, i32 addrspace(3)* %gep.2 |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_4xi32: |
| ; GCN: global_load_dwordx4 [[IN:a\[[0-9:]+\]]], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-NOT: v_accvgpr_write |
| ; GCN: v_mfma_i32_4x4x4i8 [[RES:a\[[0-9:]+\]]], v{{[0-9:]+}}, v{{[0-9:]+}}, [[IN]] |
| ; GCN-NOT: v_accvgpr_read |
| ; GCN: global_store_dwordx4 v[{{[0-9:]+}}], [[RES]], |
| define amdgpu_kernel void @test_mfma_loop_4xi32(<4 x i32> addrspace(1)* %arg) { |
| entry: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds <4 x i32>, <4 x i32> addrspace(1)* %arg, i32 %tid |
| %in = load <4 x i32>, <4 x i32> addrspace(1)* %gep |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <4 x i32> [ %in, %entry ], [ %mai.1, %for.cond.preheader ] |
| %c = phi i32 [ 0, %entry ], [ %inc, %for.cond.preheader ] |
| %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) |
| %inc = add nuw nsw i32 %c, 1 |
| %cc = icmp eq i32 %inc, 16 |
| br i1 %cc, label %exit, label %for.cond.preheader |
| |
| exit: |
| store <4 x i32> %mai.1, <4 x i32> addrspace(1)* %gep |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_32xfloat: |
| ; GCN-COUNT-8: global_load_dwordx4 a[{{[0-9:]+}}], v{{[0-9:]+}}, s[{{[0-9:]+}}] |
| ; GCN-NOT: v_accvgpr_write |
| ; GCN: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr_read |
| ; GCN-COUNT-8: global_store_dwordx4 v[{{[0-9:]+}}], a[{{[0-9:]+}}], |
| ; GCN: s_endpgm |
| define amdgpu_kernel void @test_mfma_loop_32xfloat(<32 x float> addrspace(1)* %arg) { |
| entry: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid |
| %in = load <32 x float>, <32 x float> addrspace(1)* %gep |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ %in, %entry ], [ %mai.1, %for.cond.preheader ] |
| %c = phi i32 [ 0, %entry ], [ %inc, %for.cond.preheader ] |
| %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) |
| %inc = add nuw nsw i32 %c, 1 |
| %cc = icmp eq i32 %inc, 16 |
| br i1 %cc, label %exit, label %for.cond.preheader |
| |
| exit: |
| store <32 x float> %mai.1, <32 x float> addrspace(1)* %gep |
| ret void |
| } |