Index: llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp =================================================================== --- llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp +++ llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp @@ -4092,6 +4092,14 @@ 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")]); + Error(S, "scc is not supported on this GPU"); + return false; + } + if (!(TSFlags & (SIInstrFlags::IsAtomicNoRet | SIInstrFlags::IsAtomicRet))) return true; @@ -4110,14 +4118,6 @@ } } - 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; } Index: llvm/test/MC/AMDGPU/gfx90a_asm_features.s =================================================================== --- llvm/test/MC/AMDGPU/gfx90a_asm_features.s +++ llvm/test/MC/AMDGPU/gfx90a_asm_features.s @@ -322,24 +322,6 @@ // 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] -// 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 // GFX90A: buffer_wbl2 ; encoding: [0x00,0x00,0xa0,0xe0,0x00,0x00,0x00,0x00] buffer_wbl2 @@ -564,14 +546,6 @@ // 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 -// 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 // 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 @@ -680,10 +654,6 @@ // 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 -// 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 // 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] @@ -1020,10 +990,6 @@ // NOT-GFX90A: error: instruction not supported on this GPU 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] // GFX908: error: operands are not valid for this GPU or mode // GFX1010: error: instruction not supported on this GPU @@ -1046,7 +1012,3 @@ // 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] 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 Index: llvm/test/MC/AMDGPU/gfx90a_err.s =================================================================== --- llvm/test/MC/AMDGPU/gfx90a_err.s +++ llvm/test/MC/AMDGPU/gfx90a_err.s @@ -196,55 +196,55 @@ // GFX90A: error: instruction not supported on this GPU 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 // GFX90A: error: lds_direct is not supported on this GPU Index: llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s =================================================================== --- llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s +++ llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s @@ -9,7 +9,7 @@ global_load_dwordx4 v[1:4], v[0:1], off // 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 v_mfma_f32_32x32x8f16 a[0:15], a[1:2], v[0:1], a[0:15] Index: llvm/test/MC/Disassembler/AMDGPU/gfx90a_dasm_features.txt =================================================================== --- llvm/test/MC/Disassembler/AMDGPU/gfx90a_dasm_features.txt +++ llvm/test/MC/Disassembler/AMDGPU/gfx90a_dasm_features.txt @@ -1,4 +1,3 @@ -# 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 # 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] @@ -241,22 +240,6 @@ # 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 -# 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] 0x00,0x00,0xa0,0xe0,0x00,0x00,0x00,0x00 @@ -425,14 +408,6 @@ # 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 -# 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] 0xff,0x0f,0x3c,0xdd,0x00,0x02,0x00,0x00 @@ -517,10 +492,6 @@ # 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 -# 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] 0x02,0x09,0x08,0x08 @@ -793,6 +764,3 @@ # 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 - -# 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