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

[Kirill Vedernikov via llvm-commits]

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

>From cfaa469370a4f59666ebd976e915dbf69e35f9ae Mon Sep 17 00:00:00 2001
From: Kirill Vedernikov <kvedernikov at nvidia.com>
Date: Fri, 14 Aug 2026 14:24:05 +0200
Subject: [PATCH 1/5] [NVVM][NVPTX] Support decompress_b feature for
 tcgen05.mma intrinsics

---
 llvm/include/llvm/IR/IntrinsicsNVVM.td        | 134 +++++++++
 llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp   |  28 ++
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td      | 197 ++++++++++++-
 .../tcgen05-mma-block-scale-decompress-b.ll   | 264 ++++++++++++++++++
 .../CodeGen/NVPTX/tcgen05-mma-decompress-b.ll | 260 +++++++++++++++++
 ...05-mma-disable-output-lane-decompress-b.ll | 263 +++++++++++++++++
 6 files changed, 1141 insertions(+), 5 deletions(-)
 create mode 100644 llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-decompress-b.ll
 create mode 100644 llvm/test/CodeGen/NVPTX/tcgen05-mma-decompress-b.ll
 create mode 100644 llvm/test/CodeGen/NVPTX/tcgen05-mma-disable-output-lane-decompress-b.ll

diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 6d6671e79dd24..7aaff46001e9d 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1233,6 +1233,45 @@ class NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED<string Kind, string ScaleVecSize> {
   );
 }
 
+//
+// tcgen05.mma intrinsics with decompress_b extension
+//
+
+class NVVM_TCGEN05_MMA_DECOMPRESS_BASE<string Space, string Kind>:
+    NVVM_TCGEN05_MMA_BASE<Space, /* IsSparse= */ 0, Kind> {
+  string DecompressPrefix = Prefix # SpSpaceKindStr;
+  string DecompressSuffix = ".decompress_b";
+}
+
+class NVVM_TCGEN05_MMA_DECOMPRESS<string Space, string Kind>:
+    NVVM_TCGEN05_MMA_DECOMPRESS_BASE<Space, Kind> {
+  string name = DecompressPrefix # DecompressSuffix;
+  string intr_name = IntrinsicName<name>.intr_name;
+  string record_name= IntrinsicName<name>.record_name;
+}
+
+class NVVM_TCGEN05_MMA_DECOMPRESS_DISABLE_OUTPUT_LANE<string Space,
+                                                      string Kind,
+                                                      int CtaGroup>:
+    NVVM_TCGEN05_MMA_DECOMPRESS_BASE<Space, Kind> {
+  string name = DecompressPrefix
+                # ".disable_output_lane.cg" # CtaGroup
+                # DecompressSuffix;
+  string intr_name = IntrinsicName<name>.intr_name;
+  string record_name= IntrinsicName<name>.record_name;
+}
+
+class NVVM_TCGEN05_MMA_DECOMPRESS_BLOCKSCALE<string Space, string Kind,
+                                             string ScaleVecSize>:
+    NVVM_TCGEN05_MMA_DECOMPRESS_BASE<Space, Kind> {
+  string name = DecompressPrefix
+                # ".block_scale"
+                # ScaleVecSize
+                # DecompressSuffix;
+  string intr_name = IntrinsicName<name>.intr_name;
+  string record_name= IntrinsicName<name>.record_name;
+}
+
 class TexVector<string name, list<LLVMType> types> {
   string Name = name;
   list<LLVMType> Types = types;
@@ -3563,6 +3602,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">]>,
+       ArgInfo<collector_usage_b, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]);
+
+  def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
+                          intrinsic_properties, mma.intr_name>;
+} // space
+
+//
+// tcgen05.mma decompress disable_output_lane intrinsics
+//
+foreach space = ["tensor", "shared"] in {
+  foreach cta_group = [1, 2] in {
+    defvar mma = NVVM_TCGEN05_MMA_DECOMPRESS_DISABLE_OUTPUT_LANE<
+                  space, "f8f6f4", cta_group>;
+    defvar disable_output_lane_type =
+              !if(!eq(cta_group, 1), llvm_v4i32_ty, llvm_v8i32_ty);
+    defvar args = !listconcat(
+                    mma.common_args,
+                    [llvm_tmem_ptr_ty],        // decompress_b
+                    [disable_output_lane_type] // disable_output_lane
+                  );
+    defvar flags = [llvm_i32_ty,  // collector_usage_a
+                    llvm_i32_ty]; // collector_usage_b
+    defvar nargs = !size(args);
+    defvar collector_usage_a = ArgIndex<nargs>;
+    defvar collector_usage_b = ArgIndex<!add(nargs, 1)>;
+
+    defvar intrinsic_properties = !listconcat(
+        mma.common_intr_props,
+        [Range<collector_usage_a, 0, 4>,
+         Range<collector_usage_b, 0, 4>,
+         ArgInfo<collector_usage_a, [ArgName<"collector_a">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>,
+         ArgInfo<collector_usage_b, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]);
+
+    def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
+                            intrinsic_properties, mma.intr_name>;
+  } // cta_group
+} // space
+
+//
+// tcgen05.mma decompress block_scale intrinsics
+//
+foreach space = ["tensor", "shared"] in {
+  defvar mma = NVVM_TCGEN05_MMA_DECOMPRESS_BLOCKSCALE<space, "mxf8f6f4",
+                                                      ".block32">;
+  defvar args = !listconcat(
+                  mma.common_args,
+                  [llvm_tmem_ptr_ty,  // scale_a
+                   llvm_tmem_ptr_ty,  // scale_b
+                   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">]>,
+       ArgInfo<collector_usage_b, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]);
+
+  def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
+                          intrinsic_properties, mma.intr_name>;
+} // space
+
 //
 // tensormap.replace intrinsics
 //
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 6798875f5fec5..ae6a720c88f88 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -2676,6 +2676,18 @@ static unsigned getTcgen05MMADisableOutputLane(unsigned IID) {
       nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
     return NVPTXISD::
         TCGEN05_MMA_SP_TENSOR_SCALE_D_DISABLE_OUTPUT_LANE_CG2_ASHIFT;
+  case Intrinsic::
+      nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg1_decompress_b:
+    return NVPTXISD::TCGEN05_MMA_SHARED_DISABLE_OUTPUT_LANE_CG1_DECOMPRESS_B;
+  case Intrinsic::
+      nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg2_decompress_b:
+    return NVPTXISD::TCGEN05_MMA_SHARED_DISABLE_OUTPUT_LANE_CG2_DECOMPRESS_B;
+  case Intrinsic::
+      nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg1_decompress_b:
+    return NVPTXISD::TCGEN05_MMA_TENSOR_DISABLE_OUTPUT_LANE_CG1_DECOMPRESS_B;
+  case Intrinsic::
+      nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg2_decompress_b:
+    return NVPTXISD::TCGEN05_MMA_TENSOR_DISABLE_OUTPUT_LANE_CG2_DECOMPRESS_B;
   };
   llvm_unreachable("unhandled tcgen05.mma.disable_output_lane intrinsic");
 }
@@ -2888,6 +2900,14 @@ static SDValue lowerIntrinsicVoid(SDValue Op, SelectionDAG &DAG) {
       nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift:
   case Intrinsic::
       nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
+  case Intrinsic::
+      nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg1_decompress_b:
+  case Intrinsic::
+      nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg2_decompress_b:
+  case Intrinsic::
+      nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg1_decompress_b:
+  case Intrinsic::
+      nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg2_decompress_b:
     return LowerTcgen05MMADisableOutputLane(Op, DAG);
   case Intrinsic::nvvm_tensormap_replace_elemtype:
     return lowerTensormapReplaceElemtype(Op, DAG);
@@ -5497,6 +5517,10 @@ void NVPTXTargetLowering::getTgtMemIntrinsic(
     Infos.push_back(Info);
     return;
   }
+  case Intrinsic::
+      nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg1_decompress_b:
+  case Intrinsic::
+      nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg1_decompress_b:
   case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
   case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
   case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
@@ -5522,6 +5546,10 @@ void NVPTXTargetLowering::getTgtMemIntrinsic(
     return;
   }
 
+  case Intrinsic::
+      nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg2_decompress_b:
+  case Intrinsic::
+      nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg2_decompress_b:
   case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
   case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
   case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index df5248a5bf7b5..5552e24ac9fb8 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6500,7 +6500,8 @@ foreach sparse = [0, 1] in {
 //
 
 class Tcgen05MMADisableOutputLaneTypeProfile<bit IsSparse, string ASpace,
-                                             int CtaGroup, bit IsScaleInputD>:
+                                             int CtaGroup, bit IsScaleInputD,
+                                             bit IsDecompressB>:
         SDTypeProfile<0, 0, []> {
   int DisableOutputLaneVecSize = !mul(4, CtaGroup);
 
@@ -6510,8 +6511,9 @@ class Tcgen05MMADisableOutputLaneTypeProfile<bit IsSparse, string ASpace,
     [i64, i32, i1],                             // b, idesc, enable_inp_d
     !if(IsSparse, [i32], []),                   // spmetadata
     !if(IsScaleInputD, [i64], []),              // scale_input_d
+    !if(IsDecompressB, [i32], []),              // decompress_b
     !listsplat(i32, DisableOutputLaneVecSize),  // disable_output_lane
-    [i32],                                      // kind
+    !if(IsDecompressB, [], [i32]),              // kind
     [i32],                                      // collector_usage_a
     [i32]                                       // collector_usage_b
   );
@@ -6521,14 +6523,16 @@ class Tcgen05MMADisableOutputLaneTypeProfile<bit IsSparse, string ASpace,
 
 class Tcgen05MMADisableOutputLaneSDNode<bit IsSparse, string ASpace,
                                         int CtaGroup, bit IsScaleInput,
-                                        bit IsAShift>:
+                                        bit IsAShift, bit IsDecompressB = 0>:
         SDNode<"NVPTXISD::TCGEN05_MMA"
                 # !if(IsSparse, "_SP", "")
                 # "_" # !toupper(ASpace)
                 # !if(IsScaleInput, "_SCALE_D", "")
                 # "_DISABLE_OUTPUT_LANE_CG" # CtaGroup
-                # !if(IsAShift, "_ASHIFT", ""),
-        Tcgen05MMADisableOutputLaneTypeProfile<IsSparse, ASpace, CtaGroup, IsScaleInput>,
+                # !if(IsAShift, "_ASHIFT", "")
+                # !if(IsDecompressB, "_DECOMPRESS_B", ""),
+        Tcgen05MMADisableOutputLaneTypeProfile<IsSparse, ASpace, CtaGroup,
+                                               IsScaleInput, IsDecompressB>,
                     [SDNPHasChain, SDNPSideEffect, SDNPMemOperand]>;
 
 //
@@ -6764,6 +6768,189 @@ foreach sparse = [0, 1] in {
   } // space
 } // sparse
 
+class Tcgen05MMADecompressBase<string ASpace, string Kind, int CtaGroup,
+                               string CollectorUsageA, string CollectorUsageB> :
+        Tcgen05MMABase</*IsSparse=*/ 0, ASpace, Kind, CtaGroup,
+                       CollectorUsageA, CollectorUsageB> {
+  let Predicates = [hasRubinFamilySupport];
+
+  let KindVal = !cond(
+                  !eq(Kind, "f8f6f4") : 0
+                );
+
+  dag DecompressBaseInOperandList = !con(BaseInOperandList,
+                                         (ins B32:$decompress_b));
+
+  dag DecompressBasePatternArgs = !con(BasePatternArgs, (ins i32:$decompress_b));
+
+  string DecompressBaseOperandsStr = BaseOperandsCommonStr
+                                     # ", [$decompress_b]"
+                                     # IDescStr;
+
+  string DecompressPrefix = Prefix # SpCtaKindStr;
+  string DecompressLutStr = ".decompress::lut::b";
+  string DecompressCollectorStr = ".collector::a::" # CollectorUsageA
+                                  # ".collector::b::" # CollectorUsageB;
+}
+
+class Tcgen05MMADecompressInst<string ASpace, string Kind, int CtaGroup,
+                               string CollectorUsageA,
+                               string CollectorUsageB> :
+    Tcgen05MMADecompressBase<ASpace, Kind, CtaGroup, CollectorUsageA,
+                             CollectorUsageB> {
+  Intrinsic Intrin = !cast<Intrinsic>(
+                        NVVM_TCGEN05_MMA_DECOMPRESS<ASpace, Kind>.record_name
+                     );
+
+  let InOperandList = DecompressBaseInOperandList;
+
+  let AsmString = DecompressPrefix
+                  # DecompressLutStr
+                  # DecompressCollectorStr
+                  # DecompressBaseOperandsStr
+                  # InputDStr
+                  # ";";
+
+  dag IntrinsicPattern = !foreach(tmp, DecompressBasePatternArgs, !subst(ins, Intrin, tmp));
+
+  dag FlagOperands = (Intrin (i32 CtaGroup), (i32 CollectorUsageAVal),
+                             (i32 CollectorUsageBVal));
+
+  let Pattern = [!con(IntrinsicPattern, FlagOperands)];
+}
+
+// tcgen05.mma decompress
+foreach space = ["tensor", "shared"] in {
+  foreach cta_group = [1, 2] in {
+    foreach collector_usage_a = ["discard", "lastuse", "fill", "use"] in {
+      foreach collector_usage_b = ["discard", "lastuse", "fill", "use"] in {
+        def : Tcgen05MMADecompressInst<space, "f8f6f4", cta_group,
+                                       collector_usage_a, collector_usage_b>;
+      } // collector_usage_b
+    } // collector_usage_a
+  } // cta_group
+} // space
+
+class Tcgen05MMADecompressDisableOutputLaneInst<string ASpace, string Kind,
+                                                int CtaGroup,
+                                                string CollectorUsageA,
+                                                string CollectorUsageB> :
+    Tcgen05MMADecompressBase<ASpace, Kind, CtaGroup, CollectorUsageA,
+                             CollectorUsageB> {
+  SDNode Opcode = Tcgen05MMADisableOutputLaneSDNode<0, // IsSparse
+                                                    ASpace, CtaGroup,
+                                                    0,  // IsScaleInputD
+                                                    0,  // IsAShift
+                                                    1>; // IsDecompressB
+
+  // disable output lane
+  int DisableOutputLaneVecSize = !mul(4, CtaGroup);
+
+  dag DisableOutputLaneIns = !dag(ins,
+                              !listsplat(B32, DisableOutputLaneVecSize),
+                              !foreach(x,
+                                      !range(DisableOutputLaneVecSize),
+                                      "disable_output_lane" # x));
+
+  dag DisableOutputLaneInput = !dag(Opcode,
+                                !listsplat(i32, DisableOutputLaneVecSize),
+                                !foreach(x,
+                                         !range(DisableOutputLaneVecSize),
+                                         "disable_output_lane" # x));
+
+  string DisableOutputLaneStr = "{{" #
+                                  !interleave(
+                                    !foreach(x,
+                                      !range(DisableOutputLaneVecSize),
+                                              "$disable_output_lane" # x),
+                                    ", ")
+                                # "}}";
+
+  let InOperandList = !con(DecompressBaseInOperandList,
+                           DisableOutputLaneIns);
+
+  let AsmString = DecompressPrefix
+                  # DecompressLutStr
+                  # DecompressCollectorStr
+                  # DecompressBaseOperandsStr
+                  # ", " # DisableOutputLaneStr
+                  # InputDStr
+                  # ";";
+
+  dag IntrinsicPattern = !con(!foreach(tmp, DecompressBasePatternArgs, !subst(ins, Opcode, tmp)),
+                              DisableOutputLaneInput);
+
+  dag FlagOperands = (Opcode (i32 CollectorUsageAVal),
+                             (i32 CollectorUsageBVal));
+
+  let Pattern = [!con(IntrinsicPattern, FlagOperands)];
+}
+
+// tcgen05.mma decompress disable output lane
+foreach space = ["tensor", "shared"] in {
+  foreach cta_group = [1, 2] in {
+    foreach collector_usage_a = ["discard", "lastuse", "fill", "use"] in {
+      foreach collector_usage_b = ["discard", "lastuse", "fill", "use"] in {
+        def : Tcgen05MMADecompressDisableOutputLaneInst<
+                  space, "f8f6f4", cta_group, collector_usage_a,
+                  collector_usage_b>;
+      } // collector_usage_b
+    } // collector_usage_a
+  } // cta_group
+} // space
+
+class Tcgen05MMADecompressBlockScaleInst<string ASpace, string Kind,
+                                         int CtaGroup, string CollectorUsageA,
+                                         string CollectorUsageB,
+                                         string ScaleVecSize> :
+    Tcgen05MMADecompressBase<ASpace, Kind, CtaGroup, CollectorUsageA,
+                             CollectorUsageB> {
+  Intrinsic Intrin = !cast<Intrinsic>(
+                        NVVM_TCGEN05_MMA_DECOMPRESS_BLOCKSCALE<
+                            ASpace, Kind, ScaleVecSize>.record_name
+                     );
+
+  let KindVal = !cond(
+                  !eq(Kind, "mxf8f6f4") : 0
+                );
+
+  let InOperandList = !con(BaseInOperandList,
+                           (ins B32:$scale_a, B32:$scale_b,
+                                B32:$decompress_b));
+
+  let AsmString = DecompressPrefix
+                  # ".block_scale"
+                  # DecompressLutStr
+                  # ScaleVecSize
+                  # DecompressCollectorStr
+                  # DecompressBaseOperandsStr
+                  # ", [$scale_a], [$scale_b]"
+                  # InputDStr
+                  # ";";
+
+  dag IntrinsicPattern = !con(!foreach(tmp, BasePatternArgs, !subst(ins, Intrin, tmp)),
+                              (Intrin i32:$scale_a, i32:$scale_b, i32:$decompress_b));
+
+  dag FlagOperands = (Intrin (i32 CtaGroup),
+                             (i32 CollectorUsageAVal),
+                             (i32 CollectorUsageBVal));
+
+  let Pattern = [!con(IntrinsicPattern, FlagOperands)];
+}
+
+// tcgen05.mma decompress block_scale
+foreach space = ["tensor", "shared"] in {
+  foreach cta_group = [1, 2] in {
+    foreach collector_usage_a = ["discard", "lastuse", "fill", "use"] in {
+      foreach collector_usage_b = ["discard", "lastuse", "fill", "use"] in {
+        def : Tcgen05MMADecompressBlockScaleInst<
+                  space, "mxf8f6f4", cta_group, collector_usage_a,
+                  collector_usage_b, ".block32">;
+      } // collector_usage_b
+    } // collector_usage_a
+  } // cta_group
+} // space
+
 //
 // tensormap.replace Instructions
 //
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-decompress-b.ll b/llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-decompress-b.ll
new file mode 100644
index 0000000000000..70dc720c5a7d0
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-decompress-b.ll
@@ -0,0 +1,264 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -o - -mcpu=sm_107f -march=nvptx64 -mattr=+ptx94 | FileCheck %s
+; RUN: llvm-as < %s | llvm-dis | FileCheck %s --check-prefixes=FORMAT
+; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mattr=+ptx94 -mcpu=sm_107f | %ptxas-verify -arch=sm_107f %}
+
+define void @tcgen05_mma_decompress_b_cta1_block32(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b) {
+; FORMAT-LABEL: define void @tcgen05_mma_decompress_b_cta1_block32(
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; CHECK-LABEL: tcgen05_mma_decompress_b_cta1_block32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b16 %rs<3>;
+; CHECK-NEXT:    .reg .b32 %r<7>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param::func.b8 %rs1, [tcgen05_mma_decompress_b_cta1_block32_param_5];
+; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
+; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT:    ld.param::func.b32 %r1, [tcgen05_mma_decompress_b_cta1_block32_param_0];
+; CHECK-NEXT:    ld.param::func.b64 %rd1, [tcgen05_mma_decompress_b_cta1_block32_param_2];
+; CHECK-NEXT:    ld.param::func.b64 %rd2, [tcgen05_mma_decompress_b_cta1_block32_param_3];
+; CHECK-NEXT:    ld.param::func.b32 %r2, [tcgen05_mma_decompress_b_cta1_block32_param_4];
+; CHECK-NEXT:    ld.param::func.b32 %r3, [tcgen05_mma_decompress_b_cta1_block32_param_6];
+; CHECK-NEXT:    ld.param::func.b32 %r4, [tcgen05_mma_decompress_b_cta1_block32_param_7];
+; CHECK-NEXT:    ld.param::func.b32 %r5, [tcgen05_mma_decompress_b_cta1_block32_param_8];
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::discard [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    ld.param::func.b32 %r6, [tcgen05_mma_decompress_b_cta1_block32_param_1];
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::discard [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::lastuse [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::lastuse [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::fill [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::fill [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::use [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::use [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::discard [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::discard [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::lastuse [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::lastuse [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::fill [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::fill [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::use [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::use [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::discard [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::discard [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::lastuse [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::lastuse [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::fill [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::fill [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::use [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::use [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::discard [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::discard [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::lastuse [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::lastuse [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::fill [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::fill [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::use [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::use [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    ret;
+
+  ; cta_group::1, collector::a::discard, .block32
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 3)
+
+  ; cta_group::1, collector::a::lastuse, .block32
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 3)
+
+  ; cta_group::1, collector::a::fill, .block32
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 3)
+
+  ; cta_group::1, collector::a::use, .block32
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 3)
+
+  ret void
+}
+
+define void @tcgen05_mma_decompress_b_cta2_block32(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b) {
+; FORMAT-LABEL: define void @tcgen05_mma_decompress_b_cta2_block32(
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; CHECK-LABEL: tcgen05_mma_decompress_b_cta2_block32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b16 %rs<3>;
+; CHECK-NEXT:    .reg .b32 %r<7>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param::func.b8 %rs1, [tcgen05_mma_decompress_b_cta2_block32_param_5];
+; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
+; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT:    ld.param::func.b32 %r1, [tcgen05_mma_decompress_b_cta2_block32_param_0];
+; CHECK-NEXT:    ld.param::func.b64 %rd1, [tcgen05_mma_decompress_b_cta2_block32_param_2];
+; CHECK-NEXT:    ld.param::func.b64 %rd2, [tcgen05_mma_decompress_b_cta2_block32_param_3];
+; CHECK-NEXT:    ld.param::func.b32 %r2, [tcgen05_mma_decompress_b_cta2_block32_param_4];
+; CHECK-NEXT:    ld.param::func.b32 %r3, [tcgen05_mma_decompress_b_cta2_block32_param_6];
+; CHECK-NEXT:    ld.param::func.b32 %r4, [tcgen05_mma_decompress_b_cta2_block32_param_7];
+; CHECK-NEXT:    ld.param::func.b32 %r5, [tcgen05_mma_decompress_b_cta2_block32_param_8];
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::discard [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    ld.param::func.b32 %r6, [tcgen05_mma_decompress_b_cta2_block32_param_1];
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::discard [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::lastuse [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::lastuse [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::fill [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::fill [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::use [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::discard.collector::b::use [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::discard [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::discard [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::lastuse [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::lastuse [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::fill [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::fill [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::use [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::lastuse.collector::b::use [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::discard [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::discard [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::lastuse [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::lastuse [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::fill [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::fill [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::use [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::fill.collector::b::use [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::discard [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::discard [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::lastuse [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::lastuse [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::fill [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::fill [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::use [%r1], %rd1, %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::mxf8f6f4.block_scale.decompress::lut::b.block32.collector::a::use.collector::b::use [%r1], [%r6], %rd2, [%r5], %r2, [%r3], [%r4], %p1;
+; CHECK-NEXT:    ret;
+
+  ; cta_group::2, collector::a::discard, .block32
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 3)
+
+  ; cta_group::2, collector::a::lastuse, .block32
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 3)
+
+  ; cta_group::2, collector::a::fill, .block32
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 3)
+
+  ; cta_group::2, collector::a::use, .block32
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 3)
+
+  ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-mma-decompress-b.ll b/llvm/test/CodeGen/NVPTX/tcgen05-mma-decompress-b.ll
new file mode 100644
index 0000000000000..2a4ad931308da
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-mma-decompress-b.ll
@@ -0,0 +1,260 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -o - -mcpu=sm_107f -march=nvptx64 -mattr=+ptx94 | FileCheck %s
+; RUN: llvm-as < %s | llvm-dis | FileCheck %s --check-prefixes=FORMAT
+; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mattr=+ptx94 -mcpu=sm_107f | %ptxas-verify -arch=sm_107f %}
+
+define void @tcgen05_mma_decompress_b_cta1(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b) {
+; FORMAT-LABEL: define void @tcgen05_mma_decompress_b_cta1(
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 1, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; CHECK-LABEL: tcgen05_mma_decompress_b_cta1(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b16 %rs<3>;
+; CHECK-NEXT:    .reg .b32 %r<5>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param::func.b8 %rs1, [tcgen05_mma_decompress_b_cta1_param_5];
+; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
+; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT:    ld.param::func.b32 %r1, [tcgen05_mma_decompress_b_cta1_param_0];
+; CHECK-NEXT:    ld.param::func.b64 %rd1, [tcgen05_mma_decompress_b_cta1_param_2];
+; CHECK-NEXT:    ld.param::func.b64 %rd2, [tcgen05_mma_decompress_b_cta1_param_3];
+; CHECK-NEXT:    ld.param::func.b32 %r2, [tcgen05_mma_decompress_b_cta1_param_4];
+; CHECK-NEXT:    ld.param::func.b32 %r3, [tcgen05_mma_decompress_b_cta1_param_6];
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    ld.param::func.b32 %r4, [tcgen05_mma_decompress_b_cta1_param_1];
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::discard [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::lastuse [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::fill [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::use [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::discard [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::lastuse [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::fill [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::use [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::discard [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::lastuse [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::fill [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::use [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::discard [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::lastuse [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::fill [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::use [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    ret;
+
+  ; cta_group::1, collector::a::discard
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 0, i32 3)
+
+  ; cta_group::1, collector::a::lastuse
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 1, i32 3)
+
+  ; cta_group::1, collector::a::fill
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 2, i32 3)
+
+  ; cta_group::1, collector::a::use
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 1, i32 3, i32 3)
+
+  ret void
+}
+
+define void @tcgen05_mma_decompress_b_cta2(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b) {
+; FORMAT-LABEL: define void @tcgen05_mma_decompress_b_cta2(
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, /* cta_group= */ i32 2, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; CHECK-LABEL: tcgen05_mma_decompress_b_cta2(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b16 %rs<3>;
+; CHECK-NEXT:    .reg .b32 %r<5>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param::func.b8 %rs1, [tcgen05_mma_decompress_b_cta2_param_5];
+; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
+; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT:    ld.param::func.b32 %r1, [tcgen05_mma_decompress_b_cta2_param_0];
+; CHECK-NEXT:    ld.param::func.b64 %rd1, [tcgen05_mma_decompress_b_cta2_param_2];
+; CHECK-NEXT:    ld.param::func.b64 %rd2, [tcgen05_mma_decompress_b_cta2_param_3];
+; CHECK-NEXT:    ld.param::func.b32 %r2, [tcgen05_mma_decompress_b_cta2_param_4];
+; CHECK-NEXT:    ld.param::func.b32 %r3, [tcgen05_mma_decompress_b_cta2_param_6];
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    ld.param::func.b32 %r4, [tcgen05_mma_decompress_b_cta2_param_1];
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::discard [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::discard [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::discard [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::discard [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::lastuse [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::lastuse [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::lastuse [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::lastuse [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::fill [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::fill [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::fill [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::fill [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::use [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::use [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::use [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::use [%r1], [%r4], %rd2, [%r3], %r2, %p1;
+; CHECK-NEXT:    ret;
+
+  ; cta_group::2, collector::b::discard
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 0)
+
+  ; cta_group::2, collector::b::lastuse
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 1)
+
+  ; cta_group::2, collector::b::fill
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 2)
+
+  ; cta_group::2, collector::b::use
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 0, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 1, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 2, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, i32 2, i32 3, i32 3)
+
+  ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-mma-disable-output-lane-decompress-b.ll b/llvm/test/CodeGen/NVPTX/tcgen05-mma-disable-output-lane-decompress-b.ll
new file mode 100644
index 0000000000000..88c5534af43aa
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-mma-disable-output-lane-decompress-b.ll
@@ -0,0 +1,263 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -o - -mcpu=sm_107f -march=nvptx64 -mattr=+ptx94 | FileCheck %s
+; RUN: llvm-as < %s | llvm-dis | FileCheck %s --check-prefixes=FORMAT
+; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mattr=+ptx94 -mcpu=sm_107f | %ptxas-verify -arch=sm_107f %}
+
+define void @tcgen05_mma_decompress_b_cta1(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4) {
+; FORMAT-LABEL: define void @tcgen05_mma_decompress_b_cta1(
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; CHECK-LABEL: tcgen05_mma_decompress_b_cta1(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b16 %rs<3>;
+; CHECK-NEXT:    .reg .b32 %r<9>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param::func.b8 %rs1, [tcgen05_mma_decompress_b_cta1_param_5];
+; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
+; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT:    ld.param::func.b32 %r1, [tcgen05_mma_decompress_b_cta1_param_0];
+; CHECK-NEXT:    ld.param::func.b64 %rd1, [tcgen05_mma_decompress_b_cta1_param_2];
+; CHECK-NEXT:    ld.param::func.b64 %rd2, [tcgen05_mma_decompress_b_cta1_param_3];
+; CHECK-NEXT:    ld.param::func.b32 %r2, [tcgen05_mma_decompress_b_cta1_param_4];
+; CHECK-NEXT:    ld.param::func.b32 %r3, [tcgen05_mma_decompress_b_cta1_param_6];
+; CHECK-NEXT:    ld.param::func.v4.b32 {%r4, %r5, %r6, %r7}, [tcgen05_mma_decompress_b_cta1_param_7];
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    ld.param::func.b32 %r8, [tcgen05_mma_decompress_b_cta1_param_1];
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::discard [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::lastuse [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::fill [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::use [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::discard [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::lastuse [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::fill [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::use [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::discard [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::lastuse [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::fill [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::use [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::discard [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::lastuse [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::fill [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::1.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::use [%r1], [%r8], %rd2, [%r3], %r2, {%r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    ret;
+
+  ; cta_group::1, collector::a::discard
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 0, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 0, i32 3)
+
+  ; cta_group::1, collector::a::lastuse
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 1, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 1, i32 3)
+
+  ; cta_group::1, collector::a::fill
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 2, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 2, i32 3)
+
+  ; cta_group::1, collector::a::use
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 3, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg1.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <4 x i32> %disable_output_lanev4, i32 3, i32 3)
+
+  ret void
+}
+
+define void @tcgen05_mma_decompress_b_cta2(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8) {
+; FORMAT-LABEL: define void @tcgen05_mma_decompress_b_cta2(
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=discard */ i32 0, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=discard */ i32 0, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=discard */ i32 0, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=discard */ i32 0, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=lastuse */ i32 1, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=lastuse */ i32 1, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=lastuse */ i32 1, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=lastuse */ i32 1, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=fill */ i32 2, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=fill */ i32 2, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=fill */ i32 2, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=fill */ i32 2, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=use */ i32 3, /* collector_b=discard */ i32 0)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=use */ i32 3, /* collector_b=lastuse */ i32 1)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=use */ i32 3, /* collector_b=fill */ i32 2)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; FORMAT-NEXT: call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, /* collector_a=use */ i32 3, /* collector_b=use */ i32 3)
+; CHECK-LABEL: tcgen05_mma_decompress_b_cta2(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b16 %rs<3>;
+; CHECK-NEXT:    .reg .b32 %r<13>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param::func.b8 %rs1, [tcgen05_mma_decompress_b_cta2_param_5];
+; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
+; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT:    ld.param::func.b32 %r1, [tcgen05_mma_decompress_b_cta2_param_0];
+; CHECK-NEXT:    ld.param::func.b64 %rd1, [tcgen05_mma_decompress_b_cta2_param_2];
+; CHECK-NEXT:    ld.param::func.b64 %rd2, [tcgen05_mma_decompress_b_cta2_param_3];
+; CHECK-NEXT:    ld.param::func.b32 %r2, [tcgen05_mma_decompress_b_cta2_param_4];
+; CHECK-NEXT:    ld.param::func.b32 %r3, [tcgen05_mma_decompress_b_cta2_param_6];
+; CHECK-NEXT:    ld.param::func.v4.b32 {%r4, %r5, %r6, %r7}, [tcgen05_mma_decompress_b_cta2_param_7+16];
+; CHECK-NEXT:    ld.param::func.v4.b32 {%r8, %r9, %r10, %r11}, [tcgen05_mma_decompress_b_cta2_param_7];
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    ld.param::func.b32 %r12, [tcgen05_mma_decompress_b_cta2_param_1];
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::discard [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::lastuse [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::fill [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::discard.collector::b::use [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::discard [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::lastuse [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::fill [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::lastuse.collector::b::use [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::discard [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::lastuse [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::fill [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::fill.collector::b::use [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::discard [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::discard [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::lastuse [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::lastuse [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::fill [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::fill [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::use [%r1], %rd1, %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    tcgen05.mma.cta_group::2.kind::f8f6f4.decompress::lut::b.collector::a::use.collector::b::use [%r1], [%r12], %rd2, [%r3], %r2, {%r8, %r9, %r10, %r11, %r4, %r5, %r6, %r7}, %p1;
+; CHECK-NEXT:    ret;
+
+  ; cta_group::2, collector::a::discard
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 0, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 0, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 0, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 0, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 0, i32 3)
+
+  ; cta_group::2, collector::a::lastuse
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 1, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 1, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 1, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 1, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 1, i32 3)
+
+  ; cta_group::2, collector::a::fill
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 2, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 2, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 2, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 2, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 2, i32 3)
+
+  ; cta_group::2, collector::a::use
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 3, i32 0)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 3, i32 1)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 3, i32 2)
+  call void @llvm.nvvm.tcgen05.mma.shared.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, i64 %ashared, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 3, i32 3)
+  call void @llvm.nvvm.tcgen05.mma.tensor.f8f6f4.disable_output_lane.cg2.decompress_b(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %decompress_b, <8 x i32> %disable_output_lanev8, i32 3, i32 3)
+
+  ret void
+}

>From 8e304a7b7d64a7404900c68520517aa73d372afb Mon Sep 17 00:00:00 2001
From: Kirill Vedernikov <kvedernikov at nvidia.com>
Date: Tue, 18 Aug 2026 13:31:12 +0200
Subject: [PATCH 2/5] [NVPTX] Address feedback

---
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 8 ++------
 1 file changed, 2 insertions(+), 6 deletions(-)

diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 5552e24ac9fb8..61b6a8ad58f26 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6774,9 +6774,7 @@ class Tcgen05MMADecompressBase<string ASpace, string Kind, int CtaGroup,
                        CollectorUsageA, CollectorUsageB> {
   let Predicates = [hasRubinFamilySupport];
 
-  let KindVal = !cond(
-                  !eq(Kind, "f8f6f4") : 0
-                );
+  let KindVal = !cond(!eq(Kind, "f8f6f4") : 0);
 
   dag DecompressBaseInOperandList = !con(BaseInOperandList,
                                          (ins B32:$decompress_b));
@@ -6910,9 +6908,7 @@ class Tcgen05MMADecompressBlockScaleInst<string ASpace, string Kind,
                             ASpace, Kind, ScaleVecSize>.record_name
                      );
 
-  let KindVal = !cond(
-                  !eq(Kind, "mxf8f6f4") : 0
-                );
+  let KindVal = !cond(!eq(Kind, "mxf8f6f4") : 0);
 
   let InOperandList = !con(BaseInOperandList,
                            (ins B32:$scale_a, B32:$scale_b,

>From 3417317b987fefdabc737d44a3e301d68697cdeb Mon Sep 17 00:00:00 2001
From: Kirill Vedernikov <kvedernikov at nvidia.com>
Date: Tue, 18 Aug 2026 18:41:47 +0200
Subject: [PATCH 3/5] [NVPTX] Updates according to feedback

---
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 43 +++++++-----------------
 1 file changed, 12 insertions(+), 31 deletions(-)

diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 61b6a8ad58f26..8ab3219b58bc1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6773,18 +6773,15 @@ class Tcgen05MMADecompressBase<string ASpace, string Kind, int CtaGroup,
         Tcgen05MMABase</*IsSparse=*/ 0, ASpace, Kind, CtaGroup,
                        CollectorUsageA, CollectorUsageB> {
   let Predicates = [hasRubinFamilySupport];
-
   let KindVal = !cond(!eq(Kind, "f8f6f4") : 0);
 
   dag DecompressBaseInOperandList = !con(BaseInOperandList,
                                          (ins B32:$decompress_b));
-
   dag DecompressBasePatternArgs = !con(BasePatternArgs, (ins i32:$decompress_b));
 
   string DecompressBaseOperandsStr = BaseOperandsCommonStr
                                      # ", [$decompress_b]"
                                      # IDescStr;
-
   string DecompressPrefix = Prefix # SpCtaKindStr;
   string DecompressLutStr = ".decompress::lut::b";
   string DecompressCollectorStr = ".collector::a::" # CollectorUsageA
@@ -6797,11 +6794,9 @@ class Tcgen05MMADecompressInst<string ASpace, string Kind, int CtaGroup,
     Tcgen05MMADecompressBase<ASpace, Kind, CtaGroup, CollectorUsageA,
                              CollectorUsageB> {
   Intrinsic Intrin = !cast<Intrinsic>(
-                        NVVM_TCGEN05_MMA_DECOMPRESS<ASpace, Kind>.record_name
-                     );
+                        NVVM_TCGEN05_MMA_DECOMPRESS<ASpace, Kind>.record_name);
 
   let InOperandList = DecompressBaseInOperandList;
-
   let AsmString = DecompressPrefix
                   # DecompressLutStr
                   # DecompressCollectorStr
@@ -6810,7 +6805,6 @@ class Tcgen05MMADecompressInst<string ASpace, string Kind, int CtaGroup,
                   # ";";
 
   dag IntrinsicPattern = !foreach(tmp, DecompressBasePatternArgs, !subst(ins, Intrin, tmp));
-
   dag FlagOperands = (Intrin (i32 CtaGroup), (i32 CollectorUsageAVal),
                              (i32 CollectorUsageBVal));
 
@@ -6843,30 +6837,22 @@ class Tcgen05MMADecompressDisableOutputLaneInst<string ASpace, string Kind,
 
   // disable output lane
   int DisableOutputLaneVecSize = !mul(4, CtaGroup);
+  list<string> DisableOutputLaneNames =
+      !foreach(x, !range(DisableOutputLaneVecSize),
+               "disable_output_lane" # x);
 
   dag DisableOutputLaneIns = !dag(ins,
-                              !listsplat(B32, DisableOutputLaneVecSize),
-                              !foreach(x,
-                                      !range(DisableOutputLaneVecSize),
-                                      "disable_output_lane" # x));
-
+                                  !listsplat(B32, DisableOutputLaneVecSize),
+                                  DisableOutputLaneNames);
   dag DisableOutputLaneInput = !dag(Opcode,
-                                !listsplat(i32, DisableOutputLaneVecSize),
-                                !foreach(x,
-                                         !range(DisableOutputLaneVecSize),
-                                         "disable_output_lane" # x));
+                                    !listsplat(i32, DisableOutputLaneVecSize),
+                                    DisableOutputLaneNames);
 
-  string DisableOutputLaneStr = "{{" #
-                                  !interleave(
-                                    !foreach(x,
-                                      !range(DisableOutputLaneVecSize),
-                                              "$disable_output_lane" # x),
-                                    ", ")
+  string DisableOutputLaneStr = "{{$"
+                                # !interleave(DisableOutputLaneNames, ", $")
                                 # "}}";
 
-  let InOperandList = !con(DecompressBaseInOperandList,
-                           DisableOutputLaneIns);
-
+  let InOperandList = !con(DecompressBaseInOperandList, DisableOutputLaneIns);
   let AsmString = DecompressPrefix
                   # DecompressLutStr
                   # DecompressCollectorStr
@@ -6877,7 +6863,6 @@ class Tcgen05MMADecompressDisableOutputLaneInst<string ASpace, string Kind,
 
   dag IntrinsicPattern = !con(!foreach(tmp, DecompressBasePatternArgs, !subst(ins, Opcode, tmp)),
                               DisableOutputLaneInput);
-
   dag FlagOperands = (Opcode (i32 CollectorUsageAVal),
                              (i32 CollectorUsageBVal));
 
@@ -6905,15 +6890,12 @@ class Tcgen05MMADecompressBlockScaleInst<string ASpace, string Kind,
                              CollectorUsageB> {
   Intrinsic Intrin = !cast<Intrinsic>(
                         NVVM_TCGEN05_MMA_DECOMPRESS_BLOCKSCALE<
-                            ASpace, Kind, ScaleVecSize>.record_name
-                     );
+                            ASpace, Kind, ScaleVecSize>.record_name);
 
   let KindVal = !cond(!eq(Kind, "mxf8f6f4") : 0);
-
   let InOperandList = !con(BaseInOperandList,
                            (ins B32:$scale_a, B32:$scale_b,
                                 B32:$decompress_b));
-
   let AsmString = DecompressPrefix
                   # ".block_scale"
                   # DecompressLutStr
@@ -6926,7 +6908,6 @@ class Tcgen05MMADecompressBlockScaleInst<string ASpace, string Kind,
 
   dag IntrinsicPattern = !con(!foreach(tmp, BasePatternArgs, !subst(ins, Intrin, tmp)),
                               (Intrin i32:$scale_a, i32:$scale_b, i32:$decompress_b));
-
   dag FlagOperands = (Intrin (i32 CtaGroup),
                              (i32 CollectorUsageAVal),
                              (i32 CollectorUsageBVal));

>From 59a049a2e458eef98d8592d29b6618c77ff58890 Mon Sep 17 00:00:00 2001
From: Kirill Vedernikov <kvedernikov at nvidia.com>
Date: Wed, 19 Aug 2026 18:31:11 +0200
Subject: [PATCH 4/5] [NVVM][NVPTX] Improve decompress_b according to feedback

---
 llvm/include/llvm/IR/IntrinsicsNVVM.td   | 66 +++++++++++-------------
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 36 +++++++------
 2 files changed, 48 insertions(+), 54 deletions(-)

diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 7aaff46001e9d..69905899f3007 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1157,6 +1157,12 @@ class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num,
   LLVMType type = LLVMType<!cast<ValueType>("v"#veclen#ElemType)>;
 }
 
+class NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<ArgIndex ColIdx, string Buffer> {
+  IntrinsicProperty Prop =
+      ArgInfo<ColIdx, [ArgName<"collector_" # Buffer>,
+                       ImmArgPrinter<"printTcgen05CollectorUsageOp">]>;
+}
+
 class NVVM_TCGEN05_MMA_BASE<string Space, bit IsSparse, string Kind = ""> {
   LLVMType a_operand_type = !if(!eq(Space, "tensor"),
                                   llvm_tmem_ptr_ty, llvm_i64_ty);
@@ -1237,40 +1243,27 @@ class NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED<string Kind, string ScaleVecSize> {
 // tcgen05.mma intrinsics with decompress_b extension
 //
 
-class NVVM_TCGEN05_MMA_DECOMPRESS_BASE<string Space, string Kind>:
+class NVVM_TCGEN05_MMA_DECOMPRESS_BASE<string Space, string Kind,
+                                       string Modifier = "">:
     NVVM_TCGEN05_MMA_BASE<Space, /* IsSparse= */ 0, Kind> {
-  string DecompressPrefix = Prefix # SpSpaceKindStr;
-  string DecompressSuffix = ".decompress_b";
+  string name = Prefix # SpSpaceKindStr # Modifier # ".decompress_b";
+  string intr_name = IntrinsicName<name>.intr_name;
+  string record_name = IntrinsicName<name>.record_name;
 }
 
 class NVVM_TCGEN05_MMA_DECOMPRESS<string Space, string Kind>:
-    NVVM_TCGEN05_MMA_DECOMPRESS_BASE<Space, Kind> {
-  string name = DecompressPrefix # DecompressSuffix;
-  string intr_name = IntrinsicName<name>.intr_name;
-  string record_name= IntrinsicName<name>.record_name;
-}
+    NVVM_TCGEN05_MMA_DECOMPRESS_BASE<Space, Kind>;
 
 class NVVM_TCGEN05_MMA_DECOMPRESS_DISABLE_OUTPUT_LANE<string Space,
                                                       string Kind,
                                                       int CtaGroup>:
-    NVVM_TCGEN05_MMA_DECOMPRESS_BASE<Space, Kind> {
-  string name = DecompressPrefix
-                # ".disable_output_lane.cg" # CtaGroup
-                # DecompressSuffix;
-  string intr_name = IntrinsicName<name>.intr_name;
-  string record_name= IntrinsicName<name>.record_name;
-}
+    NVVM_TCGEN05_MMA_DECOMPRESS_BASE<Space, Kind,
+                                     ".disable_output_lane.cg" # CtaGroup>;
 
 class NVVM_TCGEN05_MMA_DECOMPRESS_BLOCKSCALE<string Space, string Kind,
                                              string ScaleVecSize>:
-    NVVM_TCGEN05_MMA_DECOMPRESS_BASE<Space, Kind> {
-  string name = DecompressPrefix
-                # ".block_scale"
-                # ScaleVecSize
-                # DecompressSuffix;
-  string intr_name = IntrinsicName<name>.intr_name;
-  string record_name= IntrinsicName<name>.record_name;
-}
+    NVVM_TCGEN05_MMA_DECOMPRESS_BASE<Space, Kind,
+                                     ".block_scale" # ScaleVecSize>;
 
 class TexVector<string name, list<LLVMType> types> {
   string Name = name;
@@ -3472,8 +3465,8 @@ foreach sparse = [0, 1] in {
            Range<collector_b_idx, 0, 4>,
            ArgInfo<kind_idx, [ArgName<"kind">, ImmArgPrinter<"printTcgen05MMAKind">]>,
            ArgInfo<cta_group_idx, [ArgName<"cta_group">]>,
-           ArgInfo<collector_a_idx, [ArgName<"collector_a">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>,
-           ArgInfo<collector_b_idx, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]
+           NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_a_idx, "a">.Prop,
+           NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_b_idx, "b">.Prop]
         );
 
         def mma.record_name:
@@ -3517,8 +3510,8 @@ foreach sparse = [0, 1] in {
                      Range<collector_a_idx, 0, !if(!eq(ashift, 1), 2, 4)>,
                      Range<collector_b_idx, 0, 4>,
                      ArgInfo<kind_idx, [ArgName<"kind">, ImmArgPrinter<"printTcgen05MMAKind">]>,
-                     ArgInfo<collector_a_idx, [ArgName<"collector_a">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>,
-                     ArgInfo<collector_b_idx, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]
+                     NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_a_idx, "a">.Prop,
+                     NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_b_idx, "b">.Prop]
                   );
 
           def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
@@ -3555,8 +3548,8 @@ foreach sparse = [0, 1] in {
                          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">]>,
-                         ArgInfo<collector_usage_b, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]),
+                         NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_usage_a, "a">.Prop,
+                         NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_usage_b, "b">.Prop]),
             mma.intr_name>;
         }
       } // scale_vec_size
@@ -3591,8 +3584,7 @@ foreach sparse = [0, 1] in {
                             ImmArgPrinter<"printTcgen05MMAKind">]>,
          ArgInfo<collector_buffer_b_idx, [ArgName<"collector_b_buffer">,
                                           ImmArgPrinter<"printTcgen05MMACollectorBBuffer">]>,
-         ArgInfo<collector_usage_b_op_idx, [ArgName<"collector_b">,
-                                            ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]
+         NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_usage_b_op_idx, "b">.Prop]
       );
 
       def mma.record_name:
@@ -3625,8 +3617,8 @@ foreach space = ["tensor", "shared"] in {
        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">]>,
-       ArgInfo<collector_usage_b, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]);
+       NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_usage_a, "a">.Prop,
+       NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_usage_b, "b">.Prop]);
 
   def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
                           intrinsic_properties, mma.intr_name>;
@@ -3656,8 +3648,8 @@ foreach space = ["tensor", "shared"] in {
         mma.common_intr_props,
         [Range<collector_usage_a, 0, 4>,
          Range<collector_usage_b, 0, 4>,
-         ArgInfo<collector_usage_a, [ArgName<"collector_a">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>,
-         ArgInfo<collector_usage_b, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]);
+         NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_usage_a, "a">.Prop,
+         NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_usage_b, "b">.Prop]);
 
     def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
                             intrinsic_properties, mma.intr_name>;
@@ -3690,8 +3682,8 @@ foreach space = ["tensor", "shared"] in {
        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">]>,
-       ArgInfo<collector_usage_b, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]);
+       NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_usage_a, "a">.Prop,
+       NVVM_TCGEN05_MMA_COLLECTOR_ARGPROP<collector_usage_b, "b">.Prop]);
 
   def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
                           intrinsic_properties, mma.intr_name>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 8ab3219b58bc1..edaa2b56a833b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6366,6 +6366,8 @@ foreach dim = ["x", "y", "z"] in {
 // tcgen05.mma Helpers
 //
 
+defvar CollectorUsages = ["discard", "lastuse", "fill", "use"];
+
 class Tcgen05MMABase<bit IsSparse, string ASpace, string Kind, int CtaGroup,
                      string CollectorUsage, string CollectorUsageB = "discard"> :
         NVPTXInst<(outs), (ins), "?", []>,
@@ -6478,11 +6480,11 @@ foreach sparse = [0, 1] in {
   foreach space = ["tensor", "shared"] in {
     foreach kind = ["f16", "tf32", "f8f6f4", "i8", "ti16"] in {
       foreach cta_group = [1, 2] in {
-        foreach collector_usage_a = ["discard", "lastuse", "fill", "use"] in {
+        foreach collector_usage_a = CollectorUsages in {
           foreach scale_input_d = !if(!or(!eq(kind, "f16"),
                                           !eq(kind, "tf32")), [0, 1], [0]) in {
             foreach ashift = !if(!eq(space, "tensor"), [0, 1], [0]) in {
-              foreach collector_usage_b = ["discard", "lastuse", "fill", "use"] in {
+              foreach collector_usage_b = CollectorUsages in {
                 def : Tcgen05MMAInst<sparse, space, kind, cta_group,
                                     collector_usage_a, collector_usage_b,
                                     scale_input_d, ashift>;
@@ -6615,11 +6617,11 @@ foreach sparse = [0, 1] in {
   foreach space = ["tensor", "shared"] in {
     foreach kind = ["f16", "tf32", "f8f6f4", "i8", "ti16"] in {
       foreach cta_group = [1, 2] in {
-        foreach collector_usage_a = ["fill", "use", "lastuse", "discard"] in {
+        foreach collector_usage_a = CollectorUsages in {
           foreach scale_input_d = !if(!or(!eq(kind, "f16"),
                                           !eq(kind, "tf32")), [0, 1], [0]) in {
             foreach ashift = !if(!eq(space, "tensor"), [0, 1], [0]) in {
-              foreach collector_usage_b = ["discard", "lastuse", "fill", "use"] in {
+              foreach collector_usage_b = CollectorUsages in {
                 def :
                   Tcgen05MMADisableOutputLaneInst<sparse, space, kind, cta_group,
                                                   collector_usage_a, collector_usage_b,
@@ -6692,8 +6694,8 @@ foreach sparse = [0, 1] in {
     foreach kind = ["mxf8f6f4", "mxf4", "mxf4nvf4"] in {
       foreach scale_vec_size = ["", ".block16", ".block32"] in {
         foreach cta_group = [1, 2] in {
-          foreach collector_usage_a = ["fill", "use", "lastuse", "discard"] in {
-            foreach collector_usage_b = ["discard", "lastuse", "fill", "use"] in {
+          foreach collector_usage_a = CollectorUsages in {
+            foreach collector_usage_b = CollectorUsages in {
               if NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED<kind, scale_vec_size>.ret then {
                 def : Tcgen05MMABlockScaleInst<sparse, space, kind, cta_group,
                                                scale_vec_size, collector_usage_a,
@@ -6757,7 +6759,7 @@ foreach sparse = [0, 1] in {
   foreach space = ["shared", "tensor"] in {
     foreach kind = ["f16", "tf32", "f8f6f4", "i8", "ti16"] in {
       foreach collector_buffer_b = [0, 1, 2, 3] in {
-        foreach collector_usage_op = ["discard", "fill", "use", "lastuse"] in {
+        foreach collector_usage_op = CollectorUsages in {
           foreach zero_col_mask = [0, 1] in {
               def : Tcgen05MMAWSInst<sparse, space, kind, collector_buffer_b,
                                      collector_usage_op, zero_col_mask>;
@@ -6814,8 +6816,8 @@ class Tcgen05MMADecompressInst<string ASpace, string Kind, int CtaGroup,
 // tcgen05.mma decompress
 foreach space = ["tensor", "shared"] in {
   foreach cta_group = [1, 2] in {
-    foreach collector_usage_a = ["discard", "lastuse", "fill", "use"] in {
-      foreach collector_usage_b = ["discard", "lastuse", "fill", "use"] in {
+    foreach collector_usage_a = CollectorUsages in {
+      foreach collector_usage_b = CollectorUsages in {
         def : Tcgen05MMADecompressInst<space, "f8f6f4", cta_group,
                                        collector_usage_a, collector_usage_b>;
       } // collector_usage_b
@@ -6829,11 +6831,11 @@ class Tcgen05MMADecompressDisableOutputLaneInst<string ASpace, string Kind,
                                                 string CollectorUsageB> :
     Tcgen05MMADecompressBase<ASpace, Kind, CtaGroup, CollectorUsageA,
                              CollectorUsageB> {
-  SDNode Opcode = Tcgen05MMADisableOutputLaneSDNode<0, // IsSparse
+  SDNode Opcode = Tcgen05MMADisableOutputLaneSDNode</*IsSparse=*/ 0,
                                                     ASpace, CtaGroup,
-                                                    0,  // IsScaleInputD
-                                                    0,  // IsAShift
-                                                    1>; // IsDecompressB
+                                                    /*IsScaleInputD=*/ 0,
+                                                    /*IsAShift=*/ 0,
+                                                    /*IsDecompressB=*/ 1>;
 
   // disable output lane
   int DisableOutputLaneVecSize = !mul(4, CtaGroup);
@@ -6872,8 +6874,8 @@ class Tcgen05MMADecompressDisableOutputLaneInst<string ASpace, string Kind,
 // tcgen05.mma decompress disable output lane
 foreach space = ["tensor", "shared"] in {
   foreach cta_group = [1, 2] in {
-    foreach collector_usage_a = ["discard", "lastuse", "fill", "use"] in {
-      foreach collector_usage_b = ["discard", "lastuse", "fill", "use"] in {
+    foreach collector_usage_a = CollectorUsages in {
+      foreach collector_usage_b = CollectorUsages in {
         def : Tcgen05MMADecompressDisableOutputLaneInst<
                   space, "f8f6f4", cta_group, collector_usage_a,
                   collector_usage_b>;
@@ -6918,8 +6920,8 @@ class Tcgen05MMADecompressBlockScaleInst<string ASpace, string Kind,
 // tcgen05.mma decompress block_scale
 foreach space = ["tensor", "shared"] in {
   foreach cta_group = [1, 2] in {
-    foreach collector_usage_a = ["discard", "lastuse", "fill", "use"] in {
-      foreach collector_usage_b = ["discard", "lastuse", "fill", "use"] in {
+    foreach collector_usage_a = CollectorUsages in {
+      foreach collector_usage_b = CollectorUsages in {
         def : Tcgen05MMADecompressBlockScaleInst<
                   space, "mxf8f6f4", cta_group, collector_usage_a,
                   collector_usage_b, ".block32">;

>From 2f2c60f90d3d9cf7eeee6f916f9e9c86af9a5030 Mon Sep 17 00:00:00 2001
From: Kirill Vedernikov <kvedernikov at nvidia.com>
Date: Thu, 20 Aug 2026 09:28:28 +0200
Subject: [PATCH 5/5] [NVPTX] Combine tcgen05.mma decompress variant generation

---
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 32 +++++-------------------
 1 file changed, 6 insertions(+), 26 deletions(-)

diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index edaa2b56a833b..5e01b251ff9fd 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6813,18 +6813,6 @@ class Tcgen05MMADecompressInst<string ASpace, string Kind, int CtaGroup,
   let Pattern = [!con(IntrinsicPattern, FlagOperands)];
 }
 
-// tcgen05.mma decompress
-foreach space = ["tensor", "shared"] in {
-  foreach cta_group = [1, 2] in {
-    foreach collector_usage_a = CollectorUsages in {
-      foreach collector_usage_b = CollectorUsages in {
-        def : Tcgen05MMADecompressInst<space, "f8f6f4", cta_group,
-                                       collector_usage_a, collector_usage_b>;
-      } // collector_usage_b
-    } // collector_usage_a
-  } // cta_group
-} // space
-
 class Tcgen05MMADecompressDisableOutputLaneInst<string ASpace, string Kind,
                                                 int CtaGroup,
                                                 string CollectorUsageA,
@@ -6871,19 +6859,6 @@ class Tcgen05MMADecompressDisableOutputLaneInst<string ASpace, string Kind,
   let Pattern = [!con(IntrinsicPattern, FlagOperands)];
 }
 
-// tcgen05.mma decompress disable output lane
-foreach space = ["tensor", "shared"] in {
-  foreach cta_group = [1, 2] in {
-    foreach collector_usage_a = CollectorUsages in {
-      foreach collector_usage_b = CollectorUsages in {
-        def : Tcgen05MMADecompressDisableOutputLaneInst<
-                  space, "f8f6f4", cta_group, collector_usage_a,
-                  collector_usage_b>;
-      } // collector_usage_b
-    } // collector_usage_a
-  } // cta_group
-} // space
-
 class Tcgen05MMADecompressBlockScaleInst<string ASpace, string Kind,
                                          int CtaGroup, string CollectorUsageA,
                                          string CollectorUsageB,
@@ -6917,11 +6892,16 @@ class Tcgen05MMADecompressBlockScaleInst<string ASpace, string Kind,
   let Pattern = [!con(IntrinsicPattern, FlagOperands)];
 }
 
-// tcgen05.mma decompress block_scale
+// tcgen05.mma decompress variants
 foreach space = ["tensor", "shared"] in {
   foreach cta_group = [1, 2] in {
     foreach collector_usage_a = CollectorUsages in {
       foreach collector_usage_b = CollectorUsages in {
+        def : Tcgen05MMADecompressInst<space, "f8f6f4", cta_group,
+                                       collector_usage_a, collector_usage_b>;
+        def : Tcgen05MMADecompressDisableOutputLaneInst<
+                  space, "f8f6f4", cta_group, collector_usage_a,
+                  collector_usage_b>;
         def : Tcgen05MMADecompressBlockScaleInst<
                   space, "mxf8f6f4", cta_group, collector_usage_a,
                   collector_usage_b, ".block32">;