[llvm] b65d79f - [AMDGPU] Exclude CDNA parts from POPS exiting wave id pattern (#210892)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 29 21:15:10 PDT 2026
Author: Arseniy Obolenskiy
Date: 2026-07-30T06:15:05+02:00
New Revision: b65d79f8800e3d15d1f6cbf1a52f5426da992b59
URL: https://github.com/llvm/llvm-project/commit/b65d79f8800e3d15d1f6cbf1a52f5426da992b59
DIFF: https://github.com/llvm/llvm-project/commit/b65d79f8800e3d15d1f6cbf1a52f5426da992b59.diff
LOG: [AMDGPU] Exclude CDNA parts from POPS exiting wave id pattern (#210892)
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
Added:
Modified:
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/lib/Target/AMDGPU/AMDGPU.td
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
llvm/lib/Target/AMDGPU/SOPInstructions.td
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.pops.exiting.wave.id.ll
llvm/test/MC/AMDGPU/literals.s
llvm/test/MC/AMDGPU/mai.s
Removed:
################################################################################
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 565637b36131c..33203de2d19a3 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2871,6 +2871,7 @@ def int_amdgcn_global_load_lds : AMDGPUGlobalLoadLDS, ClangBuiltin<"__builtin_am
def int_amdgcn_global_load_async_lds : AMDGPUGlobalLoadLDS, ClangBuiltin<"__builtin_amdgcn_global_load_async_lds">;
// This is IntrHasSideEffects because it reads from a volatile hardware register.
+let TargetFeatures = "pops-exiting-wave-id" in
def int_amdgcn_pops_exiting_wave_id :
DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrHasSideEffects]>;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 09db1a9f3f009..8eba93af714dc 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -1113,6 +1113,10 @@ defm GetWaveIdInst : AMDGPUSubtargetFeature<"get-wave-id-inst",
"Has s_get_waveid_in_workgroup instruction"
>;
+defm PopsExitingWaveID : AMDGPUSubtargetFeature<"pops-exiting-wave-id",
+ "Has POPS_EXITING_WAVE_ID hardware register"
+>;
+
defm SMemTimeInst : AMDGPUSubtargetFeature<"s-memtime-inst",
"Has s_memtime instruction"
>;
@@ -1789,7 +1793,8 @@ def FeatureISAVersion9_0_Consumer_Common : FeatureSet<
[FeatureImageGather4D16Bug,
FeatureDsSrc2Insts,
FeatureExtendedImageInsts,
- FeatureGDS])>;
+ FeatureGDS,
+ FeaturePopsExitingWaveID])>;
def FeatureISAVersion9_Generic : FeatureSet<
!listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
@@ -1961,7 +1966,8 @@ def FeatureISAVersion10_Common : FeatureSet<
FeatureBackOffBarrier,
FeatureMsadInsts,
FeatureMqsadPkInsts,
- FeatureMqsadInsts]>;
+ FeatureMqsadInsts,
+ FeaturePopsExitingWaveID]>;
def FeatureISAVersion10_1_Common : FeatureSet<
!listconcat(FeatureISAVersion10_Common.Features,
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index d3027d831aaf1..27139491c2d26 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -7081,7 +7081,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 hasPopsExitingWaveID(getSTI());
case TBA:
case TBA_LO:
case TBA_HI:
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index 8fa54de570e67..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 in
+let SubtargetPredicate = HasPopsExitingWaveID in
def : GCNPat<
(int_amdgcn_pops_exiting_wave_id),
(S_MOV_B32_sideeffects (i32 SRC_POPS_EXITING_WAVE_ID))
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index 25b36d99960a1..cf1fa87d64f8e 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -2558,6 +2558,10 @@ bool isGFX9Plus(const MCSubtargetInfo &STI) {
bool isNotGFX9Plus(const MCSubtargetInfo &STI) { return !isGFX9Plus(STI); }
+bool hasPopsExitingWaveID(const MCSubtargetInfo &STI) {
+ return STI.hasFeature(AMDGPU::FeaturePopsExitingWaveID);
+}
+
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 c06ec8e5b8ab7..4f435948e1b8a 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
@@ -1524,6 +1524,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/CodeGen/AMDGPU/llvm.amdgcn.pops.exiting.wave.id.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.pops.exiting.wave.id.ll
index ba343a4eb108e..2edde15abe18c 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,21 @@
; 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 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 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 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 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 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).
+; NOPOPS: error: {{.*}} llvm.amdgcn.pops.exiting.wave.id requires target feature
declare void @foo(i32)
diff --git a/llvm/test/MC/AMDGPU/literals.s b/llvm/test/MC/AMDGPU/literals.s
index 4cf69ccac31b7..391fb93c55926 100644
--- a/llvm/test/MC/AMDGPU/literals.s
+++ b/llvm/test/MC/AMDGPU/literals.s
@@ -7,6 +7,11 @@
// RUN: not llvm-mc -triple=amdgpu12.00 -show-encoding %s | FileCheck %s --check-prefixes=GFX8PLUS,GFX12XX,GFX12
// RUN: not llvm-mc -triple=amdgpu12.50 -mattr=+real-true16 -show-encoding %s | FileCheck %s --check-prefixes=GFX8PLUS,GFX12XX,GFX1250,GFX1250-ASM
// RUN: not llvm-mc -triple=amdgpu12.50 -mattr=+real-true16 -show-encoding %s | %extract-encodings | llvm-mc -triple=amdgpu12.50 -mattr=+real-true16 -disassemble -show-encoding | FileCheck %s --check-prefixes=GFX8PLUS,GFX12XX,GFX1250,GFX1250-DIS
+// 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=amdgpu6.00 %s -filetype=null 2>&1 | FileCheck %s --check-prefixes=NOGCN,NOSICI,NOSI --implicit-check-not=error:
// RUN: not llvm-mc -triple=amdgpu7.04 %s -filetype=null 2>&1 | FileCheck %s --check-prefixes=NOGCN,NOSICI,NOCI --implicit-check-not=error:
@@ -1420,6 +1425,11 @@ 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
+// 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]
@@ -1456,6 +1466,11 @@ 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
+// 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]
diff --git a/llvm/test/MC/AMDGPU/mai.s b/llvm/test/MC/AMDGPU/mai.s
index a4f09d7cb4b6b..53e240042d1cf 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]
More information about the llvm-commits
mailing list