diff --git a/profiler/CMakeLists.txt b/profiler/CMakeLists.txt index bdd7125ac1..cf10b9997c 100644 --- a/profiler/CMakeLists.txt +++ b/profiler/CMakeLists.txt @@ -3,3 +3,4 @@ include_directories(BEFORE ) add_subdirectory(src) +add_subdirectory(ck_tile) diff --git a/profiler/ck_tile/CMakeLists.txt b/profiler/ck_tile/CMakeLists.txt new file mode 100644 index 0000000000..ee775efc03 --- /dev/null +++ b/profiler/ck_tile/CMakeLists.txt @@ -0,0 +1,5 @@ +include_directories(BEFORE + ${CMAKE_CURRENT_LIST_DIR}/include +) + +add_subdirectory(src) \ No newline at end of file diff --git a/profiler/ck_tile/include/conv_parameters.hpp b/profiler/ck_tile/include/conv_parameters.hpp new file mode 100644 index 0000000000..223526586d --- /dev/null +++ b/profiler/ck_tile/include/conv_parameters.hpp @@ -0,0 +1,304 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include +#include +#include +#include + +#include "ck/ck.hpp" +#include "ck/library/utility/numeric.hpp" +#include "ck/host_utility/io.hpp" + +namespace ck_tile { +namespace utils { +namespace conv { + +struct ConvParam +{ + ConvParam(); + ConvParam(ck::index_t n_dim, + ck::index_t group_count, + ck::index_t n_batch, + ck::index_t n_out_channels, + ck::index_t n_in_channels, + const std::vector& filters_len, + const std::vector& input_len, + const std::vector& strides, + const std::vector& dilations, + const std::vector& left_pads, + const std::vector& right_pads); + + ConvParam(ck::long_index_t n_dim, + ck::long_index_t group_count, + ck::long_index_t n_batch, + ck::long_index_t n_out_channels, + ck::long_index_t n_in_channels, + const std::vector& filters_len, + const std::vector& input_len, + const std::vector& strides, + const std::vector& dilations, + const std::vector& left_pads, + const std::vector& right_pads); + + ck::long_index_t num_dim_spatial_; + ck::long_index_t G_; + ck::long_index_t N_; + ck::long_index_t K_; + ck::long_index_t C_; + + std::vector filter_spatial_lengths_; + std::vector input_spatial_lengths_; + std::vector output_spatial_lengths_; + + std::vector conv_filter_strides_; + std::vector conv_filter_dilations_; + + std::vector input_left_pads_; + std::vector input_right_pads_; + + std::vector GetOutputSpatialLengths() const; + + std::size_t GetFlops() const; + + template + std::size_t GetInputByte() const + { + // sizeof(InDataType) * (G * N * C * ) + + return sizeof(InDataType) * + (G_ * N_ * C_ * + ck::accumulate_n( + std::begin(input_spatial_lengths_), num_dim_spatial_, 1, std::multiplies<>())); + } + + template + std::size_t GetWeightByte() const + { + // sizeof(WeiDataType) * (G * K * C * ) + + return sizeof(WeiDataType) * + (G_ * K_ * C_ * + ck::accumulate_n( + std::begin(filter_spatial_lengths_), num_dim_spatial_, 1, std::multiplies<>())); + } + + template + std::size_t GetOutputByte() const + { + // sizeof(OutDataType) * (G * N * K * ); + return sizeof(OutDataType) * (G_ * N_ * K_ * + std::accumulate(std::begin(output_spatial_lengths_), + std::end(output_spatial_lengths_), + static_cast(1), + std::multiplies())); + } + + template + std::size_t GetByte() const + { + return GetInputByte() + GetWeightByte() + + GetOutputByte(); + } +}; + +ConvParam::ConvParam(ck::index_t n_dim, + ck::index_t group_count, + ck::index_t n_batch, + ck::index_t n_out_channels, + ck::index_t n_in_channels, + const std::vector& filters_len, + const std::vector& input_len, + const std::vector& strides, + const std::vector& dilations, + const std::vector& left_pads, + const std::vector& right_pads) + : num_dim_spatial_(static_cast(n_dim)), + G_(static_cast(group_count)), + N_(static_cast(n_batch)), + K_(static_cast(n_out_channels)), + C_(static_cast(n_in_channels)), + filter_spatial_lengths_(num_dim_spatial_), + input_spatial_lengths_(num_dim_spatial_), + output_spatial_lengths_(num_dim_spatial_), + conv_filter_strides_(num_dim_spatial_), + conv_filter_dilations_(num_dim_spatial_), + input_left_pads_(num_dim_spatial_), + input_right_pads_(num_dim_spatial_) +{ + if(static_cast(filter_spatial_lengths_.size()) != num_dim_spatial_ || + static_cast(input_spatial_lengths_.size()) != num_dim_spatial_ || + static_cast(conv_filter_strides_.size()) != num_dim_spatial_ || + static_cast(conv_filter_dilations_.size()) != num_dim_spatial_ || + static_cast(input_left_pads_.size()) != num_dim_spatial_ || + static_cast(input_right_pads_.size()) != num_dim_spatial_) + { + throw( + std::runtime_error("ConvParam::ConvParam: " + "parameter size is different from number of declared dimensions!")); + } + + for(ck::index_t i = 0; i < num_dim_spatial_; ++i) + { + filter_spatial_lengths_[i] = static_cast(filters_len[i]); + input_spatial_lengths_[i] = static_cast(input_len[i]); + conv_filter_strides_[i] = static_cast(strides[i]); + conv_filter_dilations_[i] = static_cast(dilations[i]); + input_left_pads_[i] = static_cast(left_pads[i]); + input_right_pads_[i] = static_cast(right_pads[i]); + + // XEff = (X - 1) * conv_dilation_w + 1; + // Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1; + const ck::long_index_t x_eff = + (filter_spatial_lengths_[i] - 1) * conv_filter_dilations_[i] + 1; + + output_spatial_lengths_[i] = + (input_spatial_lengths_[i] + input_left_pads_[i] + input_right_pads_[i] - x_eff) / + conv_filter_strides_[i] + + 1; + } +} + +ConvParam::ConvParam(ck::long_index_t n_dim, + ck::long_index_t group_count, + ck::long_index_t n_batch, + ck::long_index_t n_out_channels, + ck::long_index_t n_in_channels, + const std::vector& filters_len, + const std::vector& input_len, + const std::vector& strides, + const std::vector& dilations, + const std::vector& left_pads, + const std::vector& right_pads) + : num_dim_spatial_(n_dim), + G_(group_count), + N_(n_batch), + K_(n_out_channels), + C_(n_in_channels), + filter_spatial_lengths_(filters_len), + input_spatial_lengths_(input_len), + output_spatial_lengths_(num_dim_spatial_), + conv_filter_strides_(strides), + conv_filter_dilations_(dilations), + input_left_pads_(left_pads), + input_right_pads_(right_pads) +{ + if(static_cast(filter_spatial_lengths_.size()) != num_dim_spatial_ || + static_cast(input_spatial_lengths_.size()) != num_dim_spatial_ || + static_cast(conv_filter_strides_.size()) != num_dim_spatial_ || + static_cast(conv_filter_dilations_.size()) != num_dim_spatial_ || + static_cast(input_left_pads_.size()) != num_dim_spatial_ || + static_cast(input_right_pads_.size()) != num_dim_spatial_) + { + throw( + std::runtime_error("ConvParam::ConvParam: " + "parameter size is different from number of declared dimensions!")); + } + + for(ck::index_t i = 0; i < num_dim_spatial_; ++i) + { + // XEff = (X - 1) * conv_dilation_w + 1; + // Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1; + const ck::long_index_t x_eff = + (filter_spatial_lengths_[i] - 1) * conv_filter_dilations_[i] + 1; + + output_spatial_lengths_[i] = + (input_spatial_lengths_[i] + input_left_pads_[i] + input_right_pads_[i] - x_eff) / + conv_filter_strides_[i] + + 1; + } +} + +ConvParam::ConvParam() + : ConvParam::ConvParam(2, 1, 128, 256, 192, {3, 3}, {71, 71}, {2, 2}, {1, 1}, {1, 1}, {1, 1}) +{ +} + +std::vector ConvParam::GetOutputSpatialLengths() const +{ + return output_spatial_lengths_; +} + +std::size_t ConvParam::GetFlops() const +{ + // 2 * G * N * K * C * * + return static_cast(2) * G_ * N_ * K_ * C_ * + ck::accumulate_n( + std::begin(output_spatial_lengths_), num_dim_spatial_, 1, std::multiplies<>()) * + ck::accumulate_n( + std::begin(filter_spatial_lengths_), num_dim_spatial_, 1, std::multiplies<>()); +} + +ck_tile::utils::conv::ConvParam parse_conv_param(int num_dim_spatial, int arg_idx, char* const argv[]) +{ + const ck::long_index_t G = std::stol(argv[arg_idx++]); + const ck::long_index_t N = std::stol(argv[arg_idx++]); + const ck::long_index_t K = std::stol(argv[arg_idx++]); + const ck::long_index_t C = std::stol(argv[arg_idx++]); + + std::vector filter_spatial_lengths(num_dim_spatial); + std::vector input_spatial_lengths(num_dim_spatial); + std::vector conv_filter_strides(num_dim_spatial); + std::vector conv_filter_dilations(num_dim_spatial); + std::vector input_left_pads(num_dim_spatial); + std::vector input_right_pads(num_dim_spatial); + + for(int i = 0; i < num_dim_spatial; ++i) + { + filter_spatial_lengths[i] = std::stol(argv[arg_idx++]); + } + + for(int i = 0; i < num_dim_spatial; ++i) + { + input_spatial_lengths[i] = std::stol(argv[arg_idx++]); + } + + for(int i = 0; i < num_dim_spatial; ++i) + { + conv_filter_strides[i] = std::stol(argv[arg_idx++]); + } + + for(int i = 0; i < num_dim_spatial; ++i) + { + conv_filter_dilations[i] = std::stol(argv[arg_idx++]); + } + + for(int i = 0; i < num_dim_spatial; ++i) + { + input_left_pads[i] = std::stol(argv[arg_idx++]); + } + + for(int i = 0; i < num_dim_spatial; ++i) + { + input_right_pads[i] = std::stol(argv[arg_idx++]); + } + + return ck_tile::utils::conv::ConvParam{num_dim_spatial, + G, + N, + K, + C, + filter_spatial_lengths, + input_spatial_lengths, + conv_filter_strides, + conv_filter_dilations, + input_left_pads, + input_right_pads}; +} +} // namespace conv +} // namespace utils +} // namespace ck_tile + +std::ostream& operator<<(std::ostream& os, const ck_tile::utils::conv::ConvParam& p) +{ + os << "ConvParam {" << "\nnum_dim_spatial: " << p.num_dim_spatial_ << "\nG: " << p.G_ + << "\nN: " << p.N_ << "\nK: " << p.K_ << "\nC: " << p.C_ + << "\nfilter_spatial_lengths: " << p.filter_spatial_lengths_ + << "\ninput_spatial_lengths: " << p.input_spatial_lengths_ + << "\nconv_filter_strides: " << p.conv_filter_strides_ + << "\nconv_filter_dilations: " << p.conv_filter_dilations_ + << "\ninput_left_pads: " << p.input_left_pads_ + << "\ninput_right_pads: " << p.input_right_pads_ << "}\n"; + + return os; +} diff --git a/profiler/ck_tile/include/tile_profile_grouped_conv_bwd_weight_impl.hpp b/profiler/ck_tile/include/tile_profile_grouped_conv_bwd_weight_impl.hpp new file mode 100644 index 0000000000..f979aaeeef --- /dev/null +++ b/profiler/ck_tile/include/tile_profile_grouped_conv_bwd_weight_impl.hpp @@ -0,0 +1,349 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include +#include +#include +#include +#include + +#include "ck/ck.hpp" +#include "conv_parameters.hpp" +// #include "ck/tensor_operation/gpu/device/tensor_layout.hpp" +// #include "ck/tensor_operation/gpu/device/impl/split_k_arg.hpp" +// #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" + +// #include "ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_weight.hpp" + +// #include "ck/library/utility/check_err.hpp" +// #include "ck/library/utility/device_memory.hpp" +// #include "ck/library/utility/host_tensor.hpp" +// #include "ck/library/utility/host_tensor_generator.hpp" +// #include "ck/library/utility/convolution_parameter.hpp" +// #include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp" +// #include "ck/library/reference_tensor_operation/cpu/reference_conv_bwd_weight.hpp" + +namespace ck_tile { +namespace profiler { + +template +bool profile_grouped_conv_bwd_weight_impl(int /*do_verification*/, + int /*init_method*/, + bool /*do_log*/, + bool /*time_kernel*/, + const ck_tile::utils::conv::ConvParam& /*conv_param*/, + const std::string& /*split_k*/) + //ck::index_t instance_index = -1) +{ + return true; + + // using InElementOp = ck::tensor_operation::element_wise::PassThrough; + // using WeiElementOp = ck::tensor_operation::element_wise::PassThrough; + // using OutElementOp = ck::tensor_operation::element_wise::PassThrough; + + // const auto in_element_op = InElementOp{}; + // const auto wei_element_op = WeiElementOp{}; + // const auto out_element_op = OutElementOp{}; + + // const auto in_g_n_c_wis_desc = + // ck::utils::conv::make_input_host_tensor_descriptor_g_n_c_wis_packed(conv_param); + + // const auto wei_g_k_c_xs_desc = + // ck::utils::conv::make_weight_host_tensor_descriptor_g_k_c_xs_packed(conv_param); + + // const auto out_g_n_k_wos_desc = + // ck::utils::conv::make_output_host_tensor_descriptor_g_n_k_wos_packed(conv_param); + + // Tensor input(in_g_n_c_wis_desc); + // Tensor weight_host_result(wei_g_k_c_xs_desc); + // Tensor weight_device_result(wei_g_k_c_xs_desc); + // Tensor output(out_g_n_k_wos_desc); + + // std::cout << "input: " << input.mDesc << std::endl; + // std::cout << "weight: " << weight_host_result.mDesc << std::endl; + // std::cout << "output: " << output.mDesc << std::endl; + + // switch(init_method) + // { + // case 0: break; + // case 1: + // input.GenerateTensorValue(GeneratorTensor_2{-5, 5}); + // output.GenerateTensorValue(GeneratorTensor_2{-5, 5}); + // break; + // default: + // input.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); + // output.GenerateTensorValue(GeneratorTensor_3{-0.5, 0.5}); + // } + + // DeviceMem in_device_buf(sizeof(InDataType) * input.mDesc.GetElementSpaceSize()); + // DeviceMem wei_device_buf(sizeof(WeiDataType) * + // weight_device_result.mDesc.GetElementSpaceSize()); + // DeviceMem out_device_buf(sizeof(OutDataType) * output.mDesc.GetElementSpaceSize()); + + // in_device_buf.ToDevice(input.mData.data()); + // out_device_buf.ToDevice(output.mData.data()); + + // float max_accumulated_value = 0; + // if(do_verification) + // { + // auto ref_conv = ck::tensor_operation::host::ReferenceConvBwdWeight{}; + // auto ref_invoker = ref_conv.MakeInvoker(); + // auto ref_argument = ref_conv.MakeArgument(input, + // weight_host_result, + // output, + // conv_param.conv_filter_strides_, + // conv_param.conv_filter_dilations_, + // conv_param.input_left_pads_, + // conv_param.input_right_pads_, + // in_element_op, + // wei_element_op, + // out_element_op, + // {}, + // {}, + // {}); + + // ref_invoker.Run(ref_argument); + // max_accumulated_value = + // *std::max_element(weight_host_result.mData.begin(), weight_host_result.mData.end()); + // } + + // using DeviceOp = ck::tensor_operation::device::DeviceGroupedConvBwdWeight; + + // // get device op instances + // const auto op_ptrs = ck::tensor_operation::device::instance::DeviceOperationInstanceFactory< + // DeviceOp>::GetInstances(); + + // std::cout << "found " << op_ptrs.size() << " instances" << std::endl; + + // std::string best_op_name; + // float best_avg_time = 0; + // float best_tflops = 0; + // float best_gb_per_sec = 0; + // std::string best_split_k("1"); + + // // profile device Conv instances + // bool all_pass = true; + + // std::array input_lengths{}; + // std::array filter_lengths{}; + // std::array output_lengths{}; + // std::array input_strides{}; + // std::array weights_strides{}; + // std::array output_strides{}; + // std::array conv_filter_strides{}; + // std::array conv_filter_dilations{}; + // std::array input_left_pads{}; + // std::array input_right_pads{}; + + // auto range_copy = [](const auto& from, auto to) { std::copy(begin(from), end(from), to); }; + + // range_copy(in_g_n_c_wis_desc.GetLengths(), begin(input_lengths)); + // range_copy(in_g_n_c_wis_desc.GetStrides(), begin(input_strides)); + // range_copy(wei_g_k_c_xs_desc.GetLengths(), begin(filter_lengths)); + // range_copy(wei_g_k_c_xs_desc.GetStrides(), begin(weights_strides)); + // range_copy(out_g_n_k_wos_desc.GetLengths(), begin(output_lengths)); + // range_copy(out_g_n_k_wos_desc.GetStrides(), begin(output_strides)); + // range_copy(conv_param.conv_filter_strides_, begin(conv_filter_strides)); + // range_copy(conv_param.conv_filter_dilations_, begin(conv_filter_dilations)); + // range_copy(conv_param.input_left_pads_, begin(input_left_pads)); + // range_copy(conv_param.input_right_pads_, begin(input_right_pads)); + + // std::vector split_k_list = {/*auto deduce value*/ -1, 1, 2, 4, 8, 16, 32, 64, 128}; + + // if(split_k != "all") + // { + // try + // { + // ck::index_t split_k_value = std::stoi(split_k); + // split_k_list = {split_k_value}; + // } + // catch(const std::exception& e) + // { + // std::cerr << e.what() << '\n'; + // exit(EXIT_FAILURE); + // } + // } + + // index_t num_kernel = 0; + // for(auto& op_ptr : op_ptrs) + // { + // for(std::size_t split_k_id = 0; split_k_id < split_k_list.size(); split_k_id++) + // { + // auto argument_ptr = op_ptr->MakeArgumentPointer( + // static_cast(in_device_buf.GetDeviceBuffer()), + // static_cast(wei_device_buf.GetDeviceBuffer()), + // static_cast(out_device_buf.GetDeviceBuffer()), + // input_lengths, + // input_strides, + // filter_lengths, + // weights_strides, + // output_lengths, + // output_strides, + // conv_filter_strides, + // conv_filter_dilations, + // input_left_pads, + // input_right_pads, + // in_element_op, + // wei_element_op, + // out_element_op, + // split_k_list[split_k_id]); + + // auto split_k_value = split_k_list[split_k_id]; + // auto split_k_param_str = std::to_string(split_k_value); + // auto* split_k_arg = + // dynamic_cast(argument_ptr.get()); + // if(split_k_arg && split_k_value < 0) + // { + // split_k_value = split_k_arg->k_batch_; + // split_k_param_str = std::to_string(split_k_value) + " (best occupancy)"; + // } + + // const std::size_t workspace_sz = op_ptr->GetWorkSpaceSize(argument_ptr.get()); + // DeviceMem workspace_dev(workspace_sz); + // op_ptr->SetWorkSpacePointer(argument_ptr.get(), workspace_dev.GetDeviceBuffer()); + + // if(op_ptr->IsSupportedArgument(argument_ptr.get())) + // { + // num_kernel++; + // if((instance_index != -1) && (instance_index + 1 != num_kernel)) + // { + // // skip test if instance_index is specified + // continue; + // } + + // 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}); + + // std::size_t flop = conv_param.GetFlops(); + // std::size_t num_btype = conv_param.GetByte(); + + // float tflops = static_cast(flop) / 1.E9 / avg_time; + // float gb_per_sec = num_btype / 1.E6 / avg_time; + + // std::cout << "Perf: " << std::setw(10) << avg_time << " ms, " << tflops + // << " TFlops, " << gb_per_sec << " GB/s, " << op_name << ", SplitK " + // << split_k_param_str << std::endl; + + // if(tflops > best_tflops) + // { + // best_op_name = op_name; + // best_tflops = tflops; + // best_avg_time = avg_time; + // best_gb_per_sec = gb_per_sec; + // best_split_k = split_k_param_str; + // } + + // if(do_verification) + // { + // wei_device_buf.FromDevice(weight_device_result.mData.data()); + + // using ComputeType = + // std::conditional_t; + // using AccDataType = + // std::conditional_t, int32_t, float>; + // const index_t num_accums = output.GetElementSize() / conv_param.K_; + // const index_t num_accums_split_k = split_k_value; + // // Calculate thresholds + // auto rtol = + // ck::utils::get_relative_threshold( + // num_accums / num_accums_split_k); + // auto atol = + // ck::utils::get_absolute_threshold( + // max_accumulated_value / num_accums_split_k, + // num_accums / num_accums_split_k); + // // Calculate error due to split_k accumulation + // auto rtol_split_k = + // ck::utils::get_relative_threshold( + // num_accums_split_k); + // auto atol_split_k = + // ck::utils::get_absolute_threshold( + // max_accumulated_value, num_accums_split_k); + // // Use higher threshold + // rtol = std::max(rtol, rtol_split_k); + // atol = std::max(atol, atol_split_k); + // // Use default atol for splitK == 1 + // bool pass = ck::utils::check_err(weight_device_result, + // weight_host_result, + // "Error: Incorrect results!", + // rtol, + // atol); + // std::cout << "Relative error threshold: " << rtol + // << " Absolute error threshold: " << atol << std::endl; + + // if(!pass) + // { + // std::cout << "Fail info: " << op_ptr->GetTypeString() << std::endl; + // } + + // all_pass &= pass; + + // if(do_log) + // { + // LogRangeAsType(std::cout << "output : ", output.mData, ",") + // << std::endl; + // LogRangeAsType( + // std::cout << "weight (device): ", weight_device_result.mData, ",") + // << std::endl; + // LogRangeAsType( + // std::cout << "weight (host): ", weight_host_result.mData, ",") + // << std::endl; + // LogRangeAsType(std::cout << "input: ", input.mData, ",") + // << std::endl; + // } + // } + // } + // else + // { + // std::cout << op_ptr->GetTypeString() << " does not support this problem" + // << std::endl; + // } + // } + // } + + // std::cout << "Best configuration parameters:" << "\nname: " << best_op_name + // << "\navg_time: " << best_avg_time << "\ntflops: " << best_tflops + // << "\nGB/s: " << best_gb_per_sec << ", SplitK " << best_split_k << std::endl; + + // if(instance_index != -1) + // { + // std::cout << "grouped_conv_bwd_weight_instance (" << instance_index << "/" << num_kernel + // << "): Passed" << std::endl; + // } + // return all_pass; +} + +} // namespace profiler +} // namespace ck_tile diff --git a/profiler/ck_tile/include/tile_profiler_operation_registry.hpp b/profiler/ck_tile/include/tile_profiler_operation_registry.hpp new file mode 100644 index 0000000000..ca2fd89f47 --- /dev/null +++ b/profiler/ck_tile/include/tile_profiler_operation_registry.hpp @@ -0,0 +1,79 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#include +#include +#include +#include +#include +#include +#include + +class ProfilerOperationRegistry final +{ + ProfilerOperationRegistry() = default; + ~ProfilerOperationRegistry() = default; + + public: + using Operation = std::function; + + private: + struct Entry final + { + explicit Entry(std::string_view description, Operation operation) noexcept + : description_(description), operation_(std::move(operation)) + { + } + + std::string_view description_; + Operation operation_; + }; + + std::map entries_; + + friend std::ostream& operator<<(std::ostream& stream, const ProfilerOperationRegistry& registry) + { + stream << "{\n"; + for(auto& [name, entry] : registry.entries_) + { + stream << "\t" << name << ": " << entry.description_ << "\n"; + } + stream << "}"; + + return stream; + } + + public: + static ProfilerOperationRegistry& GetInstance() + { + static ProfilerOperationRegistry registry; + return registry; + } + + std::optional Get(std::string_view name) const + { + const auto found = entries_.find(name); + if(found == end(entries_)) + { + return std::nullopt; + } + + return (found->second).operation_; + } + + bool Add(std::string_view name, std::string_view description, Operation operation) + { + return entries_ + .emplace(std::piecewise_construct, + std::forward_as_tuple(name), + std::forward_as_tuple(description, std::move(operation))) + .second; + } +}; + +#define PP_CONCAT(x, y) PP_CONCAT_IMPL(x, y) +#define PP_CONCAT_IMPL(x, y) x##y + +#define REGISTER_PROFILER_OPERATION(name, description, operation) \ + static const bool PP_CONCAT(operation_registration_result_, __COUNTER__) = \ + ::ProfilerOperationRegistry::GetInstance().Add(name, description, operation) diff --git a/profiler/ck_tile/src/CMakeLists.txt b/profiler/ck_tile/src/CMakeLists.txt new file mode 100644 index 0000000000..e88d58862d --- /dev/null +++ b/profiler/ck_tile/src/CMakeLists.txt @@ -0,0 +1,72 @@ +# ckTileProfiler +set(CK_PROFILER_OP_FILTER "" CACHE STRING "Filter for the operators to be profiled. Default is to include all") +set(CK_PROFILER_INSTANCE_FILTER "" CACHE STRING "Filter for the kernels instances to be profiled. Default is to be the same as the operator filter") +if (CK_PROFILER_OP_FILTER STREQUAL "") + set(CK_PROFILER_OP_FILTER ".+") +endif() +if (CK_PROFILER_INSTANCE_FILTER STREQUAL "") + set(CK_PROFILER_INSTANCE_FILTER ${CK_PROFILER_OP_FILTER}) +endif() +message(STATUS "CK_PROFILER_OP_FILTER: ${CK_PROFILER_OP_FILTER}") +message(STATUS "CK_PROFILER_INSTANCE_FILTER: ${CK_PROFILER_INSTANCE_FILTER}") + +if(SUPPORTED_GPU_TARGETS MATCHES "gfx9" OR SUPPORTED_GPU_TARGETS MATCHES "gfx1[12]") + list(APPEND PROFILER_OPS tile_profile_grouped_conv_bwd_weight.cpp) +endif() + +if(DL_KERNELS) + list(APPEND PROFILER_OPS tile_profile_grouped_conv_bwd_weight.cpp) +endif() + +set(PROFILER_SOURCES tile_profiler.cpp) +foreach(SOURCE ${PROFILER_OPS}) + string(REGEX REPLACE "tile_profile_(.+)\.cpp" "\\1" OP_NAME ${SOURCE}) + if (OP_NAME STREQUAL "") + message(FATAL_ERROR "Unexpected source file name: ${SOURCE}") + endif() + if("${OP_NAME}" MATCHES "${CK_PROFILER_OP_FILTER}") + list(APPEND PROFILER_SOURCES ${SOURCE}) + endif() +endforeach() +message(VERBOSE "ckTileProfiler sources: ${PROFILER_SOURCES}") + +set(PROFILER_EXECUTABLE ckTileProfiler) + +add_executable(${PROFILER_EXECUTABLE} ${PROFILER_SOURCES}) +#target_include_directories(${PROFILER_EXECUTABLE} PRIVATE ${CMAKE_PROJECT_DIR}/include) +target_compile_options(${PROFILER_EXECUTABLE} PRIVATE -Wno-global-constructors) +# flags to compress the library +if(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER 600241132) + message(DEBUG "Adding --offload-compress flag for ${PROFILER_EXECUTABLE}") + target_compile_options(${PROFILER_EXECUTABLE} PRIVATE --offload-compress) +endif() + +set(DEVICE_INSTANCES "") + +if(SUPPORTED_GPU_TARGETS MATCHES "gfx9" OR SUPPORTED_GPU_TARGETS MATCHES "gfx1[12]") + list(APPEND DEVICE_INSTANCES device_grouped_convnd_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv1d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv2d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv3d_bwd_weight_instance) +endif() + +if(DL_KERNELS) + list(APPEND DEVICE_INSTANCES device_grouped_conv1d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv2d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv3d_bwd_weight_instance) +endif() + +set(PROFILER_LIBS utility getopt::getopt) +foreach(LIB ${DEVICE_INSTANCES}) + string(REGEX REPLACE "device_(.+)_instance" "\\1" INSTANCE_NAME ${LIB}) + if (INSTANCE_NAME STREQUAL "") + message(FATAL_ERROR "Unexpected kernel instance name: ${LIB}") + endif() + if("${INSTANCE_NAME}" MATCHES "${CK_PROFILER_INSTANCE_FILTER}") + list(APPEND PROFILER_LIBS ${LIB}) + endif() +endforeach() +message(VERBOSE "ckTileProfiler libs: ${PROFILER_LIBS}") +target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE ${PROFILER_LIBS}) + +rocm_install(TARGETS ${PROFILER_EXECUTABLE} COMPONENT profiler) diff --git a/profiler/ck_tile/src/tile_profile_grouped_conv_bwd_weight.cpp b/profiler/ck_tile/src/tile_profile_grouped_conv_bwd_weight.cpp new file mode 100644 index 0000000000..39e5e42483 --- /dev/null +++ b/profiler/ck_tile/src/tile_profile_grouped_conv_bwd_weight.cpp @@ -0,0 +1,373 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#include +#include +#include +#include + +#include "tile_profile_grouped_conv_bwd_weight_impl.hpp" +#include "conv_parameters.hpp" +#include "tile_profiler_operation_registry.hpp" + +// Old CK library dependencies +#include "ck/library/tensor_operation_instance/device_operation_instance_factory.hpp" + +// CK Tile library dependnecies +#include "ck_tile/core/numeric/integral_constant.hpp" + +namespace { + +enum struct ConvLayout +{ + GNCHW_GKCYX_GNKHW, // 0 + GNHWC_GKYXC_GNHWK, // 1 + NHWGC_GKYXC_NHWGK, // 2 + NGCHW_GKYXC_NGKHW, // 3 + NGCHW_GKCYX_NGKHW, // 4 +}; + +enum struct ConvDataType +{ + F32_F32_F32, // 0 + F16_F16_F16, // 1 + BF16_F32_BF16, // 2 + F16_F16_F16_BF8_F8, // 3 + I8_I8_I8, // 4 + BF16_BF16_BF16, // 5 + F32_F32_F32_TF32, // 6 +}; + +#define OP_NAME "grouped_conv_bwd_weight" +#define OP_DESC "Grouped Convolution Backward Weight" + +static void print_helper_msg() +{ + std::string conv_param_parser_helper_msg; + + conv_param_parser_helper_msg += "Following arguments (depending on number of spatial dims):\n" + " Number of spatial dimensions (1=Conv1d, 2=Conv2d, 3=Conv3d)\n" + " G, N, K, C, \n" + " , (ie Y, X for 2D)\n" + " , (ie Hi, Wi for 2D)\n" + " , (ie Sy, Sx for 2D)\n" + " , (ie Dy, Dx for 2D)\n" + " , (ie LeftPy, LeftPx for 2D)\n" + " , (ie RightPy, RightPx for 2D)\n"; + + std::cout << "arg1: tensor operation (" OP_NAME ": " OP_DESC ")\n" + << "arg2: data type (0: Input fp32, Weight fp32, Output fp32\n" + << " 1: Input fp16, Weight fp16, Output fp16\n" + << " 2: Input bf16, Weight fp32, Output bf16\n" + << " 3: Input fp16, Weight fp16, Output fp16, Gemm bf8@fp8\n" + << " 4: Input int8, Weight int8, Output int8\n" + << " 5: Input bf16, Weight bf16, Output bf16\n" + << " 6: Input fp32, Weight fp32, Output fp32, Compute tf32)\n" + << "arg3: tensor layout (0: Input[G, N, C, Hi, Wi], Weight[G, K, C, Y, X], Output[G, " + "N, K, Ho, Wo]\n" + << " 1: Input[G, N, Hi, Wi, C], Weight[G, K, Y, X, C], Output[G, " + "N, Ho, Wo, K]\n" + << " 2: Input[N, Hi, Wi, G, C], Weight[G, K, Y, X, C], Output[N, " + "Ho, Wo, G, K]\n" + << " 3: Input[N, G, C, Hi, Wi], Weight[G, K, Y, X, C], Output[N, " + "G, K, Ho, Wo]\n" + << " 4: Input[N, G, C, Hi, Wi], Weight[G, K, C, Y, X], Output[N, " + "G, K, Ho, Wo]\n" + << "arg4: verification (0: no, 1: yes)\n" + << "arg5: initialization (0: no init, 1: integer value, 2: decimal value)\n" + << "arg6: print tensor value (0: no; 1: yes)\n" + << "arg7: time kernel (0: no, 1: yes)\n" + << conv_param_parser_helper_msg + << " SplitK (-1 for internally computed split-K value, positive value to set k " + "batches explicitly, or 'all' to test all internal split-K values)\n" + << std::endl; +} + +} // namespace + +int tile_profile_grouped_conv_bwd_weight(int argc, char* argv[]) +{ + // 8 for control, 1 for num_dim_spatial + if(argc < 9) + { + print_helper_msg(); + return 1; + } + + const auto data_type = static_cast(std::stoi(argv[2])); + const auto layout = static_cast(std::stoi(argv[3])); + const bool do_verification = std::stoi(argv[4]); + const int init_method = std::stoi(argv[5]); + const bool do_log = std::stoi(argv[6]); + const bool time_kernel = std::stoi(argv[7]); + const int num_dim_spatial = std::stoi(argv[8]); + + // 8 for control, 1 for num_dim_spatial, 4 for G/N/K/C, and 6 * num_dim_spatial, 1 for split-K + if(argc != 8 + 1 + 4 + 6 * num_dim_spatial + 1) + { + print_helper_msg(); + return 1; + } + + const auto params = ck_tile::utils::conv::parse_conv_param(num_dim_spatial, 9, argv); + + const auto& split_k = std::string(argv[8 + 1 + 4 + 6 * num_dim_spatial]); + + using F32 = float; + using F16 = ck::half_t; + using BF16 = ck::bhalf_t; + using F8 = ck::f8_t; + using BF8 = ck::bf8_t; +#if defined(__gfx942__) + using TF32 = ck::tf32_t; +#endif + + using namespace ck::tensor_layout::convolution; + + constexpr auto I1 = ck_tile::number<1>{}; + constexpr auto I2 = ck_tile::number<2>{}; + constexpr auto I3 = ck_tile::number<3>{}; + + auto profile = [&](auto num_dim_spatial_tmp, + auto in_layout, + auto wei_layout, + auto out_layout, + auto in_type, + auto wei_type, + auto out_type, + auto compute_type_a, + auto compute_type_b) { + constexpr ck::index_t NDimSpatial = num_dim_spatial_tmp.value; + + using InLayout = decltype(in_layout); + using WeiLayout = decltype(wei_layout); + using OutLayout = decltype(out_layout); + + using InDataType = decltype(in_type); + using WeiDataType = decltype(wei_type); + using OutDataType = decltype(out_type); + + using ComputeTypeA = decltype(compute_type_a); + using ComputeTypeB = decltype(compute_type_b); + + bool pass = ck_tile::profiler::profile_grouped_conv_bwd_weight_impl( + do_verification, init_method, do_log, time_kernel, params, split_k); + + return pass ? 0 : 1; + }; + + if(num_dim_spatial == 1 && layout == ConvLayout::GNHWC_GKYXC_GNHWK) + { + if(data_type == ConvDataType::F32_F32_F32) + { + return profile(I1, GNWC{}, GKXC{}, GNWK{}, F32{}, F32{}, F32{}, F32{}, F32{}); + } + if(data_type == ConvDataType::F16_F16_F16) + { + return profile(I1, GNWC{}, GKXC{}, GNWK{}, F16{}, F16{}, F16{}, F16{}, F16{}); + } + if(data_type == ConvDataType::BF16_F32_BF16) + { + // fp32 atomic add is used for weight tensor in bf16 kernel + return profile(I1, GNWC{}, GKXC{}, GNWK{}, BF16{}, F32{}, BF16{}, BF16{}, BF16{}); + } + else if(data_type == ConvDataType::F32_F32_F32_TF32) + { +#if defined(__gfx942__) + return profile(I1, GNWC{}, GKXC{}, GNWK{}, F32{}, F32{}, F32{}, TF32{}, TF32{}); +#endif + } + } + if(num_dim_spatial == 2 && layout == ConvLayout::GNHWC_GKYXC_GNHWK) + { + if(data_type == ConvDataType::F32_F32_F32) + { + return profile(I2, GNHWC{}, GKYXC{}, GNHWK{}, F32{}, F32{}, F32{}, F32{}, F32{}); + } + if(data_type == ConvDataType::F16_F16_F16) + { + return profile(I2, GNHWC{}, GKYXC{}, GNHWK{}, F16{}, F16{}, F16{}, F16{}, F16{}); + } + if(data_type == ConvDataType::BF16_F32_BF16) + { + // fp32 atomic add is used for weight tensor in bf16 kernel + return profile(I2, GNHWC{}, GKYXC{}, GNHWK{}, BF16{}, F32{}, BF16{}, BF16{}, BF16{}); + } + else if(data_type == ConvDataType::F32_F32_F32_TF32) + { +#if defined(__gfx942__) + return profile(I2, GNHWC{}, GKYXC{}, GNHWK{}, F32{}, F32{}, F32{}, TF32{}, TF32{}); +#endif + } + } + if(num_dim_spatial == 2 && layout == ConvLayout::NHWGC_GKYXC_NHWGK) + { + if(data_type == ConvDataType::F32_F32_F32) + { + return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, F32{}, F32{}, F32{}, F32{}, F32{}); + } + if(data_type == ConvDataType::F16_F16_F16) + { + return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, F16{}, F16{}, F16{}, F16{}, F16{}); + } + if(data_type == ConvDataType::BF16_F32_BF16) + { + // fp32 atomic add is used for weight tensor in bf16 kernel + return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, BF16{}, F32{}, BF16{}, BF16{}, BF16{}); + } + if(data_type == ConvDataType::BF16_BF16_BF16) + { + return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{}); + } + else if(data_type == ConvDataType::F32_F32_F32_TF32) + { +#if defined(__gfx942__) + return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, F32{}, F32{}, F32{}, TF32{}, TF32{}); +#endif + } + } + else if(num_dim_spatial == 2 && layout == ConvLayout::NGCHW_GKYXC_NGKHW) + { + if(data_type == ConvDataType::F16_F16_F16) + { + return profile(I2, NGCHW{}, GKYXC{}, NGKHW{}, F16{}, F16{}, F16{}, F16{}, F16{}); + } + if(data_type == ConvDataType::BF16_BF16_BF16) + { + // fp32 atomic add is used for weight tensor in bf16 kernel + return profile(I2, NGCHW{}, GKYXC{}, NGKHW{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{}); + } + } + else if(num_dim_spatial == 2 && layout == ConvLayout::NGCHW_GKCYX_NGKHW) + { + if(data_type == ConvDataType::F16_F16_F16) + { + return profile(I2, NGCHW{}, GKCYX{}, NGKHW{}, F16{}, F16{}, F16{}, F16{}, F16{}); + } + if(data_type == ConvDataType::BF16_BF16_BF16) + { + return profile(I2, NGCHW{}, GKCYX{}, NGKHW{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{}); + } + if(data_type == ConvDataType::F32_F32_F32) + { + return profile(I2, NGCHW{}, GKCYX{}, NGKHW{}, F32{}, F32{}, F32{}, F32{}, F32{}); + } + else if(data_type == ConvDataType::F32_F32_F32_TF32) + { +#if defined(__gfx942__) + return profile(I2, NGCHW{}, GKCYX{}, NGKHW{}, F32{}, F32{}, F32{}, TF32{}, TF32{}); +#endif + } + } + if(num_dim_spatial == 3 && layout == ConvLayout::GNHWC_GKYXC_GNHWK) + { + if(data_type == ConvDataType::F32_F32_F32) + { + return profile(I3, GNDHWC{}, GKZYXC{}, GNDHWK{}, F32{}, F32{}, F32{}, F32{}, F32{}); + } + if(data_type == ConvDataType::F16_F16_F16) + { + return profile(I3, GNDHWC{}, GKZYXC{}, GNDHWK{}, F16{}, F16{}, F16{}, F16{}, F16{}); + } + if(data_type == ConvDataType::BF16_F32_BF16) + { + // fp32 atomic add is used for weight tensor in bf16 kernel + return profile(I3, GNDHWC{}, GKZYXC{}, GNDHWK{}, BF16{}, F32{}, BF16{}, BF16{}, BF16{}); + } + else if(data_type == ConvDataType::I8_I8_I8) + { + return profile( + I3, GNDHWC{}, GKZYXC{}, GNDHWK{}, int8_t{}, int8_t{}, int8_t{}, int8_t{}, int8_t{}); + } + else if(data_type == ConvDataType::F32_F32_F32_TF32) + { +#if defined(__gfx942__) + return profile(I3, GNDHWC{}, GKZYXC{}, GNDHWK{}, F32{}, F32{}, F32{}, TF32{}, TF32{}); +#endif + } + } + if(num_dim_spatial == 3 && layout == ConvLayout::NHWGC_GKYXC_NHWGK) + { + if(data_type == ConvDataType::F32_F32_F32) + { + return profile(I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, F32{}, F32{}, F32{}, F32{}, F32{}); + } + if(data_type == ConvDataType::F16_F16_F16) + { + return profile(I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, F16{}, F16{}, F16{}, F16{}, F16{}); + } + if(data_type == ConvDataType::BF16_F32_BF16) + { + // fp32 atomic add is used for weight tensor in bf16 kernel + return profile(I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, BF16{}, F32{}, BF16{}, BF16{}, BF16{}); + } + if(data_type == ConvDataType::BF16_BF16_BF16) + { + return profile( + I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{}); + } + if(data_type == ConvDataType::F16_F16_F16_BF8_F8) + { + return profile(I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, F16{}, F16{}, F16{}, BF8{}, F8{}); + } + else if(data_type == ConvDataType::I8_I8_I8) + { + return profile( + I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, int8_t{}, int8_t{}, int8_t{}, int8_t{}, int8_t{}); + } + else if(data_type == ConvDataType::F32_F32_F32_TF32) + { +#if defined(__gfx942__) + return profile(I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, F32{}, F32{}, F32{}, TF32{}, TF32{}); +#endif + } + } + else if(num_dim_spatial == 3 && layout == ConvLayout::NGCHW_GKYXC_NGKHW) + { + if(data_type == ConvDataType::F16_F16_F16) + { + return profile(I3, NGCDHW{}, GKZYXC{}, NGKDHW{}, F16{}, F16{}, F16{}, F16{}, F16{}); + } + if(data_type == ConvDataType::BF16_BF16_BF16) + { + return profile( + I3, NGCDHW{}, GKZYXC{}, NGKDHW{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{}); + } + } + else if(num_dim_spatial == 3 && layout == ConvLayout::NGCHW_GKCYX_NGKHW) + { + if(data_type == ConvDataType::F16_F16_F16) + { + return profile(I3, NGCDHW{}, GKCZYX{}, NGKDHW{}, F16{}, F16{}, F16{}, F16{}, F16{}); + } + if(data_type == ConvDataType::BF16_BF16_BF16) + { + return profile( + I3, NGCDHW{}, GKCZYX{}, NGKDHW{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{}); + } + if(data_type == ConvDataType::F32_F32_F32) + { + return profile(I3, NGCDHW{}, GKCZYX{}, NGKDHW{}, F32{}, F32{}, F32{}, F32{}, F32{}); + } + else if(data_type == ConvDataType::F32_F32_F32_TF32) + { +#if defined(__gfx942__) + return profile(I3, NGCDHW{}, GKCZYX{}, NGKDHW{}, F32{}, F32{}, F32{}, TF32{}, TF32{}); +#endif + } + } + + std::cout << "this data_type & layout is not implemented" << std::endl; + + return 1; +} + +REGISTER_PROFILER_OPERATION(OP_NAME, OP_DESC, tile_profile_grouped_conv_bwd_weight); diff --git a/profiler/ck_tile/src/tile_profiler.cpp b/profiler/ck_tile/src/tile_profiler.cpp new file mode 100644 index 0000000000..deaee5709e --- /dev/null +++ b/profiler/ck_tile/src/tile_profiler.cpp @@ -0,0 +1,30 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. + +#include +#include + +#include "tile_profiler_operation_registry.hpp" + +static void print_helper_message() +{ + std::cout << "arg1: tensor operation " << ProfilerOperationRegistry::GetInstance() << std::endl; +} + +int main(int argc, char* argv[]) +{ + if(argc == 1) + { + print_helper_message(); + } + else if(const auto operation = ProfilerOperationRegistry::GetInstance().Get(argv[1]); + operation.has_value()) + { + return (*operation)(argc, argv); + } + else + { + std::cerr << "cannot find operation: " << argv[1] << std::endl; + return EXIT_FAILURE; + } +}