This is an archive of the discontinued LLVM Phabricator instance.

Differential D100069

Disable use of SCC bit from asm
ClosedPublic

Authored by rampitec on Apr 7 2021, 2:06 PM.

Details

Reviewers
t-tye
Commits
rG37878de50367: Disable use of SCC bit from asm

Diff Detail

Event Timeline

rampitec created this revision.Apr 7 2021, 2:06 PM
Herald added subscribers: kerbowa, jfb, hiraditya and 3 others. · View Herald TranscriptApr 7 2021, 2:06 PM
rampitec requested review of this revision.Apr 7 2021, 2:06 PM
Herald added a project: Restricted Project. · View Herald TranscriptApr 7 2021, 2:06 PM
arsenm added a comment.Apr 7 2021, 2:10 PM

Message / name is somewhat confusing since I assumed this meant the scc register

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
4097–4098

I don't think this should be using SMLoc to check the string content. Isn't there an already parsed token?

rampitec retitled this revision from Disable use of SCC from asm to Disable use of SCC bit from asm.Apr 7 2021, 2:14 PM
rampitec added a comment.Apr 7 2021, 2:16 PM
In D100069#2675145, @arsenm wrote:

Message / name is somewhat confusing since I assumed this meant the scc register

Retitled.

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
4097–4098

Not since I did a function to parse all these bits together and in any order. I now have to search for it. Anyhow, it is just moved from the place below.

Harbormaster completed remote builds in B97592: Diff 335923.Apr 7 2021, 2:55 PM
t-tye added inline comments.Apr 7 2021, 3:14 PM
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
4099

Could the GPU name be printed? I think that is done in other error messages?

rampitec added inline comments.Apr 7 2021, 3:17 PM
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
4099

Asm parser universally prints "on this GPU". If we want to be specific we probably need to update it everywhere.

t-tye accepted this revision.Apr 7 2021, 3:19 PM

LGTM

This revision is now accepted and ready to land.Apr 7 2021, 3:19 PM
Closed by commit rG37878de50367: Disable use of SCC bit from asm (authored by rampitec). · Explain WhyApr 7 2021, 3:44 PM
This revision was automatically updated to reflect the committed changes.
rampitec added a commit: rG37878de50367: Disable use of SCC bit from asm.

Revision Contents

PathSize
llvm/
lib/
Target/
AMDGPU/
AsmParser/
16 lines
test/
MC/
AMDGPU/
38 lines
34 lines
2 lines
Disassembler/
AMDGPU/
32 lines

Diff 335940

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 4,086 Lines • ▼ Show 20 Linesbool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst,
uint64_t TSFlags = MII.get(Inst.getOpcode()).TSFlags; uint64_t TSFlags = MII.get(Inst.getOpcode()).TSFlags;
if ((TSFlags & (SIInstrFlags::SMRD)) && if ((TSFlags & (SIInstrFlags::SMRD)) &&
(CPol & ~(AMDGPU::CPol::GLC | AMDGPU::CPol::DLC))) { (CPol & ~(AMDGPU::CPol::GLC | AMDGPU::CPol::DLC))) {
Error(IDLoc, "invalid cache policy for SMRD instruction"); Error(IDLoc, "invalid cache policy for SMRD instruction");
return false; return false;
} }
if (isGFX90A() && (CPol & CPol::SCC)) {
SMLoc S = getImmLoc(AMDGPUOperand::ImmTyCPol, Operands);
StringRef CStr(S.getPointer());
S = SMLoc::getFromPointer(&CStr.data()[CStr.find("scc")]);
arsenmUnsubmitted
Not Done
ReplyInline Actions

I don't think this should be using SMLoc to check the string content. Isn't there an already parsed token?

arsenm: I don't think this should be using SMLoc to check the string content. Isn't there an already…
rampitecAuthorUnsubmitted
Done
ReplyInline Actions

Not since I did a function to parse all these bits together and in any order. I now have to search for it. Anyhow, it is just moved from the place below.

rampitec: Not since I did a function to parse all these bits together and in any order. I now have to…
Error(S, "scc is not supported on this GPU");
t-tyeUnsubmitted
Not Done
ReplyInline Actions

Could the GPU name be printed? I think that is done in other error messages?

t-tye: Could the GPU name be printed? I think that is done in other error messages?
rampitecAuthorUnsubmitted
Done
ReplyInline Actions

Asm parser universally prints "on this GPU". If we want to be specific we probably need to update it everywhere.

rampitec: Asm parser universally prints "on this GPU". If we want to be specific we probably need to…
return false;
}
if (!(TSFlags & (SIInstrFlags::IsAtomicNoRet | SIInstrFlags::IsAtomicRet))) if (!(TSFlags & (SIInstrFlags::IsAtomicNoRet | SIInstrFlags::IsAtomicRet)))
return true; return true;
if (TSFlags & SIInstrFlags::IsAtomicRet) { if (TSFlags & SIInstrFlags::IsAtomicRet) {
if (!(TSFlags & SIInstrFlags::MIMG) && !(CPol & CPol::GLC)) { if (!(TSFlags & SIInstrFlags::MIMG) && !(CPol & CPol::GLC)) {
Error(IDLoc, "instruction must use glc"); Error(IDLoc, "instruction must use glc");
return false; return false;
} }
} else { } else {
if (CPol & CPol::GLC) { if (CPol & CPol::GLC) {
SMLoc S = getImmLoc(AMDGPUOperand::ImmTyCPol, Operands); SMLoc S = getImmLoc(AMDGPUOperand::ImmTyCPol, Operands);
StringRef CStr(S.getPointer()); StringRef CStr(S.getPointer());
S = SMLoc::getFromPointer(&CStr.data()[CStr.find("glc")]); S = SMLoc::getFromPointer(&CStr.data()[CStr.find("glc")]);
Error(S, "instruction must not use glc"); Error(S, "instruction must not use glc");
return false; return false;
} }
} }
if (isGFX90A() && (CPol & CPol::SCC) && (TSFlags & SIInstrFlags::FPAtomic)) {
SMLoc S = getImmLoc(AMDGPUOperand::ImmTyCPol, Operands);
StringRef CStr(S.getPointer());
S = SMLoc::getFromPointer(&CStr.data()[CStr.find("scc")]);
Error(S, "instruction must not use scc");
return false;
}
return true; return true;
} }
bool AMDGPUAsmParser::validateInstruction(const MCInst &Inst, bool AMDGPUAsmParser::validateInstruction(const MCInst &Inst,
const SMLoc &IDLoc, const SMLoc &IDLoc,
const OperandVector &Operands) { const OperandVector &Operands) {
if (auto ErrMsg = validateLdsDirect(Inst)) { if (auto ErrMsg = validateLdsDirect(Inst)) {
Error(getRegLoc(LDS_DIRECT, Operands), *ErrMsg); Error(getRegLoc(LDS_DIRECT, Operands), *ErrMsg);
▲ Show 20 Lines • Show All 4,121 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx90a_asm_features.s

Show First 20 Lines • Show All 316 Lines • ▼ Show 20 Lines
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,0] ; encoding: [0x00,0x48,0xb3,0xd3,0x02,0x09,0x02,0x18] // GFX90A: v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,0] ; encoding: [0x00,0x48,0xb3,0xd3,0x02,0x09,0x02,0x18]
v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,0] v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,0]
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,1] ; encoding: [0x00,0x58,0xb3,0xd3,0x02,0x09,0x02,0x18] // GFX90A: v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,1] ; encoding: [0x00,0x58,0xb3,0xd3,0x02,0x09,0x02,0x18]
v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,1] v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,1]
// GFX90A: tbuffer_load_format_xyzw v[4:7], off, s[0:3], 0 scc ; encoding: [0x00,0x80,0x09,0xe8,0x00,0x04,0x20,0x80]
// GFX1010: error: not a valid operand.
// GFX908: error: scc modifier is not supported on this GPU
tbuffer_load_format_xyzw v[4:7], off, s[0:3], dfmt:1, nfmt:0, 0 scc
// GFX90A: tbuffer_load_format_xyzw v[4:7], off, s[0:3], 0 glc scc ; encoding: [0x00,0xc0,0x09,0xe8,0x00,0x04,0x20,0x80]
// GFX1010: error: not a valid operand.
// GFX908: error: scc modifier is not supported on this GPU
tbuffer_load_format_xyzw v[4:7], off, s[0:3], dfmt:1, nfmt:0, 0 glc scc
// NOT-GFX90A: error: scc modifier is not supported on this GPU
// GFX90A: buffer_load_dword v5, off, s[8:11], s3 offset:4095 scc ; encoding: [0xff,0x8f,0x50,0xe0,0x00,0x05,0x02,0x03]
buffer_load_dword v5, off, s[8:11], s3 offset:4095 scc
// NOT-GFX90A: error: scc modifier is not supported on this GPU
// GFX90A: buffer_load_dword v5, off, s[8:11], s3 offset:4095 glc scc ; encoding: [0xff,0xcf,0x50,0xe0,0x00,0x05,0x02,0x03]
buffer_load_dword v5, off, s[8:11], s3 offset:4095 glc scc
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: buffer_wbl2 ; encoding: [0x00,0x00,0xa0,0xe0,0x00,0x00,0x00,0x00] // GFX90A: buffer_wbl2 ; encoding: [0x00,0x00,0xa0,0xe0,0x00,0x00,0x00,0x00]
buffer_wbl2 buffer_wbl2
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: buffer_invl2 ; encoding: [0x00,0x00,0xa4,0xe0,0x00,0x00,0x00,0x00] // GFX90A: buffer_invl2 ; encoding: [0x00,0x00,0xa4,0xe0,0x00,0x00,0x00,0x00]
buffer_invl2 buffer_invl2
▲ Show 20 Lines • Show All 208 Lines • ▼ Show 20 Lines
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:4 ; encoding: [0x04,0x00,0xf8,0xd8,0x01,0x02,0x00,0x04] // GFX90A: ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:4 ; encoding: [0x04,0x00,0xf8,0xd8,0x01,0x02,0x00,0x04]
ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:4 ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:4
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:65535 gds ; encoding: [0xff,0xff,0xf9,0xd8,0x01,0x02,0x00,0x04] // GFX90A: ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:65535 gds ; encoding: [0xff,0xff,0xf9,0xd8,0x01,0x02,0x00,0x04]
ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:65535 gds ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:65535 gds
// NOT-GFX90A: error: scc modifier is not supported on this GPU
// GFX90A: flat_load_dword v0, v[0:1] scc ; encoding: [0x00,0x00,0x50,0xde,0x00,0x00,0x00,0x00]
flat_load_dword v0, v[0:1] scc
// NOT-GFX90A: error: scc modifier is not supported on this GPU
// GFX90A: flat_load_dword v0, v[0:1] glc scc ; encoding: [0x00,0x00,0x51,0xde,0x00,0x00,0x00,0x00]
flat_load_dword v0, v[0:1] glc scc
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: flat_atomic_add_f64 v[0:1], v[2:3] offset:4095 ; encoding: [0xff,0x0f,0x3c,0xdd,0x00,0x02,0x00,0x00] // GFX90A: flat_atomic_add_f64 v[0:1], v[2:3] offset:4095 ; encoding: [0xff,0x0f,0x3c,0xdd,0x00,0x02,0x00,0x00]
flat_atomic_add_f64 v[0:1], v[2:3] offset:4095 flat_atomic_add_f64 v[0:1], v[2:3] offset:4095
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: flat_atomic_add_f64 v[254:255], v[2:3] offset:4095 ; encoding: [0xff,0x0f,0x3c,0xdd,0xfe,0x02,0x00,0x00] // GFX90A: flat_atomic_add_f64 v[254:255], v[2:3] offset:4095 ; encoding: [0xff,0x0f,0x3c,0xdd,0xfe,0x02,0x00,0x00]
flat_atomic_add_f64 v[254:255], v[2:3] offset:4095 flat_atomic_add_f64 v[254:255], v[2:3] offset:4095
▲ Show 20 Lines • Show All 92 Lines • ▼ Show 20 Lines
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: global_atomic_min_f64 v[0:1], v[2:3], off ; encoding: [0x00,0x80,0x40,0xdd,0x00,0x02,0x7f,0x00] // GFX90A: global_atomic_min_f64 v[0:1], v[2:3], off ; encoding: [0x00,0x80,0x40,0xdd,0x00,0x02,0x7f,0x00]
global_atomic_min_f64 v[0:1], v[2:3], off global_atomic_min_f64 v[0:1], v[2:3], off
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: global_atomic_max_f64 v[0:1], v[2:3], off ; encoding: [0x00,0x80,0x44,0xdd,0x00,0x02,0x7f,0x00] // GFX90A: global_atomic_max_f64 v[0:1], v[2:3], off ; encoding: [0x00,0x80,0x44,0xdd,0x00,0x02,0x7f,0x00]
global_atomic_max_f64 v[0:1], v[2:3], off global_atomic_max_f64 v[0:1], v[2:3], off
// NOT-GFX90A: error: scc modifier is not supported on this GPU
// GFX90A: image_load v[0:4], v2, s[0:7] dmask:0xf unorm scc ; encoding: [0x80,0x1f,0x00,0xf0,0x02,0x00,0x00,0x00]
image_load v[0:4], v2, s[0:7] dmask:0xf unorm scc
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: v_fmac_f64_e32 v[4:5], v[2:3], v[4:5] ; encoding: [0x02,0x09,0x08,0x08] // GFX90A: v_fmac_f64_e32 v[4:5], v[2:3], v[4:5] ; encoding: [0x02,0x09,0x08,0x08]
v_fmac_f64_e32 v[4:5], v[2:3], v[4:5] v_fmac_f64_e32 v[4:5], v[2:3], v[4:5]
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
// GFX90A: v_fmac_f64_e32 v[254:255], v[2:3], v[4:5] ; encoding: [0x02,0x09,0xfc,0x09] // GFX90A: v_fmac_f64_e32 v[254:255], v[2:3], v[4:5] ; encoding: [0x02,0x09,0xfc,0x09]
v_fmac_f64_e32 v[254:255], v[2:3], v[4:5] v_fmac_f64_e32 v[254:255], v[2:3], v[4:5]
▲ Show 20 Lines • Show All 320 Lines • ▼ Show 20 Lines
// GFX90A: flat_atomic_max_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x45,0xdd,0x00,0x02,0x00,0x00] // GFX90A: flat_atomic_max_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x45,0xdd,0x00,0x02,0x00,0x00]
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
flat_atomic_max_f64 v[0:1], v[0:1], v[2:3] glc flat_atomic_max_f64 v[0:1], v[0:1], v[2:3] glc
// GFX90A: flat_atomic_min_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x41,0xdd,0x00,0x02,0x00,0x00] // GFX90A: flat_atomic_min_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x41,0xdd,0x00,0x02,0x00,0x00]
// NOT-GFX90A: error: instruction not supported on this GPU // NOT-GFX90A: error: instruction not supported on this GPU
flat_atomic_min_f64 v[0:1], v[0:1], v[2:3] glc flat_atomic_min_f64 v[0:1], v[0:1], v[2:3] glc
// GFX90A: global_atomic_add v[2:3], v5, off scc ; encoding: [0x00,0x80,0x08,0xdf,0x02,0x05,0x7f,0x00]
// NOT-GFX90A: error: scc modifier is not supported on this GPU
global_atomic_add v[2:3], v5, off scc
// GFX90A: global_atomic_add_f32 v0, v[0:1], v2, off glc ; encoding: [0x00,0x80,0x35,0xdd,0x00,0x02,0x7f,0x00] // GFX90A: global_atomic_add_f32 v0, v[0:1], v2, off glc ; encoding: [0x00,0x80,0x35,0xdd,0x00,0x02,0x7f,0x00]
// GFX908: error: operands are not valid for this GPU or mode // GFX908: error: operands are not valid for this GPU or mode
// GFX1010: error: instruction not supported on this GPU // GFX1010: error: instruction not supported on this GPU
global_atomic_add_f32 v0, v[0:1], v2, off glc global_atomic_add_f32 v0, v[0:1], v2, off glc
// GFX90A: global_atomic_add_f32 v[0:1], v2, off ; encoding: [0x00,0x80,0x34,0xdd,0x00,0x02,0x7f,0x00] // GFX90A: global_atomic_add_f32 v[0:1], v2, off ; encoding: [0x00,0x80,0x34,0xdd,0x00,0x02,0x7f,0x00]
// GFX1010: error: instruction not supported on this GPU // GFX1010: error: instruction not supported on this GPU
global_atomic_add_f32 v[0:1], v2, off global_atomic_add_f32 v[0:1], v2, off
// GFX90A: global_atomic_add_f32 v0, v2, s[0:1] ; encoding: [0x00,0x80,0x34,0xdd,0x00,0x02,0x00,0x00] // GFX90A: global_atomic_add_f32 v0, v2, s[0:1] ; encoding: [0x00,0x80,0x34,0xdd,0x00,0x02,0x00,0x00]
// GFX1010: error: instruction not supported on this GPU // GFX1010: error: instruction not supported on this GPU
global_atomic_add_f32 v0, v2, s[0:1] global_atomic_add_f32 v0, v2, s[0:1]
// GFX90A: global_atomic_add_f32 v1, v0, v2, s[0:1] glc // GFX90A: global_atomic_add_f32 v1, v0, v2, s[0:1] glc
// GFX908: error: operands are not valid for this GPU or mode // GFX908: error: operands are not valid for this GPU or mode
// GFX1010: error: instruction not supported on this GPU // GFX1010: error: instruction not supported on this GPU
global_atomic_add_f32 v1, v0, v2, s[0:1] glc ; encoding: [0x00,0x80,0x35,0xdd,0x00,0x02,0x00,0x01] global_atomic_add_f32 v1, v0, v2, s[0:1] glc ; encoding: [0x00,0x80,0x35,0xdd,0x00,0x02,0x00,0x01]
// GFX908: error: operands are not valid for this GPU or mode // GFX908: error: operands are not valid for this GPU or mode
// GFX1010: error: instruction not supported on this GPU // GFX1010: error: instruction not supported on this GPU
// GFX90A: global_atomic_pk_add_f16 v0, v[0:1], v2, off glc ; encoding: [0x00,0x80,0x39,0xdd,0x00,0x02,0x7f,0x00] // GFX90A: global_atomic_pk_add_f16 v0, v[0:1], v2, off glc ; encoding: [0x00,0x80,0x39,0xdd,0x00,0x02,0x7f,0x00]
global_atomic_pk_add_f16 v0, v[0:1], v2, off glc global_atomic_pk_add_f16 v0, v[0:1], v2, off glc
// NOT-GFX90A: error: scc modifier is not supported on this GPU
// GFX90A: buffer_atomic_add v4, off, s[8:11], s3 scc ; encoding: [0x00,0x80,0x08,0xe1,0x00,0x04,0x02,0x03]
buffer_atomic_add v4, off, s[8:11], s3 scc

llvm/test/MC/AMDGPU/gfx90a_err.s

Show First 20 Lines • Show All 190 Lines • ▼ Show 20 Lines
image_sample_cd v[0:3], v[0:1], s[4:11], s[16:19] dmask:0xf image_sample_cd v[0:3], v[0:1], s[4:11], s[16:19] dmask:0xf
// GFX90A: error: instruction not supported on this GPU // GFX90A: error: instruction not supported on this GPU
image_sample_b v[0:3], v[0:1], s[4:11], s[16:19] dmask:0xf image_sample_b v[0:3], v[0:1], s[4:11], s[16:19] dmask:0xf
// GFX90A: error: instruction not supported on this GPU // GFX90A: error: instruction not supported on this GPU
global_atomic_add_f32 v0, v[0:1], v2, off glc scc global_atomic_add_f32 v0, v[0:1], v2, off glc scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
global_atomic_add_f32 v[0:1], v2, off scc global_atomic_add_f32 v[0:1], v2, off scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
global_atomic_add_f32 v0, v2, s[0:1] scc global_atomic_add_f32 v0, v2, s[0:1] scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
global_atomic_add_f32 v1, v0, v2, s[0:1] glc scc global_atomic_add_f32 v1, v0, v2, s[0:1] glc scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
global_atomic_pk_add_f16 v0, v[0:1], v2, off glc scc global_atomic_pk_add_f16 v0, v[0:1], v2, off glc scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
flat_atomic_add_f64 v[0:1], v[0:1], v[2:3] glc scc flat_atomic_add_f64 v[0:1], v[0:1], v[2:3] glc scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
flat_atomic_add_f64 v[0:1], v[2:3] scc flat_atomic_add_f64 v[0:1], v[2:3] scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
flat_atomic_min_f64 v[0:1], v[2:3] scc flat_atomic_min_f64 v[0:1], v[2:3] scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
flat_atomic_max_f64 v[0:1], v[2:3] scc flat_atomic_max_f64 v[0:1], v[2:3] scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
global_atomic_add_f64 v[0:1], v[2:3], off scc global_atomic_add_f64 v[0:1], v[2:3], off scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
global_atomic_min_f64 v[0:1], v[2:3], off scc global_atomic_min_f64 v[0:1], v[2:3], off scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
global_atomic_max_f64 v[0:1], v[2:3], off scc global_atomic_max_f64 v[0:1], v[2:3], off scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
buffer_atomic_add_f32 v4, off, s[8:11], s3 scc buffer_atomic_add_f32 v4, off, s[8:11], s3 scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
buffer_atomic_pk_add_f16 v4, off, s[8:11], s3 scc buffer_atomic_pk_add_f16 v4, off, s[8:11], s3 scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
buffer_atomic_add_f64 v[4:5], off, s[8:11], s3 scc buffer_atomic_add_f64 v[4:5], off, s[8:11], s3 scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
buffer_atomic_max_f64 v[4:5], off, s[8:11], s3 scc buffer_atomic_max_f64 v[4:5], off, s[8:11], s3 scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
buffer_atomic_min_f64 v[4:5], off, s[8:11], s3 scc buffer_atomic_min_f64 v[4:5], off, s[8:11], s3 scc
// GFX90A: error: instruction must not use scc // GFX90A: error: scc is not supported on this GPU
v_mov_b32_sdwa v1, src_lds_direct dst_sel:DWORD v_mov_b32_sdwa v1, src_lds_direct dst_sel:DWORD
// GFX90A: error: lds_direct is not supported on this GPU // GFX90A: error: lds_direct is not supported on this GPU
v_add_f32_sdwa v5, v1, lds_direct dst_sel:DWORD v_add_f32_sdwa v5, v1, lds_direct dst_sel:DWORD
// GFX90A: error: lds_direct is not supported on this GPU // GFX90A: error: lds_direct is not supported on this GPU
v_ashrrev_i16 v0, lds_direct, v0 v_ashrrev_i16 v0, lds_direct, v0
// GFX90A: error: lds_direct is not supported on this GPU // GFX90A: error: lds_direct is not supported on this GPU
v_add_f32 v5, v1, lds_direct v_add_f32 v5, v1, lds_direct
// GFX90A: error: lds_direct is not supported on this GPU // GFX90A: error: lds_direct is not supported on this GPU

llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s

// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck --check-prefixes=GFX90A --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck --check-prefixes=GFX90A --implicit-check-not=error: %s
v_add_f64 v[1:2], v[1:2], v[1:2] v_add_f64 v[1:2], v[1:2], v[1:2]
// GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
global_load_dwordx3 v[1:3], v[0:1], off global_load_dwordx3 v[1:3], v[0:1], off
// GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
global_load_dwordx4 v[1:4], v[0:1], off global_load_dwordx4 v[1:4], v[0:1], off
// GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
image_load v[1:5], v2, s[0:7] dmask:0xf unorm scc image_load v[1:5], v2, s[0:7] dmask:0xf unorm
// GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
v_mfma_f32_32x32x8f16 a[0:15], a[1:2], v[0:1], a[0:15] v_mfma_f32_32x32x8f16 a[0:15], a[1:2], v[0:1], a[0:15]
// GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
v_mfma_i32_4x4x4i8 a[1:4], a0, v1, 2 v_mfma_i32_4x4x4i8 a[1:4], a0, v1, 2
// GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[1:16] v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[1:16]
// GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned
v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32]
// GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned // GFX90A: error: invalid register class: vgpr tuples must be 64 bit aligned

llvm/test/MC/Disassembler/AMDGPU/gfx90a_dasm_features.txt

# RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -disassemble -show-encoding %s | FileCheck --check-prefix=GFX908 %s
# RUN: llvm-mc -arch=amdgcn -mcpu=gfx90a -disassemble -show-encoding %s | FileCheck --check-prefix=GFX90A %s # RUN: llvm-mc -arch=amdgcn -mcpu=gfx90a -disassemble -show-encoding %s | FileCheck --check-prefix=GFX90A %s
# GFX90A: v_pk_fma_f32 v[8:9], v[0:1], s[0:1], v[4:5] ; encoding: [0x08,0x40,0xb0,0xd3,0x00,0x01,0x10,0x1c] # GFX90A: v_pk_fma_f32 v[8:9], v[0:1], s[0:1], v[4:5] ; encoding: [0x08,0x40,0xb0,0xd3,0x00,0x01,0x10,0x1c]
0x08,0x40,0xb0,0xd3,0x00,0x01,0x10,0x1c 0x08,0x40,0xb0,0xd3,0x00,0x01,0x10,0x1c
# GFX90A: v_pk_fma_f32 v[8:9], v[0:1], s[0:1], v[4:5] ; encoding: [0x08,0x40,0xb0,0xd3,0x00,0x01,0x10,0x1c] # GFX90A: v_pk_fma_f32 v[8:9], v[0:1], s[0:1], v[4:5] ; encoding: [0x08,0x40,0xb0,0xd3,0x00,0x01,0x10,0x1c]
0x08,0x40,0xb0,0xd3,0x00,0x01,0x10,0x1c 0x08,0x40,0xb0,0xd3,0x00,0x01,0x10,0x1c
▲ Show 20 Lines • Show All 226 Lines • ▼ Show 20 Lines
0x00,0x00,0xb3,0xd3,0x02,0x09,0x02,0x10 0x00,0x00,0xb3,0xd3,0x02,0x09,0x02,0x10
# GFX90A: v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,0] ; encoding: [0x00,0x48,0xb3,0xd3,0x02,0x09,0x02,0x18] # GFX90A: v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,0] ; encoding: [0x00,0x48,0xb3,0xd3,0x02,0x09,0x02,0x18]
0x00,0x08,0xb3,0xd3,0x02,0x09,0x02,0x18 0x00,0x08,0xb3,0xd3,0x02,0x09,0x02,0x18
# GFX90A: v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,1] ; encoding: [0x00,0x58,0xb3,0xd3,0x02,0x09,0x02,0x18] # GFX90A: v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,1] ; encoding: [0x00,0x58,0xb3,0xd3,0x02,0x09,0x02,0x18]
0x00,0x18,0xb3,0xd3,0x02,0x09,0x02,0x18 0x00,0x18,0xb3,0xd3,0x02,0x09,0x02,0x18
# GFX908: tbuffer_load_format_xyzw v[4:7], off, s[0:3], 0 ; encoding: [0x00,0x80,0x09,0xe8,0x00,0x04,0x20,0x80]
# GFX90A: tbuffer_load_format_xyzw v[4:7], off, s[0:3], 0 scc ; encoding: [0x00,0x80,0x09,0xe8,0x00,0x04,0x20,0x80]
0x00,0x80,0x09,0xe8,0x00,0x04,0x20,0x80
# GFX908: tbuffer_load_format_xyzw v[4:7], off, s[0:3], 0 glc ; encoding: [0x00,0xc0,0x09,0xe8,0x00,0x04,0x20,0x80]
# GFX90A: tbuffer_load_format_xyzw v[4:7], off, s[0:3], 0 glc scc ; encoding: [0x00,0xc0,0x09,0xe8,0x00,0x04,0x20,0x80]
0x00,0xc0,0x09,0xe8,0x00,0x04,0x20,0x80
# GFX908: buffer_load_dword v5, off, s[8:11], s3 offset:4095 ; encoding: [0xff,0x8f,0x50,0xe0,0x00,0x05,0x02,0x03]
# GFX90A: buffer_load_dword v5, off, s[8:11], s3 offset:4095 scc ; encoding: [0xff,0x8f,0x50,0xe0,0x00,0x05,0x02,0x03]
0xff,0x8f,0x50,0xe0,0x00,0x05,0x02,0x03
# GFX908: buffer_load_dword v5, off, s[8:11], s3 offset:4095 glc ; encoding: [0xff,0xcf,0x50,0xe0,0x00,0x05,0x02,0x03]
# GFX90A: buffer_load_dword v5, off, s[8:11], s3 offset:4095 glc scc ; encoding: [0xff,0xcf,0x50,0xe0,0x00,0x05,0x02,0x03]
0xff,0xcf,0x50,0xe0,0x00,0x05,0x02,0x03
# GFX90A: buffer_wbl2 ; encoding: [0x00,0x00,0xa0,0xe0,0x00,0x00,0x00,0x00] # GFX90A: buffer_wbl2 ; encoding: [0x00,0x00,0xa0,0xe0,0x00,0x00,0x00,0x00]
0x00,0x00,0xa0,0xe0,0x00,0x00,0x00,0x00 0x00,0x00,0xa0,0xe0,0x00,0x00,0x00,0x00
# GFX90A: buffer_invl2 ; encoding: [0x00,0x00,0xa4,0xe0,0x00,0x00,0x00,0x00] # GFX90A: buffer_invl2 ; encoding: [0x00,0x00,0xa4,0xe0,0x00,0x00,0x00,0x00]
0x00,0x00,0xa4,0xe0,0x00,0x00,0x00,0x00 0x00,0x00,0xa4,0xe0,0x00,0x00,0x00,0x00
# GFX90A: buffer_atomic_add_f64 v[4:5], off, s[8:11], s3 offset:4095 ; encoding: [0xff,0x0f,0x3c,0xe1,0x00,0x04,0x02,0x03] # GFX90A: buffer_atomic_add_f64 v[4:5], off, s[8:11], s3 offset:4095 ; encoding: [0xff,0x0f,0x3c,0xe1,0x00,0x04,0x02,0x03]
0xff,0x0f,0x3c,0xe1,0x00,0x04,0x02,0x03 0xff,0x0f,0x3c,0xe1,0x00,0x04,0x02,0x03
▲ Show 20 Lines • Show All 152 Lines • ▼ Show 20 Lines
0x00,0x00,0xf8,0xd8,0x01,0x02,0x00,0x04 0x00,0x00,0xf8,0xd8,0x01,0x02,0x00,0x04
# GFX90A: ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:4 ; encoding: [0x04,0x00,0xf8,0xd8,0x01,0x02,0x00,0x04] # GFX90A: ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:4 ; encoding: [0x04,0x00,0xf8,0xd8,0x01,0x02,0x00,0x04]
0x04,0x00,0xf8,0xd8,0x01,0x02,0x00,0x04 0x04,0x00,0xf8,0xd8,0x01,0x02,0x00,0x04
# GFX90A: ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:65535 gds ; encoding: [0xff,0xff,0xf9,0xd8,0x01,0x02,0x00,0x04] # GFX90A: ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:65535 gds ; encoding: [0xff,0xff,0xf9,0xd8,0x01,0x02,0x00,0x04]
0xff,0xff,0xf9,0xd8,0x01,0x02,0x00,0x04 0xff,0xff,0xf9,0xd8,0x01,0x02,0x00,0x04
# GFX908: flat_load_dword v0, v[0:1] ; encoding: [0x00,0x00,0x50,0xde,0x00,0x00,0x00,0x00]
# GFX90A: flat_load_dword v0, v[0:1] scc ; encoding: [0x00,0x00,0x50,0xde,0x00,0x00,0x00,0x00]
0x00,0x00,0x50,0xde,0x00,0x00,0x00,0x00
# GFX908: flat_load_dword v0, v[0:1] glc ; encoding: [0x00,0x00,0x51,0xde,0x00,0x00,0x00,0x00]
# GFX90A: flat_load_dword v0, v[0:1] glc scc ; encoding: [0x00,0x00,0x51,0xde,0x00,0x00,0x00,0x00]
0x00,0x00,0x51,0xde,0x00,0x00,0x00,0x00
# GFX90A: flat_atomic_add_f64 v[0:1], v[2:3] offset:4095 ; encoding: [0xff,0x0f,0x3c,0xdd,0x00,0x02,0x00,0x00] # GFX90A: flat_atomic_add_f64 v[0:1], v[2:3] offset:4095 ; encoding: [0xff,0x0f,0x3c,0xdd,0x00,0x02,0x00,0x00]
0xff,0x0f,0x3c,0xdd,0x00,0x02,0x00,0x00 0xff,0x0f,0x3c,0xdd,0x00,0x02,0x00,0x00
# GFX90A: flat_atomic_add_f64 v[0:1], v[0:1], v[2:3] offset:4095 glc ; encoding: [0xff,0x0f,0x3d,0xdd,0x00,0x02,0x00,0x00] # GFX90A: flat_atomic_add_f64 v[0:1], v[0:1], v[2:3] offset:4095 glc ; encoding: [0xff,0x0f,0x3d,0xdd,0x00,0x02,0x00,0x00]
0xff,0x0f,0x3d,0xdd,0x00,0x02,0x00,0x00 0xff,0x0f,0x3d,0xdd,0x00,0x02,0x00,0x00
# GFX90A: flat_atomic_add_f64 v[254:255], v[2:3] offset:4095 ; encoding: [0xff,0x0f,0x3c,0xdd,0xfe,0x02,0x00,0x00] # GFX90A: flat_atomic_add_f64 v[254:255], v[2:3] offset:4095 ; encoding: [0xff,0x0f,0x3c,0xdd,0xfe,0x02,0x00,0x00]
0xff,0x0f,0x3c,0xdd,0xfe,0x02,0x00,0x00 0xff,0x0f,0x3c,0xdd,0xfe,0x02,0x00,0x00
▲ Show 20 Lines • Show All 68 Lines • ▼ Show 20 Lines
0x00,0x80,0x3c,0xdd,0x00,0x02,0x7f,0x00 0x00,0x80,0x3c,0xdd,0x00,0x02,0x7f,0x00
# GFX90A: global_atomic_min_f64 v[0:1], v[2:3], off ; encoding: [0x00,0x80,0x40,0xdd,0x00,0x02,0x7f,0x00] # GFX90A: global_atomic_min_f64 v[0:1], v[2:3], off ; encoding: [0x00,0x80,0x40,0xdd,0x00,0x02,0x7f,0x00]
0x00,0x80,0x40,0xdd,0x00,0x02,0x7f,0x00 0x00,0x80,0x40,0xdd,0x00,0x02,0x7f,0x00
# GFX90A: global_atomic_max_f64 v[0:1], v[2:3], off ; encoding: [0x00,0x80,0x44,0xdd,0x00,0x02,0x7f,0x00] # GFX90A: global_atomic_max_f64 v[0:1], v[2:3], off ; encoding: [0x00,0x80,0x44,0xdd,0x00,0x02,0x7f,0x00]
0x00,0x80,0x44,0xdd,0x00,0x02,0x7f,0x00 0x00,0x80,0x44,0xdd,0x00,0x02,0x7f,0x00
# GFX908: image_load v[0:4], v2, s[0:7] dmask:0xf unorm tfe ; encoding: [0x80,0x1f,0x01,0xf0,0x02,0x00,0x00,0x00]
# GFX90A: image_load a0, v2, s[0:7] dmask:0xf unorm scc ; encoding: [0x80,0x1f,0x01,0xf0,0x02,0x00,0x00,0x00]
0x80,0x1f,0x01,0xf0,0x02,0x00,0x00,0x00
# GFX90A: v_fmac_f64_e32 v[4:5], v[2:3], v[4:5] ; encoding: [0x02,0x09,0x08,0x08] # GFX90A: v_fmac_f64_e32 v[4:5], v[2:3], v[4:5] ; encoding: [0x02,0x09,0x08,0x08]
0x02,0x09,0x08,0x08 0x02,0x09,0x08,0x08
# GFX90A: v_fmac_f64_e32 v[254:255], v[2:3], v[4:5] ; encoding: [0x02,0x09,0xfc,0x09] # GFX90A: v_fmac_f64_e32 v[254:255], v[2:3], v[4:5] ; encoding: [0x02,0x09,0xfc,0x09]
0x02,0x09,0xfc,0x09 0x02,0x09,0xfc,0x09
# GFX90A: v_fmac_f64_e32 v[4:5], v[254:255], v[4:5] ; encoding: [0xfe,0x09,0x08,0x08] # GFX90A: v_fmac_f64_e32 v[4:5], v[254:255], v[4:5] ; encoding: [0xfe,0x09,0x08,0x08]
0xfe,0x09,0x08,0x08 0xfe,0x09,0x08,0x08
▲ Show 20 Lines • Show All 256 Lines • ▼ Show 20 Lines
# GFX90A: flat_atomic_add_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x3d,0xdd,0x00,0x02,0x00,0x00] # GFX90A: flat_atomic_add_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x3d,0xdd,0x00,0x02,0x00,0x00]
0x00,0x00,0x3d,0xdd,0x00,0x02,0x00,0x00 0x00,0x00,0x3d,0xdd,0x00,0x02,0x00,0x00
# GFX90A: flat_atomic_max_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x45,0xdd,0x00,0x02,0x00,0x00] # GFX90A: flat_atomic_max_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x45,0xdd,0x00,0x02,0x00,0x00]
0x00,0x00,0x45,0xdd,0x00,0x02,0x00,0x00 0x00,0x00,0x45,0xdd,0x00,0x02,0x00,0x00
# GFX90A: flat_atomic_min_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x41,0xdd,0x00,0x02,0x00,0x00] # GFX90A: flat_atomic_min_f64 v[0:1], v[0:1], v[2:3] glc ; encoding: [0x00,0x00,0x41,0xdd,0x00,0x02,0x00,0x00]
0x00,0x00,0x41,0xdd,0x00,0x02,0x00,0x00 0x00,0x00,0x41,0xdd,0x00,0x02,0x00,0x00
# GFX90A: buffer_atomic_add v4, off, s[8:11], s3 scc ; encoding: [0x00,0x80,0x08,0xe1,0x00,0x04,0x02,0x03]
0x00,0x80,0x08,0xe1,0x00,0x04,0x02,0x03