Addressed review comments

1. Moved tablegen check to Subtarget and using callSubtarget to access
it.
2. Removed AnyPred, we don't need it as of now.

Author: rajatbajpai

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

@@ -101,13 +101,6 @@ 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`.
@@ -125,14 +118,6 @@ class PTXWithAccelSMs<int PTXVersion, list<int> 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]>
-                                  ]>;
-
 def hasAtomAddF64 : Predicate<"Subtarget->hasAtomAddF64()">;
 def hasAtomScope : Predicate<"Subtarget->hasAtomScope()">;
 def hasAtomBitwise64 : Predicate<"Subtarget->hasAtomBitwise64()">;

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

@@ -5176,7 +5176,7 @@ foreach src_fmt = ["", "b6x16_p32", "b4x16_p64"] in {
 }
 } // Predicates
 
-let Predicates = [hasTcgen05ShiftSupport] in {
+let Predicates = [callSubtarget<"hasTcgen05ShiftSupport">] in {
 multiclass TCGEN05_SHIFT_INTR<string num, Intrinsic Intr> {
   def "" : BasicNVPTXInst<(outs),
              (ins ADDR:$tmem_addr),

llvm/lib/Target/NVPTX/NVPTXSubtarget.h

@@ -146,11 +146,20 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
   // - tcgen05.fence/wait
   // - tcgen05.commit
   bool hasTcgen05InstSupport() const {
+    // sm_101 renamed to sm_110 in PTX 9.0
     return hasPTXWithFamilySMs(90, {100, 110}) ||
            hasPTXWithFamilySMs(88, {100, 101}) ||
            hasPTXWithAccelSMs(86, {100, 101});
   }
 
+  // Checks tcgen05.shift instruction support.
+  bool hasTcgen05ShiftSupport() const {
+    // sm_101 renamed to sm_110 in PTX 9.0
+    return hasPTXWithAccelSMs(90, {100, 110, 103}) ||
+           hasPTXWithAccelSMs(88, {100, 101, 103}) ||
+           hasPTXWithAccelSMs(86, {100, 101});
+  }
+
   bool hasTcgen05MMAScaleInputDImm() const {
     return FullSmVersion == 1003 && PTXVersion >= 86;
   }