Hopper Grouped GEMM support for FP8 Accum (#2123)

* Add support for fp8accum, with profiler extension

* Update .gitignore

* contri

---------

Co-authored-by: Haicheng Wu <[email protected]>

Author: ANIKET-SHIVAM

.gitignore

@@ -1,3 +1,4 @@
 # PyCache files
 __pycache__/
-cutlass_library.egg-info/
+cutlass_library.egg-info/
+/build*

CONTRIBUTORS.md

@@ -1,17 +1,18 @@
 ![ALT](./media/images/gemm-hierarchy-with-epilogue-no-labels.png "CUTLASS")
 
-[README](./README.md#documentation) > **Active Developers**
+[README](./README.md#documentation) > **Contributors**
 
 # CUTLASS Developers **
 
-Andrew Kerr (CUTLASS founding member)<br />
+Andrew Kerr<br />
+Paul Springer<br />
 Dustyn Blasig<br />
 Albert Xu<br />
 Junkai Wu<br />
 Xiuxia Zhang<br />
-Haicheng Wu (CUTLASS founding member)<br />
+Haicheng Wu<br />
 Jack Yang<br />
-Pradeep Ramani (CUTLASS 3.x founding member)<br />
+Pradeep Ramani<br />
 Aditya Atluri<br />
 Han Li<br />
 Nick Zhao<br />
@@ -20,15 +21,15 @@ Yu-Jung Chen<br />
 Markus Hoehnerbach<br />
 Honghao Lu<br />
 Mihir Awatramani<br />
-Hao  Sheng<br />
+Hao Sheng<br />
 Zekun Fan<br />
-Aniket	Shivam<br />
+Aniket Shivam<br />
 Siyu Liu<br />
 Richard Cai<br />
 Vikas Gupta<br />
 Ethan Yan<br />
-Vijay Thakkar (CUTLASS 3.x and CuTe founding member)<br />
-Cris Cecka (CuTe and CUTLASS 3.x founding member)<br />
+Vijay Thakkar<br />
+Cris Cecka<br />
 Lawrence Ryan<br />
 Qun Song<br />
 Daniel Ricketts<br />
@@ -69,5 +70,61 @@ Shreya Gaur<br />
 
 ** _The list is sorted in order of the author's first contribution to the CUTLASS project._
 
+
+# CUTE Developers
+
+Cris Cecka<br />
+Vijay Thakkar<br />
+
+
 # CUTLASS Product Manager
+
 Matthew Nicely<br />
+
+
+# Former CUTLASS Developers
+
+Manish Gupta<br />
+Duane Merrill<br />
+Piotr Majcher<br />
+Naila Farooqui<br />
+Mark Hoemmen<br />
+Rawn Henry<br />
+Jin Wang<br />
+Timmy Liu<br />
+Manikandan Ananth<br />
+David Tanner<br />
+
+
+# Acknowledgements
+
+Tri Dao<br />
+Jay Shah<br />
+Timothy Costa<br />
+Julien Demouth<br />
+Brian Fahs<br />
+Michael Garland<br />
+Michael Goldfarb<br />
+Mostafa Hagog<br />
+Fei Hu<br />
+Alan Kaatz<br />
+Tina Li<br />
+Wei Liu<br />
+Tim Martin<br />
+Kevin Siu<br />
+Markus Tavenrath<br />
+John Tran<br />
+Vicki Wang<br />
+Fung Xie<br />
+Yang Xu<br />
+Scott Yokim<br />
+Girish Bharambe<br />
+Luke Durant<br />
+Carter Edwards<br />
+Olivier Giroux<br />
+Stephen Jones<br />
+Rishkul Kulkarni<br />
+Bryce Lelbach<br />
+Joel McCormack<br />
+Kyrylo Perelygin<br />
+Sean Treichler<br />

include/cutlass/gemm/collective/builders/sm90_gmma_builder.inl

@@ -234,8 +234,6 @@ struct CollectiveBuilder<
                                                                    KernelPtrArrayTmaWarpSpecializedCooperative,
                                                                    KernelPtrArrayTmaWarpSpecializedPingpong>);
   static constexpr bool IsFP8Input = detail::is_input_fp8<ElementA, ElementB>();
-  static_assert(!IsFP8Input || (IsFP8Input && !IsArrayOfPointersGemm),
-                "KernelPtrArrayTmaWarpSpecialized[Cooperative|Pingpong] is only compatible with FP8 FastAccum version right now.");
 
   // For fp32 types, map to tf32 MMA value type
   using ElementAMma = cute::conditional_t<cute::is_same_v<ElementA, float>, tfloat32_t, ElementA>;
@@ -267,12 +265,17 @@ struct CollectiveBuilder<
 
   static constexpr int PipelineStages = detail::compute_stage_count_or_override<Sm90ReducedSmemCapacityBytes,
       ElementAMma, ElementBMma, TileShape_MNK>(StageCountType{});
+  /* For FP8 use a separate mainloop compared to other datatypes */
   using DispatchPolicy = cute::conditional_t<IsArrayOfPointersGemm,
-      MainloopSm90ArrayTmaGmmaWarpSpecialized<PipelineStages, ClusterShape_MNK, KernelScheduleType>,
-      /* For FP8 use a separate mainloop compared to other datatypes */
+      cute::conditional_t<IsFP8Input,
+          MainloopSm90ArrayTmaGmmaWarpSpecializedFP8<PipelineStages, ClusterShape_MNK, KernelScheduleType>,
+          MainloopSm90ArrayTmaGmmaWarpSpecialized<PipelineStages, ClusterShape_MNK, KernelScheduleType>
+      >,
       cute::conditional_t<IsFP8Input,
           MainloopSm90TmaGmmaWarpSpecializedFP8<PipelineStages, ClusterShape_MNK, KernelScheduleType>,
-          MainloopSm90TmaGmmaWarpSpecialized<PipelineStages, ClusterShape_MNK, KernelScheduleType>>>;
+          MainloopSm90TmaGmmaWarpSpecialized<PipelineStages, ClusterShape_MNK, KernelScheduleType>
+      >
+  >;
 
   using SmemCopyAtomA = void;
   using SmemCopyAtomB = void;

include/cutlass/gemm/collective/collective_mma.hpp

@@ -46,8 +46,9 @@
 #include "cutlass/gemm/collective/sm90_sparse_mma_tma_gmma_ss_warpspecialized.hpp"
 #include "cutlass/gemm/collective/sm90_sparse_mma_tma_gmma_ss_warpspecialized_fp8.hpp"
 #include "cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized.hpp"
-#include "cutlass/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input.hpp"  
+#include "cutlass/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input.hpp"
 #include "cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8.hpp"
+#include "cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized_fp8.hpp"
 #include "cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp"
 
 #if !defined(__CUDACC_RTC__)