[llvm] [AMDGPU] Exclude CDNA parts from POPS exiting wave id pattern (PR #210892)
Arseniy Obolenskiy via llvm-commits
llvm-commits at lists.llvm.org
Tue Jul 21 01:51:12 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/210892
>From be72fc9dbfa6a1279429b61769dffa8609981b4c Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 21 Jul 2026 09:12:41 +0200
Subject: [PATCH 1/4] [AMDGPU] Exclude CDNA parts from POPS exiting wave id
pattern
POPS hardware is graphics-pipe-only and absent on compute-only CDNA targets (gfx908, gfx90a, gfx940, gfx942, gfx950), which incorrectly matched the isGFX9GFX10 predicate and selected a nonexistent register
---
llvm/lib/Target/AMDGPU/SOPInstructions.td | 2 +-
.../AMDGPU/llvm.amdgcn.pops.exiting.wave.id.ll | 16 ++++++++++++++++
2 files changed, 17 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index 8fa54de570e67..f0bfd431922ca 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -2042,7 +2042,7 @@ let SubtargetPredicate = isNotGFX9Plus in {
def : GetFPModePat<fpmode_mask_gfx6plus>;
}
-let SubtargetPredicate = isGFX9GFX10 in
+let SubtargetPredicate = isGFX9GFX10, OtherPredicates = [NotHasMAIInsts] in
def : GCNPat<
(int_amdgcn_pops_exiting_wave_id),
(S_MOV_B32_sideeffects (i32 SRC_POPS_EXITING_WAVE_ID))
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.pops.exiting.wave.id.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.pops.exiting.wave.id.ll
index ba343a4eb108e..23989f7deaee1 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.pops.exiting.wave.id.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.pops.exiting.wave.id.ll
@@ -3,6 +3,22 @@
; RUN: llc -global-isel=1 -mtriple=amdgpu9.00 < %s | FileCheck %s -check-prefixes=GFX9,GFX9-GISEL
; RUN: llc -global-isel=0 -mtriple=amdgpu10.10 < %s | FileCheck %s -check-prefixes=GFX10,SDAG,GFX10-SDAG
; RUN: llc -global-isel=1 -mtriple=amdgpu10.10 < %s | FileCheck %s -check-prefixes=GFX10,GFX10-GISEL
+; RUN: not --crash llc -global-isel=0 -mtriple=amdgpu9.08 < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+; RUN: not llc -global-isel=1 -mtriple=amdgpu9.08 < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+; RUN: not --crash llc -global-isel=0 -mtriple=amdgpu9.0a < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+; RUN: not llc -global-isel=1 -mtriple=amdgpu9.0a < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+; RUN: not --crash llc -global-isel=0 -mtriple=amdgpu9.4 < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+; RUN: not llc -global-isel=1 -mtriple=amdgpu9.4 < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+; RUN: not --crash llc -global-isel=0 -mtriple=amdgpu9.42 < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+; RUN: not llc -global-isel=1 -mtriple=amdgpu9.42 < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+; RUN: not --crash llc -global-isel=0 -mtriple=amdgpu9.50 < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+; RUN: not llc -global-isel=1 -mtriple=amdgpu9.50 < %s 2>&1 | FileCheck %s -check-prefix=NOPOPS
+
+; POPS (Primitive Ordered Pixel Shading) is a graphics-pipe-only HW feature,
+; absent on all compute-only CDNA parts (gfx908, gfx90a, gfx940, gfx942,
+; gfx950). SRC_POPS_EXITING_WAVE_ID is not a valid SGPR source on those
+; targets.
+; NOPOPS: LLVM ERROR: {{[Cc]}}annot select: {{.*}}pops.exiting.wave.id
declare void @foo(i32)
>From e14537a91abf8672e0256c3d96ae0feba7b37711 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 21 Jul 2026 09:41:38 +0200
Subject: [PATCH 2/4] Address comment
---
llvm/lib/Target/AMDGPU/AMDGPU.td | 7 +++++++
llvm/lib/Target/AMDGPU/SOPInstructions.td | 2 +-
2 files changed, 8 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index ed454806d7193..cb6f5da727092 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -2653,6 +2653,13 @@ def isGFX9GFX10 :
"Subtarget->getGeneration() == AMDGPUSubtarget::GFX10">,
AssemblerPredicate<(all_of FeatureGFX9Insts, (not FeatureGFX11Insts))>;
+def HasPopsExitingWaveID :
+ Predicate<"(Subtarget->getGeneration() == AMDGPUSubtarget::GFX9 ||"
+ " Subtarget->getGeneration() == AMDGPUSubtarget::GFX10) &&"
+ "!Subtarget->hasMAIInsts()">,
+ AssemblerPredicate<(all_of FeatureGFX9Insts, (not FeatureGFX11Insts),
+ (not FeatureMAIInsts))>;
+
def isGFX9GFX10GFX11 :
Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::GFX9 &&"
"Subtarget->getGeneration() < AMDGPUSubtarget::GFX12">,
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index f0bfd431922ca..d18e96b7621b1 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -2042,7 +2042,7 @@ let SubtargetPredicate = isNotGFX9Plus in {
def : GetFPModePat<fpmode_mask_gfx6plus>;
}
-let SubtargetPredicate = isGFX9GFX10, OtherPredicates = [NotHasMAIInsts] in
+let SubtargetPredicate = HasPopsExitingWaveID in
def : GCNPat<
(int_amdgcn_pops_exiting_wave_id),
(S_MOV_B32_sideeffects (i32 SRC_POPS_EXITING_WAVE_ID))
>From 2c7ad7e8a946a94096dd7acf75ff52d8f9dbdd46 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 21 Jul 2026 10:25:26 +0200
Subject: [PATCH 3/4] Adjust AsmParser
---
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 2 +-
llvm/test/MC/AMDGPU/literals.s | 3 +++
llvm/test/MC/AMDGPU/mai.s | 2 +-
3 files changed, 5 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index 7ead3d6f2b263..1068b500078de 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -7197,7 +7197,7 @@ bool AMDGPUAsmParser::subtargetHasRegister(const MCRegisterInfo &MRI,
case SRC_FLAT_SCRATCH_BASE_HI:
return hasGloballyAddressableScratch();
case SRC_POPS_EXITING_WAVE_ID:
- return isGFX9Plus() && !isGFX11Plus();
+ return isGFX9Plus() && !isGFX11Plus() && !hasMAIInsts(getSTI());
case TBA:
case TBA_LO:
case TBA_HI:
diff --git a/llvm/test/MC/AMDGPU/literals.s b/llvm/test/MC/AMDGPU/literals.s
index 54399319813b6..8bcb6f3514bee 100644
--- a/llvm/test/MC/AMDGPU/literals.s
+++ b/llvm/test/MC/AMDGPU/literals.s
@@ -7,6 +7,7 @@
// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1200 -show-encoding %s | FileCheck %s --check-prefixes=GFX8PLUS,GFX12XX,GFX12
// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -show-encoding %s | FileCheck %s --check-prefixes=GFX8PLUS,GFX12XX,GFX1250,GFX1250-ASM
// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -show-encoding %s | %extract-encodings | llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -disassemble -show-encoding | FileCheck %s --check-prefixes=GFX8PLUS,GFX12XX,GFX1250,GFX1250-DIS
+// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx90a -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOGFX90A
// RUN: not llvm-mc -triple=amdgcn -mcpu=tahiti %s -filetype=null 2>&1 | FileCheck %s --check-prefixes=NOGCN,NOSICI,NOSI --implicit-check-not=error:
// RUN: not llvm-mc -triple=amdgcn -mcpu=bonaire %s -filetype=null 2>&1 | FileCheck %s --check-prefixes=NOGCN,NOSICI,NOCI --implicit-check-not=error:
@@ -1420,6 +1421,7 @@ s_add_i32 s0, src_pops_exiting_wave_id, s0
// NOGFX1250: :[[@LINE-4]]:15: error: src_pops_exiting_wave_id register not available on this GPU
// NOSICI: :[[@LINE-5]]:15: error: src_pops_exiting_wave_id register not available on this GPU
// NOVI: :[[@LINE-6]]:15: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX90A: :[[@LINE-7]]:15: error: src_pops_exiting_wave_id register not available on this GPU
s_and_b64 s[0:1], s[0:1], src_shared_base
// GFX11: s_and_b64 s[0:1], s[0:1], src_shared_base ; encoding: [0x00,0xeb,0x80,0x8b]
@@ -1456,6 +1458,7 @@ s_and_b64 s[0:1], s[0:1], src_pops_exiting_wave_id
// NOGFX1250: :[[@LINE-4]]:27: error: src_pops_exiting_wave_id register not available on this GPU
// NOSICI: :[[@LINE-5]]:27: error: src_pops_exiting_wave_id register not available on this GPU
// NOVI: :[[@LINE-6]]:27: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX90A: :[[@LINE-7]]:27: error: src_pops_exiting_wave_id register not available on this GPU
v_add_u16 v0, src_shared_base, v0
// GFX9: v_add_u16_e32 v0, src_shared_base, v0 ; encoding: [0xeb,0x00,0x00,0x4c]
diff --git a/llvm/test/MC/AMDGPU/mai.s b/llvm/test/MC/AMDGPU/mai.s
index afaa1f4d5be36..290d91a00ab5f 100644
--- a/llvm/test/MC/AMDGPU/mai.s
+++ b/llvm/test/MC/AMDGPU/mai.s
@@ -41,7 +41,7 @@ v_accvgpr_write a2, shared_base
// NOGFX908: :[[@LINE-1]]:{{[0-9]+}}: error: source operand must be either a VGPR or an inline constant
v_accvgpr_write a2, pops_exiting_wave_id
-// NOGFX908: :[[@LINE-1]]:{{[0-9]+}}: error: source operand must be either a VGPR or an inline constant
+// NOGFX908: :[[@LINE-1]]:{{[0-9]+}}: error: src_pops_exiting_wave_id register not available on this GPU
v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]
// GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x02,0x04]
>From 7ae5a04d3a6cdecca2d2805cee6f8e4a5e7476fa Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 21 Jul 2026 10:50:59 +0200
Subject: [PATCH 4/4] Address comments
---
.../Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 2 +-
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp | 4 ++++
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h | 1 +
llvm/test/MC/AMDGPU/literals.s | 14 +++++++++++++-
4 files changed, 19 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index 1068b500078de..e3e18eb3e10ee 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -7197,7 +7197,7 @@ bool AMDGPUAsmParser::subtargetHasRegister(const MCRegisterInfo &MRI,
case SRC_FLAT_SCRATCH_BASE_HI:
return hasGloballyAddressableScratch();
case SRC_POPS_EXITING_WAVE_ID:
- return isGFX9Plus() && !isGFX11Plus() && !hasMAIInsts(getSTI());
+ return hasPopsExitingWaveID(getSTI());
case TBA:
case TBA_LO:
case TBA_HI:
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index e5d9c4237c2b4..2bcc17b3fb986 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -2576,6 +2576,10 @@ bool isGFX9Plus(const MCSubtargetInfo &STI) {
bool isNotGFX9Plus(const MCSubtargetInfo &STI) { return !isGFX9Plus(STI); }
+bool hasPopsExitingWaveID(const MCSubtargetInfo &STI) {
+ return isGFX9Plus(STI) && !isGFX11Plus(STI) && !hasMAIInsts(STI);
+}
+
bool isGFX10(const MCSubtargetInfo &STI) {
return STI.hasFeature(AMDGPU::FeatureGFX10);
}
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
index 25fc7d9e34744..e65134ad16b34 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
@@ -1533,6 +1533,7 @@ bool isGFX90A(const MCSubtargetInfo &STI);
bool isGFX940(const MCSubtargetInfo &STI);
bool hasArchitectedFlatScratch(const MCSubtargetInfo &STI);
bool hasMAIInsts(const MCSubtargetInfo &STI);
+bool hasPopsExitingWaveID(const MCSubtargetInfo &STI);
bool hasVOPD(const MCSubtargetInfo &STI);
bool hasDPPSrc1SGPR(const MCSubtargetInfo &STI);
diff --git a/llvm/test/MC/AMDGPU/literals.s b/llvm/test/MC/AMDGPU/literals.s
index 8bcb6f3514bee..89a160f6cfbb3 100644
--- a/llvm/test/MC/AMDGPU/literals.s
+++ b/llvm/test/MC/AMDGPU/literals.s
@@ -7,7 +7,11 @@
// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1200 -show-encoding %s | FileCheck %s --check-prefixes=GFX8PLUS,GFX12XX,GFX12
// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -show-encoding %s | FileCheck %s --check-prefixes=GFX8PLUS,GFX12XX,GFX1250,GFX1250-ASM
// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -show-encoding %s | %extract-encodings | llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -disassemble -show-encoding | FileCheck %s --check-prefixes=GFX8PLUS,GFX12XX,GFX1250,GFX1250-DIS
-// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx90a -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOGFX90A
+// RUN: not llvm-mc -triple=amdgpu9.0a -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOGFX90A
+// RUN: not llvm-mc -triple=amdgpu9.08 -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOGFX908
+// RUN: not llvm-mc -triple=amdgpu9.4 -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOGFX940
+// RUN: not llvm-mc -triple=amdgpu9.42 -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOGFX942
+// RUN: not llvm-mc -triple=amdgpu9.50 -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOGFX950
// RUN: not llvm-mc -triple=amdgcn -mcpu=tahiti %s -filetype=null 2>&1 | FileCheck %s --check-prefixes=NOGCN,NOSICI,NOSI --implicit-check-not=error:
// RUN: not llvm-mc -triple=amdgcn -mcpu=bonaire %s -filetype=null 2>&1 | FileCheck %s --check-prefixes=NOGCN,NOSICI,NOCI --implicit-check-not=error:
@@ -1422,6 +1426,10 @@ s_add_i32 s0, src_pops_exiting_wave_id, s0
// NOSICI: :[[@LINE-5]]:15: error: src_pops_exiting_wave_id register not available on this GPU
// NOVI: :[[@LINE-6]]:15: error: src_pops_exiting_wave_id register not available on this GPU
// NOGFX90A: :[[@LINE-7]]:15: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX908: :[[@LINE-8]]:15: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX940: :[[@LINE-9]]:15: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX942: :[[@LINE-10]]:15: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX950: :[[@LINE-11]]:15: error: src_pops_exiting_wave_id register not available on this GPU
s_and_b64 s[0:1], s[0:1], src_shared_base
// GFX11: s_and_b64 s[0:1], s[0:1], src_shared_base ; encoding: [0x00,0xeb,0x80,0x8b]
@@ -1459,6 +1467,10 @@ s_and_b64 s[0:1], s[0:1], src_pops_exiting_wave_id
// NOSICI: :[[@LINE-5]]:27: error: src_pops_exiting_wave_id register not available on this GPU
// NOVI: :[[@LINE-6]]:27: error: src_pops_exiting_wave_id register not available on this GPU
// NOGFX90A: :[[@LINE-7]]:27: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX908: :[[@LINE-8]]:27: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX940: :[[@LINE-9]]:27: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX942: :[[@LINE-10]]:27: error: src_pops_exiting_wave_id register not available on this GPU
+// NOGFX950: :[[@LINE-11]]:27: error: src_pops_exiting_wave_id register not available on this GPU
v_add_u16 v0, src_shared_base, v0
// GFX9: v_add_u16_e32 v0, src_shared_base, v0 ; encoding: [0xeb,0x00,0x00,0x4c]
More information about the llvm-commits
mailing list