Remove old CK Tile Stream-K implementation

The original CK Stream-K implementation was based on old CK's Stream-K
block to C tile map. However, this implementation did not align with the
original Stream-K paper. Thus, we implemented a new tile partitioner and
associated Stream-K kernel, which was placed in the reboot namespace.

Now that the new Stream-K implementation is ready, this change removes
all artifacts of the old implementation. Specifically, the following
changes were made:
- Removes old Stream-K tile partitioner from CK Tile
- Removes the reboot namespace such that the new implementation resides
  in the ck_tile namespace only.
- Adds tests for bf8 and fp8 using the new implementation
- Removes tests for the old implementation
- Remove the v2 suffix from the new CK Tile Tile Partitioner
derived classes.
- Updates Stream-K Kernel ops file to use /** commenting style.

Author: ecamartins

example/ck_tile/40_streamk_gemm/run_gemm_example.inc

@@ -71,16 +71,16 @@ 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,
-                                          reduction_strategy};
+    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,
+                                  reduction_strategy};
 
     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