[Mlir-commits] [mlir] [mlir][nvvm] Verify kind-specific target requirements of tcgen05.mma ops (PR #235628)

[Vito Secona]

================
@@ -6228,6 +6228,94 @@ mlir::NVVM::IDArgPair TensormapReplaceOp::getIntrinsicIDAndArgs(
   return {IDs[fieldIndex], args};
 }
 
+//===----------------------------------------------------------------------===//
+// NVVM tcgen05.mma target requirements
+//===----------------------------------------------------------------------===//
+
+// Full SM versions as encoded by NVVMCheckSMVersion: the SM version times 10,
+// plus 2 for the family-specific ('f') and 3 for the architecture-specific
+// ('a') variant.
+static constexpr unsigned kSM100a = 1003, kSM100f = 1002, kSM101f = 1012,
+                          kSM103a = 1033, kSM107a = 1073, kSM107f = 1072,
+                          kSM110a = 1103, kSM110f = 1102;
+
+// The requirements below mirror the predicates of the tcgen05.mma instructions
+// in the NVPTX backend.
+static constexpr unsigned kTcgen05MMABaseSM[] = {kSM100f, kSM101f, kSM110f};
+
+static NVVMCheckSMVersion makeSMRequirement(ArrayRef<unsigned> versions) {
+  NVVMCheckSMVersion requirement;
+  requirement.fullSmVersionList.assign(versions.begin(), versions.end());
+  return requirement;
+}
----------------
secona wrote:

It might be worth to make this into a constructor of NVVMCheckSMVersion; it already has the variadic constructor.

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