diff options
| author | Stanislav Mekhanoshin <Stanislav.Mekhanoshin@amd.com> | 2019-11-20 12:27:20 -0800 |
|---|---|---|
| committer | Stanislav Mekhanoshin <Stanislav.Mekhanoshin@amd.com> | 2019-11-20 12:33:12 -0800 |
| commit | 899cdf95d9ad199e6c2f87cd544455e3991626dc (patch) | |
| tree | 2056f5a6d47e34eab05b99af99f17ca72633a354 /llvm/test/CodeGen/AMDGPU | |
| parent | 6f4398d1b9950d48ead91b2b550792f2bbe4778e (diff) | |
| download | bcm5719-llvm-899cdf95d9ad199e6c2f87cd544455e3991626dc.tar.gz bcm5719-llvm-899cdf95d9ad199e6c2f87cd544455e3991626dc.zip | |
[AMDGPU] Fixed mfma test check. NFC.
Diffstat (limited to 'llvm/test/CodeGen/AMDGPU')
| -rw-r--r-- | llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll | 78 |
1 files changed, 9 insertions, 69 deletions
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll index 00199faf04a..7b34d873f7a 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll @@ -1,5 +1,5 @@ -; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck --check-prefixes=GCN,NOLIT-SRCC %s -; RUN: llc -march=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug -verify-machineinstrs < %s | FileCheck --check-prefixes=GCN,LIT-SRCC %s +; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,NOLIT-SRCC %s +; RUN: llc -march=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,LIT-SRCC %s declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32) declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32, i32, i32) @@ -1295,73 +1295,13 @@ bb: } ; 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 -; GDN-DAG: global_load_dwordx4 -; GDN-DAG: global_load_dwordx4 -; GDN-DAG: global_load_dwordx4 -; GDN-DAG: global_load_dwordx4 -; GDN-DAG: global_load_dwordx4 -; GDN-DAG: global_load_dwordx4 -; GDN-DAG: global_load_dwordx4 -; GDN-DAG: global_load_dwordx4 -; 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-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-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-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-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-DAG: 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-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: v_accvgpr_read_b32 -; GCN-DAG: global_store_dwordx4 -; GCN-DAG: global_store_dwordx4 -; GCN-DAG: global_store_dwordx4 -; GCN-DAG: global_store_dwordx4 -; GCN-DAG: global_store_dwordx4 -; GCN-DAG: global_store_dwordx4 -; GCN-DAG: global_store_dwordx4 -; GCN-DAG: 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-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 define amdgpu_kernel void @test_mfma_f32_32x32x1f32_vecarg(<32 x float> addrspace(1)* %arg) { bb: %tid = call i32 @llvm.amdgcn.workitem.id.x() |
