[llvm] [NVPTX] Update architecture support checks for tcgen05 intrinsics (PR #161519)

[Rajat Bajpai via llvm-commits]

================
@@ -99,6 +99,45 @@ 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.
+class PTXWithFamilySMs<int PTXVersion, list<int> SMVersions> :
+  Predicate<"Subtarget->getPTXVersion() >= " # PTXVersion #
+            " && Subtarget->hasFamilySpecificFeatures()" #
+            " && (" #
+            !interleave(!foreach(sm, SMVersions,
+                        "(Subtarget->getSmFamilyVersion() == " # !div(sm, 10) #
+                        " && Subtarget->getSmVersion() >= " # sm # ")"),
+                        " || ") #
+            ")">;
+
+// Checks PTX version and architecture-specific SM versions.
+// For example, sm_100{a} will match.
+class PTXWithAccelSMs<int PTXVersion, list<int> SMVersions> :
+  Predicate<"Subtarget->getPTXVersion() >= " # PTXVersion #
+            " && Subtarget->hasArchAccelFeatures()" #
+            " && (" #
+            !interleave(!foreach(sm, SMVersions,
+                        "Subtarget->getSmVersion() == " # sm),
+                        " || ") #
+            ")">;
+
+// 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<[
+                                  PTXWithAccelSMs<90, [100, 110, 103]>,
----------------
rajatbajpai wrote:

Two reasons I went with this approach: 1) it serves as a documentation and 2) it aligns with PTX ISA doc.

> Ideally we want to maintain some sort of automatic applicability where we can.

Yes, but for now I don't have better/cleaner solution.

> I think we need to make the check on the first PTX version to be ">=", so it automatically works for the newer PTX versions without us having to add new entries to the list.

PTX check is `>=`, so it'll work for newer PTX ISA versions.

https://github.com/llvm/llvm-project/pull/161519