[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 b53921f6b4c02c0fc9e0423b26a5ef23dbc885c4 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/3] [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 4357ad367d269..d7b5a16ba095d 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;
@@ -3565,6 +3604,101 @@ foreach sparse = [0, 1] in {
   } // space
 } // sparse
 
+//
+// tcgen05.mma decompress intrinsics
+//
+foreach space = ["tensor", "shared"] in {
+  defvar mma = NVVM_TCGEN05_MMA_DECOMPRESS<space, "f8f6f4">;
+  defvar args = !listconcat(
+                  mma.common_args,
+                  [llvm_tmem_ptr_ty] // decompress_b
+                );
+  defvar flags = [llvm_i32_ty,  // cta_group
+                  llvm_i32_ty,  // collector_usage_a
+                  llvm_i32_ty]; // collector_usage_b
+  defvar nargs = !size(args);
+  defvar cta_group = ArgIndex<nargs>;
+  defvar collector_usage_a = ArgIndex<!add(nargs, 1)>;
+  defvar collector_usage_b = ArgIndex<!add(nargs, 2)>;
+
+  defvar intrinsic_properties = !listconcat(
+      mma.common_intr_props,
+      [Range<cta_group, 1, 3>,
+       Range<collector_usage_a, 0, 4>,
+       Range<collector_usage_b, 0, 4>,
+       ArgInfo<cta_group, [ArgName<"cta_group">]>,
+       ArgInfo<collector_usage_a, [ArgName<"collector_a">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>,
+       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 e788b0e44041f..86cf37baeaede 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 02bcfcb6af926..828e97512ece3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6489,7 +6489,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);
 
@@ -6499,8 +6500,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
   );
@@ -6510,14 +6512,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]>;
 
 //
@@ -6753,6 +6757,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 dbdce40ab843f5316bfee1300f78535b36a2320b 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/3] [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 828e97512ece3..420a6ce8d8f97 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6763,9 +6763,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));
@@ -6899,9 +6897,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 4d99fdb16ea60cb595522cfdc81b1534a42bb2ad 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/3] [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 420a6ce8d8f97..496133b95fe03 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6762,18 +6762,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
@@ -6786,11 +6783,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
@@ -6799,7 +6794,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));
 
@@ -6832,30 +6826,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
@@ -6866,7 +6852,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));
 
@@ -6894,15 +6879,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
@@ -6915,7 +6897,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));