https://github.com/vpykhtin created https://github.com/llvm/llvm-project/pull/221989
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 <[email protected]> >From a943ffe28b415f2cb4f73c5812da4093f35f7d2b Mon Sep 17 00:00:00 2001 From: Valery Pykhtin <[email protected]> Date: Mon, 31 Aug 2026 18:48:14 +0000 Subject: [PATCH] [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 <[email protected]> --- .../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 19b08b4b77a89..db024282e4f9d 100644 --- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp +++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp @@ -449,6 +449,10 @@ class AMDGPUOperand : public MCParsedAsmOperand { bool isRegOrInlineNoMods(unsigned RCID, MVT type) const { 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); @@ -553,7 +557,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 { @@ -579,7 +583,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 { @@ -664,7 +668,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); @@ -723,91 +730,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 { @@ -815,21 +822,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 { @@ -857,29 +864,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 { @@ -887,11 +894,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 { @@ -899,29 +906,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 { @@ -929,21 +936,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 ac9a1c7c65ba8..07b196c4b3542 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 register class: vgpr tuples must be 64 bit aligned 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 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
