[llvm] [NVVM][NVPTX] Support decompress_b feature for tcgen05.mma intrinsics (PR #216312)

[Rajat Bajpai via llvm-commits]

================
@@ -3565,6 +3604,101 @@ foreach sparse = [0, 1] in {
   } // space
 } // sparse
 
+//
+// tcgen05.mma decompress intrinsics
+//
+foreach space = ["tensor", "shared"] in {
+  defvar mma = NVVM_TCGEN05_MMA_DECOMPRESS<space, "f8f6f4">;
+  defvar args = !listconcat(
+                  mma.common_args,
+                  [llvm_tmem_ptr_ty] // decompress_b
+                );
+  defvar flags = [llvm_i32_ty,  // cta_group
+                  llvm_i32_ty,  // collector_usage_a
+                  llvm_i32_ty]; // collector_usage_b
+  defvar nargs = !size(args);
+  defvar cta_group = ArgIndex<nargs>;
+  defvar collector_usage_a = ArgIndex<!add(nargs, 1)>;
+  defvar collector_usage_b = ArgIndex<!add(nargs, 2)>;
+
+  defvar intrinsic_properties = !listconcat(
+      mma.common_intr_props,
+      [Range<cta_group, 1, 3>,
+       Range<collector_usage_a, 0, 4>,
+       Range<collector_usage_b, 0, 4>,
+       ArgInfo<cta_group, [ArgName<"cta_group">]>,
+       ArgInfo<collector_usage_a, [ArgName<"collector_a">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>,
----------------
rajatbajpai wrote:

should we have a util to provide this property? I see it in multiple ops now. Something along the lines of.

```
class NVVM_COLLECTOR_ARGPROP<int col_idx, string buffer = "a">
list<IntrinsicProperty> Prop = ...

...

NVVM_COLLECTOR_ARGPROP<collector_usage_a>.Prop
NVVM_COLLECTOR_ARGPROP<collector_usage_b, "b">.Prop
```

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