This is an archive of the discontinued LLVM Phabricator instance.

Differential D65158

[AMDGPU] Add all vgpr classes to asm parser
ClosedPublic

Authored by rampitec on Jul 23 2019, 11:08 AM.

Details

Reviewers
arsenm
tpr
dp
Commits
rG5cdacea29719: [AMDGPU] Add all vgpr classes to asm parser
rL366917: [AMDGPU] Add all vgpr classes to asm parser

Diff Detail

Event Timeline

rampitec created this revision.Jul 23 2019, 11:08 AM
Herald added subscribers: t-tye, tpr, dstuttard and 5 others. · View Herald TranscriptJul 23 2019, 11:08 AM
rampitec updated this revision to Diff 211349.Jul 23 2019, 12:28 PM
rampitec retitled this revision from [AMDGPU] Add all vgpr classes to AMDGPUOperand::isVReg() to [AMDGPU] Add all vgpr classes to asm parser.
rampitec edited the summary of this revision. (Show Details)
rampitec added a reviewer: tpr.

Found where it was a problem.

rampitec added a reviewer: dp.Jul 24 2019, 8:44 AM
dp accepted this revision.Jul 24 2019, 8:54 AM

LGTM

This revision is now accepted and ready to land.Jul 24 2019, 8:54 AM
Closed by commit rL366917: [AMDGPU] Add all vgpr classes to asm parser (authored by rampitec). · Explain WhyJul 24 2019, 9:21 AM
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptJul 24 2019, 9:21 AM

Revision Contents

PathSize
lib/
Target/
AMDGPU/
AsmParser/
6 lines
test/
MC/
AMDGPU/
3 lines
4 lines

Diff 211349

lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

Show First 20 Lines • Show All 269 Lines • ▼ Show 20 Lines bool isRegOrImmWithFP64InputMods() const {
return isRegOrImmWithInputMods(AMDGPU::VS_64RegClassID, MVT::f64); return isRegOrImmWithInputMods(AMDGPU::VS_64RegClassID, MVT::f64);
} }
bool isVReg() const { bool isVReg() const {
return isRegClass(AMDGPU::VGPR_32RegClassID) || return isRegClass(AMDGPU::VGPR_32RegClassID) ||
isRegClass(AMDGPU::VReg_64RegClassID) || isRegClass(AMDGPU::VReg_64RegClassID) ||
isRegClass(AMDGPU::VReg_96RegClassID) || isRegClass(AMDGPU::VReg_96RegClassID) ||
isRegClass(AMDGPU::VReg_128RegClassID) || isRegClass(AMDGPU::VReg_128RegClassID) ||
isRegClass(AMDGPU::VReg_160RegClassID) ||
isRegClass(AMDGPU::VReg_256RegClassID) || isRegClass(AMDGPU::VReg_256RegClassID) ||
isRegClass(AMDGPU::VReg_512RegClassID); isRegClass(AMDGPU::VReg_512RegClassID) ||
isRegClass(AMDGPU::VReg_1024RegClassID);
} }
bool isVReg32() const { bool isVReg32() const {
return isRegClass(AMDGPU::VGPR_32RegClassID); return isRegClass(AMDGPU::VGPR_32RegClassID);
} }
bool isVReg32OrOff() const { bool isVReg32OrOff() const {
return isOff() || isVReg32(); return isOff() || isVReg32();
▲ Show 20 Lines • Show All 1,579 Lines • ▼ Show 20 Lines
static int getRegClass(RegisterKind Is, unsigned RegWidth) { static int getRegClass(RegisterKind Is, unsigned RegWidth) {
if (Is == IS_VGPR) { if (Is == IS_VGPR) {
switch (RegWidth) { switch (RegWidth) {
default: return -1; default: return -1;
case 1: return AMDGPU::VGPR_32RegClassID; case 1: return AMDGPU::VGPR_32RegClassID;
case 2: return AMDGPU::VReg_64RegClassID; case 2: return AMDGPU::VReg_64RegClassID;
case 3: return AMDGPU::VReg_96RegClassID; case 3: return AMDGPU::VReg_96RegClassID;
case 4: return AMDGPU::VReg_128RegClassID; case 4: return AMDGPU::VReg_128RegClassID;
case 5: return AMDGPU::VReg_160RegClassID;
case 8: return AMDGPU::VReg_256RegClassID; case 8: return AMDGPU::VReg_256RegClassID;
case 16: return AMDGPU::VReg_512RegClassID; case 16: return AMDGPU::VReg_512RegClassID;
case 32: return AMDGPU::VReg_1024RegClassID;
} }
} else if (Is == IS_TTMP) { } else if (Is == IS_TTMP) {
switch (RegWidth) { switch (RegWidth) {
default: return -1; default: return -1;
case 1: return AMDGPU::TTMP_32RegClassID; case 1: return AMDGPU::TTMP_32RegClassID;
case 2: return AMDGPU::TTMP_64RegClassID; case 2: return AMDGPU::TTMP_64RegClassID;
case 4: return AMDGPU::TTMP_128RegClassID; case 4: return AMDGPU::TTMP_128RegClassID;
case 8: return AMDGPU::TTMP_256RegClassID; case 8: return AMDGPU::TTMP_256RegClassID;
▲ Show 20 Lines • Show All 991 Lines • Show Last 20 Lines

test/MC/AMDGPU/gfx9_asm_all.s

Show First 20 Lines • Show All 991 Lines • ▼ Show 20 Lines
// CHECK: [0x00,0x01,0x02,0xf0,0x01,0x05,0x02,0x00] // CHECK: [0x00,0x01,0x02,0xf0,0x01,0x05,0x02,0x00]
image_load v5, v[1:4], s[8:15] dmask:0x1 da image_load v5, v[1:4], s[8:15] dmask:0x1 da
// CHECK: [0x00,0x41,0x00,0xf0,0x01,0x05,0x02,0x00] // CHECK: [0x00,0x41,0x00,0xf0,0x01,0x05,0x02,0x00]
image_load v5, v[1:4], s[8:15] dmask:0x1 d16 image_load v5, v[1:4], s[8:15] dmask:0x1 d16
// CHECK: [0x00,0x01,0x00,0xf0,0x01,0x05,0x02,0x80] // 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 image_load_mip v5, v[1:4], s[8:15] dmask:0x1
// CHECK: [0x00,0x01,0x04,0xf0,0x01,0x05,0x02,0x00] // CHECK: [0x00,0x01,0x04,0xf0,0x01,0x05,0x02,0x00]
image_load_mip v252, v[1:4], s[8:15] dmask:0x1 image_load_mip v252, v[1:4], s[8:15] dmask:0x1
// CHECK: [0x00,0x01,0x04,0xf0,0x01,0xfc,0x02,0x00] // CHECK: [0x00,0x01,0x04,0xf0,0x01,0xfc,0x02,0x00]
image_load_mip v5, v[252:255], s[8:15] dmask:0x1 image_load_mip v5, v[252:255], s[8:15] dmask:0x1
// CHECK: [0x00,0x01,0x04,0xf0,0xfc,0x05,0x02,0x00] // CHECK: [0x00,0x01,0x04,0xf0,0xfc,0x05,0x02,0x00]
▲ Show 20 Lines • Show All 991 Lines • Show Last 20 Lines

test/MC/AMDGPU/mai-err.s

Show All 26 Lines
v_accvgpr_write_b32 a0, 65 v_accvgpr_write_b32 a0, 65
// GFX908: error: invalid operand for instruction // GFX908: error: invalid operand for instruction
v_accvgpr_write_b32 a0, v0 v_accvgpr_write_b32 a0, v0
// GFX900: error: instruction not supported on this GPU // GFX900: error: instruction not supported on this GPU
v_mfma_f32_32x32x1f32 v[0:31], v0, v1, a[1:32] 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] 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] v_mfma_f32_32x32x1f32 a[0:31], s0, v1, a[1:32]
// GFX908: error: invalid operand for instruction // GFX908: error: invalid operand for instruction
v_mfma_f32_32x32x1f32 a[0:31], 1, v1, a[1:32] v_mfma_f32_32x32x1f32 a[0:31], 1, v1, a[1:32]
// GFX908: error: invalid operand for instruction // GFX908: error: invalid operand for instruction
v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 65 v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 65
// GFX908: error: invalid operand for instruction // GFX908: error: invalid operand for instruction
v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 0 v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 0
// GFX900: error: instruction not supported on this GPU // GFX900: error: instruction not supported on this GPU