mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-15 18:42:06 +00:00
example of conv bwd weight 1d/2d/3d fp32/fp16/bf16 xdl (#244)
* enable example of conv 1d/3d for bwd weight
* make bf16 kernel do not use atomic add
* using new gridwise gemm for bwd weight on convnd bwd weight
Co-authored-by: Chao Liu <chao.liu2@amd.com>
[ROCm/composable_kernel commit: ac543313bf]
This commit is contained in:
Shaojie WANG
2
example/20_convnd_bwd_weight_xdl/CMakeLists.txt
Normal file
2
example/20_convnd_bwd_weight_xdl/CMakeLists.txt
Normal file
@@ -0,0 +1,2 @@
|
||||
add_example_executable(example_convnd_bwd_weight_xdl convnd_bwd_weight_xdl.cpp)
|
||||
target_link_libraries(example_convnd_bwd_weight_xdl PRIVATE conv_util)
|
||||
387
example/20_convnd_bwd_weight_xdl/convnd_bwd_weight_xdl.cpp
Normal file
387
example/20_convnd_bwd_weight_xdl/convnd_bwd_weight_xdl.cpp
Normal file
@@ -0,0 +1,387 @@
|
||||
#include <iostream>
|
||||
#include <numeric>
|
||||
#include <initializer_list>
|
||||
#include <cstdlib>
|
||||
#include <stdlib.h>
|
||||
#include <half.hpp>
|
||||
|
||||
#include "check_err.hpp"
|
||||
#include "conv_util.hpp"
|
||||
#include "config.hpp"
|
||||
#include "print.hpp"
|
||||
#include "device.hpp"
|
||||
#include "host_tensor.hpp"
|
||||
#include "host_tensor_generator.hpp"
|
||||
#include "device_tensor.hpp"
|
||||
#include "tensor_layout.hpp"
|
||||
#include "element_wise_operation.hpp"
|
||||
#include "device_convnd_backward_weight_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp"
|
||||
#include "reference_conv_backward_weight.hpp"
|
||||
|
||||
using InDataType = ck::half_t;
|
||||
using WeiDataType = ck::half_t;
|
||||
using OutDataType = ck::half_t;
|
||||
using AccDataType = float;
|
||||
|
||||
template <ck::index_t... Is>
|
||||
using S = ck::Sequence<Is...>;
|
||||
|
||||
using InElementOp = ck::tensor_operation::element_wise::PassThrough;
|
||||
using WeiElementOp = ck::tensor_operation::element_wise::PassThrough;
|
||||
using OutElementOp = ck::tensor_operation::element_wise::PassThrough;
|
||||
|
||||
static constexpr auto ConvBwdWeightDefault =
|
||||
ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization::Default;
|
||||
|
||||
using DeviceConvBwdWeightBasePtr =
|
||||
ck::tensor_operation::device::DeviceConvBwdWeightPtr<InElementOp, WeiElementOp, OutElementOp>;
|
||||
|
||||
// clang-format off
|
||||
template <ck::index_t NumDimSpatial>
|
||||
using DeviceConvndBwdWeightInstance = ck::tensor_operation::device::
|
||||
DeviceConvndBwdWeightXdl_C_Shuffle_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K<
|
||||
InDataType, // InDataType
|
||||
WeiDataType, // WeiDataType
|
||||
OutDataType, // OutDataType
|
||||
AccDataType, // AccDataType
|
||||
InElementOp, // InElementwiseOperation
|
||||
WeiElementOp, // WeiElementwiseOperation
|
||||
OutElementOp, // OutElementwiseOperation
|
||||
ConvBwdWeightDefault, // ConvolutionBackwardWeightSpecialization
|
||||
NumDimSpatial, // NumDimSpatial
|
||||
256, // BlockSize
|
||||
128, // MPerBlock
|
||||
128, // NPerBlock
|
||||
4, // K0PerBlock
|
||||
8, // K1
|
||||
32, // MPerXdl
|
||||
32, // NPerXdl
|
||||
2, // MXdlPerWave
|
||||
2, // NXdlPerWave
|
||||
S<1, 4, 16, 4>, // ABlockTransferThreadClusterLengths_K0_M_K1
|
||||
S<0, 3, 1, 2>, // ABlockTransferThreadClusterArrangeOrder
|
||||
S<0, 2, 1, 3>, // ABlockTransferSrcAccessOrder
|
||||
2, // ABlockTransferSrcVectorDim
|
||||
8, // ABlockTransferSrcScalarPerVector
|
||||
2, // ABlockTransferDstScalarPerVector_K1
|
||||
true, // ABlockLdsAddExtraM
|
||||
S<1, 4, 16, 4>, // BBlockTransferThreadClusterLengths_K0_N_K1
|
||||
S<0, 3, 1, 2>, // BBlockTransferThreadClusterArrangeOrder
|
||||
S<0, 2, 1, 3>, // BBlockTransferSrcAccessOrder
|
||||
2, // BBlockTransferSrcVectorDim
|
||||
8, // BBlockTransferSrcScalarPerVector
|
||||
2, // BBlockTransferDstScalarPerVector_K1
|
||||
true, // BBlockLdsAddExtraN
|
||||
1, // CShuffleMXdlPerWavePerShuffle
|
||||
1, // CShuffleNXdlPerWavePerShuffle
|
||||
S<1, 32, 1, 4>, // CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock
|
||||
8>; // CBlockTransferScalarPerVector_NWaveNPerXdl
|
||||
// clang-format on
|
||||
|
||||
template <ck::index_t NumDimSpatial>
|
||||
using ReferenceConvBwdWeightInstance =
|
||||
ck::tensor_operation::host::ReferenceConvBwdWeight<InDataType,
|
||||
WeiDataType,
|
||||
OutDataType,
|
||||
InElementOp,
|
||||
WeiElementOp,
|
||||
OutElementOp,
|
||||
NumDimSpatial>;
|
||||
|
||||
void print_use_msg()
|
||||
{
|
||||
std::cout << "arg1: verification (0=no, 1=yes)\n"
|
||||
<< "arg2: initialization (0=no init, 1=random value, 2= init to 1 )\n"
|
||||
<< "arg3: time kernel (0=n0, 1=yes)\n"
|
||||
<< "arg4: is show log (0=no, 1=yes)\n"
|
||||
<< "arg5: split-k \n"
|
||||
<< "arg6: N spatial dimensions (default 2)\n"
|
||||
<< "Following arguments (depending on number of spatial dims):\n"
|
||||
<< " 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::endl;
|
||||
}
|
||||
|
||||
ck::utils::conv::ConvParams parse_conv_params(int num_dim_spatial, char* argv[])
|
||||
{
|
||||
// (N, K, C) + num_dim_spatial * 6 (filter, input, strides, dilations, pad left, pad right)
|
||||
ck::utils::conv::ConvParams params;
|
||||
int arg_idx = 7;
|
||||
|
||||
params.num_dim_spatial_ = num_dim_spatial;
|
||||
params.N_ = std::stoi(argv[arg_idx++]);
|
||||
params.K_ = std::stoi(argv[arg_idx++]);
|
||||
params.C_ = std::stoi(argv[arg_idx++]);
|
||||
|
||||
params.filter_spatial_lengths_.resize(num_dim_spatial);
|
||||
for(int i = 0; i < num_dim_spatial; ++i)
|
||||
{
|
||||
params.filter_spatial_lengths_[i] = std::stoi(argv[arg_idx++]);
|
||||
}
|
||||
params.input_spatial_lengths_.resize(num_dim_spatial);
|
||||
for(int i = 0; i < num_dim_spatial; ++i)
|
||||
{
|
||||
params.input_spatial_lengths_[i] = std::stoi(argv[arg_idx++]);
|
||||
}
|
||||
params.conv_filter_strides_.resize(num_dim_spatial);
|
||||
for(int i = 0; i < num_dim_spatial; ++i)
|
||||
{
|
||||
params.conv_filter_strides_[i] = std::stoi(argv[arg_idx++]);
|
||||
}
|
||||
params.conv_filter_dilations_.resize(num_dim_spatial);
|
||||
for(int i = 0; i < num_dim_spatial; ++i)
|
||||
{
|
||||
params.conv_filter_dilations_[i] = std::stoi(argv[arg_idx++]);
|
||||
}
|
||||
params.input_left_pads_.resize(num_dim_spatial);
|
||||
for(int i = 0; i < num_dim_spatial; ++i)
|
||||
{
|
||||
params.input_left_pads_[i] = std::stoi(argv[arg_idx++]);
|
||||
}
|
||||
params.input_right_pads_.resize(num_dim_spatial);
|
||||
for(int i = 0; i < num_dim_spatial; ++i)
|
||||
{
|
||||
params.input_right_pads_[i] = std::stoi(argv[arg_idx++]);
|
||||
}
|
||||
|
||||
return params;
|
||||
}
|
||||
|
||||
DeviceConvBwdWeightBasePtr get_conv_instance(int num_dim_spatial)
|
||||
{
|
||||
switch(num_dim_spatial)
|
||||
{
|
||||
case 3: {
|
||||
return std::make_unique<DeviceConvndBwdWeightInstance<3>>();
|
||||
}
|
||||
case 2: {
|
||||
return std::make_unique<DeviceConvndBwdWeightInstance<2>>();
|
||||
}
|
||||
case 1: {
|
||||
return std::make_unique<DeviceConvndBwdWeightInstance<1>>();
|
||||
}
|
||||
default: {
|
||||
throw std::runtime_error("Unsupported number of spatial dimensions provided!");
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[])
|
||||
{
|
||||
bool do_verification = true;
|
||||
int init_method = 1;
|
||||
bool time_kernel = false;
|
||||
int num_dim_spatial = 2;
|
||||
int do_log = 0;
|
||||
int split_k = 1;
|
||||
|
||||
ck::utils::conv::ConvParams params;
|
||||
params.C_ = 128;
|
||||
|
||||
if(argc == 6)
|
||||
{
|
||||
do_verification = std::stoi(argv[1]);
|
||||
init_method = std::stoi(argv[2]);
|
||||
time_kernel = std::stoi(argv[3]);
|
||||
do_log = std::stoi(argv[4]);
|
||||
split_k = std::stoi(argv[5]);
|
||||
}
|
||||
else if(argc > 6)
|
||||
{
|
||||
do_verification = std::stoi(argv[1]);
|
||||
init_method = std::stoi(argv[2]);
|
||||
time_kernel = std::stoi(argv[3]);
|
||||
do_log = std::stoi(argv[4]);
|
||||
split_k = std::stoi(argv[5]);
|
||||
num_dim_spatial = std::stoi(argv[6]);
|
||||
// check args number
|
||||
int conv_args = 3 + num_dim_spatial * 6;
|
||||
int cmdline_nargs = conv_args + 7;
|
||||
if(cmdline_nargs != argc)
|
||||
{
|
||||
print_use_msg();
|
||||
exit(1);
|
||||
}
|
||||
|
||||
params = parse_conv_params(num_dim_spatial, argv);
|
||||
}
|
||||
else if(argc != 1)
|
||||
{
|
||||
print_use_msg();
|
||||
exit(1);
|
||||
}
|
||||
|
||||
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> in_n_c_hi_wi(
|
||||
ck::utils::conv::get_input_host_tensor_descriptor(input_dims, num_dim_spatial));
|
||||
Tensor<WeiDataType> wei_k_c_y_x_host_result(
|
||||
ck::utils::conv::get_filters_host_tensor_descriptor(filter_dims, num_dim_spatial));
|
||||
Tensor<WeiDataType> wei_k_c_y_x_device_result(
|
||||
ck::utils::conv::get_filters_host_tensor_descriptor(filter_dims, num_dim_spatial));
|
||||
Tensor<OutDataType> out_n_k_ho_wo(
|
||||
ck::utils::conv::get_output_host_tensor_descriptor(output_dims, num_dim_spatial));
|
||||
|
||||
std::cout << "in_n_c_hi_wi: " << in_n_c_hi_wi.mDesc << std::endl;
|
||||
std::cout << "wei_k_c_y_x: " << wei_k_c_y_x_device_result.mDesc << std::endl;
|
||||
std::cout << "out_n_k_ho_wo: " << out_n_k_ho_wo.mDesc << std::endl;
|
||||
|
||||
std::cout << "in_n_c_hi_wi: " << in_n_c_hi_wi.mDesc << std::endl;
|
||||
std::cout << "wei_k_c_y_x: " << wei_k_c_y_x_host_result.mDesc << std::endl;
|
||||
std::cout << "out_n_k_ho_wo: " << out_n_k_ho_wo.mDesc << std::endl;
|
||||
|
||||
switch(init_method)
|
||||
{
|
||||
case 0: break;
|
||||
case 1:
|
||||
out_n_k_ho_wo.GenerateTensorValue(GeneratorTensor_2<OutDataType>{-2, 2});
|
||||
in_n_c_hi_wi.GenerateTensorValue(GeneratorTensor_2<WeiDataType>{-2, 2});
|
||||
break;
|
||||
default:
|
||||
out_n_k_ho_wo.GenerateTensorValue(GeneratorTensor_1<OutDataType>{1});
|
||||
in_n_c_hi_wi.GenerateTensorValue(GeneratorTensor_1<WeiDataType>{1});
|
||||
}
|
||||
|
||||
DeviceMem in_device_buf(sizeof(InDataType) * in_n_c_hi_wi.mDesc.GetElementSpace());
|
||||
DeviceMem wei_device_buf(sizeof(WeiDataType) *
|
||||
wei_k_c_y_x_device_result.mDesc.GetElementSpace());
|
||||
DeviceMem out_device_buf(sizeof(OutDataType) * out_n_k_ho_wo.mDesc.GetElementSpace());
|
||||
|
||||
in_device_buf.ToDevice(in_n_c_hi_wi.mData.data());
|
||||
out_device_buf.ToDevice(out_n_k_ho_wo.mData.data());
|
||||
// reset input to zero
|
||||
wei_device_buf.SetZero();
|
||||
|
||||
// do GEMM
|
||||
auto conv = get_conv_instance(num_dim_spatial);
|
||||
auto invoker = conv->MakeInvokerPointer();
|
||||
auto argument =
|
||||
conv->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_,
|
||||
InElementOp{},
|
||||
WeiElementOp{},
|
||||
OutElementOp{},
|
||||
split_k);
|
||||
|
||||
if(!conv->IsSupportedArgument(argument.get()))
|
||||
{
|
||||
std::cout << "wrong! device_conv with the specified compilation parameters does "
|
||||
"not support this Conv problem"
|
||||
<< std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
float ave_time = invoker->Run(argument.get(), StreamConfig{nullptr, time_kernel});
|
||||
|
||||
std::size_t flop = ck::utils::conv::get_flops(
|
||||
params.N_, params.C_, params.K_, params.filter_spatial_lengths_, output_spatial_lengths);
|
||||
std::size_t num_btype = ck::utils::conv::get_btype<InDataType, WeiDataType, OutDataType>(
|
||||
params.N_,
|
||||
params.C_,
|
||||
params.K_,
|
||||
params.input_spatial_lengths_,
|
||||
params.filter_spatial_lengths_,
|
||||
output_spatial_lengths);
|
||||
|
||||
float tflops = static_cast<float>(flop) / 1.E9 / ave_time;
|
||||
|
||||
float gb_per_sec = num_btype / 1.E6 / ave_time;
|
||||
|
||||
std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s"
|
||||
<< std::endl;
|
||||
|
||||
if(do_verification)
|
||||
{
|
||||
auto verify_f = [&](const auto& ref_conv) {
|
||||
auto ref_invoker = ref_conv.MakeInvoker();
|
||||
|
||||
auto ref_argument = ref_conv.MakeArgument(in_n_c_hi_wi,
|
||||
wei_k_c_y_x_host_result,
|
||||
out_n_k_ho_wo,
|
||||
params.conv_filter_strides_,
|
||||
params.conv_filter_dilations_,
|
||||
params.input_left_pads_,
|
||||
params.input_right_pads_,
|
||||
InElementOp{},
|
||||
WeiElementOp{},
|
||||
OutElementOp{});
|
||||
|
||||
ref_invoker.Run(ref_argument);
|
||||
|
||||
wei_device_buf.FromDevice(wei_k_c_y_x_device_result.mData.data());
|
||||
|
||||
if(do_log)
|
||||
{
|
||||
LogRangeAsType<float>(std::cout << "out: ", out_n_k_ho_wo.mData, ",") << std::endl;
|
||||
LogRangeAsType<float>(std::cout << "in : ", in_n_c_hi_wi.mData, ",") << std::endl;
|
||||
LogRangeAsType<float>(
|
||||
std::cout << "wei_device(after): ", wei_k_c_y_x_device_result.mData, ",")
|
||||
<< std::endl;
|
||||
LogRangeAsType<float>(
|
||||
std::cout << "wei_host : ", wei_k_c_y_x_host_result.mData, ",")
|
||||
<< std::endl;
|
||||
}
|
||||
|
||||
return ck::utils::check_err(wei_k_c_y_x_device_result.mData,
|
||||
wei_k_c_y_x_host_result.mData)
|
||||
? 0
|
||||
: 1;
|
||||
};
|
||||
|
||||
switch(num_dim_spatial)
|
||||
{
|
||||
case 3: {
|
||||
auto ref_conv = ReferenceConvBwdWeightInstance<3>();
|
||||
verify_f(ref_conv);
|
||||
break;
|
||||
}
|
||||
case 2: {
|
||||
auto ref_conv = ReferenceConvBwdWeightInstance<2>();
|
||||
verify_f(ref_conv);
|
||||
break;
|
||||
}
|
||||
case 1: {
|
||||
auto ref_conv = ReferenceConvBwdWeightInstance<1>();
|
||||
verify_f(ref_conv);
|
||||
break;
|
||||
}
|
||||
default: {
|
||||
throw std::runtime_error("Unsupported number of spatial dimensions provided!");
|
||||
}
|
||||
}
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
@@ -52,3 +52,4 @@ add_subdirectory(15_grouped_gemm)
|
||||
add_subdirectory(16_gemm_reduce)
|
||||
add_subdirectory(18_batched_gemm_reduce)
|
||||
add_subdirectory(19_binary_elementwise)
|
||||
add_subdirectory(20_convnd_bwd_weight_xdl)
|
||||
|
||||
@@ -0,0 +1,17 @@
|
||||
#pragma once
|
||||
|
||||
namespace ck {
|
||||
namespace tensor_operation {
|
||||
namespace device {
|
||||
|
||||
enum struct ConvolutionBackwardWeightSpecialization
|
||||
{
|
||||
Default,
|
||||
Filter1x1Stride1Pad0,
|
||||
Filter1x1Pad0,
|
||||
OddC,
|
||||
};
|
||||
|
||||
} // namespace device
|
||||
} // namespace tensor_operation
|
||||
} // namespace ck
|
||||
File diff suppressed because it is too large
Load Diff
@@ -1,5 +1,4 @@
|
||||
#ifndef REFERENCE_CONV_WRW_HPP
|
||||
#define REFERENCE_CONV_WRW_HPP
|
||||
#pragma once
|
||||
|
||||
#include <iostream>
|
||||
#include <sstream>
|
||||
@@ -16,7 +15,9 @@ template <typename InDataType,
|
||||
typename OutDataType,
|
||||
typename InElementwiseOperation,
|
||||
typename WeiElementwiseOperation,
|
||||
typename OutElementwiseOperation>
|
||||
typename OutElementwiseOperation,
|
||||
ck::index_t NumDimSpatial = 2,
|
||||
typename ck::enable_if<NumDimSpatial >= 1 && NumDimSpatial <= 3, bool>::type = false>
|
||||
struct ReferenceConvBwdWeight : public device::BaseOperator
|
||||
{
|
||||
// Argument
|
||||
@@ -32,9 +33,9 @@ struct ReferenceConvBwdWeight : public device::BaseOperator
|
||||
InElementwiseOperation in_element_op,
|
||||
WeiElementwiseOperation wei_element_op,
|
||||
OutElementwiseOperation out_element_op)
|
||||
: in_n_c_hi_wi_{in_n_c_hi_wi},
|
||||
wei_k_c_y_x_{wei_k_c_y_x},
|
||||
out_n_k_ho_wo_{out_n_k_ho_wo},
|
||||
: input_{in_n_c_hi_wi},
|
||||
weight_{wei_k_c_y_x},
|
||||
output_{out_n_k_ho_wo},
|
||||
conv_strides_{conv_filter_strides},
|
||||
conv_dilations_{conv_filter_dilations},
|
||||
in_left_pads_{input_left_pads},
|
||||
@@ -45,9 +46,9 @@ struct ReferenceConvBwdWeight : public device::BaseOperator
|
||||
{
|
||||
}
|
||||
|
||||
const Tensor<InDataType>& in_n_c_hi_wi_;
|
||||
Tensor<WeiDataType>& wei_k_c_y_x_;
|
||||
const Tensor<OutDataType>& out_n_k_ho_wo_;
|
||||
const Tensor<InDataType>& input_;
|
||||
Tensor<WeiDataType>& weight_;
|
||||
const Tensor<OutDataType>& output_;
|
||||
|
||||
std::vector<index_t> conv_strides_;
|
||||
std::vector<index_t> conv_dilations_;
|
||||
@@ -66,59 +67,180 @@ struct ReferenceConvBwdWeight : public device::BaseOperator
|
||||
|
||||
float Run(const Argument& arg)
|
||||
{
|
||||
constexpr auto I0 = Number<0>{};
|
||||
constexpr auto I1 = Number<1>{};
|
||||
auto f_kcyx = [&](auto k, auto c, auto y, auto x) {
|
||||
float v_acc = 0;
|
||||
for(std::size_t n = 0; n < arg.out_n_k_ho_wo_.mDesc.GetLengths()[0]; ++n)
|
||||
{
|
||||
for(std::size_t ho = 0; ho < arg.out_n_k_ho_wo_.mDesc.GetLengths()[2]; ++ho)
|
||||
if constexpr(NumDimSpatial == 1)
|
||||
{
|
||||
constexpr auto I0 = Number<0>{};
|
||||
auto f_kcx = [&](auto k, auto c, auto x) {
|
||||
float v_acc = 0;
|
||||
for(std::size_t n = 0; n < arg.output_.mDesc.GetLengths()[0]; ++n)
|
||||
{
|
||||
auto hi = ck::type_convert<ck::long_index_t>(ho * arg.conv_strides_[I0]) +
|
||||
ck::type_convert<ck::long_index_t>(y * arg.conv_dilations_[I0]) -
|
||||
ck::type_convert<ck::long_index_t>(arg.in_left_pads_[I0]);
|
||||
for(std::size_t wo = 0; wo < arg.out_n_k_ho_wo_.mDesc.GetLengths()[3]; ++wo)
|
||||
for(std::size_t wo = 0; wo < arg.output_.mDesc.GetLengths()[2]; ++wo)
|
||||
{
|
||||
auto wi =
|
||||
ck::type_convert<ck::long_index_t>(wo * arg.conv_strides_[I1]) +
|
||||
ck::type_convert<ck::long_index_t>(x * arg.conv_dilations_[I1]) -
|
||||
ck::type_convert<ck::long_index_t>(arg.in_left_pads_[I1]);
|
||||
if(hi >= 0 &&
|
||||
ck::type_convert<std::size_t>(hi) <
|
||||
arg.in_n_c_hi_wi_.mDesc.GetLengths()[2] &&
|
||||
wi >= 0 &&
|
||||
ck::type_convert<std::size_t>(wi) <
|
||||
arg.in_n_c_hi_wi_.mDesc.GetLengths()[3])
|
||||
ck::type_convert<ck::long_index_t>(wo * arg.conv_strides_[I0]) +
|
||||
ck::type_convert<ck::long_index_t>(x * arg.conv_dilations_[I0]) -
|
||||
ck::type_convert<ck::long_index_t>(arg.in_left_pads_[I0]);
|
||||
if(wi >= 0 &&
|
||||
ck::type_convert<std::size_t>(wi) < arg.input_.mDesc.GetLengths()[2])
|
||||
{
|
||||
float v_out;
|
||||
float v_in;
|
||||
|
||||
arg.out_element_op_(
|
||||
v_out,
|
||||
ck::type_convert<float>(arg.out_n_k_ho_wo_(n, k, ho, wo)));
|
||||
arg.in_element_op_(
|
||||
v_in, ck::type_convert<float>(arg.in_n_c_hi_wi_(n, c, hi, wi)));
|
||||
arg.out_element_op_(v_out,
|
||||
ck::type_convert<float>(arg.output_(n, k, wo)));
|
||||
arg.in_element_op_(v_in,
|
||||
ck::type_convert<float>(arg.input_(n, c, wi)));
|
||||
|
||||
v_acc += v_out * v_in;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
float v_wei;
|
||||
float v_wei;
|
||||
|
||||
arg.wei_element_op_(v_wei, v_acc);
|
||||
arg.wei_element_op_(v_wei, v_acc);
|
||||
|
||||
arg.wei_k_c_y_x_(k, c, y, x) = ck::type_convert<OutDataType>(v_wei);
|
||||
};
|
||||
arg.weight_(k, c, x) = ck::type_convert<WeiDataType>(v_wei);
|
||||
};
|
||||
|
||||
make_ParallelTensorFunctor(f_kcyx,
|
||||
arg.wei_k_c_y_x_.mDesc.GetLengths()[0],
|
||||
arg.wei_k_c_y_x_.mDesc.GetLengths()[1],
|
||||
arg.wei_k_c_y_x_.mDesc.GetLengths()[2],
|
||||
arg.wei_k_c_y_x_.mDesc.GetLengths()[3])(
|
||||
std::thread::hardware_concurrency());
|
||||
make_ParallelTensorFunctor(f_kcx,
|
||||
arg.weight_.mDesc.GetLengths()[0],
|
||||
arg.weight_.mDesc.GetLengths()[1],
|
||||
arg.weight_.mDesc.GetLengths()[2])(
|
||||
std::thread::hardware_concurrency());
|
||||
|
||||
return 0;
|
||||
return 0;
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 2)
|
||||
{
|
||||
constexpr auto I0 = Number<0>{};
|
||||
constexpr auto I1 = Number<1>{};
|
||||
auto f_kcyx = [&](auto k, auto c, auto y, auto x) {
|
||||
float v_acc = 0;
|
||||
for(std::size_t n = 0; n < arg.output_.mDesc.GetLengths()[0]; ++n)
|
||||
{
|
||||
for(std::size_t ho = 0; ho < arg.output_.mDesc.GetLengths()[2]; ++ho)
|
||||
{
|
||||
auto hi =
|
||||
ck::type_convert<ck::long_index_t>(ho * arg.conv_strides_[I0]) +
|
||||
ck::type_convert<ck::long_index_t>(y * arg.conv_dilations_[I0]) -
|
||||
ck::type_convert<ck::long_index_t>(arg.in_left_pads_[I0]);
|
||||
for(std::size_t wo = 0; wo < arg.output_.mDesc.GetLengths()[3]; ++wo)
|
||||
{
|
||||
auto wi =
|
||||
ck::type_convert<ck::long_index_t>(wo * arg.conv_strides_[I1]) +
|
||||
ck::type_convert<ck::long_index_t>(x *
|
||||
arg.conv_dilations_[I1]) -
|
||||
ck::type_convert<ck::long_index_t>(arg.in_left_pads_[I1]);
|
||||
if(hi >= 0 &&
|
||||
ck::type_convert<std::size_t>(hi) <
|
||||
arg.input_.mDesc.GetLengths()[2] &&
|
||||
wi >= 0 &&
|
||||
ck::type_convert<std::size_t>(wi) <
|
||||
arg.input_.mDesc.GetLengths()[3])
|
||||
{
|
||||
float v_out;
|
||||
float v_in;
|
||||
|
||||
arg.out_element_op_(
|
||||
v_out, ck::type_convert<float>(arg.output_(n, k, ho, wo)));
|
||||
arg.in_element_op_(
|
||||
v_in, ck::type_convert<float>(arg.input_(n, c, hi, wi)));
|
||||
|
||||
v_acc += v_out * v_in;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
float v_wei;
|
||||
|
||||
arg.wei_element_op_(v_wei, v_acc);
|
||||
|
||||
arg.weight_(k, c, y, x) = ck::type_convert<WeiDataType>(v_wei);
|
||||
};
|
||||
|
||||
make_ParallelTensorFunctor(f_kcyx,
|
||||
arg.weight_.mDesc.GetLengths()[0],
|
||||
arg.weight_.mDesc.GetLengths()[1],
|
||||
arg.weight_.mDesc.GetLengths()[2],
|
||||
arg.weight_.mDesc.GetLengths()[3])(
|
||||
std::thread::hardware_concurrency());
|
||||
|
||||
return 0;
|
||||
}
|
||||
else if constexpr(NumDimSpatial == 3)
|
||||
{
|
||||
constexpr auto I0 = Number<0>{};
|
||||
constexpr auto I1 = Number<1>{};
|
||||
constexpr auto I2 = Number<2>{};
|
||||
auto f_kczyx = [&](auto k, auto c, auto z, auto y, auto x) {
|
||||
float v_acc = 0;
|
||||
for(std::size_t n = 0; n < arg.output_.mDesc.GetLengths()[0]; ++n)
|
||||
{
|
||||
for(std::size_t do_ = 0; do_ < arg.output_.mDesc.GetLengths()[2]; ++do_)
|
||||
{
|
||||
auto di =
|
||||
ck::type_convert<ck::long_index_t>(do_ * arg.conv_strides_[I0]) +
|
||||
ck::type_convert<ck::long_index_t>(z * arg.conv_dilations_[I0]) -
|
||||
ck::type_convert<ck::long_index_t>(arg.in_left_pads_[I0]);
|
||||
for(std::size_t ho = 0; ho < arg.output_.mDesc.GetLengths()[3]; ++ho)
|
||||
{
|
||||
auto hi =
|
||||
ck::type_convert<ck::long_index_t>(ho * arg.conv_strides_[I1]) +
|
||||
ck::type_convert<ck::long_index_t>(y *
|
||||
arg.conv_dilations_[I1]) -
|
||||
ck::type_convert<ck::long_index_t>(arg.in_left_pads_[I1]);
|
||||
for(std::size_t wo = 0; wo < arg.output_.mDesc.GetLengths()[4];
|
||||
++wo)
|
||||
{
|
||||
auto wi =
|
||||
ck::type_convert<ck::long_index_t>(wo *
|
||||
arg.conv_strides_[I2]) +
|
||||
ck::type_convert<ck::long_index_t>(
|
||||
x * arg.conv_dilations_[I2]) -
|
||||
ck::type_convert<ck::long_index_t>(arg.in_left_pads_[I2]);
|
||||
if(di >= 0 &&
|
||||
ck::type_convert<std::size_t>(di) <
|
||||
arg.input_.mDesc.GetLengths()[2] &&
|
||||
hi >= 0 &&
|
||||
ck::type_convert<std::size_t>(hi) <
|
||||
arg.input_.mDesc.GetLengths()[3] &&
|
||||
wi >= 0 &&
|
||||
ck::type_convert<std::size_t>(wi) <
|
||||
arg.input_.mDesc.GetLengths()[4])
|
||||
{
|
||||
float v_out;
|
||||
float v_in;
|
||||
|
||||
arg.out_element_op_(v_out,
|
||||
ck::type_convert<float>(
|
||||
arg.output_(n, k, do_, ho, wo)));
|
||||
arg.in_element_op_(
|
||||
v_in,
|
||||
ck::type_convert<float>(arg.input_(n, c, di, hi, wi)));
|
||||
|
||||
v_acc += v_out * v_in;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
float v_wei;
|
||||
|
||||
arg.wei_element_op_(v_wei, v_acc);
|
||||
|
||||
arg.weight_(k, c, z, y, x) = ck::type_convert<WeiDataType>(v_wei);
|
||||
};
|
||||
|
||||
make_ParallelTensorFunctor(f_kczyx,
|
||||
arg.weight_.mDesc.GetLengths()[0],
|
||||
arg.weight_.mDesc.GetLengths()[1],
|
||||
arg.weight_.mDesc.GetLengths()[2],
|
||||
arg.weight_.mDesc.GetLengths()[3],
|
||||
arg.weight_.mDesc.GetLengths()[4])(
|
||||
std::thread::hardware_concurrency());
|
||||
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
|
||||
float Run(const device::BaseArgument* p_arg,
|
||||
@@ -182,4 +304,3 @@ struct ReferenceConvBwdWeight : public device::BaseOperator
|
||||
} // namespace host
|
||||
} // namespace tensor_operation
|
||||
} // namespace ck
|
||||
#endif
|
||||
|
||||
Reference in New Issue
Block a user