Author: Changpeng Fang Date: 2026-10-30T13:49:58-07:00 New Revision: 11c03ef9119fede49ab2822a43e7b31438c70bb0
URL: https://github.com/llvm/llvm-project/commit/11c03ef9119fede49ab2822a43e7b31438c70bb0 DIFF: https://github.com/llvm/llvm-project/commit/11c03ef9119fede49ab2822a43e7b31438c70bb0.diff LOG: [AMDGPU][NFC] Add a target feature for load monitor instructions (#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". Co-authored-by: Claude Opus 5 <[email protected]> Added: Modified: clang/include/clang/Basic/BuiltinsAMDGPU.td llvm/lib/Target/AMDGPU/FLATInstructions.td llvm/lib/Target/AMDGPU/GCNFeatures.td Removed: ################################################################################ diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 5a0b71fa50c19..e5c61f7e7916a 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 8aeb279487389..5178da814fcfc 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 65a5644c5ab62..313f67ea31e22 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" >; @@ -2530,6 +2534,7 @@ def FeatureISAVersion12_50_Common : FeatureSet< FeatureTensorCvtLutInsts, FeatureTDMInsts, FeatureVPermPk16Hazard, + FeatureLoadMonitorInsts, FeatureTransposeLoadF4F6Insts, FeatureBF16TransInsts, FeatureBF16ConversionInsts, @@ -2834,7 +2839,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
