[llvm] [NVPTX] Update architecture support checks for tcgen05 intrinsics (PR #161519)
Durgadoss R via llvm-commits
llvm-commits at lists.llvm.org
Wed Oct 8 00:52:22 PDT 2025
================
@@ -101,6 +101,37 @@ def PrmtMode : Operand<i32> {
// NVPTX Instruction Predicate Definitions
//===----------------------------------------------------------------------===//
+// Helper predicate to compose multiple predicates.
+class AnyPred<list<Predicate> Preds>
+ : Predicate<"(" #
+ !interleave(!foreach(pred, Preds, pred.CondString),
+ ") || (") #
+ ")">;
+
+// Checks PTX version and family-specific and architecture-specific SM versions.
+// For example, sm_100{f/a} and any future variants in the same family will match
+// for any PTX version greater than or equal to `PTXVersion`.
+class PTXWithFamilySMs<int PTXVersion, list<int> SMVersions> :
+ Predicate<"Subtarget->hasPTXWithFamilySMs(" # PTXVersion # ", {" #
+ !interleave(SMVersions, ", ") # "})">;
+
+// Checks PTX version and architecture-specific SM versions.
+// For example, sm_100{a} will match for any PTX version
+// greater than or equal to `PTXVersion`.
+class PTXWithAccelSMs<int PTXVersion, list<int> SMVersions> :
+ Predicate<"Subtarget->hasPTXWithAccelSMs(" # PTXVersion # ", {" #
+ !interleave(SMVersions, ", ") # "})">;
+
+// Helper predicate to call a subtarget method.
+class callSubtarget<string SubtargetMethod> : Predicate<"Subtarget->" # SubtargetMethod # "()">;
+
+// Composed predicate to check tcgen05.shift instructions support.
+def hasTcgen05ShiftSupport : AnyPred<[
+ // sm_101 renamed to sm_110 in PTX 9.0
+ PTXWithAccelSMs<90, [100, 110, 103]>,
+ PTXWithAccelSMs<88, [100, 101, 103]>,
+ PTXWithAccelSMs<86, [100, 101]>
+ ]>;
----------------
durga4github wrote:
ok, thanks!
https://github.com/llvm/llvm-project/pull/161519
More information about the llvm-commits
mailing list