Remove old CK Tile Stream-K implementation

The original CK Stream-K implementation used a Tile Partitioner that was
based on old CK's Stream-K block to C tile map. However, old CK's
implementation did not align with the original Stream-K paper. Thus, we
implemented a new Tile Partitioner and associated Stream-K Kernel. The
kernel implementation was placed in the reboot namespace. Now that all
functionality for the new implementation is in place, this change makes
the following changes:
- Removes all uses of the old CK Tile Stream-K implementation.
- Removes the reboot namespace such that the new implementation is in
  the ck_tile namespace only.
- Adds tests for fp8 and bf8 for the new implementation as these were
  only in place for the old implementation.
- Removes the old CK Tile Stream-K Tile partitioner
- Remove the v2 suffix from the new CK Tile Tile Partitioner derived
  classes.

Author: ecamartins

example/ck_tile/40_streamk_gemm/run_gemm_example.inc

@@ -71,15 +71,15 @@ invoke_gemm(ck_tile::DeviceMem& a_m_k_dev_buf,
             bool flush_cache,
             ck_tile::StreamKReductionStrategy reduction_strategy)
 {
-    ck_tile::reboot::StreamKHostArgs args{a_m_k_dev_buf.GetDeviceBuffer(),
-                                          b_k_n_dev_buf.GetDeviceBuffer(),
-                                          c_m_n_dev_buf.GetDeviceBuffer(),
-                                          M,
-                                          N,
-                                          K,
-                                          stride_A,
-                                          stride_B,
-                                          stride_C};
+    ck_tile::StreamKHostArgs args{a_m_k_dev_buf.GetDeviceBuffer(),
+                                  b_k_n_dev_buf.GetDeviceBuffer(),
+                                  c_m_n_dev_buf.GetDeviceBuffer(),
+                                  M,
+                                  N,
+                                  K,
+                                  stride_A,
+                                  stride_B,
+                                  stride_C};
 
     std::tuple<float, ck_tile::index_t> ave_time_and_batch;
 

example/ck_tile/40_streamk_gemm/streamk_gemm_basic.cpp

@@ -16,7 +16,7 @@ template <typename GemmConfig,
           typename ELayout,
           typename CDEElementWise,
           ck_tile::StreamKReductionStrategy ReductionStrategy>
-std::tuple<float, ck_tile::index_t> gemm(const ck_tile::reboot::StreamKHostArgs& args,
+std::tuple<float, ck_tile::index_t> gemm(const ck_tile::StreamKHostArgs& args,
                                          const ck_tile::stream_config& s)
 {
     using GemmShape = ck_tile::TileGemmShape<
@@ -28,7 +28,7 @@ std::tuple<float, ck_tile::index_t> gemm(const ck_tile::reboot::StreamKHostArgs&
         GemmConfig::PermuteB>;
 
     using TilePartitioner =
-        ck_tile::StreamKTilePartitioner_v2<GemmShape, ReductionStrategy, GemmConfig::Persistent>;
+        ck_tile::StreamKTilePartitioner<GemmShape, ReductionStrategy, GemmConfig::Persistent>;
 
     using GemmUniversalTraits = ck_tile::TileGemmUniversalTraits<GemmConfig::kPadM,
                                                                  GemmConfig::kPadN,
@@ -77,7 +77,7 @@ std::tuple<float, ck_tile::index_t> gemm(const ck_tile::reboot::StreamKHostArgs&
                                              memory_operation.value,
                                              GemmConfig::NumWaveGroups>>;
 
-        using Kernel = ck_tile::reboot::StreamKKernel<TilePartitioner, GemmPipeline, GemmEpilogue>;
+        using Kernel = ck_tile::StreamKKernel<TilePartitioner, GemmPipeline, GemmEpilogue>;
 
         auto kargs                = Kernel::MakeKernelArgs(args);
         const auto workspace_size = Kernel::GetWorkSpaceSize(kargs);

include/ck_tile/ops/common/streamk_common.hpp

@@ -11,33 +11,4 @@ enum StreamKReductionStrategy : uint32_t
     Atomic    = 0u,
     Reduction = 1u
 };
-
-/**
- * @brief Estimates the number of Stream-K workgroups per macro tile in the C tensor.
- *
- * @param sk_ctas           Number of Stream-K workgroups.
- * @param iters_per_sk_cta  Number of iterations per Stream-K workgroup.
- * @param iters_per_tile    Number of iterations per tile (i.e., the number of macro tiles in the K
- * dimension).
- * @return ck_tile::index_t An estimate of the number of workgroups per macro tile in the C tensor.
- * @note It is assumed that `iters_per_sk_cta` > 0.
- */
-template <ck_tile::StreamKReductionStrategy ReductionStrategy>
-ck_tile::index_t
-estimate_num_wgs_per_tile(index_t sk_ctas, index_t iters_per_sk_cta, index_t iters_per_tile)
-{
-    // In the case of non-atomic reduction or data-parallel only, there will always be 1 workgroup
-    // writing final results to a given macro tile in C.
-    int num_wgs_per_tile = 1;
-
-    // Otherwise, for atomics, multiple workgroups may be writing to the same macro tile in C.
-    if(sk_ctas > 0 && ReductionStrategy == ck_tile::StreamKReductionStrategy::Atomic)
-    {
-        // Estimate the number of workgroups per macro tile.
-        num_wgs_per_tile =
-            (iters_per_tile / iters_per_sk_cta) + ((iters_per_tile % iters_per_sk_cta) != 0);
-    }
-
-    return std::max(num_wgs_per_tile, 1);
-}
 } // namespace ck_tile