maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[AMDGPU] Add all vgpr classes to asm parser 

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

llvm-svn: 366917
This commit is contained in:
Stanislav Mekhanoshin 2019-07-24 16:21:18 +00:00
parent 0b7f226311
commit 5cdacea297
3 changed files with 10 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
View File

@ -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) {

3
llvm/test/MC/AMDGPU/gfx9_asm_all.s
View File

@ -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]

4
llvm/test/MC/AMDGPU/mai-err.s
View File

@ -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