https://github.com/changpeng created
https://github.com/llvm/llvm-project/pull/237115
Introduce FeatureLoadMonitorInsts ("load-monitor-insts") and use it to gate the
flat_load_monitor_b* and global_load_monitor_b* pseudos and their selection
patterns, instead of keying off isGFX125xOnly directly.
The feature is added to FeatureISAVersion12_50_Common, which covers exactly the
same set of targets as isGFX125xOnly (gfx1250, gfx1250s, gfx1251 and
gfx12.5-generic), so there is no functional change.
The feature is also made frontend visible, and the corresponding
__builtin_amdgcn_{flat,global}_load_monitor_b* builtins now require
"load-monitor-insts" rather than "gfx1250-insts".
>From f8fa405823d373f685d42320d1cf35ab03938e91 Mon Sep 17 00:00:00 2001
From: Changpeng Fang <[email protected]>
Date: Thu, 29 Oct 2026 23:50:24 -0700
Subject: [PATCH] [AMDGPU][NFC] Add a target feature for load monitor
instructions
Introduce FeatureLoadMonitorInsts ("load-monitor-insts") and use it to
gate the flat_load_monitor_b* and global_load_monitor_b* pseudos and
their selection patterns, instead of keying off isGFX125xOnly directly.
The feature is added to FeatureISAVersion12_50_Common, which covers
exactly the same set of targets as isGFX125xOnly (gfx1250, gfx1250s,
gfx1251 and gfx12.5-generic), so there is no functional change.
The feature is also made frontend visible, and the corresponding
__builtin_amdgcn_{flat,global}_load_monitor_b* builtins now require
"load-monitor-insts" rather than "gfx1250-insts".
Co-Authored-By: Claude Opus 5 <[email protected]>
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 12 ++++++------
llvm/lib/Target/AMDGPU/FLATInstructions.td | 8 ++++----
llvm/lib/Target/AMDGPU/GCNFeatures.td | 7 ++++++-
3 files changed, 16 insertions(+), 11 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td
b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 5a0b71fa50c19b..e5c61f7e7916a5 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -1015,12 +1015,12 @@ def __builtin_amdgcn_s_cluster_barrier :
AMDGPUBuiltin<"void()", [], "gfx1250-in
def __builtin_amdgcn_flat_prefetch : AMDGPUBuiltin<"void(void const
address_space<0> *, _Constant int)", [Const], "vmem-pref-insts">;
def __builtin_amdgcn_global_prefetch : AMDGPUBuiltin<"void(void const
address_space<1> *, _Constant int)", [Const], "vmem-pref-insts">;
-def __builtin_amdgcn_global_load_monitor_b32 : AMDGPUBuiltin<"int(int
address_space<1> *, _Constant int, _Constant int)", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_global_load_monitor_b64 : AMDGPUBuiltin<"_ExtVector<2,
int>(_ExtVector<2, int> address_space<1> *, _Constant int, _Constant int)",
[Const], "gfx1250-insts">;
-def __builtin_amdgcn_global_load_monitor_b128 : AMDGPUBuiltin<"_ExtVector<4,
int>(_ExtVector<4, int> address_space<1> *, _Constant int, _Constant int)",
[Const], "gfx1250-insts">;
-def __builtin_amdgcn_flat_load_monitor_b32 : AMDGPUBuiltin<"int(int
address_space<0> *, _Constant int, _Constant int)", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_flat_load_monitor_b64 : AMDGPUBuiltin<"_ExtVector<2,
int>(_ExtVector<2, int> address_space<0> *, _Constant int, _Constant int)",
[Const], "gfx1250-insts">;
-def __builtin_amdgcn_flat_load_monitor_b128 : AMDGPUBuiltin<"_ExtVector<4,
int>(_ExtVector<4, int> address_space<0> *, _Constant int, _Constant int)",
[Const], "gfx1250-insts">;
+def __builtin_amdgcn_global_load_monitor_b32 : AMDGPUBuiltin<"int(int
address_space<1> *, _Constant int, _Constant int)", [Const],
"load-monitor-insts">;
+def __builtin_amdgcn_global_load_monitor_b64 : AMDGPUBuiltin<"_ExtVector<2,
int>(_ExtVector<2, int> address_space<1> *, _Constant int, _Constant int)",
[Const], "load-monitor-insts">;
+def __builtin_amdgcn_global_load_monitor_b128 : AMDGPUBuiltin<"_ExtVector<4,
int>(_ExtVector<4, int> address_space<1> *, _Constant int, _Constant int)",
[Const], "load-monitor-insts">;
+def __builtin_amdgcn_flat_load_monitor_b32 : AMDGPUBuiltin<"int(int
address_space<0> *, _Constant int, _Constant int)", [Const],
"load-monitor-insts">;
+def __builtin_amdgcn_flat_load_monitor_b64 : AMDGPUBuiltin<"_ExtVector<2,
int>(_ExtVector<2, int> address_space<0> *, _Constant int, _Constant int)",
[Const], "load-monitor-insts">;
+def __builtin_amdgcn_flat_load_monitor_b128 : AMDGPUBuiltin<"_ExtVector<4,
int>(_ExtVector<4, int> address_space<0> *, _Constant int, _Constant int)",
[Const], "load-monitor-insts">;
def __builtin_amdgcn_cluster_load_b32 : AMDGPUBuiltin<"int(int
address_space<1> *, _Constant int, int)", [Const],
"mcast-load-insts,wavefrontsize32">;
def __builtin_amdgcn_cluster_load_b64 : AMDGPUBuiltin<"_ExtVector<2,
int>(_ExtVector<2, int> address_space<1> *, _Constant int, int)", [Const],
"mcast-load-insts,wavefrontsize32">;
def __builtin_amdgcn_cluster_load_b128 : AMDGPUBuiltin<"_ExtVector<4,
int>(_ExtVector<4, int> address_space<1> *, _Constant int, int)", [Const],
"mcast-load-insts,wavefrontsize32">;
diff --git a/llvm/lib/Target/AMDGPU/FLATInstructions.td
b/llvm/lib/Target/AMDGPU/FLATInstructions.td
index 8aeb2794873890..5178da814fcfcd 100644
--- a/llvm/lib/Target/AMDGPU/FLATInstructions.td
+++ b/llvm/lib/Target/AMDGPU/FLATInstructions.td
@@ -1297,7 +1297,7 @@ defm SCRATCH_LOAD_LDS_USHORT :
FLAT_Scratch_Load_LDS_Pseudo <"scratch_load_lds_u
defm SCRATCH_LOAD_LDS_SSHORT : FLAT_Scratch_Load_LDS_Pseudo
<"scratch_load_lds_sshort">;
defm SCRATCH_LOAD_LDS_DWORD : FLAT_Scratch_Load_LDS_Pseudo
<"scratch_load_lds_dword">;
-let SubtargetPredicate = isGFX125xOnly in {
+let SubtargetPredicate = HasLoadMonitorInsts in {
defm FLAT_LOAD_MONITOR_B32 : FLAT_Flat_Load_Pseudo
<"flat_load_monitor_b32", VGPROp_32>;
defm FLAT_LOAD_MONITOR_B64 : FLAT_Flat_Load_Pseudo
<"flat_load_monitor_b64", VGPROp_64>;
defm FLAT_LOAD_MONITOR_B128 : FLAT_Flat_Load_Pseudo
<"flat_load_monitor_b128", VGPROp_128>;
@@ -1305,7 +1305,7 @@ defm FLAT_LOAD_MONITOR_B128 : FLAT_Flat_Load_Pseudo
<"flat_load_monitor_b128",
defm GLOBAL_LOAD_MONITOR_B32 : FLAT_Global_Load_Pseudo
<"global_load_monitor_b32", VGPROp_32>;
defm GLOBAL_LOAD_MONITOR_B64 : FLAT_Global_Load_Pseudo
<"global_load_monitor_b64", VGPROp_64>;
defm GLOBAL_LOAD_MONITOR_B128 : FLAT_Global_Load_Pseudo
<"global_load_monitor_b128", VGPROp_128>;
-} // End SubtargetPredicate = isGFX125xOnly
+} // End SubtargetPredicate = HasLoadMonitorInsts
let SubtargetPredicate = isGFX1250Plus, WaveSizePredicate = isWave32 in {
let Uses = [M0, EXEC] in { // Use M0 for broadcast workgroup mask.
@@ -2371,7 +2371,7 @@ let WaveSizePredicate = isWave32, OtherPredicates =
[HasTransposeLoadF4F6Insts]
defm : GlobalFLATLoadPats <GLOBAL_LOAD_TR6_B96,
int_amdgcn_global_load_tr6_b96, v3i32>;
}
-let OtherPredicates = [isGFX125xOnly] in {
+let SubtargetPredicate = HasLoadMonitorInsts in {
def : FlatLoadPat <FLAT_LOAD_MONITOR_B32, AMDGPUflat_load_monitor, i32>;
def : FlatLoadPat <FLAT_LOAD_MONITOR_B64, AMDGPUflat_load_monitor, v2i32>;
def : FlatLoadPat <FLAT_LOAD_MONITOR_B128, AMDGPUflat_load_monitor, v4i32>;
@@ -2379,7 +2379,7 @@ let OtherPredicates = [isGFX125xOnly] in {
defm : GlobalFLATLoadPats <GLOBAL_LOAD_MONITOR_B32,
AMDGPUglobal_load_monitor, i32>;
defm : GlobalFLATLoadPats <GLOBAL_LOAD_MONITOR_B64,
AMDGPUglobal_load_monitor, v2i32>;
defm : GlobalFLATLoadPats <GLOBAL_LOAD_MONITOR_B128,
AMDGPUglobal_load_monitor, v4i32>;
-} // End SubtargetPredicate = isGFX125xOnly
+} // End SubtargetPredicate = HasLoadMonitorInsts
let OtherPredicates = [isGFX1250Plus] in {
defm : GlobalFLATLoadPats_M0 <CLUSTER_LOAD_B32,
int_amdgcn_cluster_load_b32, i32>;
diff --git a/llvm/lib/Target/AMDGPU/GCNFeatures.td
b/llvm/lib/Target/AMDGPU/GCNFeatures.td
index 47e3155d3f51a0..b666f9796f6667 100644
--- a/llvm/lib/Target/AMDGPU/GCNFeatures.td
+++ b/llvm/lib/Target/AMDGPU/GCNFeatures.td
@@ -572,6 +572,10 @@ defm McastLoadInsts :
AMDGPUSubtargetFeature<"mcast-load-insts",
"Has multicast load instructions"
>;
+defm LoadMonitorInsts : AMDGPUSubtargetFeature<"load-monitor-insts",
+ "Has flat_load_monitor_b* and global_load_monitor_b* instructions"
+>;
+
defm AsyncLoadToLDSInsts : AMDGPUSubtargetFeature<"async-load-to-lds-insts",
"Has global async load to LDS instructions"
>;
@@ -2525,6 +2529,7 @@ def FeatureISAVersion12_50_Common : FeatureSet<
FeatureTensorCvtLutInsts,
FeatureTDMInsts,
FeatureVPermPk16Hazard,
+ FeatureLoadMonitorInsts,
FeatureTransposeLoadF4F6Insts,
FeatureBF16TransInsts,
FeatureBF16ConversionInsts,
@@ -2827,7 +2832,7 @@ def AMDGPUFrontendVisibleFeatures {
FeatureGFX13Insts, FeatureGFX8Insts, FeatureGFX90AInsts,
FeatureGFX940Insts, FeatureGFX950Insts, FeatureGFX9Insts,
FeatureGWS, FeatureImageInsts, FeatureLdsBarrierArriveAtomic,
- FeatureLerpInst, FeatureMAIInsts, FeatureMcastLoadInsts,
+ FeatureLerpInst, FeatureLoadMonitorInsts, FeatureMAIInsts,
FeatureMcastLoadInsts,
FeatureMqsadInsts, FeatureMqsadPkInsts, FeatureMsadInsts,
FeatureOCPFP8ConversionInsts, FeaturePermlane16Insts, FeaturePermlane16Swap,
FeaturePermlane32Swap, FeaturePkAddMinMaxInsts, FeaturePrngInst,
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
