https://github.com/vpykhtin updated https://github.com/llvm/llvm-project/pull/221989
>From 7ff1ca05d4e6ae271aff2f0e1d1ca98d6f97f904 Mon Sep 17 00:00:00 2001 From: Valery Pykhtin <[email protected]> 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 <[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 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 fdd26915b3709662a5d4bd66105bfe3998f358dc Mon Sep 17 00:00:00 2001 From: Valery Pykhtin <[email protected]> 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 afb42e8b441f59486a7bcd323685a86995eca6d8 Mon Sep 17 00:00:00 2001 From: Valery Pykhtin <[email protected]> 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 { _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
