[llvm] 37878de - Disable use of SCC bit from asm
Stanislav Mekhanoshin via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 7 15:44:46 PDT 2021
Author: Stanislav Mekhanoshin
Date: 2021-04-07T15:32:17-07:00
New Revision: 37878de5036718481e13d5067a17d65eb85c3388
URL: https://github.com/llvm/llvm-project/commit/37878de5036718481e13d5067a17d65eb85c3388
DIFF: https://github.com/llvm/llvm-project/commit/37878de5036718481e13d5067a17d65eb85c3388.diff
LOG: Disable use of SCC bit from asm
Differential Revision: https://reviews.llvm.org/D100069
Added:
Modified:
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
llvm/test/MC/AMDGPU/gfx90a_asm_features.s
llvm/test/MC/AMDGPU/gfx90a_err.s
llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
llvm/test/MC/Disassembler/AMDGPU/gfx90a_dasm_features.txt
Removed:
################################################################################
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index b2a79ca2ce7b..ef17b64278c9 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -4092,6 +4092,14 @@ bool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst,
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 @@ bool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst,
}
}
- 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;
}
diff --git a/llvm/test/MC/AMDGPU/gfx90a_asm_features.s b/llvm/test/MC/AMDGPU/gfx90a_asm_features.s
index 14e6f7ef0003..026fa2521cc1 100644
--- a/llvm/test/MC/AMDGPU/gfx90a_asm_features.s
+++ b/llvm/test/MC/AMDGPU/gfx90a_asm_features.s
@@ -322,24 +322,6 @@ v_pk_mov_b32 v[0:1], v[2:3], v[4:5] op_sel:[1,0]
// 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 @@ ds_add_rtn_f64 v[4:5], v1, v[2:3] offset:4
// 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 @@ global_atomic_min_f64 v[0:1], v[2:3], off
// 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 @@ flat_atomic_max_f64 v[0:1], v[0:1], v[2:3] glc
// 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 @@ global_atomic_add_f32 v1, v0, v2, s[0:1] glc ; encoding: [0x00,0x80,0x35,0xdd,0x
// 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
diff --git a/llvm/test/MC/AMDGPU/gfx90a_err.s b/llvm/test/MC/AMDGPU/gfx90a_err.s
index 44c48595ca17..31947510f873 100644
--- a/llvm/test/MC/AMDGPU/gfx90a_err.s
+++ b/llvm/test/MC/AMDGPU/gfx90a_err.s
@@ -196,55 +196,55 @@ 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
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
diff --git a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
index 0370e30a8c08..d2cd4383bb41 100644
--- a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -9,7 +9,7 @@ global_load_dwordx3 v[1:3], 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
-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]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx90a_dasm_features.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx90a_dasm_features.txt
index ec5dbbe19535..adf15e899468 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx90a_dasm_features.txt
+++ b/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
More information about the llvm-commits
mailing list