diff --git a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_bwd_data_xdl_nhwc_kyxc_nhwk.hpp b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_bwd_data_xdl_nhwc_kyxc_nhwk.hpp index 9b68c4de43..42bb8db613 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_bwd_data_xdl_nhwc_kyxc_nhwk.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_bwd_data_xdl_nhwc_kyxc_nhwk.hpp @@ -16,6 +16,7 @@ #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp" #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#include "ck/library/utility/numeric.hpp" namespace ck { namespace tensor_operation { @@ -492,6 +493,10 @@ struct DeviceConv2dBwdDataXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K c_grid_desc_m_n_container_.push_back(descs[I2]); } } + c_space_size_bytes = + ck::accumulate_n( + input_spatial_lengths.begin(), NDimSpatial, 1, std::multiplies<>()) * + Conv_N_ * Conv_C_ * sizeof(CDataType); } const ADataType* p_a_grid_; @@ -512,6 +517,8 @@ struct DeviceConv2dBwdDataXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K std::vector conv_filter_dilations_; std::vector input_left_pads_; std::vector input_right_pads_; + + long_index_t c_space_size_bytes; }; // Invoker @@ -571,18 +578,47 @@ struct DeviceConv2dBwdDataXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K DeviceOp::BGridDesc_K0_N_K1, DeviceOp::CGridDesc_M_N, true>; - - ave_time += launch_and_time_kernel(stream_config, - kernel, - dim3(gdx, gdy, gdz), - dim3(BlockSize), - 0, - arg.p_a_grid_, - arg.p_b_grid_, - arg.p_c_grid_, - arg.a_grid_desc_k0_m_k1_container_[i], - arg.b_grid_desc_k0_n_k1_container_[i], - arg.c_grid_desc_m_n_container_[i]); + if(stream_config.flush_cache) + { + // Clear input only for perf measurement. + // For non-grouped solver user has to clear input on his own. + const auto clear_input = [&]() { + if(i == 0) + { + hip_check_error(hipMemsetAsync(arg.p_c_grid_, + 0, + arg.c_space_size_bytes, + stream_config.stream_id_)); + } + }; + ave_time += launch_and_time_kernel_with_preprocess_flush_cache( + stream_config, + clear_input, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + arg.p_a_grid_, + arg.p_b_grid_, + arg.p_c_grid_, + arg.a_grid_desc_k0_m_k1_container_[i], + arg.b_grid_desc_k0_n_k1_container_[i], + arg.c_grid_desc_m_n_container_[i]); + } + else + { + ave_time += launch_and_time_kernel(stream_config, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + arg.p_a_grid_, + arg.p_b_grid_, + arg.p_c_grid_, + arg.a_grid_desc_k0_m_k1_container_[i], + arg.b_grid_desc_k0_n_k1_container_[i], + arg.c_grid_desc_m_n_container_[i]); + } } else { @@ -594,18 +630,47 @@ struct DeviceConv2dBwdDataXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K DeviceOp::BGridDesc_K0_N_K1, DeviceOp::CGridDesc_M_N, false>; - - ave_time += launch_and_time_kernel(stream_config, - kernel, - dim3(gdx, gdy, gdz), - dim3(BlockSize), - 0, - arg.p_a_grid_, - arg.p_b_grid_, - arg.p_c_grid_, - arg.a_grid_desc_k0_m_k1_container_[i], - arg.b_grid_desc_k0_n_k1_container_[i], - arg.c_grid_desc_m_n_container_[i]); + if(stream_config.flush_cache) + { + // Clear input only for perf measurement. + // For non-grouped solver user has to clear input on his own. + const auto clear_input = [&]() { + if(i == 0) + { + hip_check_error(hipMemsetAsync(arg.p_c_grid_, + 0, + arg.c_space_size_bytes, + stream_config.stream_id_)); + } + }; + ave_time += launch_and_time_kernel_with_preprocess_flush_cache( + stream_config, + clear_input, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + arg.p_a_grid_, + arg.p_b_grid_, + arg.p_c_grid_, + arg.a_grid_desc_k0_m_k1_container_[i], + arg.b_grid_desc_k0_n_k1_container_[i], + arg.c_grid_desc_m_n_container_[i]); + } + else + { + ave_time += launch_and_time_kernel(stream_config, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + arg.p_a_grid_, + arg.p_b_grid_, + arg.p_c_grid_, + arg.a_grid_desc_k0_m_k1_container_[i], + arg.b_grid_desc_k0_n_k1_container_[i], + arg.c_grid_desc_m_n_container_[i]); + } } } return ave_time; diff --git a/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp index 022bda3ed0..ebddd90381 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp @@ -16,6 +16,7 @@ #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp" #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#include "ck/library/utility/numeric.hpp" namespace ck { namespace tensor_operation { @@ -1050,6 +1051,10 @@ struct DeviceConvNdBwdDataNwcKxcNwk_Xdl input_right_pads_{input_right_pads} { CreateABCDesc(); + c_space_size_bytes = + ck::accumulate_n( + input_spatial_lengths.begin(), NDimSpatial, 1, std::multiplies<>()) * + Conv_N_ * Conv_C_ * sizeof(CDataType); } template ::type = false> @@ -1216,6 +1221,8 @@ struct DeviceConvNdBwdDataNwcKxcNwk_Xdl std::vector conv_filter_dilations_; std::vector input_left_pads_; std::vector input_right_pads_; + + long_index_t c_space_size_bytes; }; // Invoker @@ -1273,18 +1280,47 @@ struct DeviceConvNdBwdDataNwcKxcNwk_Xdl DeviceOp::BGridDesc_K0_N_K1, DeviceOp::CGridDesc_M_N, true>; - - ave_time += launch_and_time_kernel(stream_config, - kernel, - dim3(gdx, gdy, gdz), - dim3(BlockSize), - 0, - arg.p_a_grid_, - arg.p_b_grid_, - arg.p_c_grid_, - arg.a_grid_desc_k0_m_k1_container_[i], - arg.b_grid_desc_k0_n_k1_container_[i], - arg.c_grid_desc_m_n_container_[i]); + if(stream_config.flush_cache) + { + // Clear input only for perf measurement. + // For non-grouped solver user has to clear input on his own. + const auto clear_input = [&]() { + if(i == 0) + { + hip_check_error(hipMemsetAsync(arg.p_c_grid_, + 0, + arg.c_space_size_bytes, + stream_config.stream_id_)); + } + }; + ave_time += launch_and_time_kernel_with_preprocess_flush_cache( + stream_config, + clear_input, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + arg.p_a_grid_, + arg.p_b_grid_, + arg.p_c_grid_, + arg.a_grid_desc_k0_m_k1_container_[i], + arg.b_grid_desc_k0_n_k1_container_[i], + arg.c_grid_desc_m_n_container_[i]); + } + else + { + ave_time += launch_and_time_kernel(stream_config, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + arg.p_a_grid_, + arg.p_b_grid_, + arg.p_c_grid_, + arg.a_grid_desc_k0_m_k1_container_[i], + arg.b_grid_desc_k0_n_k1_container_[i], + arg.c_grid_desc_m_n_container_[i]); + } } else { @@ -1296,18 +1332,47 @@ struct DeviceConvNdBwdDataNwcKxcNwk_Xdl DeviceOp::BGridDesc_K0_N_K1, DeviceOp::CGridDesc_M_N, false>; - - ave_time += launch_and_time_kernel(stream_config, - kernel, - dim3(gdx, gdy, gdz), - dim3(BlockSize), - 0, - arg.p_a_grid_, - arg.p_b_grid_, - arg.p_c_grid_, - arg.a_grid_desc_k0_m_k1_container_[i], - arg.b_grid_desc_k0_n_k1_container_[i], - arg.c_grid_desc_m_n_container_[i]); + if(stream_config.flush_cache) + { + // Clear input only for perf measurement. + // For non-grouped solver user has to clear input on his own. + const auto clear_input = [&]() { + if(i == 0) + { + hip_check_error(hipMemsetAsync(arg.p_c_grid_, + 0, + arg.c_space_size_bytes, + stream_config.stream_id_)); + } + }; + ave_time += launch_and_time_kernel_with_preprocess_flush_cache( + stream_config, + clear_input, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + arg.p_a_grid_, + arg.p_b_grid_, + arg.p_c_grid_, + arg.a_grid_desc_k0_m_k1_container_[i], + arg.b_grid_desc_k0_n_k1_container_[i], + arg.c_grid_desc_m_n_container_[i]); + } + else + { + ave_time += launch_and_time_kernel(stream_config, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + arg.p_a_grid_, + arg.p_b_grid_, + arg.p_c_grid_, + arg.a_grid_desc_k0_m_k1_container_[i], + arg.b_grid_desc_k0_n_k1_container_[i], + arg.c_grid_desc_m_n_container_[i]); + } } } return ave_time; diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp index 8b71a4fa40..825a3f8b5c 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp @@ -1225,26 +1225,50 @@ struct DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 has_main_loop, no_main_loop, CTranspose>; - - return launch_and_time_kernel_with_preprocess( - stream_config, - clear_workspace, - kernel, - dim3(gdx, gdy, gdz), - dim3(BlockSize), - 0, - p_b_grid, - p_a_grid, - arg.p_ds_grid_, - p_e_grid, - gemm_kernel_args, - gemms_count_for_set, - arg.b_element_op_, - arg.a_element_op_, - arg.cde_element_op_, - arg.compute_ptr_offset_of_batch_, - arg.compute_ptr_offset_of_n_, - arg.k_batch_); + if(stream_config.flush_cache) + { + return launch_and_time_kernel_with_preprocess_flush_cache( + stream_config, + clear_workspace, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + p_b_grid, + p_a_grid, + arg.p_ds_grid_, + p_e_grid, + gemm_kernel_args, + gemms_count_for_set, + arg.b_element_op_, + arg.a_element_op_, + arg.cde_element_op_, + arg.compute_ptr_offset_of_batch_, + arg.compute_ptr_offset_of_n_, + arg.k_batch_); + } + else + { + return launch_and_time_kernel_with_preprocess( + stream_config, + clear_workspace, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + p_b_grid, + p_a_grid, + arg.p_ds_grid_, + p_e_grid, + gemm_kernel_args, + gemms_count_for_set, + arg.b_element_op_, + arg.a_element_op_, + arg.cde_element_op_, + arg.compute_ptr_offset_of_batch_, + arg.compute_ptr_offset_of_n_, + arg.k_batch_); + } } else { @@ -1264,26 +1288,50 @@ struct DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 has_main_loop, no_main_loop, CTranspose>; - - return launch_and_time_kernel_with_preprocess( - stream_config, - clear_workspace, - kernel, - dim3(gdx, gdy, gdz), - dim3(BlockSize), - 0, - p_a_grid, - p_b_grid, - arg.p_ds_grid_, - p_e_grid, - gemm_kernel_args, - gemms_count_for_set, - arg.a_element_op_, - arg.b_element_op_, - arg.cde_element_op_, - arg.compute_ptr_offset_of_batch_, - arg.compute_ptr_offset_of_n_, - arg.k_batch_); + if(stream_config.flush_cache) + { + return launch_and_time_kernel_with_preprocess_flush_cache( + stream_config, + clear_workspace, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + p_a_grid, + p_b_grid, + arg.p_ds_grid_, + p_e_grid, + gemm_kernel_args, + gemms_count_for_set, + arg.a_element_op_, + arg.b_element_op_, + arg.cde_element_op_, + arg.compute_ptr_offset_of_batch_, + arg.compute_ptr_offset_of_n_, + arg.k_batch_); + } + else + { + return launch_and_time_kernel_with_preprocess( + stream_config, + clear_workspace, + kernel, + dim3(gdx, gdy, gdz), + dim3(BlockSize), + 0, + p_a_grid, + p_b_grid, + arg.p_ds_grid_, + p_e_grid, + gemm_kernel_args, + gemms_count_for_set, + arg.a_element_op_, + arg.b_element_op_, + arg.cde_element_op_, + arg.compute_ptr_offset_of_batch_, + arg.compute_ptr_offset_of_n_, + arg.k_batch_); + } } }; if(has_loop_in_all_gemm) diff --git a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp index 801207106b..e353dc8b54 100644 --- a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp +++ b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp @@ -903,13 +903,11 @@ struct GroupedConvolutionBackwardDataKernel const auto& d_block_window = MakeDBlockWindows(ds_ptr, kargs, group_id, block_idx_m, block_idx_n); - const index_t num_loop = amd_wave_read_first_lane(TilePartitioner::GetLoopNum(splitted_k)); - const bool has_hot_loop = GemmPipeline::BlockHasHotloop(num_loop); - const TailNumber tail_num = GemmPipeline::GetBlockLoopTailNum(num_loop); + const index_t num_loop = amd_wave_read_first_lane(TilePartitioner::GetLoopNum(splitted_k)); // Run GEMM cooperatively by whole workgroup. const auto& c_block_tile = GemmPipeline{}.template operator()( - a_block_window, b_block_window, num_loop, has_hot_loop, tail_num, smem_ptr_0); + a_block_window, b_block_window, num_loop, smem_ptr_0); const index_t k_batch = amd_wave_read_first_lane(kargs.k_batch); diff --git a/profiler/include/profiler/grouped_convolution_backward_data_tile_algs.hpp b/profiler/include/profiler/grouped_convolution_backward_data_tile_algs.hpp index 2fa2019b07..4bbf3eb6db 100644 --- a/profiler/include/profiler/grouped_convolution_backward_data_tile_algs.hpp +++ b/profiler/include/profiler/grouped_convolution_backward_data_tile_algs.hpp @@ -65,7 +65,9 @@ run_grouped_conv_backward_data_tile_algs(const ckt::Args& args, const ckt::Outputs& outputs, const ck_tile::stream_config& s_conf) { - float best_avg_time = std::numeric_limits::max(); + // Run first instance as dummy to get proper time from the first instance + bool dummy_run_executed = false; + float best_avg_time = std::numeric_limits::max(); std::string best_op_name, op_name; int best_split_k = 0; ck::index_t best_instance_index = -1; @@ -121,6 +123,13 @@ run_grouped_conv_backward_data_tile_algs(const ckt::Args& args, run_alg_func(args_k_batch, inputs, outputs, s_conf); if(is_supported) { + if((s_conf.time_kernel_ || s_conf.flush_cache_) && !dummy_run_executed) + { + // Run first instance twice + std::tie(is_supported, avg_time, op_name) = + run_alg_func(args_k_batch, inputs, outputs, s_conf); + dummy_run_executed = true; + } ckt::ValidationReport report; auto&& [rtol, atol] = get_rtol_atol(num_accums, k_batch, max_accumulated_value); diff --git a/profiler/include/profiler/grouped_convolution_backward_weight_tile_algs.hpp b/profiler/include/profiler/grouped_convolution_backward_weight_tile_algs.hpp index e79fc44e8d..697849a019 100644 --- a/profiler/include/profiler/grouped_convolution_backward_weight_tile_algs.hpp +++ b/profiler/include/profiler/grouped_convolution_backward_weight_tile_algs.hpp @@ -106,17 +106,17 @@ run_grouped_conv_backward_weight_tile_algs(const ckt::Args& args, { ckt::Args args_k_batch = args; args_k_batch.k_batch = k_batch; - if((s_conf.time_kernel_ || s_conf.flush_cache_) && !dummy_run_executed) - { - // Run first instance twice when profiling to stabilize timing - std::tie(is_supported, avg_time, op_name) = - run_alg_func(args_k_batch, inputs, outputs, s_conf); - dummy_run_executed = true; - } std::tie(is_supported, avg_time, op_name) = run_alg_func(args_k_batch, inputs, outputs, s_conf); if(is_supported) { + if((s_conf.time_kernel_ || s_conf.flush_cache_) && !dummy_run_executed) + { + // Run first instance twice when profiling to stabilize timing + std::tie(is_supported, avg_time, op_name) = + run_alg_func(args_k_batch, inputs, outputs, s_conf); + dummy_run_executed = true; + } ckt::ValidationReport report; auto&& [rtol, atol] = get_rtol_atol(num_accums, k_batch, max_accumulated_value); diff --git a/profiler/include/profiler/grouped_convolution_forward_tile_algs.hpp b/profiler/include/profiler/grouped_convolution_forward_tile_algs.hpp index 054da8057a..a5e79706be 100644 --- a/profiler/include/profiler/grouped_convolution_forward_tile_algs.hpp +++ b/profiler/include/profiler/grouped_convolution_forward_tile_algs.hpp @@ -86,15 +86,16 @@ run_grouped_conv_forward_tile_algs(const ckt::Args& args, auto ref_conv = ReferenceInstance{}; auto ref_result = ckt::run(ref_conv, args, inputs, reference.get()); auto run_alg = [&](auto&& run_alg_func) { - if(!dummy_run_executed) - { - // Run first instance twice - std::tie(is_supported, avg_time, op_name) = run_alg_func(args, inputs, outputs, s_conf); - dummy_run_executed = true; - } std::tie(is_supported, avg_time, op_name) = run_alg_func(args, inputs, outputs, s_conf); if(is_supported) { + if((s_conf.time_kernel_ || s_conf.flush_cache_) && !dummy_run_executed) + { + // Run first instance twice + std::tie(is_supported, avg_time, op_name) = + run_alg_func(args, inputs, outputs, s_conf); + dummy_run_executed = true; + } best_avg_time = std::min(best_avg_time, avg_time); best_op_name = best_avg_time < avg_time ? best_op_name : op_name; std::cout << "Perf: " << std::setw(10) << avg_time << " ms," << " " << op_name diff --git a/profiler/include/profiler/profile_conv_bwd_data_impl.hpp b/profiler/include/profiler/profile_conv_bwd_data_impl.hpp index 937fb24f5a..2d2f9982e9 100644 --- a/profiler/include/profiler/profile_conv_bwd_data_impl.hpp +++ b/profiler/include/profiler/profile_conv_bwd_data_impl.hpp @@ -197,10 +197,11 @@ bool profile_conv_bwd_data_impl(int do_verification, } std::string best_op_name; - float best_avg_time = 0; - float best_tflops = 0; - float best_gb_per_sec = 0; - int num_kernel = 0; + float best_avg_time = 0; + float best_tflops = 0; + float best_gb_per_sec = 0; + int num_kernel = 0; + bool dummy_run_executed = false; for(auto& op_ptr : op_ptrs) { @@ -230,16 +231,38 @@ bool profile_conv_bwd_data_impl(int do_verification, // skip test if instance_index is specified continue; } - // for conv bwd data, some input tensor element are zero, but not written by kernel, - // need to set zero - in_device_buf.SetZero(); + if(!time_kernel) + { + // Don't clear for perf measurement. + // For non-grouped solver user has to clear input on his own. + // for conv bwd data, some input tensor element are zero, but not written by kernel, + // need to set zero + in_device_buf.SetZero(); + } std::string op_name = op_ptr->GetTypeString(); auto invoker_ptr = op_ptr->MakeInvokerPointer(); - float avg_time = - invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, time_kernel}); + // Run first instance twice to get proper time + if(time_kernel && !dummy_run_executed) + { + invoker_ptr->Run(argument_ptr.get(), + StreamConfig{nullptr, + time_kernel, + 0 /*log_level*/, + 5 /*cold_iters*/, + 50 /*nrepeat_*/, + time_kernel /*flush_cache*/}); + dummy_run_executed = true; + } + float avg_time = invoker_ptr->Run(argument_ptr.get(), + StreamConfig{nullptr, + time_kernel, + 0 /*log_level*/, + 5 /*cold_iters*/, + 50 /*nrepeat_*/, + time_kernel /*flush_cache*/}); std::size_t flop = conv_param.GetFlops(); std::size_t num_btype = conv_param.GetByte(); diff --git a/profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp b/profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp index 8a5bf966b7..3d053b1dc1 100644 --- a/profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp +++ b/profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp @@ -287,6 +287,8 @@ bool profile_grouped_conv_bwd_data_impl(int do_verification, bool pass = true; index_t num_kernel = 0; index_t valid_instances = 0; + bool dummy_run_executed = false; + auto run_impl = [&](auto& op_ptr, auto& argument_ptr, const index_t& split_k_for_run) { // workspace_sz will be equal to 0 for other layout than NGCHW const std::size_t workspace_sz = op_ptr->GetWorkSpaceSize(argument_ptr.get()); @@ -317,8 +319,25 @@ bool profile_grouped_conv_bwd_data_impl(int do_verification, auto invoker_ptr = op_ptr->MakeInvokerPointer(); - float avg_time = - invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, time_kernel}); + // Run first instance twice to get proper time + if(time_kernel && !dummy_run_executed) + { + invoker_ptr->Run(argument_ptr.get(), + StreamConfig{nullptr, + time_kernel, + 0 /*log_level*/, + 5 /*cold_iters*/, + 50 /*nrepeat_*/, + time_kernel /*flush_cache*/}); + dummy_run_executed = true; + } + float avg_time = invoker_ptr->Run(argument_ptr.get(), + StreamConfig{nullptr, + time_kernel, + 0 /*log_level*/, + 5 /*cold_iters*/, + 50 /*nrepeat_*/, + time_kernel /*flush_cache*/}); std::size_t flop = conv_param.GetFlops(); std::size_t num_btype = conv_param.GetByte(); @@ -495,7 +514,6 @@ bool profile_grouped_conv_bwd_data_impl(int do_verification, { std::cout << "\nValid instances for this problem:" << std::endl; } - for(auto& op_ptr : op_ptrs) { for(std::size_t split_k_id = 0; split_k_id < split_k_list.size(); split_k_id++) diff --git a/profiler/include/profiler/profile_grouped_conv_fwd_impl.hpp b/profiler/include/profiler/profile_grouped_conv_fwd_impl.hpp index 44d000f8c6..24bc67a647 100644 --- a/profiler/include/profiler/profile_grouped_conv_fwd_impl.hpp +++ b/profiler/include/profiler/profile_grouped_conv_fwd_impl.hpp @@ -296,7 +296,8 @@ bool profile_grouped_conv_fwd_impl(int do_verification, index_t best_instance_index = 0; // profile device op instances - bool pass = true; + bool pass = true; + bool dummy_run_executed = false; auto run_impl = [&](auto& op_ptr, auto& argument_ptr) { // workspace_sz will be equal to 0 for other layout than NGCHW @@ -331,6 +332,19 @@ bool profile_grouped_conv_fwd_impl(int do_verification, auto invoker_ptr = op_ptr->MakeInvokerPointer(); + // Run first instance twice to get proper time + if(time_kernel && !dummy_run_executed) + { + invoker_ptr->Run(argument_ptr.get(), + StreamConfig{nullptr, + time_kernel, + 0 /*log_level*/, + 5 /*cold_iters*/, + 50 /*nrepeat_*/, + time_kernel /*flush_cache*/}); + dummy_run_executed = true; + } + float avg_time = invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, time_kernel, @@ -437,30 +451,6 @@ bool profile_grouped_conv_fwd_impl(int do_verification, std::cout << "\nValid instances for this problem:" << std::endl; } - // Run first instance twice to get proper time - { - auto argument_ptr = op_ptrs[0]->MakeArgumentPointer(in_device_buf.GetDeviceBuffer(), - wei_device_buf.GetDeviceBuffer(), - {}, - out_device_buf.GetDeviceBuffer(), - a_g_n_c_wis_lengths, - a_g_n_c_wis_strides, - b_g_k_c_xs_lengths, - b_g_k_c_xs_strides, - {}, - {}, - e_g_n_k_wos_lengths, - e_g_n_k_wos_strides, - conv_filter_strides, - conv_filter_dilations, - input_left_pads, - input_right_pads, - in_element_op, - wei_element_op, - out_element_op); - - run_impl(op_ptrs[0], argument_ptr); - } for(auto& op_ptr : op_ptrs) { auto argument_ptr = op_ptr->MakeArgumentPointer(in_device_buf.GetDeviceBuffer(), diff --git a/profiler/src/profile_grouped_conv_bwd_data_tile.cpp b/profiler/src/profile_grouped_conv_bwd_data_tile.cpp index fe51056805..280af3fe00 100644 --- a/profiler/src/profile_grouped_conv_bwd_data_tile.cpp +++ b/profiler/src/profile_grouped_conv_bwd_data_tile.cpp @@ -89,7 +89,8 @@ int call_profiler(const ckt::Args& args, 0 /*log_level*/, 5 /*cold_iters*/, 50 /*nrepeat_*/, - true /*is_gpu_timer_*/}); + true /*is_gpu_timer_*/, + time_kernel /*flush_cache*/}); if(time_kernel) { std::cout << "\nBest configuration parameters:" << "\n\tname: " << op_name << " (instance "