From e67cc380a800d91297bae9e82ea3357ff39e379d Mon Sep 17 00:00:00 2001 From: Stanislav Mekhanoshin Date: Thu, 11 Jul 2019 21:19:33 +0000 Subject: [AMDGPU] gfx908 mfma support Differential Revision: https://reviews.llvm.org/D64584 llvm-svn: 365824 --- llvm/test/CodeGen/AMDGPU/agpr-register-count.ll | 15 +++++++++++++++ 1 file changed, 15 insertions(+) create mode 100644 llvm/test/CodeGen/AMDGPU/agpr-register-count.ll (limited to 'llvm/test/CodeGen/AMDGPU/agpr-register-count.ll') diff --git a/llvm/test/CodeGen/AMDGPU/agpr-register-count.ll b/llvm/test/CodeGen/AMDGPU/agpr-register-count.ll new file mode 100644 index 00000000000..ab4fcc54f65 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/agpr-register-count.ll @@ -0,0 +1,15 @@ +; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s + +declare <32 x i32> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x i32>, i32, i32, i32) + +; GCN-LABEL: {{^}}test_32_agprs: +; GCN: v_mfma_f32_32x32x1f32 a[0:31], {{v[0-9]+}}, {{v[0-9]+}}, 0 +; GCN-NOT: v28 +; GCN: NumVgprs: 32 +; GCN: VGPRBlocks: 7 +define amdgpu_kernel void @test_32_agprs(<32 x i32> addrspace(1)* %arg) { +bb: + %mai.1 = tail call <32 x i32> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x i32> , i32 0, i32 0, i32 0) + store <32 x i32> %mai.1, <32 x i32> addrspace(1)* %arg + ret void +} -- cgit v1.2.3