| ; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX908,GFX908_A %s |
| ; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX90A,GFX908_A %s |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_zeroinit: |
| |
| ; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX90A: v_accvgpr_write_b32 [[LEAD:a[0-9]+]], 0 |
| ; GFX90A-COUNT-31: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]] |
| |
| ; Check that we do not copy agprs to vgprs and back inside the loop. |
| |
| ; GCN: [[LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[LOOP]] |
| |
| ; Final result should be read only once after the loop. |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_loop_zeroinit(<32 x float> addrspace(1)* %arg) { |
| entry: |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ zeroinitializer, %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)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_unfoldable_splat: |
| |
| ; Check that we do not use 32 temp vgprs, but rotate 3 vgprs only. |
| ; 3 vgprs are needed to avoid wait states between writes. |
| ; Check that we do not use 32 temp sgprs as well. |
| |
| ; GFX908_A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x42f60000 |
| ; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_accvgpr_write_b32 [[LEAD:a[0-9]+]], [[TMP]] |
| ; GFX90A-COUNT-31: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]] |
| |
| ; GCN: [[LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[LOOP]] |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_loop_unfoldable_splat(<32 x float> addrspace(1)* %arg) { |
| entry: |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ <float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0, float 123.0>, %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)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_non_splat: |
| |
| ; GCN: v_accvgpr_write_b32 [[LEAD:a[0-9]+]], 0{{$}} |
| ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0{{$}} |
| ; GFX908-COUNT-30: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX90A-COUNT-30: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]]{{$}} |
| |
| ; GCN: [[LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[LOOP]] |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_loop_non_splat(<32 x float> addrspace(1)* %arg) { |
| entry: |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ <float 0.0, float 1.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0>, %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)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_unfoldable_seq: |
| |
| ; Check that we do not use 32 temp vgprs, but rotate 3 vgprs only. |
| ; 3 vgprs are needed to avoid wait states between writes. |
| |
| ; GFX908: v_mov_b32_e32 [[TMP1:v[0-9]+]], 0x42f60000 |
| ; GFX908: v_mov_b32_e32 [[TMP2:v[0-9]+]], 0x42f80000 |
| ; GFX908: v_mov_b32_e32 [[TMP3:v[0-9]+]], 0x42fe0000 |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| ; GFX908: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}} |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]] |
| ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]] |
| |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x4{{[0-9a-f]+}} |
| ; GFX90A: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| |
| |
| ; GCN: [[LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[LOOP]] |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_loop_unfoldable_seq(<32 x float> addrspace(1)* %arg) { |
| entry: |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ <float 123.0, float 124.0, float 125.0, float 126.0, float 127.0, float 128.0, float 129.0, float 130.0, float 131.0, float 132.0, float 133.0, float 134.0, float 135.0, float 136.0, float 137.0, float 138.0, float 139.0, float 140.0, float 141.0, float 142.0, float 143.0, float 144.0, float 145.0, float 146.0, float 147.0, float 148.0, float 149.0, float 150.0, float 151.0, float 152.0, float 153.0, float 154.0>, %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)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_vgpr_init: |
| |
| ; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, v0{{$}} |
| |
| ; GCN: [[LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[LOOP]] |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_loop_vgpr_init(<32 x float> addrspace(1)* %arg) { |
| entry: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %init = bitcast i32 %tid to float |
| %tmp0 = insertelement <32 x float> undef, float %init, i32 0 |
| %tmp1 = insertelement <32 x float> %tmp0, float %init, i32 1 |
| %tmp2 = insertelement <32 x float> %tmp1, float %init, i32 2 |
| %tmp3 = insertelement <32 x float> %tmp2, float %init, i32 3 |
| %tmp4 = insertelement <32 x float> %tmp3, float %init, i32 4 |
| %tmp5 = insertelement <32 x float> %tmp4, float %init, i32 5 |
| %tmp6 = insertelement <32 x float> %tmp5, float %init, i32 6 |
| %tmp7 = insertelement <32 x float> %tmp6, float %init, i32 7 |
| %tmp8 = insertelement <32 x float> %tmp7, float %init, i32 8 |
| %tmp9 = insertelement <32 x float> %tmp8, float %init, i32 9 |
| %tmp10 = insertelement <32 x float> %tmp9, float %init, i32 10 |
| %tmp11 = insertelement <32 x float> %tmp10, float %init, i32 11 |
| %tmp12 = insertelement <32 x float> %tmp11, float %init, i32 12 |
| %tmp13 = insertelement <32 x float> %tmp12, float %init, i32 13 |
| %tmp14 = insertelement <32 x float> %tmp13, float %init, i32 14 |
| %tmp15 = insertelement <32 x float> %tmp14, float %init, i32 15 |
| %tmp16 = insertelement <32 x float> %tmp15, float %init, i32 16 |
| %tmp17 = insertelement <32 x float> %tmp16, float %init, i32 17 |
| %tmp18 = insertelement <32 x float> %tmp17, float %init, i32 18 |
| %tmp19 = insertelement <32 x float> %tmp18, float %init, i32 19 |
| %tmp20 = insertelement <32 x float> %tmp19, float %init, i32 20 |
| %tmp21 = insertelement <32 x float> %tmp20, float %init, i32 21 |
| %tmp22 = insertelement <32 x float> %tmp21, float %init, i32 22 |
| %tmp23 = insertelement <32 x float> %tmp22, float %init, i32 23 |
| %tmp24 = insertelement <32 x float> %tmp23, float %init, i32 24 |
| %tmp25 = insertelement <32 x float> %tmp24, float %init, i32 25 |
| %tmp26 = insertelement <32 x float> %tmp25, float %init, i32 26 |
| %tmp27 = insertelement <32 x float> %tmp26, float %init, i32 27 |
| %tmp28 = insertelement <32 x float> %tmp27, float %init, i32 28 |
| %tmp29 = insertelement <32 x float> %tmp28, float %init, i32 29 |
| %tmp30 = insertelement <32 x float> %tmp29, float %init, i32 30 |
| %tmp31 = insertelement <32 x float> %tmp30, float %init, i32 31 |
| |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ %tmp31, %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)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_sgpr_init: |
| |
| ; GFX908_A: v_mov_b32_e32 [[TMP:v[0-9]+]], s{{[0-9]+}} |
| ; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX90A: v_accvgpr_write_b32 [[LEAD:a[0-9]+]], [[TMP]] |
| ; GFX90A-COUNT-31: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]] |
| |
| ; GCN: [[LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[LOOP]] |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_loop_sgpr_init(<32 x float> addrspace(1)* %arg, float %init) { |
| entry: |
| %tmp0 = insertelement <32 x float> undef, float %init, i32 0 |
| %tmp1 = insertelement <32 x float> %tmp0, float %init, i32 1 |
| %tmp2 = insertelement <32 x float> %tmp1, float %init, i32 2 |
| %tmp3 = insertelement <32 x float> %tmp2, float %init, i32 3 |
| %tmp4 = insertelement <32 x float> %tmp3, float %init, i32 4 |
| %tmp5 = insertelement <32 x float> %tmp4, float %init, i32 5 |
| %tmp6 = insertelement <32 x float> %tmp5, float %init, i32 6 |
| %tmp7 = insertelement <32 x float> %tmp6, float %init, i32 7 |
| %tmp8 = insertelement <32 x float> %tmp7, float %init, i32 8 |
| %tmp9 = insertelement <32 x float> %tmp8, float %init, i32 9 |
| %tmp10 = insertelement <32 x float> %tmp9, float %init, i32 10 |
| %tmp11 = insertelement <32 x float> %tmp10, float %init, i32 11 |
| %tmp12 = insertelement <32 x float> %tmp11, float %init, i32 12 |
| %tmp13 = insertelement <32 x float> %tmp12, float %init, i32 13 |
| %tmp14 = insertelement <32 x float> %tmp13, float %init, i32 14 |
| %tmp15 = insertelement <32 x float> %tmp14, float %init, i32 15 |
| %tmp16 = insertelement <32 x float> %tmp15, float %init, i32 16 |
| %tmp17 = insertelement <32 x float> %tmp16, float %init, i32 17 |
| %tmp18 = insertelement <32 x float> %tmp17, float %init, i32 18 |
| %tmp19 = insertelement <32 x float> %tmp18, float %init, i32 19 |
| %tmp20 = insertelement <32 x float> %tmp19, float %init, i32 20 |
| %tmp21 = insertelement <32 x float> %tmp20, float %init, i32 21 |
| %tmp22 = insertelement <32 x float> %tmp21, float %init, i32 22 |
| %tmp23 = insertelement <32 x float> %tmp22, float %init, i32 23 |
| %tmp24 = insertelement <32 x float> %tmp23, float %init, i32 24 |
| %tmp25 = insertelement <32 x float> %tmp24, float %init, i32 25 |
| %tmp26 = insertelement <32 x float> %tmp25, float %init, i32 26 |
| %tmp27 = insertelement <32 x float> %tmp26, float %init, i32 27 |
| %tmp28 = insertelement <32 x float> %tmp27, float %init, i32 28 |
| %tmp29 = insertelement <32 x float> %tmp28, float %init, i32 29 |
| %tmp30 = insertelement <32 x float> %tmp29, float %init, i32 30 |
| %tmp31 = insertelement <32 x float> %tmp30, float %init, i32 31 |
| |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ %tmp31, %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)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_mixed_init: |
| |
| ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v0 |
| ; GFX908_A-DAG: v_mov_b32_e32 [[TMP:v[0-9]+]], s{{[0-9]+}} |
| ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]] |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX908-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| |
| ; GFX90A-DAG: v_accvgpr_write_b32 [[LEAD:a[0-9]+]], 0 |
| ; GFX90A-COUNT-28: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]] |
| |
| ; GCN: [[LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[LOOP]] |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_loop_mixed_init(<32 x float> addrspace(1)* %arg, float %x) { |
| entry: |
| %tid = call i32 @llvm.amdgcn.workitem.id.x() |
| %init = bitcast i32 %tid to float |
| %tmp0 = insertelement <32 x float> zeroinitializer, float %init, i32 0 |
| %tmp1 = insertelement <32 x float> %tmp0, float %x, i32 1 |
| |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ %tmp1, %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)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_mfma_forward_init: |
| |
| ; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0 |
| ; GFX908: v_mfma_f32_32x32x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9:]+}}] |
| ; GFX90A-NOT: v_accvgpr |
| ; GFX90A: v_mfma_f32_32x32x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, 0{{$}} |
| ; GFX90A-NOT: v_accvgpr |
| ; GCN-NOT: v_accvgpr |
| |
| ; GCN: [[LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[LOOP]] |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_loop_mfma_forward_init(<32 x float> addrspace(1)* %arg) { |
| entry: |
| %mai.0 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> zeroinitializer, i32 0, i32 0, i32 0) |
| |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ %mai.0, %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)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_loop_agpr_init: |
| |
| ; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0 |
| ; GFX908: v_mfma_f32_32x32x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9:]+}}] |
| ; GFX90A-NOT: v_accvgpr |
| ; GFX90A: v_mfma_f32_32x32x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, 0{{$}} |
| ; GFX90A-NOT: v_accvgpr |
| |
| ; Check that we are using only one tmp VGPR. |
| |
| ; GCN: v_accvgpr_read_b32 [[TMP:v[0-9]+]], a{{[0-9]+}} |
| ; GFX908-COUNT-31: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]{{$}} |
| ; GFX90A: v_accvgpr_write_b32 [[LEAD:a[0-9]+]], [[TMP]]{{$}} |
| ; GFX90A-COUNT-29: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]] |
| |
| ; GCN: [[LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[LOOP]] |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_loop_agpr_init(<32 x float> addrspace(1)* %arg) { |
| entry: |
| %mai.0 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> zeroinitializer, i32 0, i32 0, i32 0) |
| %init = extractelement <32 x float> %mai.0, i32 0 |
| %tmp0 = insertelement <32 x float> undef, float %init, i32 0 |
| %tmp1 = insertelement <32 x float> %tmp0, float %init, i32 1 |
| %tmp2 = insertelement <32 x float> %tmp1, float %init, i32 2 |
| %tmp3 = insertelement <32 x float> %tmp2, float %init, i32 3 |
| %tmp4 = insertelement <32 x float> %tmp3, float %init, i32 4 |
| %tmp5 = insertelement <32 x float> %tmp4, float %init, i32 5 |
| %tmp6 = insertelement <32 x float> %tmp5, float %init, i32 6 |
| %tmp7 = insertelement <32 x float> %tmp6, float %init, i32 7 |
| %tmp8 = insertelement <32 x float> %tmp7, float %init, i32 8 |
| %tmp9 = insertelement <32 x float> %tmp8, float %init, i32 9 |
| %tmp10 = insertelement <32 x float> %tmp9, float %init, i32 10 |
| %tmp11 = insertelement <32 x float> %tmp10, float %init, i32 11 |
| %tmp12 = insertelement <32 x float> %tmp11, float %init, i32 12 |
| %tmp13 = insertelement <32 x float> %tmp12, float %init, i32 13 |
| %tmp14 = insertelement <32 x float> %tmp13, float %init, i32 14 |
| %tmp15 = insertelement <32 x float> %tmp14, float %init, i32 15 |
| %tmp16 = insertelement <32 x float> %tmp15, float %init, i32 16 |
| %tmp17 = insertelement <32 x float> %tmp16, float %init, i32 17 |
| %tmp18 = insertelement <32 x float> %tmp17, float %init, i32 18 |
| %tmp19 = insertelement <32 x float> %tmp18, float %init, i32 19 |
| %tmp20 = insertelement <32 x float> %tmp19, float %init, i32 20 |
| %tmp21 = insertelement <32 x float> %tmp20, float %init, i32 21 |
| %tmp22 = insertelement <32 x float> %tmp21, float %init, i32 22 |
| %tmp23 = insertelement <32 x float> %tmp22, float %init, i32 23 |
| %tmp24 = insertelement <32 x float> %tmp23, float %init, i32 24 |
| %tmp25 = insertelement <32 x float> %tmp24, float %init, i32 25 |
| %tmp26 = insertelement <32 x float> %tmp25, float %init, i32 26 |
| %tmp27 = insertelement <32 x float> %tmp26, float %init, i32 27 |
| %tmp28 = insertelement <32 x float> %tmp27, float %init, i32 28 |
| %tmp29 = insertelement <32 x float> %tmp28, float %init, i32 29 |
| %tmp30 = insertelement <32 x float> %tmp29, float %init, i32 30 |
| %tmp31 = insertelement <32 x float> %tmp30, float %init, i32 31 |
| |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi = phi <32 x float> [ %tmp31, %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)* %arg |
| ret void |
| } |
| |
| ; GCN-LABEL: {{^}}test_mfma_nested_loop_zeroinit: |
| |
| ; GFX908-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}} |
| ; GFX90A: v_accvgpr_write_b32 [[LEAD:a[0-9]+]], 0 |
| ; GFX90A-COUNT-31: v_accvgpr_mov_b32 a{{[0-9]+}}, [[LEAD]] |
| |
| ; Check that we do not copy agprs to vgprs and back in an outer loop. |
| |
| ; GCN: [[OUTER_LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GCN: [[INNER_LOOP:.LBB[0-9_]+]]: |
| ; GCN-NOT: v_accvgpr |
| ; GFX908_A: v_mfma_f32_32x32x1f32 |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[INNER_LOOP]] |
| ; GCN-NOT: v_accvgpr |
| ; GCN: s_cbranch_scc1 [[OUTER_LOOP]] |
| |
| ; Final result should be read only once after the loop. |
| |
| ; GFX908-COUNT-32: v_accvgpr_read_b32 |
| ; GFX90A-NOT: v_accvgpr_read_b32 |
| ; GFX908-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}] |
| ; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}] |
| |
| define amdgpu_kernel void @test_mfma_nested_loop_zeroinit(<32 x float> addrspace(1)* %arg) { |
| entry: |
| br label %for.cond.preheader |
| |
| for.cond.preheader: |
| %phi.0 = phi <32 x float> [ zeroinitializer, %entry ], [ %mai.1, %inner.exit ] |
| %c.0 = phi i32 [ 0, %entry ], [ %inc.0, %inner.exit ] |
| br label %inner.for.cond.preheader |
| |
| inner.for.cond.preheader: |
| %phi = phi <32 x float> [ %phi.0, %for.cond.preheader ], [ %mai.1, %inner.for.cond.preheader ] |
| %c = phi i32 [ 0, %for.cond.preheader ], [ %inc, %inner.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 %inner.exit, label %inner.for.cond.preheader |
| |
| inner.exit: |
| %inc.0 = add nuw nsw i32 %c.0, 1 |
| %cc.0 = icmp eq i32 %inc.0, 16 |
| br i1 %cc.0, label %exit, label %for.cond.preheader |
| |
| exit: |
| store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg |
| ret void |
| } |
| |
| declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32) |
| declare i32 @llvm.amdgcn.workitem.id.x() |