[llvm-branch-commits] [llvm] [AMDGPU] Route no-modifier reg-or-inline AsmParser operands through HwMode predicate (PR #221989)
Valery Pykhtin via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Wed Sep 9 23:06:41 PDT 2026
https://github.com/vpykhtin updated https://github.com/llvm/llvm-project/pull/221989
>From f800b8983e4392ac2fafaeda9c62a5cb07fb535e Mon Sep 17 00:00:00 2001
From: Valery Pykhtin <valery.pykhtin at amd.com>
Date: Mon, 31 Aug 2026 18:48:14 +0000
Subject: [PATCH 1/3] [AMDGPU] Route no-modifier reg-or-inline AsmParser
operands through HwMode predicate
Convert the reg-or-inline operands with no modifiers (MFMA VGPR/AGPR
sources, VCSrc, v_pk_mov_b32, VOP scalar f64) from the fixed-class
isRegOrInlineNoMods to the HwMode-aware isRegOrInlineNoModsByHwMode, so an
odd-aligned tuple is rejected at the offending operand column instead of by
the validateVGPRAlign catch-all.
Co-Authored-By: Claude <noreply at anthropic.com>
---
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 99 ++++++++++---------
.../MC/AMDGPU/misaligned-vgpr-tuples-err.s | 14 +--
2 files changed, 60 insertions(+), 53 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index db5dd3a38b2a9..8ebd3ef095dd8 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -434,6 +434,10 @@ class AMDGPUOperand : public MCParsedAsmOperand {
return isRegOrInline(RCID, type) && !hasModifiers();
}
+ bool isRegOrInlineNoModsByHwMode(unsigned RCByHwModeIdx, MVT type) const {
+ return isRegOrInlineByHwMode(RCByHwModeIdx, type) && !hasModifiers();
+ }
+
bool isSCSrcB16() const {
return isRegOrInlineNoMods(AMDGPU::SReg_32RegClassID, MVT::i16);
}
@@ -539,7 +543,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVCSrc_b64() const {
- return isRegOrInlineNoMods(AMDGPU::VS_64RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::i64);
}
bool isVCSrcT_b16() const {
@@ -565,7 +569,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVCSrc_f64() const {
- return isRegOrInlineNoMods(AMDGPU::VS_64RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64);
}
bool isVCSrcTBF16() const {
@@ -650,7 +654,10 @@ class AMDGPUOperand : public MCParsedAsmOperand {
return isVCSrc_f32() || isLiteralImm(MVT::f32) || isExpr();
}
- bool isVSrc_f64() const { return isVCSrc_f64() || isLiteralImm(MVT::f64); }
+ bool isVSrc_f64() const {
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64) ||
+ isLiteralImm(MVT::f64);
+ }
bool isVSrcT_bf16() const {
return isVCSrcTBF16() || isLiteralImm(MVT::bf16);
@@ -709,91 +716,91 @@ class AMDGPUOperand : public MCParsedAsmOperand {
bool isVISrcV2F16() const { return isVISrcF16() || isVISrcB32(); }
bool isVISrc_64_bf16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::bf16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::bf16);
}
bool isVISrc_64_f16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::f16);
}
bool isVISrc_64_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::i32);
}
bool isVISrc_64B64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::i64);
}
bool isVISrc_64_f64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::f64);
}
bool isVISrc_64V2FP32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::f32);
}
bool isVISrc_64V2INT32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::i32);
}
bool isVISrc_256_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::i32);
}
bool isVISrc_256_f32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::f32);
}
bool isVISrc_256B64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::i64);
}
bool isVISrc_256_f64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::f64);
}
bool isVISrc_512_f64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::f64);
}
bool isVISrc_128B16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::i16);
}
bool isVISrc_128V2B16() const { return isVISrc_128B16(); }
bool isVISrc_128_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::i32);
}
bool isVISrc_128_f32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::f32);
}
bool isVISrc_256V2FP32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::f32);
}
bool isVISrc_256V2INT32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::i32);
}
bool isVISrc_512_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::i32);
}
bool isVISrc_512B16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::i16);
}
bool isVISrc_512V2B16() const { return isVISrc_512B16(); }
bool isVISrc_512_f32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::f32);
}
bool isVISrc_512F16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::f16);
}
bool isVISrc_512V2F16() const {
@@ -801,21 +808,21 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVISrc_1024_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_1024RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::i32);
}
bool isVISrc_1024B16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_1024RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::i16);
}
bool isVISrc_1024V2B16() const { return isVISrc_1024B16(); }
bool isVISrc_1024_f32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_1024RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::f32);
}
bool isVISrc_1024F16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_1024RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::f16);
}
bool isVISrc_1024V2F16() const {
@@ -843,29 +850,29 @@ class AMDGPUOperand : public MCParsedAsmOperand {
bool isAISrcV2F16() const { return isAISrcF16() || isAISrcB32(); }
bool isAISrc_64B64() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_64RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_64_AlignTarget, MVT::i64);
}
bool isAISrc_64_f64() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_64RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_64_AlignTarget, MVT::f64);
}
bool isAISrc_128_b32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_128RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::i32);
}
bool isAISrc_128B16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_128RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::i16);
}
bool isAISrc_128V2B16() const { return isAISrc_128B16(); }
bool isAISrc_128_f32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_128RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::f32);
}
bool isAISrc_128F16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_128RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::f16);
}
bool isAISrc_128V2F16() const {
@@ -873,11 +880,11 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVISrc_128_bf16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::bf16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::bf16);
}
bool isVISrc_128_f16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::f16);
}
bool isVISrc_128V2F16() const {
@@ -885,29 +892,29 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isAISrc_256B64() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_256RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_256_AlignTarget, MVT::i64);
}
bool isAISrc_256_f64() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_256RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_256_AlignTarget, MVT::f64);
}
bool isAISrc_512_b32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_512RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::i32);
}
bool isAISrc_512B16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_512RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::i16);
}
bool isAISrc_512V2B16() const { return isAISrc_512B16(); }
bool isAISrc_512_f32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_512RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::f32);
}
bool isAISrc_512F16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_512RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::f16);
}
bool isAISrc_512V2F16() const {
@@ -915,21 +922,21 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isAISrc_1024_b32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_1024RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::i32);
}
bool isAISrc_1024B16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_1024RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::i16);
}
bool isAISrc_1024V2B16() const { return isAISrc_1024B16(); }
bool isAISrc_1024_f32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_1024RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::f32);
}
bool isAISrc_1024F16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_1024RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::f16);
}
bool isAISrc_1024V2F16() const {
diff --git a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
index 9297e98edacf5..3149a14e99b66 100644
--- a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -94,10 +94,10 @@ v_add_f64 v[1:2], v[1:2], v[1:2]
// GFX90A: :[[@LINE-1]]:11: error: invalid operand for instruction
v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[17:32]
-// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: error: invalid register class: vgpr tuples must be 64 bit aligned
v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[33:64]
-// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: 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]
@@ -107,19 +107,19 @@ v_mfma_i32_4x4x4i8 a[1:4], a0, v1, 2
// GFX90A: :[[@LINE-1]]:20: error: invalid operand for instruction
v_mfma_f32_4x4x1f32 v[0:3], v0, v1, v[5:8]
-// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:37: error: invalid register class: vgpr tuples must be 64 bit aligned
v_mfma_f32_16x16x1f32 v[0:15], v0, v1, v[17:32]
-// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: error: invalid register class: vgpr tuples must be 64 bit aligned
v_mfma_f32_32x32x1f32 v[0:31], v0, v1, v[33:64]
-// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: error: invalid register class: vgpr tuples must be 64 bit aligned
v_ceil_f64_dpp v[0:1], v[3:4] row_newbcast:1 row_mask:0xf bank_mask:0xf
// GFX90A: :[[@LINE-1]]:24: error: invalid operand for instruction
v_pk_mov_b32 v[0:1], v[3:4], v[4:5]
-// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:22: error: invalid register class: vgpr tuples must be 64 bit aligned
v_pk_mov_b32 v[0:1], v[2:3], v[5:6]
-// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:30: error: invalid register class: vgpr tuples must be 64 bit aligned
>From ac364412a6876b3901fc086238e7dc88e21b4191 Mon Sep 17 00:00:00 2001
From: Valery Pykhtin <valery.pykhtin at amd.com>
Date: Tue, 8 Sep 2026 17:01:47 +0000
Subject: [PATCH 2/3] [AMDGPU] Update no-modifier operand tests for the dropped
align diagnostic
The no-modifier reg-or-inline operands routed through the HwMode
predicate now report a misaligned tuple as a plain invalid operand,
matching the diagnostic dropped earlier in the stack.
---
llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s | 14 +++++++-------
llvm/test/MC/AMDGPU/out-of-range-registers.s | 2 +-
2 files changed, 8 insertions(+), 8 deletions(-)
diff --git a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
index 3149a14e99b66..a8871740203d4 100644
--- a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -94,10 +94,10 @@ v_add_f64 v[1:2], v[1:2], v[1:2]
// GFX90A: :[[@LINE-1]]:11: error: invalid operand for instruction
v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[17:32]
-// GFX90A: :[[@LINE-1]]:40: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: error: invalid operand for instruction
v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[33:64]
-// GFX90A: :[[@LINE-1]]:40: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: error: invalid operand for instruction
v_mfma_f32_32x32x8f16 a[0:15], a[1:2], v[0:1], a[0:15]
@@ -107,19 +107,19 @@ v_mfma_i32_4x4x4i8 a[1:4], a0, v1, 2
// GFX90A: :[[@LINE-1]]:20: error: invalid operand for instruction
v_mfma_f32_4x4x1f32 v[0:3], v0, v1, v[5:8]
-// GFX90A: :[[@LINE-1]]:37: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:37: error: invalid operand for instruction
v_mfma_f32_16x16x1f32 v[0:15], v0, v1, v[17:32]
-// GFX90A: :[[@LINE-1]]:40: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: error: invalid operand for instruction
v_mfma_f32_32x32x1f32 v[0:31], v0, v1, v[33:64]
-// GFX90A: :[[@LINE-1]]:40: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: error: invalid operand for instruction
v_ceil_f64_dpp v[0:1], v[3:4] row_newbcast:1 row_mask:0xf bank_mask:0xf
// GFX90A: :[[@LINE-1]]:24: error: invalid operand for instruction
v_pk_mov_b32 v[0:1], v[3:4], v[4:5]
-// GFX90A: :[[@LINE-1]]:22: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:22: error: invalid operand for instruction
v_pk_mov_b32 v[0:1], v[2:3], v[5:6]
-// GFX90A: :[[@LINE-1]]:30: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:30: error: invalid operand for instruction
diff --git a/llvm/test/MC/AMDGPU/out-of-range-registers.s b/llvm/test/MC/AMDGPU/out-of-range-registers.s
index 619071da6663d..2f8013660e9bf 100644
--- a/llvm/test/MC/AMDGPU/out-of-range-registers.s
+++ b/llvm/test/MC/AMDGPU/out-of-range-registers.s
@@ -132,7 +132,7 @@ s_mov_b32 tba_lo, 0
v_cvt_f32_f64 v0, v[255:256]
// SICIVI9-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: register index is out of range
// GFX10-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: register index is out of range
-// GFX1250-ERR: :[[@LINE-3]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX1250-ERR: :[[@LINE-3]]:{{[0-9]+}}: error: invalid operand for instruction
v_mqsad_u32_u8 v[254:257], v[0:1], -4.0, v[2:5]
// SI-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
>From 5b699d59d7fa6fd0dedf5f0ce9428ee9c7017c19 Mon Sep 17 00:00:00 2001
From: Valery Pykhtin <valery.pykhtin at amd.com>
Date: Wed, 9 Sep 2026 11:19:58 +0000
Subject: [PATCH 3/3] [AMDGPU] Rename the no-modifier operand predicate to
*Target, NFC
Match the *Target naming used for the other alignment-aware operand
predicates.
---
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 96 +++++++++----------
1 file changed, 48 insertions(+), 48 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index 8ebd3ef095dd8..beaa407f56cb7 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -434,8 +434,8 @@ class AMDGPUOperand : public MCParsedAsmOperand {
return isRegOrInline(RCID, type) && !hasModifiers();
}
- bool isRegOrInlineNoModsByHwMode(unsigned RCByHwModeIdx, MVT type) const {
- return isRegOrInlineByHwMode(RCByHwModeIdx, type) && !hasModifiers();
+ bool isRegOrInlineNoModsTarget(unsigned TargetRCIdx, MVT type) const {
+ return isRegOrInlineTarget(TargetRCIdx, type) && !hasModifiers();
}
bool isSCSrcB16() const {
@@ -543,7 +543,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVCSrc_b64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::i64);
+ return isRegOrInlineNoModsTarget(AMDGPU::VS_64_AlignTarget, MVT::i64);
}
bool isVCSrcT_b16() const {
@@ -569,7 +569,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVCSrc_f64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64);
+ return isRegOrInlineNoModsTarget(AMDGPU::VS_64_AlignTarget, MVT::f64);
}
bool isVCSrcTBF16() const {
@@ -655,7 +655,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVSrc_f64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64) ||
+ return isRegOrInlineNoModsTarget(AMDGPU::VS_64_AlignTarget, MVT::f64) ||
isLiteralImm(MVT::f64);
}
@@ -716,91 +716,91 @@ class AMDGPUOperand : public MCParsedAsmOperand {
bool isVISrcV2F16() const { return isVISrcF16() || isVISrcB32(); }
bool isVISrc_64_bf16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::bf16);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_64_AlignTarget, MVT::bf16);
}
bool isVISrc_64_f16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::f16);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_64_AlignTarget, MVT::f16);
}
bool isVISrc_64_b32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_64_AlignTarget, MVT::i32);
}
bool isVISrc_64B64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::i64);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_64_AlignTarget, MVT::i64);
}
bool isVISrc_64_f64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::f64);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_64_AlignTarget, MVT::f64);
}
bool isVISrc_64V2FP32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::f32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_64_AlignTarget, MVT::f32);
}
bool isVISrc_64V2INT32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_64_AlignTarget, MVT::i32);
}
bool isVISrc_256_b32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_256_AlignTarget, MVT::i32);
}
bool isVISrc_256_f32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::f32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_256_AlignTarget, MVT::f32);
}
bool isVISrc_256B64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::i64);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_256_AlignTarget, MVT::i64);
}
bool isVISrc_256_f64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::f64);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_256_AlignTarget, MVT::f64);
}
bool isVISrc_512_f64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::f64);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_512_AlignTarget, MVT::f64);
}
bool isVISrc_128B16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::i16);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_128_AlignTarget, MVT::i16);
}
bool isVISrc_128V2B16() const { return isVISrc_128B16(); }
bool isVISrc_128_b32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_128_AlignTarget, MVT::i32);
}
bool isVISrc_128_f32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::f32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_128_AlignTarget, MVT::f32);
}
bool isVISrc_256V2FP32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::f32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_256_AlignTarget, MVT::f32);
}
bool isVISrc_256V2INT32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_256_AlignTarget, MVT::i32);
}
bool isVISrc_512_b32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_512_AlignTarget, MVT::i32);
}
bool isVISrc_512B16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::i16);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_512_AlignTarget, MVT::i16);
}
bool isVISrc_512V2B16() const { return isVISrc_512B16(); }
bool isVISrc_512_f32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::f32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_512_AlignTarget, MVT::f32);
}
bool isVISrc_512F16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::f16);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_512_AlignTarget, MVT::f16);
}
bool isVISrc_512V2F16() const {
@@ -808,21 +808,21 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVISrc_1024_b32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_1024_AlignTarget, MVT::i32);
}
bool isVISrc_1024B16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::i16);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_1024_AlignTarget, MVT::i16);
}
bool isVISrc_1024V2B16() const { return isVISrc_1024B16(); }
bool isVISrc_1024_f32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::f32);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_1024_AlignTarget, MVT::f32);
}
bool isVISrc_1024F16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::f16);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_1024_AlignTarget, MVT::f16);
}
bool isVISrc_1024V2F16() const {
@@ -850,29 +850,29 @@ class AMDGPUOperand : public MCParsedAsmOperand {
bool isAISrcV2F16() const { return isAISrcF16() || isAISrcB32(); }
bool isAISrc_64B64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_64_AlignTarget, MVT::i64);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_64_AlignTarget, MVT::i64);
}
bool isAISrc_64_f64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_64_AlignTarget, MVT::f64);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_64_AlignTarget, MVT::f64);
}
bool isAISrc_128_b32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_128_AlignTarget, MVT::i32);
}
bool isAISrc_128B16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::i16);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_128_AlignTarget, MVT::i16);
}
bool isAISrc_128V2B16() const { return isAISrc_128B16(); }
bool isAISrc_128_f32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::f32);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_128_AlignTarget, MVT::f32);
}
bool isAISrc_128F16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::f16);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_128_AlignTarget, MVT::f16);
}
bool isAISrc_128V2F16() const {
@@ -880,11 +880,11 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVISrc_128_bf16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::bf16);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_128_AlignTarget, MVT::bf16);
}
bool isVISrc_128_f16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::f16);
+ return isRegOrInlineNoModsTarget(AMDGPU::VReg_128_AlignTarget, MVT::f16);
}
bool isVISrc_128V2F16() const {
@@ -892,29 +892,29 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isAISrc_256B64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_256_AlignTarget, MVT::i64);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_256_AlignTarget, MVT::i64);
}
bool isAISrc_256_f64() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_256_AlignTarget, MVT::f64);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_256_AlignTarget, MVT::f64);
}
bool isAISrc_512_b32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_512_AlignTarget, MVT::i32);
}
bool isAISrc_512B16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::i16);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_512_AlignTarget, MVT::i16);
}
bool isAISrc_512V2B16() const { return isAISrc_512B16(); }
bool isAISrc_512_f32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::f32);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_512_AlignTarget, MVT::f32);
}
bool isAISrc_512F16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::f16);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_512_AlignTarget, MVT::f16);
}
bool isAISrc_512V2F16() const {
@@ -922,21 +922,21 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isAISrc_1024_b32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::i32);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_1024_AlignTarget, MVT::i32);
}
bool isAISrc_1024B16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::i16);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_1024_AlignTarget, MVT::i16);
}
bool isAISrc_1024V2B16() const { return isAISrc_1024B16(); }
bool isAISrc_1024_f32() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::f32);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_1024_AlignTarget, MVT::f32);
}
bool isAISrc_1024F16() const {
- return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::f16);
+ return isRegOrInlineNoModsTarget(AMDGPU::AReg_1024_AlignTarget, MVT::f16);
}
bool isAISrc_1024V2F16() const {
More information about the llvm-branch-commits
mailing list