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

Skeleton for the ckTileProfiler.

Browse Source
This commit is contained in:
Ville Pietilä
2025-10-13 11:40:31 +00:00
parent 94569f3991
commit d62f34348a
8 changed files with 1213 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
profiler/CMakeLists.txt
View File

@@ -3,3 +3,4 @@ include_directories(BEFORE
)
add_subdirectory(src)
add_subdirectory(ck_tile)

5
profiler/ck_tile/CMakeLists.txt Normal file
View File

@@ -0,0 +1,5 @@
include_directories(BEFORE
${CMAKE_CURRENT_LIST_DIR}/include
)
add_subdirectory(src)

304
profiler/ck_tile/include/conv_parameters.hpp Normal file
View File

@@ -0,0 +1,304 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include <cstdlib>
#include <numeric>
#include <iterator>
#include <vector>
#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<ck::index_t>& filters_len,
const std::vector<ck::index_t>& input_len,
const std::vector<ck::index_t>& strides,
const std::vector<ck::index_t>& dilations,
const std::vector<ck::index_t>& left_pads,
const std::vector<ck::index_t>& 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<ck::long_index_t>& filters_len,
const std::vector<ck::long_index_t>& input_len,
const std::vector<ck::long_index_t>& strides,
const std::vector<ck::long_index_t>& dilations,
const std::vector<ck::long_index_t>& left_pads,
const std::vector<ck::long_index_t>& 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<ck::long_index_t> filter_spatial_lengths_;
std::vector<ck::long_index_t> input_spatial_lengths_;
std::vector<ck::long_index_t> output_spatial_lengths_;
std::vector<ck::long_index_t> conv_filter_strides_;
std::vector<ck::long_index_t> conv_filter_dilations_;
std::vector<ck::long_index_t> input_left_pads_;
std::vector<ck::long_index_t> input_right_pads_;
std::vector<ck::long_index_t> GetOutputSpatialLengths() const;
std::size_t GetFlops() const;
template <typename InDataType>
std::size_t GetInputByte() const
{
// sizeof(InDataType) * (G * N * C * <input spatial lengths product>) +
return sizeof(InDataType) *
(G_ * N_ * C_ *
ck::accumulate_n<std::size_t>(
std::begin(input_spatial_lengths_), num_dim_spatial_, 1, std::multiplies<>()));
}
template <typename WeiDataType>
std::size_t GetWeightByte() const
{
// sizeof(WeiDataType) * (G * K * C * <filter spatial lengths product>) +
return sizeof(WeiDataType) *
(G_ * K_ * C_ *
ck::accumulate_n<std::size_t>(
std::begin(filter_spatial_lengths_), num_dim_spatial_, 1, std::multiplies<>()));
}
template <typename OutDataType>
std::size_t GetOutputByte() const
{
// sizeof(OutDataType) * (G * N * K * <output spatial lengths product>);
return sizeof(OutDataType) * (G_ * N_ * K_ *
std::accumulate(std::begin(output_spatial_lengths_),
std::end(output_spatial_lengths_),
static_cast<std::size_t>(1),
std::multiplies<std::size_t>()));
}
template <typename InDataType, typename WeiDataType, typename OutDataType>
std::size_t GetByte() const
{
return GetInputByte<InDataType>() + GetWeightByte<WeiDataType>() +
GetOutputByte<OutDataType>();
}
};
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<ck::index_t>& filters_len,
const std::vector<ck::index_t>& input_len,
const std::vector<ck::index_t>& strides,
const std::vector<ck::index_t>& dilations,
const std::vector<ck::index_t>& left_pads,
const std::vector<ck::index_t>& right_pads)
: num_dim_spatial_(static_cast<ck::long_index_t>(n_dim)),
G_(static_cast<ck::long_index_t>(group_count)),
N_(static_cast<ck::long_index_t>(n_batch)),
K_(static_cast<ck::long_index_t>(n_out_channels)),
C_(static_cast<ck::long_index_t>(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<ck::index_t>(filter_spatial_lengths_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(input_spatial_lengths_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(conv_filter_strides_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(conv_filter_dilations_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(input_left_pads_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(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<ck::long_index_t>(filters_len[i]);
input_spatial_lengths_[i] = static_cast<ck::long_index_t>(input_len[i]);
conv_filter_strides_[i] = static_cast<ck::long_index_t>(strides[i]);
conv_filter_dilations_[i] = static_cast<ck::long_index_t>(dilations[i]);
input_left_pads_[i] = static_cast<ck::long_index_t>(left_pads[i]);
input_right_pads_[i] = static_cast<ck::long_index_t>(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<ck::long_index_t>& filters_len,
const std::vector<ck::long_index_t>& input_len,
const std::vector<ck::long_index_t>& strides,
const std::vector<ck::long_index_t>& dilations,
const std::vector<ck::long_index_t>& left_pads,
const std::vector<ck::long_index_t>& 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<ck::index_t>(filter_spatial_lengths_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(input_spatial_lengths_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(conv_filter_strides_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(conv_filter_dilations_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(input_left_pads_.size()) != num_dim_spatial_ ||
static_cast<ck::index_t>(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<ck::long_index_t> ConvParam::GetOutputSpatialLengths() const
{
return output_spatial_lengths_;
}
std::size_t ConvParam::GetFlops() const
{
// 2 * G * N * K * C * <output spatial lengths product> * <filter spatial lengths product>
return static_cast<std::size_t>(2) * G_ * N_ * K_ * C_ *
ck::accumulate_n<std::size_t>(
std::begin(output_spatial_lengths_), num_dim_spatial_, 1, std::multiplies<>()) *
ck::accumulate_n<std::size_t>(
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<ck::long_index_t> filter_spatial_lengths(num_dim_spatial);
std::vector<ck::long_index_t> input_spatial_lengths(num_dim_spatial);
std::vector<ck::long_index_t> conv_filter_strides(num_dim_spatial);
std::vector<ck::long_index_t> conv_filter_dilations(num_dim_spatial);
std::vector<ck::long_index_t> input_left_pads(num_dim_spatial);
std::vector<ck::long_index_t> 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;
}

349
profiler/ck_tile/include/tile_profile_grouped_conv_bwd_weight_impl.hpp Normal file
View File

@@ -0,0 +1,349 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include <algorithm>
#include <iomanip>
#include <iostream>
#include <iterator>
#include <typeinfo>
#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 <ck::index_t NDimSpatial,
typename InLayout,
typename WeiLayout,
typename OutLayout,
typename InDataType,
typename WeiDataType,
typename OutDataType,
typename ComputeTypeA = InDataType,
typename ComputeTypeB = ComputeTypeA>
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<InLayout>(conv_param);
// const auto wei_g_k_c_xs_desc =
// ck::utils::conv::make_weight_host_tensor_descriptor_g_k_c_xs_packed<WeiLayout>(conv_param);
// const auto out_g_n_k_wos_desc =
// ck::utils::conv::make_output_host_tensor_descriptor_g_n_k_wos_packed<OutLayout>(conv_param);
// Tensor<InDataType> input(in_g_n_c_wis_desc);
// Tensor<WeiDataType> weight_host_result(wei_g_k_c_xs_desc);
// Tensor<WeiDataType> weight_device_result(wei_g_k_c_xs_desc);
// Tensor<OutDataType> 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<InDataType>{-5, 5});
// output.GenerateTensorValue(GeneratorTensor_2<OutDataType>{-5, 5});
// break;
// default:
// input.GenerateTensorValue(GeneratorTensor_3<InDataType>{0.0, 1.0});
// output.GenerateTensorValue(GeneratorTensor_3<OutDataType>{-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<NDimSpatial,
// InDataType,
// WeiDataType,
// OutDataType,
// InElementOp,
// WeiElementOp,
// OutElementOp>{};
// 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<NDimSpatial,
// InLayout,
// WeiLayout,
// OutLayout,
// InDataType,
// WeiDataType,
// OutDataType,
// InElementOp,
// WeiElementOp,
// OutElementOp,
// ComputeTypeA,
// ComputeTypeB>;
// // 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<ck::index_t, NDimSpatial + 3> input_lengths{};
// std::array<ck::index_t, NDimSpatial + 3> filter_lengths{};
// std::array<ck::index_t, NDimSpatial + 3> output_lengths{};
// std::array<ck::index_t, NDimSpatial + 3> input_strides{};
// std::array<ck::index_t, NDimSpatial + 3> weights_strides{};
// std::array<ck::index_t, NDimSpatial + 3> output_strides{};
// std::array<ck::index_t, NDimSpatial> conv_filter_strides{};
// std::array<ck::index_t, NDimSpatial> conv_filter_dilations{};
// std::array<ck::index_t, NDimSpatial> input_left_pads{};
// std::array<ck::index_t, NDimSpatial> 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<ck::index_t> 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<InDataType*>(in_device_buf.GetDeviceBuffer()),
// static_cast<WeiDataType*>(wei_device_buf.GetDeviceBuffer()),
// static_cast<OutDataType*>(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<ck::tensor_operation::device::ArgumentSplitK*>(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<InDataType, WeiDataType, OutDataType>();
// float tflops = static_cast<float>(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<sizeof(ComputeTypeA) < sizeof(ComputeTypeB),
// ComputeTypeA,
// ComputeTypeB>;
// using AccDataType =
// std::conditional_t<std::is_same_v<ComputeType, int8_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<ComputeType, WeiDataType, AccDataType>(
// num_accums / num_accums_split_k);
// auto atol =
// ck::utils::get_absolute_threshold<ComputeType, WeiDataType, AccDataType>(
// 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<WeiDataType, WeiDataType, WeiDataType>(
// num_accums_split_k);
// auto atol_split_k =
// ck::utils::get_absolute_threshold<WeiDataType, WeiDataType, WeiDataType>(
// 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<float>(std::cout << "output : ", output.mData, ",")
// << std::endl;
// LogRangeAsType<float>(
// std::cout << "weight (device): ", weight_device_result.mData, ",")
// << std::endl;
// LogRangeAsType<float>(
// std::cout << "weight (host): ", weight_host_result.mData, ",")
// << std::endl;
// LogRangeAsType<float>(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

79
profiler/ck_tile/include/tile_profiler_operation_registry.hpp Normal file
View File

@@ -0,0 +1,79 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <functional>
#include <iostream>
#include <iterator>
#include <map>
#include <optional>
#include <string_view>
#include <utility>
class ProfilerOperationRegistry final
{
ProfilerOperationRegistry() = default;
~ProfilerOperationRegistry() = default;
public:
using Operation = std::function<int(int, char*[])>;
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<std::string_view, Entry> 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<Operation> 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)

72
profiler/ck_tile/src/CMakeLists.txt Normal file
View File

@@ -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)

373
profiler/ck_tile/src/tile_profile_grouped_conv_bwd_weight.cpp Normal file
View File

@@ -0,0 +1,373 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <initializer_list>
#include <iostream>
#include <numeric>
#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"
" <filter spatial dimensions>, (ie Y, X for 2D)\n"
" <input image spatial dimensions>, (ie Hi, Wi for 2D)\n"
" <strides>, (ie Sy, Sx for 2D)\n"
" <dilations>, (ie Dy, Dx for 2D)\n"
" <left padding>, (ie LeftPy, LeftPx for 2D)\n"
" <right padding>, (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<ConvDataType>(std::stoi(argv[2]));
const auto layout = static_cast<ConvLayout>(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<NDimSpatial,
InLayout,
WeiLayout,
OutLayout,
InDataType,
WeiDataType,
OutDataType,
ComputeTypeA,
ComputeTypeB>(
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);

30
profiler/ck_tile/src/tile_profiler.cpp Normal file
View File

@@ -0,0 +1,30 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iostream>
#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;
}
}