ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-12 02:05:50 +00:00
Code Issues Packages Projects Releases Wiki Activity

Remove oversubscription.

Browse Source
This commit is contained in:
Ville Pietilä
2025-07-10 15:30:30 +00:00
parent 3ad1ec2981
commit 727d578ca9
4 changed files with 9 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp
View File

@@ -655,8 +655,8 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle
const auto gemmM = a_grid_desc_kbatch_k0_m_k1.GetLength(I1);
const auto gemmN = b_grid_desc_kbatch_k0_n_k1.GetLength(I1);
const auto grid_size_mn = GridwiseGemm::Block2CTileMap::CalculateGridSize(gemmM, gemmN);
k_batch_ = get_best_occupancy_k_batch_value(max_occupancy.value_, grid_size_mn, Conv_G_);
const auto grid_size = GridwiseGemm::Block2CTileMap::CalculateGridSize(gemmM, gemmN) * Conv_G_;
k_batch_ = get_best_occupancy_k_batch_value(max_occupancy.value_, grid_size);
// Ensure that k_batch_ does not exceed the maximum value
// for the GEMM pipeline

4
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp
View File

@@ -551,8 +551,8 @@ struct DeviceGroupedConvBwdWeight_Xdl_CShuffle
const auto& c_grid_desc_m_n = descs_initial[I2];
const auto& block_2_ctile_map = GridwiseGemm::MakeCBlockClusterAdaptor(c_grid_desc_m_n, M01, N01, k_batch_initial);
const auto grid_size_mn = block_2_ctile_map.CalculateGridSize(c_grid_desc_m_n);
k_batch_ = get_best_occupancy_k_batch_value(max_occupancy.value_, grid_size_mn, Conv_G_);
const auto grid_size = block_2_ctile_map.CalculateGridSize(c_grid_desc_m_n) * Conv_G_;
k_batch_ = get_best_occupancy_k_batch_value(max_occupancy.value_, grid_size);
}
else
{

4
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle_v3.hpp
View File

@@ -520,8 +520,8 @@ struct DeviceGroupedConvBwdWeight_Xdl_CShuffleV3
const auto gemmM = a_grid_desc_kbatch_k0_m_k1.GetLength(I1);
const auto gemmN = b_grid_desc_kbatch_k0_n_k1.GetLength(I1);
const auto grid_size_mn = GridwiseGemm::Block2CTileMap::CalculateGridSize(gemmM, gemmN);
k_batch_ = get_best_occupancy_k_batch_value(max_occupancy.value_, grid_size_mn, Conv_G_);
const auto grid_size = GridwiseGemm::Block2CTileMap::CalculateGridSize(gemmM, gemmN) * Conv_G_;
k_batch_ = get_best_occupancy_k_batch_value(max_occupancy.value_, grid_size);
// Ensure that k_batch_ does not exceed the maximum value
// for the GEMM pipeline

9
include/ck/tensor_operation/gpu/device/impl/split_k_utils.hpp
View File

@@ -27,15 +27,13 @@ struct DeviceProperties
int num_cu_;
};
inline ck::index_t get_best_occupancy_k_batch_value(int max_occupancy, ck::index_t grid_size_mn, ck::index_t num_conv_groups)
inline ck::index_t get_best_occupancy_k_batch_value(int max_occupancy, ck::index_t grid_size)
{
static DeviceProperties device_properties;
const int max_capacity = max_occupancy * device_properties.num_cu_;
// Empirically, when we have more than one group, we can oversubscribe by the number of groups.
// This is because the groups are independent and can be processed in parallel.
ck::index_t k_batch = 1;
const auto optimal_split = static_cast<ck::index_t>(std::floor((1.0 * max_capacity) / (grid_size_mn)));
const auto optimal_split = static_cast<ck::index_t>(std::floor((1.0 * max_capacity) / grid_size));
if (optimal_split > 1)
{
k_batch = optimal_split;
@@ -44,8 +42,7 @@ inline ck::index_t get_best_occupancy_k_batch_value(int max_occupancy, ck::index
if (ck::EnvIsEnabled(CK_ENV(CK_LOGGING)))
{
std::cout << "[SPLIT-K AUTODEDUCE] Max active thread blocks per CU for GEMM kernel: " << max_occupancy << std::endl;
std::cout << "[SPLIT-K AUTODEDUCE] Output grid size: " << grid_size_mn << std::endl;
std::cout << "[SPLIT-K AUTODEDUCE] Number of conv groups: " << num_conv_groups << std::endl;
std::cout << "[SPLIT-K AUTODEDUCE] Output grid size: " << grid_size << std::endl;
std::cout << "[SPLIT-K AUTODEDUCE] Optimal split-k value " << k_batch << std::endl;
}
return k_batch;