mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-26 08:00:13 +00:00
Convolution FWD profiler refactor. (#183)
* Convolution ND * Code unification across dimensions for generating tensor descriptors. * Example * Instances * Move convnd f32 instance file to comply with repo structure. * Conv 1D tensor layouts. * Formatting and use ReferenceConv * Reference ConvFwd supporting 1D and 2D convolution. * Debug printing TensorLayout name. * Conv fwd 1D instance f32 * Refactor conv ND example. Needed to support various conv dimensio. Needed to support various conv dimensions * Rename conv nd example director to prevent conflicts. * Refactor some common utility to single file. Plus some tests. * Refactor GetHostTensorDescriptor + UT. * Add 1D test case. * Test reference convolution 1d/2d * Remove some leftovers. * Fix convolution example error for 1D * Refactor test check errors utility function. * Test Conv2D Fwd XDL * More UT for 1D case. * Parameterize input & weight initializers. * Rename example to prevent conflicts. * Split convnd instance into separate files for 1d/2d * Address review comments. * Fix data type for flops/gbytes calculations. * Assign example number 11. * 3D cases for convolution utility functions. * 3D reference convolution. * Add support for 3D convolution. * Check for inputs bigger than 2GB. * Formatting * Support for bf16/f16/f32/i8 - conv instances + UT. * Use check_err from test_util.hpp. * Split convnd test into separate files for each dim. * Fix data generation and use proper instances. * Formatting * Skip tensor initialization if not necessary. * Fix CMakefiles. * Remove redundant conv2d_fwd test. * Lower problem size for conv3D UT. * 3D case for convnd example. * Remove leftovers after merge. * Add Conv Specialization string to GetTypeString * Skip instance causing numerical errors. * Small fixes. * Remove redundant includes. * Fix namespace name error. * Script for automatic testing and logging convolution fwd UTs * Comment out numactl cmd. * Refine weights initalization and relax rtol for fp16 * Move test_util.hpp to check_err.hpp * Refine weights initalization and relax rtol for fp16 * Refactor common part of test conv utils. * Move utility function to single common place. * Add additional common functions to utility. * Refactor convnd_fwd_xdl examples. * Remove redundant files. * Unify structure. * Add constructor to ConvParams. * And add input parameters validation. * Modify conv examples to use single utility file. * Remove check_error from host_tensor.hpp * Get rid of check_indices function. * Remove bf16_to_f32 function overload for scalars. * Fix namespace. * Add half_float::half for check_err. * Fix conv params size in UT. * Fix weights initialization for int8. * Fix weights initialization for int8. * Add type_convert when store output in ref conv 1D. * Get back old conv2d_fwd_xdl operation. * Silence conv debug print. * format * clean * clean * Fix merge. * Fix namespace for check_err * Formatting. * Fix merge artifacts. * Remove deleted header. * Fix some includes and use ck::utils::check_err. * Remove unused check_indices restored by previous merge. * Fix namespaces after merge. * Fix compilation error. * Small fixes. * Use common functions. * Fix filename * Fix namespaces. * Fix merge artifact - retrieve removed by accident fun. * Fix ConvForwardSpecialization. * Working example of OpInstanceRunEngine for conv2dfwd UT. * Adhere to coding style rules. * Formatting and adhere to coding style rules. * Fix merge artifacts. * Utility for collecting conv fwd instances. + Plus commmon part for parsing cmdline params. * Refactor FillUniform because of segfault for int8_t. * Naming convention. * Elegant version of device mem allocation. * Use OpInstanceRunEngine in conv fwd nd tests. * Multiple refinements. * conditional init * don't run reference op if not provided. * Use OpInstanceRunEngine for ckProfiler conv_fwd * Refactor common tensor fill function to separate file. * Clean up unused functions. * Support different init methods. * Create CMake target for conv_fwd_util. * Add header for profile_convnd_fwd.cpp * Fix CMakefiles to link with conv_fwd_util where needed. * Fix some clutter. Co-authored-by: Adam Osewski <aosewski@amd.com> Co-authored-by: Chao Liu <chao.liu2@amd.com>
This commit is contained in:
Adam Osewski
@@ -1,13 +1,10 @@
|
||||
#ifndef CONV_FWD_UTIL_HPP
|
||||
#define CONV_FWD_UTIL_HPP
|
||||
#pragma once
|
||||
|
||||
#include <algorithm>
|
||||
#include <cstdlib>
|
||||
#include <functional>
|
||||
#include <iterator>
|
||||
#include <numeric>
|
||||
#include <sstream>
|
||||
#include <random>
|
||||
#include <tuple>
|
||||
#include <type_traits>
|
||||
#include <vector>
|
||||
@@ -18,10 +15,50 @@
|
||||
#include "device_conv_fwd.hpp"
|
||||
#include "device_tensor.hpp"
|
||||
#include "element_wise_operation.hpp"
|
||||
#include "fill.hpp"
|
||||
#include "host_tensor.hpp"
|
||||
#include "op_instance_engine.hpp"
|
||||
#include "reference_conv_fwd.hpp"
|
||||
#include "tensor_layout.hpp"
|
||||
|
||||
namespace ck {
|
||||
namespace tensor_operation {
|
||||
namespace device {
|
||||
|
||||
using DeviceConvFwdNoOpPtr = DeviceConvFwdPtr<element_wise::PassThrough,
|
||||
element_wise::PassThrough,
|
||||
element_wise::PassThrough>;
|
||||
namespace device_conv1d_fwd_instance {
|
||||
|
||||
void add_device_conv1d_fwd_xdl_nwc_kxc_nwk_bf16_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv1d_fwd_xdl_nwc_kxc_nwk_f16_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv1d_fwd_xdl_nwc_kxc_nwk_f32_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv1d_fwd_xdl_nwc_kxc_nwk_int8_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
|
||||
} // namespace device_conv1d_fwd_instance
|
||||
namespace device_conv2d_fwd_instance {
|
||||
|
||||
void add_device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_bf16_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_f16_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk_f16_instances(
|
||||
std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_f32_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_int8_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
|
||||
} // namespace device_conv2d_fwd_instance
|
||||
namespace device_conv3d_fwd_instance {
|
||||
|
||||
void add_device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk_bf16_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk_f16_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk_f32_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
void add_device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk_int8_instances(std::vector<DeviceConvFwdNoOpPtr>&);
|
||||
|
||||
} // namespace device_conv3d_fwd_instance
|
||||
|
||||
} // namespace device
|
||||
} // namespace tensor_operation
|
||||
} // namespace ck
|
||||
|
||||
namespace ck {
|
||||
namespace utils {
|
||||
namespace conv {
|
||||
@@ -47,20 +84,7 @@ std::size_t get_flops(ck::index_t N,
|
||||
ck::index_t C,
|
||||
ck::index_t K,
|
||||
const std::vector<ck::index_t>& filter_spatial_lengths,
|
||||
const std::vector<ck::index_t>& output_spatial_lengths)
|
||||
{
|
||||
// 2 * N * K * <output spatial lengths product> * C * <filter spatial lengths product>
|
||||
return static_cast<std::size_t>(2) * 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>()) *
|
||||
C *
|
||||
std::accumulate(std::begin(filter_spatial_lengths),
|
||||
std::end(filter_spatial_lengths),
|
||||
static_cast<std::size_t>(1),
|
||||
std::multiplies<std::size_t>());
|
||||
}
|
||||
const std::vector<ck::index_t>& output_spatial_lengths);
|
||||
|
||||
/**
|
||||
* @brief Calculate number of bytes read/write by convolution algorithm.
|
||||
@@ -110,20 +134,7 @@ std::size_t get_btype(ck::index_t N,
|
||||
|
||||
struct ConvParams
|
||||
{
|
||||
ConvParams()
|
||||
: num_dim_spatial(2),
|
||||
N(128),
|
||||
K(256),
|
||||
C(192),
|
||||
filter_spatial_lengths(2, 3),
|
||||
input_spatial_lengths(2, 71),
|
||||
conv_filter_strides(2, 2),
|
||||
conv_filter_dilations(2, 1),
|
||||
input_left_pads(2, 1),
|
||||
input_right_pads(2, 1)
|
||||
{
|
||||
}
|
||||
|
||||
ConvParams();
|
||||
ConvParams(ck::index_t n_dim,
|
||||
ck::index_t n_batch,
|
||||
ck::index_t n_out_channels,
|
||||
@@ -133,29 +144,7 @@ struct ConvParams
|
||||
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(n_dim),
|
||||
N(n_batch),
|
||||
K(n_out_channels),
|
||||
C(n_in_channels),
|
||||
filter_spatial_lengths(filters_len),
|
||||
input_spatial_lengths(input_len),
|
||||
conv_filter_strides(strides),
|
||||
conv_filter_dilations(dilations),
|
||||
input_left_pads(left_pads),
|
||||
input_right_pads(right_pads)
|
||||
{
|
||||
if(filter_spatial_lengths.size() != num_dim_spatial ||
|
||||
input_spatial_lengths.size() != num_dim_spatial ||
|
||||
conv_filter_strides.size() != num_dim_spatial ||
|
||||
conv_filter_dilations.size() != num_dim_spatial ||
|
||||
input_left_pads.size() != num_dim_spatial || input_right_pads.size() != num_dim_spatial)
|
||||
{
|
||||
throw(std::runtime_error(
|
||||
"ConvParams::GetOutputSpatialLengths: "
|
||||
"parameter size is different from number of declared dimensions!"));
|
||||
}
|
||||
}
|
||||
const std::vector<ck::index_t>& right_pads);
|
||||
|
||||
ck::index_t num_dim_spatial;
|
||||
ck::index_t N;
|
||||
@@ -171,35 +160,11 @@ struct ConvParams
|
||||
std::vector<ck::index_t> input_left_pads;
|
||||
std::vector<ck::index_t> input_right_pads;
|
||||
|
||||
std::vector<ck::index_t> GetOutputSpatialLengths() const
|
||||
{
|
||||
if(filter_spatial_lengths.size() != num_dim_spatial ||
|
||||
input_spatial_lengths.size() != num_dim_spatial ||
|
||||
conv_filter_strides.size() != num_dim_spatial ||
|
||||
conv_filter_dilations.size() != num_dim_spatial ||
|
||||
input_left_pads.size() != num_dim_spatial || input_right_pads.size() != num_dim_spatial)
|
||||
{
|
||||
throw(std::runtime_error(
|
||||
"ConvParams::GetOutputSpatialLengths: "
|
||||
"parameter size is different from number of declared dimensions!"));
|
||||
}
|
||||
|
||||
std::vector<ck::index_t> out_spatial_len(num_dim_spatial, 0);
|
||||
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::index_t idx_eff =
|
||||
(filter_spatial_lengths[i] - 1) * conv_filter_dilations[i] + 1;
|
||||
out_spatial_len[i] =
|
||||
(input_spatial_lengths[i] + input_left_pads[i] + input_right_pads[i] - idx_eff) /
|
||||
conv_filter_strides[i] +
|
||||
1;
|
||||
}
|
||||
return out_spatial_len;
|
||||
}
|
||||
std::vector<ck::index_t> GetOutputSpatialLengths() const;
|
||||
};
|
||||
|
||||
ConvParams parse_conv_params(int num_dim_spatial, int arg_idx, char* const argv[]);
|
||||
|
||||
/**
|
||||
* @brief Gets the host tensor descriptor.
|
||||
*
|
||||
@@ -221,13 +186,13 @@ HostTensorDescriptor get_host_tensor_descriptor(const std::vector<std::size_t>&
|
||||
std::is_same<TensorLayout, ck::tensor_layout::convolution::NKW>::value)
|
||||
{
|
||||
|
||||
return HostTensorDescriptor(dims, std::vector<std::size_t>({C * dims[2], dims[2], 1}));
|
||||
return HostTensorDescriptor(dims, std::vector<std::size_t>{C * dims[2], dims[2], 1});
|
||||
}
|
||||
else if constexpr(std::is_same<TensorLayout, ck::tensor_layout::convolution::NWC>::value ||
|
||||
std::is_same<TensorLayout, ck::tensor_layout::convolution::KXC>::value ||
|
||||
std::is_same<TensorLayout, ck::tensor_layout::convolution::NWK>::value)
|
||||
{
|
||||
return HostTensorDescriptor(dims, std::vector<std::size_t>({C * dims[2], 1, C}));
|
||||
return HostTensorDescriptor(dims, std::vector<std::size_t>{C * dims[2], 1, C});
|
||||
}
|
||||
// 2D
|
||||
else if constexpr(std::is_same<TensorLayout, ck::tensor_layout::convolution::NCHW>::value ||
|
||||
@@ -273,132 +238,14 @@ HostTensorDescriptor get_host_tensor_descriptor(const std::vector<std::size_t>&
|
||||
throw std::runtime_error(err_msg.str());
|
||||
}
|
||||
|
||||
template <typename InDataType = float,
|
||||
typename WeiDataType = float,
|
||||
typename OutDataType = float,
|
||||
typename InLayout = ck::tensor_layout::convolution::NHWC,
|
||||
typename WeiLayout = ck::tensor_layout::convolution::KYXC,
|
||||
typename OutLayout = ck::tensor_layout::convolution::NHWK>
|
||||
auto get_host_tensors(const ConvParams& params, bool init = true)
|
||||
{
|
||||
std::vector<std::size_t> input_dims{static_cast<std::size_t>(params.N),
|
||||
static_cast<std::size_t>(params.C)};
|
||||
input_dims.insert(std::end(input_dims),
|
||||
std::begin(params.input_spatial_lengths),
|
||||
std::end(params.input_spatial_lengths));
|
||||
|
||||
std::vector<std::size_t> filter_dims{static_cast<std::size_t>(params.K),
|
||||
static_cast<std::size_t>(params.C)};
|
||||
filter_dims.insert(std::end(filter_dims),
|
||||
std::begin(params.filter_spatial_lengths),
|
||||
std::end(params.filter_spatial_lengths));
|
||||
|
||||
const std::vector<ck::index_t>& output_spatial_lengths = params.GetOutputSpatialLengths();
|
||||
std::vector<std::size_t> output_dims{static_cast<std::size_t>(params.N),
|
||||
static_cast<std::size_t>(params.K)};
|
||||
output_dims.insert(std::end(output_dims),
|
||||
std::begin(output_spatial_lengths),
|
||||
std::end(output_spatial_lengths));
|
||||
|
||||
Tensor<InDataType> input(ck::utils::conv::get_host_tensor_descriptor(input_dims, InLayout{}));
|
||||
Tensor<WeiDataType> weights(
|
||||
ck::utils::conv::get_host_tensor_descriptor(filter_dims, WeiLayout{}));
|
||||
Tensor<OutDataType> host_output(
|
||||
ck::utils::conv::get_host_tensor_descriptor(output_dims, OutLayout{}));
|
||||
Tensor<OutDataType> device_output(
|
||||
ck::utils::conv::get_host_tensor_descriptor(output_dims, OutLayout{}));
|
||||
|
||||
if(init)
|
||||
{
|
||||
std::mt19937 gen(11939);
|
||||
if constexpr(std::is_same<InDataType, uint8_t>::value)
|
||||
{
|
||||
std::uniform_int_distribution<> dis(-5, 5);
|
||||
std::generate(
|
||||
input.begin(), input.end(), [&dis, &gen]() { return InDataType(dis(gen)); });
|
||||
std::generate(
|
||||
weights.begin(), weights.end(), [&dis, &gen]() { return WeiDataType(dis(gen)); });
|
||||
}
|
||||
else
|
||||
{
|
||||
std::uniform_real_distribution<> dis(0.f, 1.f);
|
||||
std::generate(
|
||||
input.begin(), input.end(), [&dis, &gen]() { return InDataType(dis(gen)); });
|
||||
std::generate(
|
||||
weights.begin(), weights.end(), [&dis, &gen]() { return WeiDataType(dis(gen)); });
|
||||
}
|
||||
std::fill(host_output.begin(), host_output.end(), OutDataType(0.f));
|
||||
std::fill(device_output.begin(), device_output.end(), OutDataType(0.f));
|
||||
}
|
||||
|
||||
return std::make_tuple(input, weights, host_output, device_output);
|
||||
}
|
||||
|
||||
HostTensorDescriptor get_output_host_tensor_descriptor(const std::vector<std::size_t>& dims,
|
||||
int num_dim_spatial = 2)
|
||||
{
|
||||
namespace tl = ck::tensor_layout::convolution;
|
||||
|
||||
switch(num_dim_spatial)
|
||||
{
|
||||
case 3: {
|
||||
return ck::utils::conv::get_host_tensor_descriptor(dims, tl::NDHWK{});
|
||||
}
|
||||
case 2: {
|
||||
return ck::utils::conv::get_host_tensor_descriptor(dims, tl::NHWK{});
|
||||
}
|
||||
case 1: {
|
||||
return ck::utils::conv::get_host_tensor_descriptor(dims, tl::NWK{});
|
||||
}
|
||||
default: {
|
||||
throw std::runtime_error("Unsupported number of spatial dimensions provided!");
|
||||
}
|
||||
}
|
||||
}
|
||||
int num_dim_spatial = 2);
|
||||
|
||||
HostTensorDescriptor get_filters_host_tensor_descriptor(const std::vector<std::size_t>& dims,
|
||||
int num_dim_spatial = 2)
|
||||
{
|
||||
namespace tl = ck::tensor_layout::convolution;
|
||||
|
||||
switch(num_dim_spatial)
|
||||
{
|
||||
case 3: {
|
||||
return ck::utils::conv::get_host_tensor_descriptor(dims, tl::KZYXC{});
|
||||
}
|
||||
case 2: {
|
||||
return ck::utils::conv::get_host_tensor_descriptor(dims, tl::KYXC{});
|
||||
}
|
||||
case 1: {
|
||||
return ck::utils::conv::get_host_tensor_descriptor(dims, tl::KXC{});
|
||||
}
|
||||
default: {
|
||||
throw std::runtime_error("Unsupported number of spatial dimensions provided!");
|
||||
}
|
||||
}
|
||||
}
|
||||
int num_dim_spatial = 2);
|
||||
|
||||
HostTensorDescriptor get_input_host_tensor_descriptor(const std::vector<std::size_t>& dims,
|
||||
int num_dim_spatial = 2)
|
||||
{
|
||||
namespace tl = ck::tensor_layout::convolution;
|
||||
|
||||
switch(num_dim_spatial)
|
||||
{
|
||||
case 3: {
|
||||
return ck::utils::conv::get_host_tensor_descriptor(dims, tl::NDHWC{});
|
||||
}
|
||||
case 2: {
|
||||
return ck::utils::conv::get_host_tensor_descriptor(dims, tl::NHWC{});
|
||||
}
|
||||
case 1: {
|
||||
return ck::utils::conv::get_host_tensor_descriptor(dims, tl::NWC{});
|
||||
}
|
||||
default: {
|
||||
throw std::runtime_error("Unsupported number of spatial dimensions provided!");
|
||||
}
|
||||
}
|
||||
}
|
||||
int num_dim_spatial = 2);
|
||||
|
||||
template <ck::index_t NDim,
|
||||
typename InDataType = float,
|
||||
@@ -432,123 +279,293 @@ void run_reference_convolution_forward(const ConvParams& params,
|
||||
ref_invoker.Run(ref_argument);
|
||||
}
|
||||
|
||||
template <ck::index_t NDim,
|
||||
typename InDataType = float,
|
||||
typename WeiDataType = float,
|
||||
typename OutDataType = float,
|
||||
template <ck::index_t, typename, typename, typename>
|
||||
class DeviceConvNDFwdInstance>
|
||||
void run_convolution_forward(const ConvParams& params,
|
||||
const Tensor<InDataType>& input,
|
||||
const Tensor<WeiDataType>& weights,
|
||||
Tensor<OutDataType>& output)
|
||||
template <typename InDataType, typename WeiDataType, typename OutDataType>
|
||||
struct ConvolutionFwdInstances;
|
||||
|
||||
template <>
|
||||
struct ConvolutionFwdInstances<float, float, float>
|
||||
{
|
||||
using PassThrough = ck::tensor_operation::element_wise::PassThrough;
|
||||
|
||||
DeviceMem in_device_buf(sizeof(InDataType) * input.mDesc.GetElementSpace());
|
||||
DeviceMem wei_device_buf(sizeof(WeiDataType) * weights.mDesc.GetElementSpace());
|
||||
DeviceMem out_device_buf(sizeof(OutDataType) * output.mDesc.GetElementSpace());
|
||||
|
||||
in_device_buf.ToDevice(input.mData.data());
|
||||
wei_device_buf.ToDevice(weights.mData.data());
|
||||
const std::vector<ck::index_t>& output_spatial_lengths = params.GetOutputSpatialLengths();
|
||||
|
||||
auto conv = DeviceConvNDFwdInstance<NDim, InDataType, WeiDataType, OutDataType>();
|
||||
auto invoker = conv.MakeInvoker();
|
||||
auto argument = conv.MakeArgument(static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
|
||||
static_cast<WeiDataType*>(wei_device_buf.GetDeviceBuffer()),
|
||||
static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
|
||||
params.N,
|
||||
params.K,
|
||||
params.C,
|
||||
params.input_spatial_lengths,
|
||||
params.filter_spatial_lengths,
|
||||
output_spatial_lengths,
|
||||
params.conv_filter_strides,
|
||||
params.conv_filter_dilations,
|
||||
params.input_left_pads,
|
||||
params.input_right_pads,
|
||||
PassThrough{},
|
||||
PassThrough{},
|
||||
PassThrough{});
|
||||
|
||||
if(!conv.IsSupportedArgument(argument))
|
||||
template <int NumDimSpatial,
|
||||
typename std::enable_if<NumDimSpatial >= 1 && NumDimSpatial <= 3, bool>::type = false>
|
||||
static std::vector<DeviceConvFwdNoOpPtr> Get()
|
||||
{
|
||||
throw std::runtime_error(
|
||||
"Error! device_conv with the specified compilation parameters does "
|
||||
"not support this Conv problem");
|
||||
}
|
||||
|
||||
invoker.Run(argument);
|
||||
out_device_buf.FromDevice(output.mData.data());
|
||||
}
|
||||
|
||||
template <ck::index_t NDim,
|
||||
typename InDataType = float,
|
||||
typename WeiDataType = float,
|
||||
typename OutDataType = float>
|
||||
bool run_convolution_forward_instances(const ConvParams& params,
|
||||
const std::vector<DeviceConvFwdNoOpPtr>& conv_ptrs,
|
||||
const Tensor<InDataType>& input,
|
||||
const Tensor<WeiDataType>& weights,
|
||||
Tensor<OutDataType>& output,
|
||||
const Tensor<OutDataType>& host_output)
|
||||
{
|
||||
using PassThrough = ck::tensor_operation::element_wise::PassThrough;
|
||||
|
||||
DeviceMem in_device_buf(sizeof(InDataType) * input.mDesc.GetElementSpace());
|
||||
DeviceMem wei_device_buf(sizeof(WeiDataType) * weights.mDesc.GetElementSpace());
|
||||
DeviceMem out_device_buf(sizeof(OutDataType) * output.mDesc.GetElementSpace());
|
||||
|
||||
in_device_buf.ToDevice(input.mData.data());
|
||||
wei_device_buf.ToDevice(weights.mData.data());
|
||||
const std::vector<ck::index_t>& output_spatial_lengths = params.GetOutputSpatialLengths();
|
||||
|
||||
bool res{true};
|
||||
for(auto& conv_ptr : conv_ptrs)
|
||||
{
|
||||
auto invoker = conv_ptr->MakeInvokerPointer();
|
||||
auto argument = conv_ptr->MakeArgumentPointer(
|
||||
static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
|
||||
static_cast<WeiDataType*>(wei_device_buf.GetDeviceBuffer()),
|
||||
static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
|
||||
params.N,
|
||||
params.K,
|
||||
params.C,
|
||||
params.input_spatial_lengths,
|
||||
params.filter_spatial_lengths,
|
||||
output_spatial_lengths,
|
||||
params.conv_filter_strides,
|
||||
params.conv_filter_dilations,
|
||||
params.input_left_pads,
|
||||
params.input_right_pads,
|
||||
PassThrough{},
|
||||
PassThrough{},
|
||||
PassThrough{});
|
||||
|
||||
if(conv_ptr->IsSupportedArgument(argument.get()))
|
||||
std::vector<DeviceConvFwdNoOpPtr> conv_ptrs;
|
||||
if constexpr(NumDimSpatial == 1)
|
||||
{
|
||||
float atol{1e-5f};
|
||||
float rtol{1e-4f};
|
||||
if constexpr(std::is_same_v<InDataType, ck::half_t>)
|
||||
{
|
||||
atol = 1e-4f;
|
||||
rtol = 2.5e-3f;
|
||||
}
|
||||
invoker->Run(argument.get());
|
||||
out_device_buf.FromDevice(output.mData.data());
|
||||
res = res &&
|
||||
ck::utils::check_err(
|
||||
output.mData, host_output.mData, "Error: incorrect results!", atol, rtol);
|
||||
hipGetErrorString(
|
||||
hipMemset(out_device_buf.GetDeviceBuffer(), 0, out_device_buf.mMemSize));
|
||||
ck::tensor_operation::device::device_conv1d_fwd_instance::
|
||||
add_device_conv1d_fwd_xdl_nwc_kxc_nwk_f32_instances(conv_ptrs);
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 2)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv2d_fwd_instance::
|
||||
add_device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_f32_instances(conv_ptrs);
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 3)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv3d_fwd_instance::
|
||||
add_device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk_f32_instances(conv_ptrs);
|
||||
}
|
||||
return conv_ptrs;
|
||||
}
|
||||
return res;
|
||||
}
|
||||
};
|
||||
|
||||
template <>
|
||||
struct ConvolutionFwdInstances<half_t, half_t, half_t>
|
||||
{
|
||||
template <int NumDimSpatial,
|
||||
typename std::enable_if<NumDimSpatial >= 1 && NumDimSpatial <= 3, bool>::type = false>
|
||||
static std::vector<DeviceConvFwdNoOpPtr> Get()
|
||||
{
|
||||
std::vector<DeviceConvFwdNoOpPtr> conv_ptrs;
|
||||
if constexpr(NumDimSpatial == 1)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv1d_fwd_instance::
|
||||
add_device_conv1d_fwd_xdl_nwc_kxc_nwk_f16_instances(conv_ptrs);
|
||||
return conv_ptrs;
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 2)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv2d_fwd_instance::
|
||||
add_device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_f16_instances(conv_ptrs);
|
||||
ck::tensor_operation::device::device_conv2d_fwd_instance::
|
||||
add_device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk_f16_instances(conv_ptrs);
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 3)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv3d_fwd_instance::
|
||||
add_device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk_f16_instances(conv_ptrs);
|
||||
}
|
||||
return conv_ptrs;
|
||||
}
|
||||
};
|
||||
|
||||
template <>
|
||||
struct ConvolutionFwdInstances<bhalf_t, bhalf_t, bhalf_t>
|
||||
{
|
||||
template <int NumDimSpatial,
|
||||
typename std::enable_if<NumDimSpatial >= 1 && NumDimSpatial <= 3, bool>::type = false>
|
||||
static std::vector<DeviceConvFwdNoOpPtr> Get()
|
||||
{
|
||||
std::vector<DeviceConvFwdNoOpPtr> conv_ptrs;
|
||||
if constexpr(NumDimSpatial == 1)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv1d_fwd_instance::
|
||||
add_device_conv1d_fwd_xdl_nwc_kxc_nwk_bf16_instances(conv_ptrs);
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 2)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv2d_fwd_instance::
|
||||
add_device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_bf16_instances(conv_ptrs);
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 3)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv3d_fwd_instance::
|
||||
add_device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk_bf16_instances(conv_ptrs);
|
||||
}
|
||||
return conv_ptrs;
|
||||
}
|
||||
};
|
||||
|
||||
template <>
|
||||
struct ConvolutionFwdInstances<int8_t, int8_t, int8_t>
|
||||
{
|
||||
template <int NumDimSpatial,
|
||||
typename std::enable_if<NumDimSpatial >= 1 && NumDimSpatial <= 3, bool>::type = false>
|
||||
static std::vector<DeviceConvFwdNoOpPtr> Get()
|
||||
{
|
||||
std::vector<DeviceConvFwdNoOpPtr> conv_ptrs;
|
||||
if constexpr(NumDimSpatial == 1)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv1d_fwd_instance::
|
||||
add_device_conv1d_fwd_xdl_nwc_kxc_nwk_int8_instances(conv_ptrs);
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 2)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv2d_fwd_instance::
|
||||
add_device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_int8_instances(conv_ptrs);
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 3)
|
||||
{
|
||||
ck::tensor_operation::device::device_conv3d_fwd_instance::
|
||||
add_device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk_int8_instances(conv_ptrs);
|
||||
}
|
||||
return conv_ptrs;
|
||||
}
|
||||
};
|
||||
|
||||
template <typename InDataType,
|
||||
typename WeiDataType,
|
||||
typename OutDataType,
|
||||
typename InLayout = ck::tensor_layout::convolution::NHWC,
|
||||
typename WeiLayout = ck::tensor_layout::convolution::KYXC,
|
||||
typename OutLayout = ck::tensor_layout::convolution::NHWK,
|
||||
typename InElementwiseOp = ck::tensor_operation::element_wise::PassThrough,
|
||||
typename WeiElementwiseOp = ck::tensor_operation::element_wise::PassThrough,
|
||||
typename OutElementwiseOp = ck::tensor_operation::element_wise::PassThrough,
|
||||
typename InputInitFun = FillUniform<InDataType>,
|
||||
typename WeightsInitFun = FillUniform<WeiDataType>>
|
||||
class ConvFwdOpInstance : public ck::utils::OpInstance<OutDataType, InDataType, WeiDataType>
|
||||
{
|
||||
using DeviceConvFwdOp = tensor_operation::device::
|
||||
DeviceConvFwd<InElementwiseOp, WeiElementwiseOp, OutElementwiseOp>;
|
||||
using DeviceMemPtr = std::unique_ptr<DeviceMem>;
|
||||
using DeviceBuffers = std::vector<DeviceMemPtr>;
|
||||
using BaseType = ck::utils::OpInstance<OutDataType, InDataType, WeiDataType>;
|
||||
template <typename T>
|
||||
using TensorPtr = std::unique_ptr<Tensor<T>>;
|
||||
using InTensorsTuple = std::tuple<TensorPtr<InDataType>, TensorPtr<WeiDataType>>;
|
||||
|
||||
public:
|
||||
ConvFwdOpInstance() = delete;
|
||||
ConvFwdOpInstance(const ConvFwdOpInstance&) = default;
|
||||
ConvFwdOpInstance& operator=(const ConvFwdOpInstance&) = default;
|
||||
|
||||
ConvFwdOpInstance(const ConvParams& params,
|
||||
bool do_init = true,
|
||||
const InputInitFun& input_init_f = InputInitFun{},
|
||||
const WeightsInitFun& weights_init_f = WeightsInitFun{})
|
||||
: BaseType(),
|
||||
params_{params},
|
||||
output_spatial_lengths_{params.GetOutputSpatialLengths()},
|
||||
do_init_{do_init},
|
||||
input_init_f_{input_init_f},
|
||||
weights_init_f_{weights_init_f}
|
||||
{
|
||||
}
|
||||
|
||||
virtual ~ConvFwdOpInstance() override{};
|
||||
|
||||
virtual InTensorsTuple GetInputTensors() const override
|
||||
{
|
||||
std::vector<std::size_t> input_dims{static_cast<std::size_t>(params_.N),
|
||||
static_cast<std::size_t>(params_.C)};
|
||||
input_dims.insert(std::end(input_dims),
|
||||
std::begin(params_.input_spatial_lengths),
|
||||
std::end(params_.input_spatial_lengths));
|
||||
|
||||
std::vector<std::size_t> filter_dims{static_cast<std::size_t>(params_.K),
|
||||
static_cast<std::size_t>(params_.C)};
|
||||
filter_dims.insert(std::end(filter_dims),
|
||||
std::begin(params_.filter_spatial_lengths),
|
||||
std::end(params_.filter_spatial_lengths));
|
||||
|
||||
auto input = std::make_unique<Tensor<InDataType>>(
|
||||
get_host_tensor_descriptor(input_dims, InLayout{}));
|
||||
auto weights = std::make_unique<Tensor<WeiDataType>>(
|
||||
get_host_tensor_descriptor(filter_dims, WeiLayout{}));
|
||||
|
||||
if(do_init_)
|
||||
{
|
||||
input_init_f_(input->begin(), input->end());
|
||||
weights_init_f_(weights->begin(), weights->end());
|
||||
}
|
||||
|
||||
return std::make_tuple(std::move(input), std::move(weights));
|
||||
}
|
||||
|
||||
virtual TensorPtr<OutDataType> GetOutputTensor() const override
|
||||
{
|
||||
std::vector<std::size_t> output_dims{static_cast<std::size_t>(params_.N),
|
||||
static_cast<std::size_t>(params_.K)};
|
||||
output_dims.insert(std::end(output_dims),
|
||||
std::begin(output_spatial_lengths_),
|
||||
std::end(output_spatial_lengths_));
|
||||
auto output = std::make_unique<Tensor<OutDataType>>(
|
||||
get_host_tensor_descriptor(output_dims, OutLayout{}));
|
||||
|
||||
if(do_init_)
|
||||
{
|
||||
std::fill(output->begin(), output->end(), OutDataType(0.f));
|
||||
}
|
||||
return output;
|
||||
}
|
||||
|
||||
virtual std::unique_ptr<tensor_operation::device::BaseInvoker>
|
||||
MakeInvokerPointer(tensor_operation::device::BaseOperator* op_ptr) const override
|
||||
{
|
||||
static_assert(
|
||||
std::is_same_v<InElementwiseOp, ck::tensor_operation::element_wise::PassThrough>);
|
||||
static_assert(
|
||||
std::is_same_v<OutElementwiseOp, ck::tensor_operation::element_wise::PassThrough>);
|
||||
static_assert(
|
||||
std::is_same_v<WeiElementwiseOp, ck::tensor_operation::element_wise::PassThrough>);
|
||||
|
||||
auto conv_ptr = dynamic_cast<DeviceConvFwdOp*>(op_ptr);
|
||||
if(!conv_ptr)
|
||||
{
|
||||
throw std::runtime_error(
|
||||
"[ConvFwdOpInstance]: couldn't cast op_ptr to DeviceConvFwdNoOpPtr type!");
|
||||
}
|
||||
return conv_ptr->MakeInvokerPointer();
|
||||
}
|
||||
|
||||
virtual std::unique_ptr<tensor_operation::device::BaseArgument>
|
||||
MakeArgumentPointer(tensor_operation::device::BaseOperator* op_ptr,
|
||||
const DeviceBuffers& in_device_buffers,
|
||||
const DeviceMemPtr& out_device_buffer) const override
|
||||
{
|
||||
static_assert(
|
||||
std::is_same_v<InElementwiseOp, ck::tensor_operation::element_wise::PassThrough>);
|
||||
static_assert(
|
||||
std::is_same_v<OutElementwiseOp, ck::tensor_operation::element_wise::PassThrough>);
|
||||
static_assert(
|
||||
std::is_same_v<WeiElementwiseOp, ck::tensor_operation::element_wise::PassThrough>);
|
||||
|
||||
auto conv_ptr = dynamic_cast<DeviceConvFwdOp*>(op_ptr);
|
||||
if(!conv_ptr)
|
||||
{
|
||||
throw std::runtime_error(
|
||||
"[ConvFwdOpInstance]: couldn't cast op_ptr to DeviceConvFwdNoOpPtr type!");
|
||||
}
|
||||
|
||||
return conv_ptr->MakeArgumentPointer(
|
||||
static_cast<InDataType*>(in_device_buffers[0]->GetDeviceBuffer()),
|
||||
static_cast<WeiDataType*>(in_device_buffers[1]->GetDeviceBuffer()),
|
||||
static_cast<OutDataType*>(out_device_buffer->GetDeviceBuffer()),
|
||||
params_.N,
|
||||
params_.K,
|
||||
params_.C,
|
||||
params_.input_spatial_lengths,
|
||||
params_.filter_spatial_lengths,
|
||||
output_spatial_lengths_,
|
||||
params_.conv_filter_strides,
|
||||
params_.conv_filter_dilations,
|
||||
params_.input_left_pads,
|
||||
params_.input_right_pads,
|
||||
InElementwiseOp{},
|
||||
WeiElementwiseOp{},
|
||||
OutElementwiseOp{});
|
||||
}
|
||||
|
||||
virtual std::size_t GetFlops() const override
|
||||
{
|
||||
return get_flops(params_.N,
|
||||
params_.C,
|
||||
params_.K,
|
||||
params_.filter_spatial_lengths,
|
||||
output_spatial_lengths_);
|
||||
}
|
||||
|
||||
virtual std::size_t GetBtype() const override
|
||||
{
|
||||
return get_btype<InDataType, WeiDataType, OutDataType>(params_.N,
|
||||
params_.C,
|
||||
params_.K,
|
||||
params_.input_spatial_lengths,
|
||||
params_.filter_spatial_lengths,
|
||||
output_spatial_lengths_);
|
||||
}
|
||||
|
||||
private:
|
||||
const ConvParams& params_;
|
||||
const std::vector<ck::index_t> output_spatial_lengths_;
|
||||
const bool do_init_;
|
||||
const InputInitFun& input_init_f_;
|
||||
const WeightsInitFun& weights_init_f_;
|
||||
};
|
||||
|
||||
} // namespace conv
|
||||
} // namespace utils
|
||||
} // namespace ck
|
||||
|
||||
#endif
|
||||
std::ostream& operator<<(std::ostream& os, const ck::utils::conv::ConvParams& p);
|
||||
|
||||
81
library/include/ck/library/utility/fill.hpp
Normal file
81
library/include/ck/library/utility/fill.hpp
Normal file
@@ -0,0 +1,81 @@
|
||||
#pragma once
|
||||
|
||||
#include <algorithm>
|
||||
#include <random>
|
||||
|
||||
#include "data_type.hpp"
|
||||
|
||||
namespace ck {
|
||||
namespace utils {
|
||||
|
||||
// template <typename T, class Enable = void>
|
||||
// struct FillUniform;
|
||||
|
||||
// TODO: what's wrong with this specialization???
|
||||
// err: segmentation fault in mt19937 - infinite loop like.
|
||||
// template <typename T>
|
||||
// struct FillUniform<T, typename std::enable_if<std::is_integral<T>::value &&
|
||||
// !std::is_same<T, bhalf_t>::value>::type>
|
||||
// {
|
||||
// int a_{0};
|
||||
// int b_{5};
|
||||
// // T a_ = T{0};
|
||||
// // T b_ = T{5};
|
||||
|
||||
// template <typename ForwardIter>
|
||||
// void operator()(ForwardIter first, ForwardIter last) const
|
||||
// {
|
||||
// std::mt19937 gen{11939};
|
||||
// std::uniform_int_distribution<int> dis(a_, b_);
|
||||
// std::generate(first, last, [&dis, &gen]() { return ck::type_convert<T>(dis(gen)); });
|
||||
// }
|
||||
// };
|
||||
|
||||
// struct FillUniform<T, typename std::enable_if<std::is_floating_point<T>::value ||
|
||||
// std::is_same<T, bhalf_t>::value>::type>
|
||||
template <typename T>
|
||||
struct FillUniform
|
||||
{
|
||||
float a_{0};
|
||||
float b_{5};
|
||||
|
||||
template <typename ForwardIter>
|
||||
void operator()(ForwardIter first, ForwardIter last) const
|
||||
{
|
||||
std::mt19937 gen{11939};
|
||||
std::uniform_real_distribution<> dis(a_, b_);
|
||||
std::generate(first, last, [&dis, &gen]() { return ck::type_convert<T>(dis(gen)); });
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct FillMonotonicSeq
|
||||
{
|
||||
T init_value_{0};
|
||||
T step_{1};
|
||||
|
||||
template <typename ForwardIter>
|
||||
void operator()(ForwardIter first, ForwardIter last) const
|
||||
{
|
||||
std::generate(first, last, [=, n = init_value_]() mutable {
|
||||
auto tmp = n;
|
||||
n += step_;
|
||||
return tmp;
|
||||
});
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct FillConstant
|
||||
{
|
||||
T value_{0};
|
||||
|
||||
template <typename ForwardIter>
|
||||
void operator()(ForwardIter first, ForwardIter last) const
|
||||
{
|
||||
std::fill(first, last, value_);
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace utils
|
||||
} // namespace ck
|
||||
231
library/include/ck/library/utility/op_instance_engine.hpp
Normal file
231
library/include/ck/library/utility/op_instance_engine.hpp
Normal file
@@ -0,0 +1,231 @@
|
||||
#pragma once
|
||||
|
||||
#include <cstdlib>
|
||||
#include <limits>
|
||||
#include <memory>
|
||||
#include <stdexcept>
|
||||
#include <tuple>
|
||||
#include <utility>
|
||||
#include <vector>
|
||||
|
||||
#include "check_err.hpp"
|
||||
#include "device_base.hpp"
|
||||
#include "functional2.hpp"
|
||||
|
||||
namespace ck {
|
||||
namespace utils {
|
||||
|
||||
struct ProfileBestConfig
|
||||
{
|
||||
std::string best_op_name;
|
||||
float best_avg_time = std::numeric_limits<float>::max();
|
||||
float best_tflops = std::numeric_limits<float>::max();
|
||||
float best_gb_per_sec = std::numeric_limits<float>::max();
|
||||
};
|
||||
|
||||
/**
|
||||
* @brief This class describes an operation instance(s).
|
||||
*
|
||||
* Op instance defines a particular specializations of operator
|
||||
* template. Thanks to this specific input/output data types, data
|
||||
* layouts and modifying elementwise operations it is able to create
|
||||
* it's input/output tensors, provide pointers to instances which
|
||||
* can execute it and all operation specific parameters.
|
||||
*/
|
||||
template <typename OutDataType, typename... InArgTypes>
|
||||
class OpInstance
|
||||
{
|
||||
public:
|
||||
template <typename T>
|
||||
using TensorPtr = std::unique_ptr<Tensor<T>>;
|
||||
using InTensorsTuple = std::tuple<TensorPtr<InArgTypes>...>;
|
||||
using DeviceMemPtr = std::unique_ptr<DeviceMem>;
|
||||
using DeviceBuffers = std::vector<DeviceMemPtr>;
|
||||
|
||||
OpInstance() = default;
|
||||
OpInstance(const OpInstance&) = default;
|
||||
OpInstance& operator=(const OpInstance&) = default;
|
||||
virtual ~OpInstance(){};
|
||||
|
||||
virtual InTensorsTuple GetInputTensors() const = 0;
|
||||
virtual TensorPtr<OutDataType> GetOutputTensor() const = 0;
|
||||
virtual std::unique_ptr<tensor_operation::device::BaseInvoker>
|
||||
MakeInvokerPointer(tensor_operation::device::BaseOperator*) const = 0;
|
||||
virtual std::unique_ptr<tensor_operation::device::BaseArgument>
|
||||
MakeArgumentPointer(tensor_operation::device::BaseOperator*,
|
||||
const DeviceBuffers&,
|
||||
const DeviceMemPtr&) const = 0;
|
||||
virtual std::size_t GetFlops() const = 0;
|
||||
virtual std::size_t GetBtype() const = 0;
|
||||
};
|
||||
|
||||
/**
|
||||
* @brief A generic operation instance run engine.
|
||||
*/
|
||||
template <typename OutDataType, typename... InArgTypes>
|
||||
class OpInstanceRunEngine
|
||||
{
|
||||
public:
|
||||
using OpInstanceT = OpInstance<InArgTypes..., OutDataType>;
|
||||
template <typename T>
|
||||
using TensorPtr = std::unique_ptr<Tensor<T>>;
|
||||
using DeviceMemPtr = std::unique_ptr<DeviceMem>;
|
||||
using InTensorsTuple = std::tuple<TensorPtr<InArgTypes>...>;
|
||||
using DeviceBuffers = std::vector<DeviceMemPtr>;
|
||||
using InArgsTypesTuple = std::tuple<InArgTypes...>;
|
||||
|
||||
OpInstanceRunEngine() = delete;
|
||||
|
||||
template <typename ReferenceOp = std::function<void()>>
|
||||
OpInstanceRunEngine(const OpInstanceT& op_instance,
|
||||
const ReferenceOp& reference_op = ReferenceOp{})
|
||||
: op_instance_{op_instance}
|
||||
{
|
||||
in_tensors_ = op_instance_.GetInputTensors();
|
||||
out_tensor_ = op_instance_.GetOutputTensor();
|
||||
|
||||
if constexpr(std::is_invocable_v<ReferenceOp,
|
||||
const Tensor<InArgTypes>&...,
|
||||
Tensor<OutDataType>&>)
|
||||
{
|
||||
ref_output_ = op_instance_.GetOutputTensor();
|
||||
CallRefOpUnpackArgs(reference_op, std::make_index_sequence<kNInArgs_>{});
|
||||
}
|
||||
AllocateDeviceInputTensors(std::make_index_sequence<kNInArgs_>{});
|
||||
out_device_buffer_ =
|
||||
std::make_unique<DeviceMem>(sizeof(OutDataType) * out_tensor_->mDesc.GetElementSpace());
|
||||
out_device_buffer_->SetZero();
|
||||
}
|
||||
|
||||
virtual ~OpInstanceRunEngine(){};
|
||||
|
||||
template <typename OpInstancePtr>
|
||||
bool Test(const std::vector<OpInstancePtr>& op_ptrs)
|
||||
{
|
||||
bool res{true};
|
||||
for(auto& op_ptr : op_ptrs)
|
||||
{
|
||||
auto invoker = op_instance_.MakeInvokerPointer(op_ptr.get());
|
||||
auto argument = op_instance_.MakeArgumentPointer(
|
||||
op_ptr.get(), in_device_buffers_, out_device_buffer_);
|
||||
if(op_ptr->IsSupportedArgument(argument.get()))
|
||||
{
|
||||
invoker->Run(argument.get());
|
||||
out_device_buffer_->FromDevice(out_tensor_->mData.data());
|
||||
if(!ref_output_)
|
||||
{
|
||||
throw std::runtime_error(
|
||||
"OpInstanceRunEngine::Test: Reference value not availabe."
|
||||
" You have to provide reference function.");
|
||||
}
|
||||
// TODO: enable flexible use of custom check_error functions
|
||||
res = res && check_err(out_tensor_->mData, ref_output_->mData);
|
||||
out_device_buffer_->SetZero();
|
||||
}
|
||||
}
|
||||
return res;
|
||||
}
|
||||
|
||||
template <typename OpInstancePtr>
|
||||
ProfileBestConfig Profile(const std::vector<OpInstancePtr>& op_ptrs,
|
||||
int nrepeat = 100,
|
||||
bool do_verification = false,
|
||||
bool do_log = false)
|
||||
{
|
||||
bool res{true};
|
||||
ProfileBestConfig best_config;
|
||||
|
||||
for(auto& op_ptr : op_ptrs)
|
||||
{
|
||||
auto invoker = op_instance_.MakeInvokerPointer(op_ptr.get());
|
||||
auto argument = op_instance_.MakeArgumentPointer(
|
||||
op_ptr.get(), in_device_buffers_, out_device_buffer_);
|
||||
if(op_ptr->IsSupportedArgument(argument.get()))
|
||||
{
|
||||
std::string op_name = op_ptr->GetTypeString();
|
||||
float avg_time = invoker->Run(argument.get(), nrepeat);
|
||||
|
||||
std::size_t flops = op_instance_.GetFlops();
|
||||
std::size_t num_btype = op_instance_.GetBtype();
|
||||
float tflops = static_cast<float>(flops) / 1.E9 / avg_time;
|
||||
float gb_per_sec = num_btype / 1.E6 / avg_time;
|
||||
|
||||
std::cout << "Perf: " << avg_time << " ms, " << tflops << " TFlops, " << gb_per_sec
|
||||
<< " GB/s, " << op_name << std::endl;
|
||||
|
||||
if(tflops < best_config.best_tflops)
|
||||
{
|
||||
best_config.best_op_name = op_name;
|
||||
best_config.best_tflops = tflops;
|
||||
best_config.best_gb_per_sec = gb_per_sec;
|
||||
best_config.best_avg_time = avg_time;
|
||||
}
|
||||
|
||||
if(do_verification)
|
||||
{
|
||||
out_device_buffer_->FromDevice(out_tensor_->mData.data());
|
||||
if(!ref_output_)
|
||||
{
|
||||
throw std::runtime_error(
|
||||
"OpInstanceRunEngine::Profile: Reference value not availabe."
|
||||
" You have to provide reference function.");
|
||||
}
|
||||
// TODO: enable flexible use of custom check_error functions
|
||||
res = res && CheckErr(out_tensor_->mData, ref_output_->mData);
|
||||
|
||||
if(do_log) {}
|
||||
}
|
||||
out_device_buffer_->SetZero();
|
||||
}
|
||||
}
|
||||
return best_config;
|
||||
}
|
||||
|
||||
void SetAtol(double a) { atol_ = a; }
|
||||
void SetRtol(double r) { rtol_ = r; }
|
||||
|
||||
private:
|
||||
template <typename F, std::size_t... Is>
|
||||
void CallRefOpUnpackArgs(const F& f, std::index_sequence<Is...>) const
|
||||
{
|
||||
f(*std::get<Is>(in_tensors_)..., *ref_output_);
|
||||
}
|
||||
|
||||
template <std::size_t... Is>
|
||||
void AllocateDeviceInputTensors(std::index_sequence<Is...>)
|
||||
{
|
||||
(AllocateDeviceInputTensorsImpl<Is>(), ...);
|
||||
}
|
||||
|
||||
template <std::size_t Index>
|
||||
void AllocateDeviceInputTensorsImpl()
|
||||
{
|
||||
const auto& ts = std::get<Index>(in_tensors_);
|
||||
in_device_buffers_
|
||||
.emplace_back(
|
||||
std::make_unique<DeviceMem>(sizeof(std::tuple_element_t<Index, InArgsTypesTuple>) *
|
||||
ts->mDesc.GetElementSpace()))
|
||||
->ToDevice(ts->mData.data());
|
||||
}
|
||||
|
||||
static constexpr std::size_t kNInArgs_ = std::tuple_size_v<InTensorsTuple>;
|
||||
const OpInstanceT& op_instance_;
|
||||
double rtol_{1e-5};
|
||||
double atol_{1e-8};
|
||||
|
||||
InTensorsTuple in_tensors_;
|
||||
TensorPtr<OutDataType> out_tensor_;
|
||||
TensorPtr<OutDataType> ref_output_;
|
||||
|
||||
DeviceBuffers in_device_buffers_;
|
||||
DeviceMemPtr out_device_buffer_;
|
||||
|
||||
template <typename T>
|
||||
bool CheckErr(const std::vector<T>& dev_out, const std::vector<T>& ref_out) const
|
||||
{
|
||||
return ck::utils::check_err(dev_out, ref_out, "Error: incorrect results!", atol_, rtol_);
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace utils
|
||||
} // namespace ck
|
||||
Reference in New Issue
Block a user