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

Fix grid size calculation.

Browse Source
This commit is contained in:
Ville Pietilä
2025-06-10 12:02:41 +00:00
parent 692d71d466
commit 3e3a73eee3
3 changed files with 89 additions and 43 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -451,7 +451,7 @@ struct DeviceGroupedConvBwdWeight_Xdl_CShuffle
true>, // TODO: Do we need to test both true/false for HasMainKBlockLoop?
BlockSize,
dynSharedMemPerBlk));
value_ = max_occupancy;
value_ = std::max(1, max_occupancy);
}
int value_;
};
@@ -504,13 +504,6 @@ struct DeviceGroupedConvBwdWeight_Xdl_CShuffle
{
static MaximumActiveBlocksPerMultiprocessor max_occupancy;
k_batch_ = get_k_batch_value<MPerBlock, NPerBlock>(
split_k,
max_occupancy.value_,
M01,
N01,
Conv_G_);
constexpr index_t spatial_offset = 3;
std::copy(begin(b_g_n_c_wis_lengths) + spatial_offset,
end(b_g_n_c_wis_lengths),
@@ -531,6 +524,40 @@ struct DeviceGroupedConvBwdWeight_Xdl_CShuffle
std::array<index_t, NDimSpatial + 3> e_g_k_c_xs_strides_transposed =
conv_ngchw_to_nhwgc_transformer.TransposeWeiStrides(e_g_k_c_xs_lengths,
e_g_k_c_xs_strides);
if (split_k < 0)
{
constexpr int k_batch_initial = 1;
const auto descs_initial =
conv_to_gemm_transformer
.template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N<NDimSpatial>(
Conv_N_,
Conv_K_,
Conv_C_,
input_spatial_lengths_,
filter_spatial_lengths_,
output_spatial_lengths_,
b_g_n_c_wis_strides_transposed,
e_g_k_c_xs_strides_transposed,
a_g_n_k_wos_strides_transposed,
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads,
k_batch_initial);
const auto& a_grid_desc_kbatch_k0_m_k1 = descs_initial[I0];
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 = block_2_ctile_map.CalculateGridSize(c_grid_desc_m_n);
const auto k_size = a_grid_desc_kbatch_k0_m_k1.GetLength(I0) * a_grid_desc_kbatch_k0_m_k1.GetLength(I1);
k_batch_ = get_k_batch_value(max_occupancy.value_, grid_size, k_size, Conv_G_);
}
else {
k_batch_ = split_k;
}
const auto descs =
conv_to_gemm_transformer
.template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N<NDimSpatial>(
@@ -701,7 +728,7 @@ struct DeviceGroupedConvBwdWeight_Xdl_CShuffle
// Invoker
struct Invoker : public BaseInvoker
{
using Argument = DeviceOp::Argument; // Refers to the argument defined above.
using Argument = DeviceOp::Argument;
void ShowInfo(const Argument& arg)
{

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

@@ -25,28 +25,26 @@ struct DeviceProperties
int num_cu_;
};
template<
ck::index_t MPerBlock,
ck::index_t NPerBlock>
ck::index_t get_k_batch_value(ck::index_t split_k, int max_occupancy, ck::index_t M, ck::index_t N, ck::index_t conv_G)
inline ck::index_t get_k_batch_value(int max_occupancy, ck::index_t grid_size, ck::index_t K_size, ck::index_t conv_G)
{
static DeviceProperties device_properties;
// For now, assume that negative value signals automatic computation of the split_k value.
if(split_k <= 0)
constexpr ck::index_t k_batch_min = 1;
constexpr ck::index_t batch_size_min = 16;
const int num_cu = device_properties.num_cu_;
const auto k_batch_max = math::integer_divide_ceil(K_size, batch_size_min);
auto k_batch = static_cast<ck::index_t>(std::ceil((max_occupancy * num_cu) / (1.0 * grid_size))); // Exclude th egrid size from the occupancy calculation
k_batch = std::min(std::max(k_batch_min, k_batch), k_batch_max);
if (ck::EnvIsEnabled(CK_ENV(CK_LOGGING)))
{
const int num_cu = device_properties.num_cu_;
const auto M0 = math::integer_divide_ceil(M, MPerBlock);
const auto N0 = math::integer_divide_ceil(N, NPerBlock);
const auto n_output_tiles = M0 * N0;
const auto k_batch = std::ceil((max_occupancy * num_cu) / (1.0 * n_output_tiles * conv_G));
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] Using optimized split-k value " << k_batch << " for K-batch."<< std::endl;
}
return k_batch;
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 (M tiles x N tiles x Conv groups): " << grid_size << std::endl;
std::cout << "[SPLIT-K AUTODEDUCE] K-dim size: " << K_size << std::endl;
std::cout << "[SPLIT-K AUTODEDUCE] Conv groups: " << conv_G << std::endl;
std::cout << "[SPLIT-K AUTODEDUCE] Maximum k_batch value: " << k_batch_max << std::endl;
std::cout << "[SPLIT-K AUTODEDUCE] Optimal split-k value " << k_batch << " for K-batch."<< std::endl;
}
return split_k;
return k_batch;
}
} // namespace device

53
profiler/include/profiler/profile_grouped_conv_bwd_weight_impl.hpp
View File

@@ -139,24 +139,41 @@ void write_perf_results_to_file(const PerfResults& perf_results_global,
const std::vector<PerfResults>& perf_results_list)
{
const auto& results_file = ck::EnvGetString(CK_ENV(CK_PROFILER_OUTPUT_FILE));
const std::string separator(";");
const auto& write_to_file = [&](const PerfResults res, std::ofstream& file, bool only_one_op = false) {
auto best_split_k = res.best_split_k_ > 0 ? res.best_split_k_ : res.best_split_k_arg_;
ck::index_t rank, total_num;
std::tie(rank, total_num) = res.get_ranking(res.opt_split_k_best_op_name_, res.opt_split_k_best_arg_);
file << res.best_op_name_ << separator
<< res.best_avg_time_ << separator
<< best_split_k << separator;
if (!only_one_op)
{
file << res.opt_split_k_best_op_name_ << separator;
}
file << res.opt_split_k_best_arg_ << separator
<< rank << separator
<< total_num;
};
if(!results_file.empty())
{
std::ofstream file(results_file, std::ios::out | std::ios::app);
if(file.is_open())
{
file << "Best configuration parameters:"
<< perf_results_global.print_best_op() << std::endl;
// First the global results
write_to_file(perf_results_global, file);
file << separator;
if (!perf_results_list.empty())
// Then the local results - one set for each op
const auto size = perf_results_list.size();
for (size_t i = 0; i < size; ++i)
{
file << "Optimized split-K results:"
<< perf_results_global.print_best_split_k() << std::endl;
}
for (const auto& res : perf_results_list)
{
file << res.print_best_op() << std::endl;
write_to_file(perf_results_list[i], file, true);
if (i < size - 2) file << separator;
}
file << std::endl;
file.close();
}
else
@@ -313,13 +330,13 @@ bool profile_grouped_conv_bwd_weight_impl(int do_verification,
profile_all = false;
}
std::vector<PerfResults> perf_results_list;
PerfResults perf_results_global;
std::vector<PerfResults> perf_results_list;
const auto& disabled_ops = get_disabled_ops();
for(auto& op_ptr : op_ptrs)
{
std::string op_name = op_ptr->GetTypeString();
// Skip disabled ops
@@ -333,6 +350,7 @@ bool profile_grouped_conv_bwd_weight_impl(int do_verification,
PerfResults perf_results_local;
bool supports_split_k_optimization = false;
bool is_supported = false;
for(std::size_t split_k_id = 0; split_k_id < split_k_list.size(); split_k_id++)
{
@@ -375,6 +393,7 @@ bool profile_grouped_conv_bwd_weight_impl(int do_verification,
if(op_ptr->IsSupportedArgument(argument_ptr.get()))
{
is_supported = true;
// using atomic add, so need to reset input
wei_device_buf.SetZero();
@@ -493,7 +512,7 @@ bool profile_grouped_conv_bwd_weight_impl(int do_verification,
}
}
if (supports_split_k_optimization)
if (supports_split_k_optimization && is_supported)
{
perf_results_list.push_back(perf_results_local);
}
@@ -506,16 +525,18 @@ bool profile_grouped_conv_bwd_weight_impl(int do_verification,
{
std::cout << "Optimized split-K results:"
<< perf_results_global.print_best_split_k() << std::endl;
const auto& local_per_result = std::find_if(perf_results_list.begin(), perf_results_list.end(),
const auto& local_perf_result = std::find_if(perf_results_list.begin(), perf_results_list.end(),
[&](const PerfResults& res) { return res.opt_split_k_best_op_name_ == perf_results_global.opt_split_k_best_op_name_; });
std::cout << "Global ranking: "
<< std::get<0>(perf_results_global.get_ranking(perf_results_global.opt_split_k_best_op_name_, perf_results_global.opt_split_k_best_arg_))
<< " / " << std::get<1>(perf_results_global.get_ranking(perf_results_global.opt_split_k_best_op_name_, perf_results_global.opt_split_k_best_arg_))
<< std::endl;
std::cout << "Local ranking: "
<< std::get<0>(local_per_result->get_ranking(perf_results_global.opt_split_k_best_op_name_, perf_results_global.opt_split_k_best_arg_))
<< " / " << std::get<1>(local_per_result->get_ranking(perf_results_global.opt_split_k_best_op_name_, perf_results_global.opt_split_k_best_arg_))
<< std::get<0>(local_perf_result->get_ranking(perf_results_global.opt_split_k_best_op_name_, perf_results_global.opt_split_k_best_arg_))
<< " / " << std::get<1>(local_perf_result->get_ranking(perf_results_global.opt_split_k_best_op_name_, perf_results_global.opt_split_k_best_arg_))
<< std::endl;
write_perf_results_to_file(perf_results_global, perf_results_list);
}
return all_pass;