AMDGPU: Fix lit test checks with dag option

[dstuttard]

Summary

I was seeing some failures on a test with slightly different instruction
ordering. Adding in some DAG directives solved the issue.

Change-Id: If5a3d3969055fb19279943bd45161bb70a3dabce

[rampitec]

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll

1300

AFAIK DAG-COUNT does not work.

[dstuttard]

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll

1300

Yes, you're right. It just wasn't triggering the check at all - hence passing.
I'll re-submit with lots of GCN-DAG checks instead, unless you can think of a better way to do it.

[dstuttard]

Replace COUNT-n with n DAG versions (DAG-COUNT-n doesn't work)

[arsenm]

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll

1300

Repeating -DAG doesn’t work either as far as I know. Might as well just generate this test

[dstuttard]

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll

1300

Doesn't it? The other tests in this file do this.

Problem is that I'm seeing different ordering of the mov instructions - hence requiring the DAG variants.
Doesn't the -DAG here mean that FileCheck is expecting to see the starting LABEL, then 2 mov's 8 global_load_dwordx4's and 16 v_accvgpr_write_b32's in any order, followed by the v_mfma instruction?

[rampitec]

LGTM. Repeating -DAG statements do work, just verified.

[dstuttard]

Matt - are you happy for this to proceed?

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll

 define amdgpu_kernel void @test_mfma_f32_4x4x1f32_lit_splat(<4 x float> addrspace(1)* %arg) {
 bb:
   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> <float 123.0, float 123.0, float 123.0, float 123.0>, i32 0, i32 0, i32 0)
   store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_vecarg:
-; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
-; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
-; GCN-COUNT-8:     global_load_dwordx4
-; GCN-COUNT-16:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
-; GCN:             v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
-; GCN-COUNT-32:    v_accvgpr_read_b32
-; GCN-COUNT-8:     global_store_dwordx4
+; GCN-DAG:             v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
+; GCN-DAG:             v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
+; GCN-DAG-COUNT-8:     global_load_dwordx4
+; GCN-DAG-COUNT-16:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; GCN:                 v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
+; GCN-COUNT-32:        v_accvgpr_read_b32
+; GCN-COUNT-8:         global_store_dwordx4
 define amdgpu_kernel void @test_mfma_f32_32x32x1f32_vecarg(<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
 }