[QST] Why did I get a wrong result from GemmGrouped?

I'm using GemmGrouped in this way:

using GemmKernel = typename cutlass::gemm::kernel::DefaultGemmGrouped<
    cutlass_t,                                      // Element A
    cutlass::layout::RowMajor,                      // Layout A
    cutlass::ComplexTransform::kNone,               //
    1,                                              // Granularity A
    cutlass_t,                                      // Element B
    cutlass::layout::RowMajor,                      // Layout B
    cutlass::ComplexTransform::kNone,               //
    1,                                              // Granularity B
    cutlass_t,                                      // Element C&D
    cutlass::layout::RowMajor,                      // Layout C&D
    float,                                          // Element Accumulator
    cutlass::arch::OpClassTensorOp,                 // Operator Class Tag
    cutlass::arch::Sm80,                            // Architecture
    cutlass::gemm::GemmShape<16, 64, 64>,           // Thread Block Shape
    cutlass::gemm::GemmShape<16, 16, 64>,           // Warp Shape
    cutlass::gemm::GemmShape<16, 8, 16>,            // Instruction Shape
    LinearCombination<cutlass_t, 1, float, float>,  // Epilogue
    GemmIdentityThreadblockSwizzle<>,              // Swizzling Operator
    2                                               // Stages
    >::GemmKernel;

using EpilogueOutputOp = typename GemmKernel::Epilogue::OutputOp;
typename EpilogueOutputOp::Params epilogue_op(1.0, 0.0);

using GemmGrouped = cutlass::gemm::device::GemmGrouped<GemmKernel>;
typename GemmGrouped::Arguments args_(
    all_problems, num_problems, 512, epilogue_op, ptr_X, ptr_W, ptr_Y,
    ptr_Y, ld_X, ld_W, ld_Y, ld_Y);

GemmGrouped gemm;

gemm.initialize(args, nullptr, stream);
gemm.run(stream);

With num_problems = 1, M = 2035, N = 48, K = 3584
all_problems[0] = cutlass::gemm::GemmCoord(2035, 48, 3584)
ptr_X[0] = matrix A, ptr_W[0] = matrix B and ptr_Y[0] = matrix C
ld_X[0] = K, ld_W[0] = N, ld_Y[0] = N
cutlass_t = cutlass::half_t
and the GPU I used is A100
cutlass version: 3.5.0
nvcc version: 12.4

The calculation I'm expecting is C = matmul(A, B) so I set alpha = 1.0 and beta = 0.0 for epilogue_op.
The shape of C is (2035, 48), however only elements in the first 16 columns were correct, all other elements of C were incorrect.

I spent a lot of time on tracing the execution procedure with cuda-gdb, and I found something is wrong in loading warp fragments of matrix B.
The loading procedure is done by code below:
https://github.com/NVIDIA/cutlass/blob/main/include/cutlass/gemm/warp/mma_tensor_op_tile_iterator.h#L395-L415

for (int s = 0; s < Policy::LdsmIterations::kStrided; ++s) {
  for (int c = 0; c < Policy::LdsmIterations::kContiguous; ++c) {
    int access_idx = c + s * Policy::LdsmIterations::kContiguous;

    AccessType const *source_ptr =
        pointer_[c % kPointerCount] +
        Layout::TileShape::kContiguous * (c / kPointerCount) +
        Policy::kLdsmOpInner * Policy::LdsmShape::kStrided * s * stride_;

    char const *source_byte_ptr = reinterpret_cast<char const *>(source_ptr) + byte_offset + byte_offset_;

    cutlass::arch::ldsm<layout::ColumnMajor, Policy::LdsmShape::kCount>(
      fetch_ptr[access_idx],
      source_byte_ptr
    );
  }
}

Before the 'ldsm' here, I printed the values in global memory from address 'source_byte_ptr' of block 0 thread 0, and the values and 'source_byte_ptr' of some threads are shown below:

Looks like cutlass has automatically done padding and swizzle for matrix B, and I think the memory layout of B looks correct. The ''source_byte_ptr' of 32 threads in Warp 0(thread 0 - 31) were all pointed to B[threadIdx.x][0:8], however for Warp 1, thread 34 got the address of 8 zeros...and many threads in Warp 1-3 got wrong memory addresses in my point of view.
As a result, after the 'ldsm' here, thread 0 got 4 32-bit values consisted of {B[0][0], B[1][0], B[8][0], B[9][0], B[0][8], B[1][8], B[8][8], B[9][8]}.

According to the figure above showing the element layout of an m16n8k16 mma instruction, threads in Warp 0 got correct fragments. But threads in other Warps would get wrong fragments, for example, thread 33 got all zeros.
I think that's the reason why only the part belonged to Warp 0 has correct values in matrix C, but I still don't know why the fragment loading procedure would go wrong.

Reproduce the A100 case with CUTLASS 3.5.0, CUDA 12.4, num_problems=1, and the provided M/N/K values. Start in include/cutlass/gemm/warp/mma_tensor_op_tile_iterator.h around lines 395-415 and inspect the GemmGrouped fragment-loading path. Done means explaining the incorrect non-Warp-0 fragments and identifying the change or configuration needed for all 48 output columns to match the expected matmul.