smem_atom_layoutSFB branches on majorSFA instead of majorSFB

Description

Sm1xxBlockwiseScaleConfig::smem_atom_layoutSFB (include/cutlass/detail/blockwise_scale_layout.hpp:110) branches on majorSFA instead of majorSFB:

template<typename CtaShape_MNK>
static constexpr auto
smem_atom_layoutSFB(CtaShape_MNK cta_shape_mnk) {
  static_assert(cute::is_static_v<CtaShape_MNK>, "Expect static CTA shape");
  auto strides = [&]() CUTLASS_LAMBDA_FUNC_INLINE {
    if constexpr (majorSFA == UMMA::Major::MN) {          // <-- should be majorSFB
      return make_stride(make_stride(_0{}, _1{}), make_stride(_0{}, Int<cute::ceil_div(size<1>(CtaShape_MNK{}), SFVecSizeN)>{}));
    }
    else {
      return make_stride(make_stride(_0{}, Int<cute::ceil_div(size<2>(CtaShape_MNK{}), SFVecSizeK)>{}), make_stride(_0{}, _1{}));
    }
  }();
  ...
}

The SFB shared-memory atom is therefore built from the A operand's scale-factor major. Every other SFB quantity in the same struct uses majorSFB: the global layout tile_atom_to_shape_SFB (line 163) branches correctly on majorSFB, and the sibling smem_atom_layoutSFA (line 81) branches on majorSFA. Whenever a kernel configures different majors for A and B scales, the SFB smem layout and the SFB global layout describe different orderings.

All four consumers derive the two majors independently from user-provided scale layouts and call this function:

• sm100_mma_warpspecialized_blockwise_scaling.hpp:163
• sm100_mma_array_warpspecialized_blockwise_scaling.hpp:179
• sm90_mma_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp:166
• sm90_mma_array_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp:166

Reproduction

Host program against the current header (g++, CTA shape (_128,_256,_512)):

(MN,MN) smem SFB strides: ((_0,_1),(_0,_2))
(K, MN) smem SFB strides: ((_0,_4),(_0,_1))   // majorSFB == MN, but got the K-major pattern
(K, K ) smem SFB strides: ((_0,_4),(_0,_1))
(MN,K ) smem SFB strides: ((_0,_1),(_0,_2))   // majorSFB == K,  but got the MN-major pattern

Both mixed-major configurations produce an SFB smem atom whose stride pattern contradicts tile_atom_to_shape_SFB for the same config (that one follows majorSFB).

Impact

• (SFA=MN, SFB=K) appears in shipped examples (e.g. example 67 uses <..., Major::MN, Major::K>); there the scalar cp.async copies happen to place values coordinate-correctly, which masks the wrong atom.
• The mirror case (SFA=K, SFB=MN) does not compile today: downstream vectorized-copy selection built from the true N-major global layout fails partitioning onto the wrongly K-major smem atom with Copy_Traits: dst failed to vectorize into registers. Layout is incompatible with this CopyOp.

So mixed-major blockwise kernels are either silently using an inconsistent smem descriptor or cannot be instantiated at all, depending on combination.

Suggested fix

One token: test majorSFB == UMMA::Major::MN in smem_atom_layoutSFB.

Start at include/cutlass/detail/blockwise_scale_layout.hpp:110 and compare smem_atom_layoutSFB with tile_atom_to_shape_SFB and smem_atom_layoutSFA. Check the four listed consumers and reproduce the mixed-major cases with the host program; done means SFB shared-memory strides follow majorSFB and the reported configurations agree with the global layout.