NVIDIA / NVIDIA/cutlass

[BUG] SM120 sparse block-scaled GEMM: 256x128x256 and 128x128x128 tiles fail with a raw template backtrace instead of a static_assert

Open
#3,393 1 comment 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

inactive-30d
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Describe the bug

The SM120 block-scaled 2:4-sparse collective (MainloopSm120TmaWarpSpecializedSparseBlockScaled, NVFP4 nv_float4_t A/B) rejects two tile shapes at compile time that are natural choices for large training-style GEMMs:

  • 256x128x256 (larger M): fails with a ~200-line template error rooted at cute/numeric/arithmetic_tuple.hpp(372): error: no operator "==" matches these operands (const int32_t == cute::R<1, 16>).
  • 128x128x128 (narrower K): fails with cute/layout.hpp(1792): static assertion failed: "tile_to_shape: block shape does not divide the target shape." plus sm120_gemm_tma_warpspecialized_cooperative_asymmetric_dma.hpp(181): static assertion failed: "SMEM usage exceeded capacity."

The failure surfaces as a raw template backtrace rather than a static_assert naming the supported tile set. The stock 128x128x256 builds and runs fine. Reproduced on main @ d01487db and identically on tags v4.4.0 and v4.5.3.

Steps/Code to reproduce bug

One-line tile substitution in the in-tree example:

cd cutlass
sed 's|Shape<_128,_128,_256>|Shape<_256,_128,_256>|' \
  examples/80_blackwell_geforce_sparse_gemm/80b_blackwell_geforce_nvfp4_nvfp4_sparse_gemm.cu > /tmp/probe.cu
nvcc -O3 -DNDEBUG -std=c++17 --generate-code=arch=compute_121a,code=[sm_121a] \
  -DCUTLASS_ENABLE_TENSOR_CORE_MMA=1 --expt-relaxed-constexpr \
  -Icutlass/include -Icutlass/examples/common -Icutlass/tools/util/include \
  -c /tmp/probe.cu -o /tmp/probe.o
# same with Shape<_128,_128,_128>

compute_120a behaves the same; I ran on sm_121.

Expected behavior

Either the tile shapes compile, or a first-line static_assert states the supported SM120 sparse block-scaled tile set, so the unsupported-tile case is diagnosed in one line instead of a long cute backtrace.

Environment details

  • CUTLASS main @ d01487db (also verified on tags v4.4.0 / v4.5.3)
  • CUDA 13.1, nvcc from NGC PyTorch 26.01 container (bare-metal Docker)
  • GPU: GB10 (DGX Spark), sm_121a. Driver 580.95.05.

Additional context

Measured motivation, from profiling NVFP4 dense vs 2:4-sparse block-scaled GEMMs on this GPU across large training shapes (M up to 28672, N/K 4096 to 16384, bf16 D): the best dense config for fat-M/N shapes is the 256x128 tile family, which is inexpressible for the sparse collective, so the sparse kernel is limited to a smaller tile than the best dense config. With the one legal big tile (128x128x256) the sparse kernel reaches only about 44-54% of its 2x sparse tensor-op roof (ncu: 1 CTA/SM both arms, SMEM 99-101 KB/block), and measured sparse:dense speedups land at 1.1-1.5x instead of near 2x. Tile-shape freedom is the main in-library lever against that, which is what is locked. Side note: 128x64x256 (narrow N) compiles on current main but not on v4.4.0, so narrow-N support was added between 4.4 and 4.5.

Questions:

  1. Are 256-row and K=128 tiles intended-unsupported for the SM120 sparse block-scaled collective (metadata/TMA layout constraints?), or is this a gap that could be closed?
  2. If intended: would you take a PR adding an early static_assert naming the supported tile set for this collective?
  3. If not intended: is larger-M sparse tile support on any roadmap? Happy to test candidate branches on sm_121 hardware.

Trimmed compile logs for both failing shapes:

256x128x256 compile error (tail) 
      return bool_constant<((Ns == Ms) && ...)>{} && t.value() == u.value();
                                                               ^
/nvfp4/upstream/clones/cutlass-main/include/cute/numeric/arithmetic_tuple.hpp(372): note #3328-D: built-in operator==(<nullptr>, <nullptr>) does not match because argument #1 does not match parameter
      return bool_constant<((Ns == Ms) && ...)>{} && t.value() == u.value();
                                                               ^
          detected during:
            instantiation of "auto cute::operator==(const cute::ScaledBasis<T, Ns...> &, const cute::ScaledBasis<U, Ms...> &) [with T=int, Ns=<1>, U=cute::R<1, 16>, Ms=<0>]" at line 801 of /nvfp4/upstream/clones/cutlass-main/include/cute/layout.hpp
            instantiation of "auto cute::detail::bw_coalesce<I,OldShape,OldStride,NewShape,NewStride>(const OldShape &, const OldStride &, const NewShape &, const NewStride &) [with I=1, OldShape=cute::tuple<cute::_128, cute::_2, cute::_256>, OldStride=cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>, cute::ScaledBasis<cute::R<1, 16>, 0>>, NewShape=cute::_256, NewStride=cute::ScaledBasis<cute::R<1, 16>, 0>]" at line 835 of /nvfp4/upstream/clones/cutlass-main/include/cute/layout.hpp
            instantiation of "auto cute::detail::coalesce_x(const cute::Layout<Shape, Stride> &) [with Shape=cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256>, Stride=cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>>]" at line 851 of /nvfp4/upstream/clones/cutlass-main/include/cute/layout.hpp
            instantiation of "auto cute::detail::coalesce_x(const cute::Layout<Shape, Stride> &, const IntTuple &) [with Shape=cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256>, Stride=cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>>, IntTuple=cute::C<0>]" at line 1139 of /nvfp4/upstream/clones/cutlass-main/include/cute/layout.hpp
            instantiation of "auto cute::composition(const cute::Layout<LShape, LStride> &, const cute::Layout<RShape, RStride> &) [with LShape=cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256>, LStride=cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>>, RShape=cute::tuple<cute::_1, cute::tuple<cute::tuple<cute::C<256>, cute::C<256>>, cute::_1>>, RStride=cute::tuple<cute::C<0>, cute::tuple<cute::tuple<cute::_256, cute::_1>, cute::C<0>>>]" at line 1152 of /nvfp4/upstream/clones/cutlass-main/include/cute/layout.hpp
            instantiation of function "lambda [](const auto &, const auto &)->auto [with <auto-1>=cute::Layout<cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256>, cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>>>, <auto-2>=cute::Layout<cute::tuple<cute::_1, cute::tuple<cute::tuple<cute::C<256>, cute::C<256>>, cute::_1>>, cute::tuple<cute::C<0>, cute::tuple<cute::tuple<cute::_256, cute::_1>, cute::C<0>>>>]" at line 743 of /nvfp4/upstream/clones/cutlass-main/include/cute/layout.hpp
            instantiation of "auto cute::detail::transform_layout(const Tuple0 &, const Tuple1 &, F &&, cute::seq<I...>, cute::seq<I0...>, cute::seq<I1...>) [with Tuple0=cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256>, cute::tuple<cute::_1, cute::_1, int32_t>>, cute::tuple<cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>>, cute::tuple<cute::C<0>, cute::C<0>, cute::ScaledBasis<cute::C<1>, 2>>>>, Tuple1=cute::tuple<cute::Layout<cute::tuple<cute::_1, cute::tuple<cute::tuple<cute::C<256>, cute::C<256>>, cute::_1>>, cute::tuple<cute::C<0>, cute::tuple<cute::tuple<cute::_256, cute::_1>, cute::C<0>>>>, cute::Underscore>, F=lambda [](const auto &, const auto &)->auto, I=<0, 1>, I0=<>, I1=<>]" at line 1152 of /nvfp4/upstream/clones/cutlass-main/include/cute/layout.hpp
            instantiation of "auto cute::composition(const cute::Layout<LShape, LStride> &, const Tiler &) [with LShape=cute::tuple<cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256>, cute::tuple<cute::_1, cute::_1, int32_t>>, LStride=cute::tuple<cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>>, cute::tuple<cute::C<0>, cute::C<0>, cute::ScaledBasis<cute::C<1>, 2>>>, Tiler=cute::tuple<cute::Layout<cute::tuple<cute::_1, cute::tuple<cute::tuple<cute::C<256>, cute::C<256>>, cute::_1>>, cute::tuple<cute::C<0>, cute::tuple<cute::tuple<cute::_256, cute::_1>, cute::C<0>>>>, cute::Underscore>]" at line 200 of /nvfp4/upstream/clones/cutlass-main/include/cute/layout.hpp
            instantiation of "auto cute::Layout<Shape, Stride>::compose(const Layouts &...) const [with Shape=cute::tuple<cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256>, cute::tuple<cute::_1, cute::_1, int32_t>>, Stride=cute::tuple<cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>>, cute::tuple<cute::C<0>, cute::C<0>, cute::ScaledBasis<cute::C<1>, 2>>>, Layouts=<cute::Layout<cute::tuple<cute::_1, cute::tuple<cute::tuple<cute::C<256>, cute::C<256>>, cute::_1>>, cute::tuple<cute::C<0>, cute::tuple<cute::tuple<cute::_256, cute::_1>, cute::C<0>>>>, cute::Underscore>]" at line 276 of /nvfp4/upstream/clones/cutlass-main/include/cute/atom/copy_atom.hpp
            instantiation of "auto cute::TiledCopy<Copy_Atom, LayoutCopy_TV, ShapeTiler_MN>::tile2thrfrg(Tensor &&, const Ref2TrgLayout &) [with Copy_Atom=cute::Copy_Atom<cute::Copy_Traits<cute::SM90_TMA_LOAD, cute::C<32768>, cute::AuxTmaParams<cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::tuple<cute::ScaledBasis<cute::R</* etc... */>, 0>, cute::ScaledBasis<cute::C</* etc... */>, 2>>, cute::ScaledBasis<cute::C<1>, 3>>, const cute::Layout<cute::tuple<cute::_16, cute::tuple<cute::_128, cute::_2>, cute::_1, cute::_1>, cute::tuple<cute::ScaledBasis<cute::C</* etc... */>, 1, 0>, cute::tuple<cute::ScaledBasis</* etc... */>, cute::ScaledBasis</* etc... */>>, cute::ScaledBasis<cute::C</* etc... */>, 1, 1>, cute::ScaledBasis<cute::C</* etc... */>, 2>>> &, const cute::Swizzle<0, 4, 3> &>>, cute::sparse_elem<16, uint8_t>>, LayoutCopy_TV=cute::Layout<cute::tuple<cute::_1, cute::tuple<cute::tuple<cute::tuple<cute::C<256>, cute::C<256>>, cute::_1>>>, cute::tuple<cute::C<0>, cute::tuple<cute::tuple<cute::tuple<cute::_256, cute::_1>, cute::C<0>>>>>, ShapeTiler_MN=cute::tuple<cute::C<256>, cute::C<256>>, Tensor=cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256>, cute::tuple<cute::_1, cute::_1, int32_t>>, cute::tuple<cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>>, cute::tuple<cute::C<0>, cute::C<0>, cute::ScaledBasis<cute::C<1>, 2>>>>, Ref2TrgLayout=cute::Layout<cute::tuple<cute::_1, cute::C<65536>>, cute::tuple<cute::C<0>, cute::C<1>>>]" at line 226 of /nvfp4/upstream/clones/cutlass-main/include/cute/atom/copy_atom.hpp
            instantiation of "auto cute::TiledCopy<Copy_Atom, LayoutCopy_TV, ShapeTiler_MN>::tidfrg_S(STensor &&) [with Copy_Atom=cute::Copy_Atom<cute::Copy_Traits<cute::SM90_TMA_LOAD, cute::C<32768>, cute::AuxTmaParams<cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::tuple<cute::ScaledBasis<cute::R</* etc... */>, 0>, cute::ScaledBasis<cute::C</* etc... */>, 2>>, cute::ScaledBasis<cute::C<1>, 3>>, const cute::Layout<cute::tuple<cute::_16, cute::tuple<cute::_128, cute::_2>, cute::_1, cute::_1>, cute::tuple<cute::ScaledBasis<cute::C</* etc... */>, 1, 0>, cute::tuple<cute::ScaledBasis</* etc... */>, cute::ScaledBasis</* etc... */>>, cute::ScaledBasis<cute::C</* etc... */>, 1, 1>, cute::ScaledBasis<cute::C</* etc... */>, 2>>> &, const cute::Swizzle<0, 4, 3> &>>, cute::sparse_elem<16, uint8_t>>, LayoutCopy_TV=cute::Layout<cute::tuple<cute::_1, cute::tuple<cute::tuple<cute::tuple<cute::C<256>, cute::C<256>>, cute::_1>>>, cute::tuple<cute::C<0>, cute::tuple<cute::tuple<cute::tuple<cute::_256, cute::_1>, cute::C<0>>>>>, ShapeTiler_MN=cute::tuple<cute::C<256>, cute::C<256>>, STensor=const cute::Layout<cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256, int32_t>, cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>, cute::ScaledBasis<cute::C<1>, 2>>> &]" at line 371 of /nvfp4/upstream/clones/cutlass-main/include/cute/atom/copy_atom.hpp
            instantiation of "auto cute::ThrCopy<TiledCopy, ThrIdx>::partition_S(STensor &&) const [with TiledCopy=cute::TiledCopy<cute::Copy_Atom<cute::Copy_Traits<cute::SM90_TMA_LOAD, cute::C<32768>, cute::AuxTmaParams<cute::tuple<cute::tuple<cute::ScaledBasis</* etc... */>, cute::ScaledBasis</* etc... */>>, cute::tuple<cute::ScaledBasis</* etc... */>, cute::ScaledBasis</* etc... */>>, cute::ScaledBasis<cute::C</* etc... */>, 3>>, const cute::Layout<cute::tuple<cute::_16, cute::tuple</* etc... */>, cute::_1, cute::_1>, cute::tuple<cute::ScaledBasis</* etc... */>, cute::tuple</* etc... */>, cute::ScaledBasis</* etc... */>, cute::ScaledBasis</* etc... */>>> &, const cute::Swizzle<0, 4, 3> &>>, cute::sparse_elem<16, uint8_t>>, cute::Layout<cute::tuple<cute::_1, cute::tuple<cute::tuple<cute::tuple<cute::C</* etc... */>, cute::C</* etc... */>>, cute::_1>>>, cute::tuple<cute::C<0>, cute::tuple<cute::tuple<cute::tuple<cute::_256, cute::_1>, cute::C<0>>>>>, cute::tuple<cute::C<256>, cute::C<256>>>, ThrIdx=int, STensor=cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<cute::C<0>, int, cute::C<0>, int>>>, cute::Layout<cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256, int32_t>, cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>, cute::ScaledBasis<cute::C<1>, 2>>>> &]" at line 796 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/gemm/collective/sm120_blockscaled_sparse_mma_tma.hpp
            instantiation of "void cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120TmaWarpSpecializedSparseBlockScaled<StagesA, StagesB, StagesE, SchedulerPipelineStageCount, ClusterShape>, TileShape_, ElementPairA_, LayoutPairsA_, ElementPairB_, StridePairB_, TiledMma_, GmemTiledCopyPairA_, SmemLayoutAtomsA_, SmemCopyAtomsA_, TransformA_, GmemTiledCopyPairB_, SmemLayoutAtomsB_, SmemCopyAtomsB_, TransformB_>::load_MK(const cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120TmaWarpSpecializedSparseBlockScaled<StagesA, StagesB, StagesE, SchedulerPipelineStageCount, ClusterShape>, TileShape_, ElementPairA_, LayoutPairsA_, ElementPairB_, StridePairB_, TiledMma_, GmemTiledCopyPairA_, SmemLayoutAtomsA_, SmemCopyAtomsA_, TransformA_, GmemTiledCopyPairB_, SmemLayoutAtomsB_, SmemCopyAtomsB_, TransformB_>::Params &, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120TmaWarpSpecializedSparseBlockScaled<StagesA, StagesB, StagesE, SchedulerPipelineStageCount, ClusterShape>, TileShape_, ElementPairA_, LayoutPairsA_, ElementPairB_, StridePairB_, TiledMma_, GmemTiledCopyPairA_, SmemLayoutAtomsA_, SmemCopyAtomsA_, TransformA_, GmemTiledCopyPairB_, SmemLayoutAtomsB_, SmemCopyAtomsB_, TransformB_>::MainloopPipelineMK, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120TmaWarpSpecializedSparseBlockScaled<StagesA, StagesB, StagesE, SchedulerPipelineStageCount, ClusterShape>, TileShape_, ElementPairA_, LayoutPairsA_, ElementPairB_, StridePairB_, TiledMma_, GmemTiledCopyPairA_, SmemLayoutAtomsA_, SmemCopyAtomsA_, TransformA_, GmemTiledCopyPairB_, SmemLayoutAtomsB_, SmemCopyAtomsB_, TransformB_>::PipelineStateMK, const cute::tuple<TensorA, TensorB, TensorE, TensorSFA, TensorSFB> &, const BlockCoord &, KTileIterator, int, int, uint32_t, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120TmaWarpSpecializedSparseBlockScaled<StagesA, StagesB, StagesE, SchedulerPipelineStageCount, ClusterShape>, TileShape_, ElementPairA_, LayoutPairsA_, ElementPairB_, StridePairB_, TiledMma_, GmemTiledCopyPairA_, SmemLayoutAtomsA_, SmemCopyAtomsA_, TransformA_, GmemTiledCopyPairB_, SmemLayoutAtomsB_, SmemCopyAtomsB_, TransformB_>::TensorStorage &) [with StagesA=2, StagesB=2, StagesE=2, SchedulerPipelineStageCount=2, ClusterShape=ClusterShape, TileShape_=ThreadBlockShape, ElementPairA_=cute::tuple<cutlass::float_e2m1_t, cutlass::float_ue4m3_t>, LayoutPairsA_=cute::tuple<cute::Layout<cute::tuple<int32_t, cute::tuple<cute::_2, int32_t>, int32_t>, cute::tuple<int64_t, cute::tuple<cute::_1, cute::_2>, int64_t>>, cute::Layout<cute::tuple<cute::tuple<cute::_128, int32_t>, cute::tuple<cute::_256, int32_t>, int32_t>, cute::tuple<cute::tuple<cute::_256, cute::C<32768>>, cute::tuple<cute::C<1>, int64_t>, int64_t>>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, int32_t>, cute::tuple<cute::tuple<cute::_32, cute::_4>, int32_t>, cute::tuple<cute::_1, int32_t>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, int32_t>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_512>, cute::tuple<cute::_0, int32_t>>>, cute::tuple<int64_t, cute::C<1>, int64_t>>, ElementPairB_=cute::tuple<cutlass::float_e2m1_t, cutlass::float_ue4m3_t>, StridePairB_=cute::tuple<cute::tuple<int64_t, cute::C<1>, int64_t>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, int32_t>, cute::tuple<cute::tuple<cute::_32, cute::_4>, int32_t>, cute::tuple<cute::_1, int32_t>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, int32_t>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_512>, cute::tuple<cute::_0, int32_t>>>>, TiledMma_=cute::TiledMMA<cute::MMA_Atom<cute::SM120::BLOCKSCALED::SPARSE::SM120_SPARSE_16x8x128_TN_VS<cutlass::float_e2m1_t, cutlass::float_e2m1_t, ElementAccumulator, cutlass::float_ue4m3_t, 32>>, cute::Layout<cute::tuple<cute::_4, cute::_2, cute::_1>, cute::tuple<cute::_1, cute::_4, cute::C<0>>>, cute::tuple<cute::C<128>, cute::C<32>, cute::C<128>>>, GmemTiledCopyPairA_=cute::tuple<cute::SM90_TMA_LOAD, cute::SM90_TMA_LOAD>, SmemLayoutAtomsA_=cute::tuple<cute::ComposedLayout<cute::Swizzle<2, 4, 3>, cute::smem_sparse_ptr_flag_bits<4, 8>, cute::Layout<cute::tuple<cute::tuple<cute::_1, cute::_8>, cute::tuple<cute::_4, cute::C<64>>>, cute::tuple<cute::tuple<cute::C<0>, cute::_256>, cute::tuple<cute::_1, cute::_4>>>>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::C<2>>, cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::_1, cute::_2>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, cute::C<512>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_4, cute::_1024>>>>, SmemCopyAtomsA_=cute::tuple<cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, cute::sparse_elem<4, uint8_t>>, cute::Copy_Atom<cute::UniversalCopy<uint64_t, uint64_t>, cute::sparse_elem<16, uint8_t>>, cute::Copy_Atom<cute::UniversalCopy<cutlass::float_ue4m3_t, cutlass::float_ue4m3_t>, cutlass::float_ue4m3_t>>, TransformA_=cute::identity, GmemTiledCopyPairB_=cute::tuple<cute::SM90_TMA_LOAD, cute::SM90_TMA_LOAD>, SmemLayoutAtomsB_=cute::tuple<cute::ComposedLayout<cute::Swizzle<3, 4, 3>, cute::smem_ptr_flag_bits<4>, cute::Layout<cute::tuple<cute::_8, cute::_256>, cute::tuple<cute::_256, cute::_1>>>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::C<1>>, cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::_1, cute::_2>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, cute::C<512>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_4, cute::_512>>>>, SmemCopyAtomsB_=cute::tuple<cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, cute::uint4_t>, cute::Copy_Atom<cute::UniversalCopy<cutlass::float_ue4m3_t, cutlass::float_ue4m3_t>, cutlass::float_ue4m3_t>>, TransformB_=cute::identity, TensorA=cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<cute::C<0>, cute::C<0>, cute::C<0>>>>, cute::Layout<cute::tuple<cute::_256, cute::tuple<cute::_2, cute::C<128>>, int32_t, int32_t, int32_t>, cute::tuple<cute::ScaledBasis<cute::C<1>, 1>, cute::tuple<cute::C<0>, cute::ScaledBasis<cute::C<1>, 0>>, cute::ScaledBasis<cute::C<256>, 1>, cute::ScaledBasis<cute::C<128>, 0>, cute::ScaledBasis<cute::C<1>, 2>>>>, TensorB=cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<cute::C<0>, cute::C<0>, cute::C<0>>>>, cute::Layout<cute::tuple<cute::_128, cute::_256, int32_t, int32_t, int32_t>, cute::tuple<cute::ScaledBasis<cute::C<1>, 1>, cute::ScaledBasis<cute::C<1>, 0>, cute::ScaledBasis<cute::C<128>, 1>, cute::ScaledBasis<cute::C<256>, 0>, cute::ScaledBasis<cute::C<1>, 2>>>>, TensorE=cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<cute::C<0>, cute::C<0>, cute::C<0>, cute::C<0>>>>, cute::Layout<cute::tuple<cute::tuple<cute::_128, cute::_2>, cute::_256, int32_t, int32_t, int32_t>, cute::tuple<cute::tuple<cute::ScaledBasis<int, 1>, cute::ScaledBasis<int, 1>>, cute::ScaledBasis<cute::R<1, 16>, 0>, cute::ScaledBasis<int, 1>, cute::ScaledBasis<cute::C<1>, 2>, cute::ScaledBasis<cute::C<1>, 3>>>>, TensorSFA=cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<cute::C<0>, cute::C<0>, cute::C<0>, cute::C<0>>>>, cute::Layout<cute::tuple<cute::tuple<cute::C<32>, cute::_4, cute::C<2>>, cute::tuple<cute::C<32>, cute::_4, cute::C<2>>, int32_t, int32_t, cute::tuple<cute::_1, int32_t>>, cute::tuple<cute::tuple<cute::ScaledBasis<cute::C<8>, 0>, cute::ScaledBasis<cute::C<2>, 0>, cute::ScaledBasis<cute::C<1>, 1>>, cute::tuple<cute::C<0>, cute::ScaledBasis<cute::R<1, 2>, 0>, cute::ScaledBasis<cute::C<1>, 2>>, cute::ScaledBasis<cute::C<2>, 1>, cute::ScaledBasis<cute::C<2>, 2>, cute::tuple<cute::C<0>, cute::ScaledBasis<cute::C<1>, 3>>>>>, TensorSFB=cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<cute::C<0>, cute::C<0>, cute::C<0>, cute::C<0>>>>, cute::Layout<cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::tuple<cute::C<32>, cute::_4, cute::C<2>>, int32_t, int32_t, cute::tuple<cute::_1, int32_t>>, cute::tuple<cute::tuple<cute::ScaledBasis<cute::C<8>, 0>, cute::ScaledBasis<cute::C<2>, 0>>, cute::tuple<cute::C<0>, cute::ScaledBasis<cute::R<1, 2>, 0>, cute::ScaledBasis<cute::C<1>, 1>>, cute::ScaledBasis<cute::C<1>, 2>, cute::ScaledBasis<cute::C<2>, 1>, cute::tuple<cute::C<0>, cute::ScaledBasis<cute::C<1>, 3>>>>>, KTileIterator=cute::ForwardCoordIterator<uint32_t, int32_t, cute::tuple<cute::ScaledBasis<cute::C<1>>>>, BlockCoord=cute::tuple<int32_t, int32_t, cute::Underscore, int32_t>]" at line 630 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/gemm/kernel/sm120_gemm_tma_warpspecialized_cooperative_asymmetric_dma.hpp
            instantiation of "void cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileSchedulerTag_, std::enable_if_t<<expression>, void>>::operator()(const cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileSchedulerTag_, std::enable_if_t<<expression>, void>>::Params &, char *) [with ProblemShape_=cute::tuple<int, int, int, int>, CollectiveMainloop_=CollectiveMainloop, CollectiveEpilogue_=CollectiveEpilogue, TileSchedulerTag_=void]" at line 123 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/device_kernel.h
            instantiation of "void cutlass::device_kernel<Operator>(Operator::Params) [with Operator=GemmKernel]" at line 347 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/gemm/device/gemm_universal_adapter.h
            instantiation of "cutlass::Status cutlass::gemm::device::GemmUniversalAdapter<GemmKernel_, std::enable_if_t<cutlass::gemm::detail::IsCutlass3GemmKernel<cutlass::GetUnderlyingKernel_t<GemmKernel_>, void>::value, void>>::initialize(const cutlass::gemm::device::GemmUniversalAdapter<GemmKernel_, std::enable_if_t<cutlass::gemm::detail::IsCutlass3GemmKernel<cutlass::GetUnderlyingKernel_t<GemmKernel_>, void>::value, void>>::Arguments &, void *, cudaStream_t, cutlass::CudaHostAdapter *) [with GemmKernel_=GemmKernel]" at line 503 of /tmp/mainprobe_m256.cu
            instantiation of "int run<Gemm>(Options &) [with Gemm=Gemm]" at line 577 of /tmp/mainprobe_m256.cu

1 error detected in the compilation of "/tmp/mainprobe_m256.cu".
128x128x128 compile error (tail) 

/nvfp4/upstream/clones/cutlass-main/include/cute/atom/copy_traits_sm90_tma.hpp(745): error: static assertion failed with "TMA requires CTA_Tile and SLayout top-level size equivalence."
    static_assert(decltype(size(slayout) == size(cta_v_map))::value, "TMA requires CTA_Tile and SLayout top-level size equivalence.")
    ^
          detected during:
            instantiation of "auto cute::detail::construct_tma_gbasis<TmaInternalType,GEngine,GLayout,SShape,SStride,VShape,VStride>(const cute::Tensor<GEngine, GLayout> &, const cute::Layout<SShape, SStride> &, const cute::Layout<VShape, VStride> &) [with TmaInternalType=uint8_t, GEngine=cute::ViewEngine<cute::sparse_ptr<16, cute::sparse_elem<16, uint8_t> *>>, GLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, int32_t>, cute::tuple<cute::_256, int32_t>, int32_t>, cute::tuple<cute::tuple<cute::_256, cute::C<32768>>, cute::tuple<cute::C<1>, int64_t>, int64_t>>, SShape=cute::tuple<cute::tuple<cute::_128, cute::_1>, cute::tuple<cute::_256, cute::_1>>, SStride=cute::tuple<cute::tuple<cute::_256, cute::_0>, cute::tuple<cute::_1, cute::_0>>, VShape=cute::tuple<cute::_128, cute::_128>, VStride=cute::tuple<cute::ScaledBasis<cute::C<1>, 0, 0>, cute::ScaledBasis<cute::C<1>, 1, 0>>]" at line 1152
            instantiation of "auto cute::detail::make_tma_copy_atom<TmaInternalType,CopyOp,GEngine,GLayout,SLayout,VShape,VStride>(CopyOp, const cute::Tensor<GEngine, GLayout> &, const SLayout &, const uint32_t &, const cute::Layout<VShape, VStride> &) [with TmaInternalType=uint8_t, CopyOp=cute::SM90_TMA_LOAD, GEngine=cute::ViewEngine<cute::sparse_ptr<16, cute::sparse_elem<16, uint8_t> *>>, GLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, int32_t>, cute::tuple<cute::_256, int32_t>, int32_t>, cute::tuple<cute::tuple<cute::_256, cute::C<32768>>, cute::tuple<cute::C<1>, int64_t>, int64_t>>, SLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, cute::_1>, cute::tuple<cute::_256, cute::_1>>, cute::tuple<cute::tuple<cute::_256, cute::_0>, cute::tuple<cute::_1, cute::_0>>>, VShape=cute::tuple<cute::_128, cute::_128>, VStride=cute::tuple<cute::ScaledBasis<cute::C<1>, 0, 0>, cute::ScaledBasis<cute::C<1>, 1, 0>>]" at line 1215
            instantiation of "auto cute::detail::make_tma_copy_tiled<TmaInternalType,CopyOp,GEngine,GLayout,SLayout,TShape,TStride,VShape,VStride>(const CopyOp &, const cute::Tensor<GEngine, GLayout> &, const SLayout &, const cute::Layout<TShape, TStride> &, const cute::Layout<VShape, VStride> &) [with TmaInternalType=uint8_t, CopyOp=cute::SM90_TMA_LOAD, GEngine=cute::ViewEngine<cute::sparse_ptr<16, cute::sparse_elem<16, uint8_t> *>>, GLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, int32_t>, cute::tuple<cute::_256, int32_t>, int32_t>, cute::tuple<cute::tuple<cute::_256, cute::C<32768>>, cute::tuple<cute::C<1>, int64_t>, int64_t>>, SLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, cute::_1>, cute::tuple<cute::_256, cute::_1>>, cute::tuple<cute::tuple<cute::_256, cute::_0>, cute::tuple<cute::_1, cute::_0>>>, TShape=cute::_1, TStride=cute::_0, VShape=cute::tuple<cute::_128, cute::_128>, VStride=cute::tuple<cute::ScaledBasis<cute::C<1>, 0, 0>, cute::ScaledBasis<cute::C<1>, 1, 0>>]" at line 1352
            instantiation of "auto cute::make_tma_copy(const CopyOp &, const cute::Tensor<GEngine, GLayout> &, const SLayout &, const CTA_Tiler &, const Cluster_Size &) [with TmaInternalType=uint8_t, CopyOp=cute::SM90_TMA_LOAD, GEngine=cute::ViewEngine<cute::sparse_ptr<16, cute::sparse_elem<16, uint8_t> *>>, GLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, int32_t>, cute::tuple<cute::_256, int32_t>, int32_t>, cute::tuple<cute::tuple<cute::_256, cute::C<32768>>, cute::tuple<cute::C<1>, int64_t>, int64_t>>, SLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, cute::_1>, cute::tuple<cute::_256, cute::_1>>, cute::tuple<cute::tuple<cute::_256, cute::_0>, cute::tuple<cute::_1, cute::_0>>>, CTA_Tiler=cute::tuple<cute::_128, cute::_128>, Cluster_Size=cute::_1]" at line 347 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/gemm/collective/sm120_blockscaled_sparse_mma_tma.hpp
            instantiation of class "cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120TmaWarpSpecializedSparseBlockScaled<StagesA, StagesB, StagesE, SchedulerPipelineStageCount, ClusterShape>, TileShape_, ElementPairA_, LayoutPairsA_, ElementPairB_, StridePairB_, TiledMma_, GmemTiledCopyPairA_, SmemLayoutAtomsA_, SmemCopyAtomsA_, TransformA_, GmemTiledCopyPairB_, SmemLayoutAtomsB_, SmemCopyAtomsB_, TransformB_>::Params [with StagesA=6, StagesB=6, StagesE=6, SchedulerPipelineStageCount=2, ClusterShape=ClusterShape, TileShape_=ThreadBlockShape, ElementPairA_=cute::tuple<cutlass::float_e2m1_t, cutlass::float_ue4m3_t>, LayoutPairsA_=cute::tuple<cute::Layout<cute::tuple<int32_t, cute::tuple<cute::_2, int32_t>, int32_t>, cute::tuple<int64_t, cute::tuple<cute::_1, cute::_2>, int64_t>>, cute::Layout<cute::tuple<cute::tuple<cute::_128, int32_t>, cute::tuple<cute::_256, int32_t>, int32_t>, cute::tuple<cute::tuple<cute::_256, cute::C<32768>>, cute::tuple<cute::C<1>, int64_t>, int64_t>>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, int>, cute::tuple<cute::tuple<cute::_32, cute::_4>, int>, cute::tuple<cute::_1, int>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, int>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_512>, cute::tuple<cute::_0, int32_t>>>, cute::tuple<int64_t, cute::C<1>, int64_t>>, ElementPairB_=cute::tuple<cutlass::float_e2m1_t, cutlass::float_ue4m3_t>, StridePairB_=cute::tuple<cute::tuple<int64_t, cute::C<1>, int64_t>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, int>, cute::tuple<cute::tuple<cute::_32, cute::_4>, int>, cute::tuple<cute::_1, int>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, int>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_512>, cute::tuple<cute::_0, int32_t>>>>, TiledMma_=cute::TiledMMA<cute::MMA_Atom<cute::SM120::BLOCKSCALED::SPARSE::SM120_SPARSE_16x8x128_TN_VS<cutlass::float_e2m1_t, cutlass::float_e2m1_t, ElementAccumulator, cutlass::float_ue4m3_t, 32>>, cute::Layout<cute::tuple<cute::_4, cute::_2, cute::_1>, cute::tuple<cute::_1, cute::_4, cute::C<0>>>, cute::tuple<cute::C<128>, cute::C<32>, cute::C<128>>>, GmemTiledCopyPairA_=cute::tuple<cute::SM90_TMA_LOAD, cute::SM90_TMA_LOAD>, SmemLayoutAtomsA_=cute::tuple<cute::ComposedLayout<cute::Swizzle<1, 4, 3>, cute::smem_sparse_ptr_flag_bits<4, 8>, cute::Layout<cute::tuple<cute::tuple<cute::_1, cute::_8>, cute::tuple<cute::_4, cute::C<32>>>, cute::tuple<cute::tuple<cute::_0, cute::_128>, cute::tuple<cute::_1, cute::_4>>>>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::C<1>>, cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::_1, cute::_1>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, cute::C<512>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_4, cute::_512>>>>, SmemCopyAtomsA_=cute::tuple<cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, cute::sparse_elem<4, uint8_t>>, cute::Copy_Atom<cute::UniversalCopy<uint64_t, uint64_t>, cute::sparse_elem<16, uint8_t>>, cute::Copy_Atom<cute::UniversalCopy<cutlass::float_ue4m3_t, cutlass::float_ue4m3_t>, cutlass::float_ue4m3_t>>, TransformA_=cute::identity, GmemTiledCopyPairB_=cute::tuple<cute::SM90_TMA_LOAD, cute::SM90_TMA_LOAD>, SmemLayoutAtomsB_=cute::tuple<cute::ComposedLayout<cute::Swizzle<2, 4, 3>, cute::smem_ptr_flag_bits<4>, cute::Layout<cute::tuple<cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_1>>>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::C<1>>, cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::_1, cute::_1>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, cute::C<512>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_4, cute::_512>>>>, SmemCopyAtomsB_=cute::tuple<cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, cute::uint4_t>, cute::Copy_Atom<cute::UniversalCopy<cutlass::float_ue4m3_t, cutlass::float_ue4m3_t>, cutlass::float_ue4m3_t>>, TransformB_=cute::identity]" at line 197 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/gemm/kernel/sm120_gemm_tma_warpspecialized_cooperative_asymmetric_dma.hpp
            instantiation of class "cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileSchedulerTag_, std::enable_if_t<<expression>, void>>::Params [with ProblemShape_=cute::tuple<int32_t, int32_t, int32_t, int>, CollectiveMainloop_=CollectiveMainloop, CollectiveEpilogue_=CollectiveEpilogue, TileSchedulerTag_=void]" at line 221 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/gemm/device/gemm_universal_adapter.h
            instantiation of class "cutlass::gemm::device::GemmUniversalAdapter<GemmKernel_, std::enable_if_t<cutlass::gemm::detail::IsCutlass3GemmKernel<cutlass::GetUnderlyingKernel_t<GemmKernel_>, void>::value, void>> [with GemmKernel_=GemmKernel]" at line 140 of /tmp/mainprobe_k128.cu

/nvfp4/upstream/clones/cutlass-main/include/cute/layout.hpp(1792): error: static assertion failed with "tile_to_shape: block shape does not divide the target shape."
      static_assert(decltype(evenly_divides(target_shape, block_shape))::value, "tile_to_shape: block shape does not divide the target shape.")
      ^
          detected during:
            instantiation of "auto cute::tile_to_shape(const cute::Layout<Shape, Stride> &, const TrgShape &, const ModeOrder &) [with Shape=cute::tuple<cute::tuple<cute::_256, cute::_128>>, Stride=cute::tuple<cute::tuple<cute::_128, cute::_1>>, TrgShape=cute::C<16384>, ModeOrder=cute::LayoutLeft]" at line 1230 of /nvfp4/upstream/clones/cutlass-main/include/cute/atom/copy_traits_sm90_tma.hpp
            instantiation of "auto cute::detail::make_tma_copy_tiled<TmaInternalType,CopyOp,GEngine,GLayout,SLayout,TShape,TStride,VShape,VStride>(const CopyOp &, const cute::Tensor<GEngine, GLayout> &, const SLayout &, const cute::Layout<TShape, TStride> &, const cute::Layout<VShape, VStride> &) [with TmaInternalType=uint8_t, CopyOp=cute::SM90_TMA_LOAD, GEngine=cute::ViewEngine<cute::sparse_ptr<16, cute::sparse_elem<16, uint8_t> *>>, GLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, int32_t>, cute::tuple<cute::_256, int32_t>, int32_t>, cute::tuple<cute::tuple<cute::_256, cute::C<32768>>, cute::tuple<cute::C<1>, int64_t>, int64_t>>, SLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, cute::_1>, cute::tuple<cute::_256, cute::_1>>, cute::tuple<cute::tuple<cute::_256, cute::_0>, cute::tuple<cute::_1, cute::_0>>>, TShape=cute::_1, TStride=cute::_0, VShape=cute::tuple<cute::_128, cute::_128>, VStride=cute::tuple<cute::ScaledBasis<cute::C<1>, 0, 0>, cute::ScaledBasis<cute::C<1>, 1, 0>>]" at line 1352 of /nvfp4/upstream/clones/cutlass-main/include/cute/atom/copy_traits_sm90_tma.hpp
            instantiation of "auto cute::make_tma_copy(const CopyOp &, const cute::Tensor<GEngine, GLayout> &, const SLayout &, const CTA_Tiler &, const Cluster_Size &) [with TmaInternalType=uint8_t, CopyOp=cute::SM90_TMA_LOAD, GEngine=cute::ViewEngine<cute::sparse_ptr<16, cute::sparse_elem<16, uint8_t> *>>, GLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, int32_t>, cute::tuple<cute::_256, int32_t>, int32_t>, cute::tuple<cute::tuple<cute::_256, cute::C<32768>>, cute::tuple<cute::C<1>, int64_t>, int64_t>>, SLayout=cute::Layout<cute::tuple<cute::tuple<cute::_128, cute::_1>, cute::tuple<cute::_256, cute::_1>>, cute::tuple<cute::tuple<cute::_256, cute::_0>, cute::tuple<cute::_1, cute::_0>>>, CTA_Tiler=cute::tuple<cute::_128, cute::_128>, Cluster_Size=cute::_1]" at line 347 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/gemm/collective/sm120_blockscaled_sparse_mma_tma.hpp
            instantiation of class "cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120TmaWarpSpecializedSparseBlockScaled<StagesA, StagesB, StagesE, SchedulerPipelineStageCount, ClusterShape>, TileShape_, ElementPairA_, LayoutPairsA_, ElementPairB_, StridePairB_, TiledMma_, GmemTiledCopyPairA_, SmemLayoutAtomsA_, SmemCopyAtomsA_, TransformA_, GmemTiledCopyPairB_, SmemLayoutAtomsB_, SmemCopyAtomsB_, TransformB_>::Params [with StagesA=6, StagesB=6, StagesE=6, SchedulerPipelineStageCount=2, ClusterShape=ClusterShape, TileShape_=ThreadBlockShape, ElementPairA_=cute::tuple<cutlass::float_e2m1_t, cutlass::float_ue4m3_t>, LayoutPairsA_=cute::tuple<cute::Layout<cute::tuple<int32_t, cute::tuple<cute::_2, int32_t>, int32_t>, cute::tuple<int64_t, cute::tuple<cute::_1, cute::_2>, int64_t>>, cute::Layout<cute::tuple<cute::tuple<cute::_128, int32_t>, cute::tuple<cute::_256, int32_t>, int32_t>, cute::tuple<cute::tuple<cute::_256, cute::C<32768>>, cute::tuple<cute::C<1>, int64_t>, int64_t>>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, int>, cute::tuple<cute::tuple<cute::_32, cute::_4>, int>, cute::tuple<cute::_1, int>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, int>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_512>, cute::tuple<cute::_0, int32_t>>>, cute::tuple<int64_t, cute::C<1>, int64_t>>, ElementPairB_=cute::tuple<cutlass::float_e2m1_t, cutlass::float_ue4m3_t>, StridePairB_=cute::tuple<cute::tuple<int64_t, cute::C<1>, int64_t>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, int>, cute::tuple<cute::tuple<cute::_32, cute::_4>, int>, cute::tuple<cute::_1, int>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, int>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_512>, cute::tuple<cute::_0, int32_t>>>>, TiledMma_=cute::TiledMMA<cute::MMA_Atom<cute::SM120::BLOCKSCALED::SPARSE::SM120_SPARSE_16x8x128_TN_VS<cutlass::float_e2m1_t, cutlass::float_e2m1_t, ElementAccumulator, cutlass::float_ue4m3_t, 32>>, cute::Layout<cute::tuple<cute::_4, cute::_2, cute::_1>, cute::tuple<cute::_1, cute::_4, cute::C<0>>>, cute::tuple<cute::C<128>, cute::C<32>, cute::C<128>>>, GmemTiledCopyPairA_=cute::tuple<cute::SM90_TMA_LOAD, cute::SM90_TMA_LOAD>, SmemLayoutAtomsA_=cute::tuple<cute::ComposedLayout<cute::Swizzle<1, 4, 3>, cute::smem_sparse_ptr_flag_bits<4, 8>, cute::Layout<cute::tuple<cute::tuple<cute::_1, cute::_8>, cute::tuple<cute::_4, cute::C<32>>>, cute::tuple<cute::tuple<cute::_0, cute::_128>, cute::tuple<cute::_1, cute::_4>>>>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::C<1>>, cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::_1, cute::_1>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, cute::C<512>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_4, cute::_512>>>>, SmemCopyAtomsA_=cute::tuple<cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, cute::sparse_elem<4, uint8_t>>, cute::Copy_Atom<cute::UniversalCopy<uint64_t, uint64_t>, cute::sparse_elem<16, uint8_t>>, cute::Copy_Atom<cute::UniversalCopy<cutlass::float_ue4m3_t, cutlass::float_ue4m3_t>, cutlass::float_ue4m3_t>>, TransformA_=cute::identity, GmemTiledCopyPairB_=cute::tuple<cute::SM90_TMA_LOAD, cute::SM90_TMA_LOAD>, SmemLayoutAtomsB_=cute::tuple<cute::ComposedLayout<cute::Swizzle<2, 4, 3>, cute::smem_ptr_flag_bits<4>, cute::Layout<cute::tuple<cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_1>>>, cute::Layout<cute::tuple<cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::C<1>>, cute::tuple<cute::tuple<cute::_32, cute::_4>, cute::_1, cute::_1>>, cute::tuple<cute::tuple<cute::tuple<cute::_16, cute::_4>, cute::C<512>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::_4, cute::_512>>>>, SmemCopyAtomsB_=cute::tuple<cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, cute::uint4_t>, cute::Copy_Atom<cute::UniversalCopy<cutlass::float_ue4m3_t, cutlass::float_ue4m3_t>, cutlass::float_ue4m3_t>>, TransformB_=cute::identity]" at line 197 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/gemm/kernel/sm120_gemm_tma_warpspecialized_cooperative_asymmetric_dma.hpp
            instantiation of class "cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileSchedulerTag_, std::enable_if_t<<expression>, void>>::Params [with ProblemShape_=cute::tuple<int32_t, int32_t, int32_t, int>, CollectiveMainloop_=CollectiveMainloop, CollectiveEpilogue_=CollectiveEpilogue, TileSchedulerTag_=void]" at line 221 of /nvfp4/upstream/clones/cutlass-main/include/cutlass/gemm/device/gemm_universal_adapter.h
            instantiation of class "cutlass::gemm::device::GemmUniversalAdapter<GemmKernel_, std::enable_if_t<cutlass::gemm::detail::IsCutlass3GemmKernel<cutlass::GetUnderlyingKernel_t<GemmKernel_>, void>::value, void>> [with GemmKernel_=GemmKernel]" at line 140 of /tmp/mainprobe_k128.cu

4 errors detected in the compilation of "/tmp/mainprobe_k128.cu".

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start with examples/80_blackwell_geforce_sparse_gemm/80b_blackwell_geforce_nvfp4_nvfp4_sparse_gemm.cu and reproduce both tile substitutions using the shown nvcc command. Trace the SM120 sparse block-scaled collective through the reported cute/layout.hpp and sm120_gemm_tma_warpspecialized_cooperative_asymmetric_dma.hpp failures. Done means the intended tile support is established and unsupported shapes receive a concise static_assert naming the supported set, or the shapes compile if they are intended to work.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
hpc
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Active
Clarity
Mostly clear
Newbie friendliness
55/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.