[BUG] Incorrect results with Hopper Mixed-input Kernel using 8x1x1 cluster

I have created a reproducer for the failure I am seeing for the below template instantiation that you can get using the cmake.

Reproducer by changing the example 55

Reproducer is in the branch here.

./examples/55_hopper_mixed_dtype_gemm/55_hopper_mixed_dtype_gemm --mode=0
Running in no scale mode.
  Problem Size: 5120x4096x4096x1
  Avg runtime: 0.637257 ms
  GFLOPS: 269591
  Disposition: Failed

We originally found the issue while using the below instance in our stack that can be generated using the below cmake.

cmake

cmake -DCMAKE_BUILD_TYPE:STRING=Release -DCUTLASS_NVCC_ARCHS:STRING=90a -DCUTLASS_NVCC_KEEP:STRING=OFF -DCUTLASS_ENABLE_F16C:STRING=ON -DCUTLASS_LIBRARY_INSTANTIATION_LEVEL:STRING=max -DCUTLASS_LIBRARY_KERNELS:STRING=s64x32x16gemm_e4m3_bf16_f32_bf16_bf16*tnn*align16,s64x64x16gemm_e4m3_bf16_f32_bf16_bf16*tnn*align16,s64x128x16gemm_e4m3_bf16_f32_bf16_bf16*tnn*align16 -DCUTLASS_LIBRARY_IGNORE_KERNELS:STRING=gemm_grouped*,gemm_planar*,shfl,scl,sclzr,

Failing e4m3_bf16 cvt-only template instance from the above cmake.

using cutlass3x_sm90_tensorop_s64x32x16gemm_e4m3_bf16_f32_bf16_bf16_cvt_128x32x128_8x1x1_0_tnn_align16_warpspecialized_cooperative_epi_tma_epilogue =
  typename cutlass::epilogue::collective::CollectiveBuilder<
    cutlass::arch::Sm90, cutlass::arch::OpClassTensorOp,
    cute::Shape<cute::_128, cute::_32, cute::_128>,
    cute::Shape<cute::_8, cute::_1, cute::_1>,
    cutlass::epilogue::collective::EpilogueTileAuto,
    float, float,
    cutlass::bfloat16_t, cutlass::layout::ColumnMajor, 8,
    cutlass::bfloat16_t, cutlass::layout::ColumnMajor, 8,
    cutlass::epilogue::TmaWarpSpecializedCooperative,
    cutlass::epilogue::fusion::LinearCombination<
      cutlass::bfloat16_t,
      float,
      cutlass::bfloat16_t,
      float
    >
  >::CollectiveOp;

using cutlass3x_sm90_tensorop_s64x32x16gemm_e4m3_bf16_f32_bf16_bf16_cvt_128x32x128_8x1x1_0_tnn_align16_warpspecialized_cooperative_epi_tma_mainloop =
  typename cutlass::gemm::collective::CollectiveBuilder<
    cutlass::arch::Sm90, cutlass::arch::OpClassTensorOp,
    cutlass::float_e4m3_t, cutlass::layout::RowMajor, 16,
    cutlass::bfloat16_t, cutlass::layout::ColumnMajor, 8,
    float,
    cute::Shape<cute::_128, cute::_32, cute::_128>,
    cute::Shape<cute::_8, cute::_1, cute::_1>,
    cutlass::gemm::collective::StageCountAutoCarveout<static_cast<int>(sizeof(typename cutlass3x_sm90_tensorop_s64x32x16gemm_e4m3_bf16_f32_bf16_bf16_cvt_128x32x128_8x1x1_0_tnn_align16_warpspecialized_cooperative_epi_tma_epilogue::SharedStorage))>,
    cutlass::gemm::KernelTmaWarpSpecializedCooperative
  >::CollectiveOp;

// Gemm operator cutlass3x_sm90_tensorop_s64x32x16gemm_e4m3_bf16_f32_bf16_bf16_cvt_128x32x128_8x1x1_0_tnn_align16_warpspecialized_cooperative_epi_tma
using cutlass3x_sm90_tensorop_s64x32x16gemm_e4m3_bf16_f32_bf16_bf16_cvt_128x32x128_8x1x1_0_tnn_align16_warpspecialized_cooperative_epi_tma_base = cutlass::gemm::kernel::GemmUniversal<
    cute::Shape<int,int,int,int>,
    cutlass3x_sm90_tensorop_s64x32x16gemm_e4m3_bf16_f32_bf16_bf16_cvt_128x32x128_8x1x1_0_tnn_align16_warpspecialized_cooperative_epi_tma_mainloop,
    cutlass3x_sm90_tensorop_s64x32x16gemm_e4m3_bf16_f32_bf16_bf16_cvt_128x32x128_8x1x1_0_tnn_align16_warpspecialized_cooperative_epi_tma_epilogue,
    void>;

Why correctness isn't checked during the profiling runs?

../build/tools/profiler/cutlass_profiler --kernels=cutlass3x_sm90_tensorop_s64x32x16gemm_e4m3_bf16_f32_bf16_bf16_cvt_128x32x128_8x1x1_0_tnn_align16_warpspecialized_cooperative_epi_tma --verification-enabled=true



=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: gemm
       Operation: cutlass3x_sm90_tensorop_s64x32x16gemm_e4m3_bf16_f32_bf16_bf16_cvt_128x32x128_8x1x1_0_tnn_align16_warpspecialized_cooperative_epi_tma

          Status: Success
    Verification: ON
     Disposition: Not verified

reference_device: Not run
          cuBLAS: Not run
           cuDNN: Not run

       Arguments: --gemm_kind=universal --m=1024 --n=1024 --k=1024 --A=fe4m3:row --B=bf16:column --C=bf16:column --D=bf16:column  \
                  --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1 --batch_count=1 --raster_order=heuristic  \
                  --runtime_input_datatype_a=invalid --runtime_input_datatype_b=invalid --use_pdl=false --enable_sm90_mixed_dtype_shuffle_test=false  \
                  --swizzle_size=1 --op_class=tensorop --accum=f32 --cta_m=128 --cta_n=32 --cta_k=128 --cluster_m=1 --cluster_n=1  \
                  --cluster_k=1 --cluster_m_fallback=0 --cluster_n_fallback=0 --cluster_k_fallback=0 --stages=8 --warps_m=4  \
                  --warps_n=2 --warps_k=1 --inst_m=64 --inst_n=32 --inst_k=16 --min_cc=90 --max_cc=90

           Bytes: 5242880  bytes
           FLOPs: 2149580800  flops
           FLOPs/Byte: 410

         Runtime: 0.0215533  ms
          Memory: 226.546 GiB/s

            Math: 99733.3 GFLOP/s


=============================

Start with the reproducer in examples/55_hopper_mixed_dtype_gemm and the linked sm90_mixed_input_incorrect_results branch, then run the provided CMake configuration and example command. Compare the 8x1x1 mixed-input result with the expected output and inspect the profiler command's verification status. Done means the reproducer reports correct results and profiler verification runs for this template.