Fix MMA promotion interval assertions (#1641)

NVIDIA · Sep 16, 2024 · 1ebda1c · 1ebda1c
1 parent 9f68995
commit 1ebda1c
Showing 1 changed file with 2 additions and 2 deletions.
diff --git a/include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8.hpp b/include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8.hpp
@@ -246,8 +246,8 @@ struct CollectiveMma<
  implementable = implementable && cutlass::detail::check_alignment<min_tma_aligned_elements_A>(cute::make_shape(M,K,L), StrideA{});
  constexpr int min_tma_aligned_elements_B = tma_alignment_bits / cutlass::sizeof_bits<ElementB>::value;
  implementable = implementable && cutlass::detail::check_alignment<min_tma_aligned_elements_B>(cute::make_shape(N,K,L), StrideB{});
- /* MMA promotion interval should be a multiple of 4, since each mainloop iteration would issue 4 MMA instructions. */
- implementable = implementable && (args.mma_promotion_interval % 4 == 0);
+ /* MMA promotion interval should be a multiple of the number of MMA instructions issued by each mainloop iteration. */
+ implementable = implementable && (args.mma_promotion_interval % (size<2>(TileShape{})() / TiledMma().template tile_size_mnk<2>()()) == 0);
 
  if (!implementable) {
  CUTLASS_TRACE_HOST(" CAN IMPLEMENT: Problem Size doesn't meet the minimum alignment requirements for TMA.\n");