From b90efd9c2b3dbb303545e09be18615450785a41c Mon Sep 17 00:00:00 2001 From: Stanislav Mekhanoshin Date: Wed, 24 Jul 2019 16:21:18 +0000 Subject: [PATCH] [AMDGPU] Add all vgpr classes to asm parser Differential Revision: https://reviews.llvm.org/D65158 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@366917 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 6 +++++- test/MC/AMDGPU/gfx9_asm_all.s | 3 +++ test/MC/AMDGPU/mai-err.s | 4 ++-- 3 files changed, 10 insertions(+), 3 deletions(-) diff --git a/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp index cb2d37ad819..b51858e4f7c 100644 --- a/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp +++ b/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp @@ -275,8 +275,10 @@ public: isRegClass(AMDGPU::VReg_64RegClassID) || isRegClass(AMDGPU::VReg_96RegClassID) || isRegClass(AMDGPU::VReg_128RegClassID) || + isRegClass(AMDGPU::VReg_160RegClassID) || isRegClass(AMDGPU::VReg_256RegClassID) || - isRegClass(AMDGPU::VReg_512RegClassID); + isRegClass(AMDGPU::VReg_512RegClassID) || + isRegClass(AMDGPU::VReg_1024RegClassID); } bool isVReg32() const { @@ -1872,8 +1874,10 @@ static int getRegClass(RegisterKind Is, unsigned RegWidth) { case 2: return AMDGPU::VReg_64RegClassID; case 3: return AMDGPU::VReg_96RegClassID; case 4: return AMDGPU::VReg_128RegClassID; + case 5: return AMDGPU::VReg_160RegClassID; case 8: return AMDGPU::VReg_256RegClassID; case 16: return AMDGPU::VReg_512RegClassID; + case 32: return AMDGPU::VReg_1024RegClassID; } } else if (Is == IS_TTMP) { switch (RegWidth) { diff --git a/test/MC/AMDGPU/gfx9_asm_all.s b/test/MC/AMDGPU/gfx9_asm_all.s index fc4d8f438cc..cb800da9ca6 100644 --- a/test/MC/AMDGPU/gfx9_asm_all.s +++ b/test/MC/AMDGPU/gfx9_asm_all.s @@ -5977,6 +5977,9 @@ image_load v5, v[1:4], s[8:15] dmask:0x1 da image_load v5, v[1:4], s[8:15] dmask:0x1 d16 // CHECK: [0x00,0x01,0x00,0xf0,0x01,0x05,0x02,0x80] +image_load v[0:4], v5, s[0:7] dmask:0xf unorm tfe +// CHECK: [0x00,0x1f,0x01,0xf0,0x05,0x00,0x00,0x00] + image_load_mip v5, v[1:4], s[8:15] dmask:0x1 // CHECK: [0x00,0x01,0x04,0xf0,0x01,0x05,0x02,0x00] diff --git a/test/MC/AMDGPU/mai-err.s b/test/MC/AMDGPU/mai-err.s index 4834b895e37..24d90d61441 100644 --- a/test/MC/AMDGPU/mai-err.s +++ b/test/MC/AMDGPU/mai-err.s @@ -32,10 +32,10 @@ v_accvgpr_write_b32 a0, v0 // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x1f32 v[0:31], v0, v1, a[1:32] -// GFX908: error: not a valid operand +// GFX908: error: invalid operand for instruction v_mfma_f32_32x32x1f32 a[0:31], v0, v1, v[1:32] -// GFX908: error: not a valid operand +// GFX908: error: invalid operand for instruction v_mfma_f32_32x32x1f32 a[0:31], s0, v1, a[1:32] // GFX908: error: invalid operand for instruction -- 2.40.0