Renaming the get_grid function and grid_ parameter to max_active_wgs to avoid confusion with grid_size

Author: arai713

projects/composablekernel/include/ck_tile/ops/gemm/kernel/streamk_gemm/streamk_gemm_kernel.hpp

@@ -790,7 +790,7 @@ struct StreamKKernel
 
         // Data-parallel section
         for(index_t tile_idx = block_idx; tile_idx < kargs.tile_partitioner.get_dp_tiles();
-            tile_idx += kargs.tile_partitioner.get_grid())
+            tile_idx += kargs.tile_partitioner.get_max_active_wgs())
         {
             BaseGemm(kargs, tile_idx, dp_num_loop, 0, 0, kargs.K, smem_ptr_0);
             block_sync_lds();

projects/composablekernel/include/ck_tile/ops/gemm/kernel/streamk_gemm/streamk_gemm_tile_partitioner.hpp

@@ -156,7 +156,7 @@ struct StreamKTilePartitionerBase
      * @brief Returns the maximum number of active workgroups; this is assumed to be number of CUs *
      * occupancy.
      */
-    CK_TILE_HOST_DEVICE index_t get_grid() const noexcept;
+    CK_TILE_HOST_DEVICE index_t get_max_active_wgs() const noexcept;
 
     /**
      * @brief Returns the number of tiles in the C tensor that will use the data-parallel (DP)
@@ -215,7 +215,7 @@ struct StreamKTilePartitionerBase
 
     protected:
     index_t num_tiles_;
-    index_t grid_;
+    index_t max_active_wgs_;
     index_t dp_tiles_;
 
     private:
@@ -290,7 +290,7 @@ struct StreamKTilePartitioner<BlockGemmShapeType, ReductionStrategyType, true>
 
     /**
      * @brief Returns the total number of DP tiles left over when `dp_tiles_` is not evenly
-     * divisible by `grid_`.
+     * divisible by `max_active_wgs_`.
      */
     CK_TILE_HOST_DEVICE index_t get_extra_dp_tiles() const noexcept;
 

projects/composablekernel/include/ck_tile/ops/gemm/kernel/streamk_gemm/streamk_gemm_tile_partitioner_impl.hpp

@@ -8,23 +8,24 @@ namespace ck_tile {
 template <typename BlockGemmShapeType, StreamKReductionStrategy ReductionStrategyType>
 StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::StreamKTilePartitionerBase(
     index_t m, index_t n, index_t k, index_t grid)
-    : grid_{grid}, n_{n}
+    : max_active_wgs_{grid}, n_{n}
 {
     iters_per_tile_ = integer_divide_ceil(k, KPerBlock);
     num_tiles_      = integer_divide_ceil(m, MPerBlock) * integer_divide_ceil(n_, NPerBlock);
 
-    bool big_enough         = num_tiles_ > grid_;
-    index_t remainder_tiles = num_tiles_ % grid_;
+    bool big_enough         = num_tiles_ > max_active_wgs_;
+    index_t remainder_tiles = num_tiles_ % max_active_wgs_;
 
     if(remainder_tiles)
     {
-        sk_tiles_       = big_enough ? full_tiles_ * grid_ + (num_tiles_ % grid_) : num_tiles_;
-        sk_tiles_       = min(num_tiles_, sk_tiles_);
-        sk_ctas_        = grid_;
+        sk_tiles_ = big_enough ? full_tiles_ * max_active_wgs_ + (num_tiles_ % max_active_wgs_)
+                               : num_tiles_;
+        sk_tiles_ = min(num_tiles_, sk_tiles_);
+        sk_ctas_  = max_active_wgs_;
         total_sk_iters_ = sk_tiles_ * iters_per_tile_;
 
         // If there still isn't enough work to saturate all CUs, then just revert to DP only.
-        if(total_sk_iters_ < grid_)
+        if(total_sk_iters_ < max_active_wgs_)
         {
             sk_tiles_       = 0;
             sk_ctas_        = 0;
@@ -175,9 +176,10 @@ StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::get_num_t
 
 template <typename BlockGemmShapeType, StreamKReductionStrategy ReductionStrategyType>
 CK_TILE_HOST_DEVICE index_t
-StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::get_grid() const noexcept
+StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>::get_max_active_wgs()
+    const noexcept
 {
-    return grid_;
+    return max_active_wgs_;
 }
 
 template <typename BlockGemmShapeType, StreamKReductionStrategy ReductionStrategyType>
@@ -290,8 +292,8 @@ StreamKTilePartitioner<BlockGemmShapeType, ReductionStrategyType, true>::StreamK
     ck_tile::index_t m, ck_tile::index_t n, ck_tile::index_t k, ck_tile::index_t grid)
     : StreamKTilePartitionerBase<BlockGemmShapeType, ReductionStrategyType>(m, n, k, grid)
 { // inherit from base constructor
-    dp_tiles_per_cta_ = this->dp_tiles_ / this->grid_;
-    extra_dp_tiles_   = this->dp_tiles_ % this->grid_;
+    dp_tiles_per_cta_ = this->dp_tiles_ / this->max_active_wgs_;
+    extra_dp_tiles_   = this->dp_tiles_ % this->max_active_wgs_;
 }
 
 template <typename BlockGemmShapeType, StreamKReductionStrategy ReductionStrategyType>
@@ -301,7 +303,7 @@ StreamKTilePartitioner<BlockGemmShapeType, ReductionStrategyType, true>::grid_si
 {
     if(extra_dp_tiles_ == 0)
     {
-        return dim3(this->grid_, 1, 1);
+        return dim3(this->max_active_wgs_, 1, 1);
     }
     else
     {

projects/composablekernel/test/ck_tile/gemm_streamk/test_streamk_tile_partitioner_common.hpp

@@ -183,7 +183,7 @@ void validate_streamk_base_constructor(
     EXPECT_EQ(tile_partitioner.get_iters_per_tile(), expected_values.iters_per_tile_);
     EXPECT_EQ(tile_partitioner.get_total_dp_iters(), expected_values.total_dp_iters_);
     EXPECT_EQ(tile_partitioner.get_num_tiles(), expected_values.num_tiles_);
-    EXPECT_EQ(tile_partitioner.get_grid(), expected_values.grid_);
+    EXPECT_EQ(tile_partitioner.get_max_active_wgs(), expected_values.grid_);
     EXPECT_EQ(tile_partitioner.get_n(), expected_values.n_);
 }
 
@@ -446,7 +446,7 @@ void validate_streamk_persistent(
 {
     EXPECT_EQ(tile_partitioner.get_dp_tiles_per_cta(), expected_values.dp_tiles_per_cta_);
     EXPECT_EQ(tile_partitioner.get_extra_dp_tiles(), expected_values.extra_dp_tiles_);
-    EXPECT_EQ(tile_partitioner.get_grid(), expected_values.grid_);
+    EXPECT_EQ(tile_partitioner.get_max_active_wgs(), expected_values.grid_);
 }
 
 // Non-Persistent
@@ -459,5 +459,5 @@ void validate_streamk_nonpersistent(
     EXPECT_EQ(tile_partitioner.get_dp_ctas(), expected_values.dp_ctas_);
     EXPECT_EQ(tile_partitioner.get_dp_start_block_idx(), expected_values.dp_start_block_idx_);
     EXPECT_EQ(tile_partitioner.get_sk_start_block_idx(), expected_values.sk_start_block_idx_);
-    EXPECT_EQ(tile_partitioner.get_grid(), expected_values.grid_);
+    EXPECT_EQ(tile_partitioner.get_max_active_wgs(), expected_values.grid_);
 }