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

[Artem Belevich 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 # ")"),
+                        " || ") #
+            ")">;
----------------
Artem-B wrote:

I think this is a bit more complicated than it should be.
We can probably make it somewhat more readable.
E.g. we can extract comparison logic to a lambda, and just use tablegen to populate the values to compare:

```
     // the values are populated by tablegen
     std::array values {1, 2, 3, 4, 5};
     int key = 2;
     // But the logic is separate and readable.
     return llvm::any_of(values, [key](auto x) { return x == key; });

```

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