[llvm] r366917 - [AMDGPU] Add all vgpr classes to asm parser

[Stanislav Mekhanoshin via llvm-commits]

Author: rampitec
Date: Wed Jul 24 09:21:18 2019
New Revision: 366917

URL: http://llvm.org/viewvc/llvm-project?rev=366917&view=rev
Log:
[AMDGPU] Add all vgpr classes to asm parser

Differential Revision: https://reviews.llvm.org/D65158

Modified:
    llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
    llvm/trunk/test/MC/AMDGPU/gfx9_asm_all.s
    llvm/trunk/test/MC/AMDGPU/mai-err.s

Modified: llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp?rev=366917&r1=366916&r2=366917&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp Wed Jul 24 09:21:18 2019
@@ -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,
       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) {

Modified: llvm/trunk/test/MC/AMDGPU/gfx9_asm_all.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/gfx9_asm_all.s?rev=366917&r1=366916&r2=366917&view=diff
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/gfx9_asm_all.s (original)
+++ llvm/trunk/test/MC/AMDGPU/gfx9_asm_all.s Wed Jul 24 09:21:18 2019
@@ -5977,6 +5977,9 @@ image_load v5, v[1:4], s[8:15] dmask:0x1
 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]
 

Modified: llvm/trunk/test/MC/AMDGPU/mai-err.s
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/mai-err.s?rev=366917&r1=366916&r2=366917&view=diff
==============================================================================
--- llvm/trunk/test/MC/AMDGPU/mai-err.s (original)
+++ llvm/trunk/test/MC/AMDGPU/mai-err.s Wed Jul 24 09:21:18 2019
@@ -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