Skip to content

N
ngraph
Commit 1ec803ee authored by Sergey Shalnov's avatar Sergey Shalnov Committed by Scott Cyphers
Browse files
Options

IntelGPU backend: GroupConvolution implementation (#2835) 

* IntelGPU backend: GroupConvolution implementation

* IntelGPU backend: GroupConvolution with clDNN support added

* use better ctor in cldnn
parent cdb2f98e
Hide whitespace changes
Inline Side-by-side
Showing with 141 additions and 13 deletions
src/ngraph/runtime/intelgpu/intelgpu_backend.cpp
View file @ 1ec803ee
...@@ -79,6 +79,7 @@ ...@@ -79,6 +79,7 @@
#include "ngraph/op/fused/conv_fused.hpp" #include "ngraph/op/fused/conv_fused.hpp"
#include "ngraph/op/fused/depth_to_space.hpp" #include "ngraph/op/fused/depth_to_space.hpp"
#include "ngraph/op/fused/elu.hpp" #include "ngraph/op/fused/elu.hpp"
#include "ngraph/op/fused/group_conv.hpp"
#include "ngraph/op/fused/space_to_depth.hpp" #include "ngraph/op/fused/space_to_depth.hpp"
#include "ngraph/op/get_output_element.hpp" #include "ngraph/op/get_output_element.hpp"
#include "ngraph/op/greater.hpp" #include "ngraph/op/greater.hpp"
...@@ -1414,6 +1415,7 @@ shared_ptr<runtime::Executable> ...@@ -1414,6 +1415,7 @@ shared_ptr<runtime::Executable>
case OP_TYPEID::Convolution: case OP_TYPEID::Convolution:
case OP_TYPEID::ConvolutionBias: case OP_TYPEID::ConvolutionBias:
case OP_TYPEID::ConvolutionBiasAdd: case OP_TYPEID::ConvolutionBiasAdd:
case OP_TYPEID::GroupConvolution:
{ {
// since bad inheritance design of these classes // since bad inheritance design of these classes
Strides win_stride; Strides win_stride;
...@@ -1446,6 +1448,18 @@ shared_ptr<runtime::Executable> ...@@ -1446,6 +1448,18 @@ shared_ptr<runtime::Executable>
pad_below = conv_op->get_padding_below(); pad_below = conv_op->get_padding_below();
pad_above = conv_op->get_padding_above(); pad_above = conv_op->get_padding_above();
} }
else if (op_type_id == OP_TYPEID::GroupConvolution)
{
arguments_check(op, 2, 1);
const shared_ptr<op::GroupConvolution> conv_op =
static_pointer_cast<op::GroupConvolution>(op);
win_stride = conv_op->get_window_movement_strides();
win_dilation = conv_op->get_window_dilation_strides();
data_dilation = conv_op->get_data_dilation_strides();
pad_below = conv_op->get_padding_below();
pad_above = conv_op->get_padding_above();
}
else else
{ {
arguments_check(op, 2, 1); arguments_check(op, 2, 1);
...@@ -1475,6 +1489,10 @@ shared_ptr<runtime::Executable> ...@@ -1475,6 +1489,10 @@ shared_ptr<runtime::Executable>
kern.emit<op::ConvolutionBiasAdd>( kern.emit<op::ConvolutionBiasAdd>(
static_pointer_cast<op::ConvolutionBiasAdd>(op)); static_pointer_cast<op::ConvolutionBiasAdd>(op));
} }
else if (op_type_id == OP_TYPEID::GroupConvolution)
{
kern.emit<op::GroupConvolution>(static_pointer_cast<op::GroupConvolution>(op));
}
else else
{ {
kern.emit<op::Convolution>(static_pointer_cast<op::Convolution>(op)); kern.emit<op::Convolution>(static_pointer_cast<op::Convolution>(op));
...@@ -1542,6 +1560,20 @@ shared_ptr<runtime::Executable> ...@@ -1542,6 +1560,20 @@ shared_ptr<runtime::Executable>
topology.add(cldnn_conv_bias_add); topology.add(cldnn_conv_bias_add);
} }
else if (op_type_id == OP_TYPEID::GroupConvolution)
{
const shared_ptr<op::GroupConvolution> conv_op =
static_pointer_cast<op::GroupConvolution>(op);
const cldnn::convolution cldnn_conv(op->get_output_tensor_name(0),
op_input_name,
{op->get_input_tensor_name(1)},
conv_op->get_groups(),
strides,
input_offset,
dilation);
topology.add(cldnn_conv);
}
else else
{ {
const cldnn::convolution cldnn_conv(op->get_output_tensor_name(0), const cldnn::convolution cldnn_conv(op->get_output_tensor_name(0),
...@@ -1946,7 +1978,6 @@ shared_ptr<runtime::Executable> ...@@ -1946,7 +1978,6 @@ shared_ptr<runtime::Executable>
case OP_TYPEID::SpaceToDepth: case OP_TYPEID::SpaceToDepth:
case OP_TYPEID::StopGradient: case OP_TYPEID::StopGradient:
case OP_TYPEID::Transpose: case OP_TYPEID::Transpose:
case OP_TYPEID::GroupConvolution:
default: default:
{ {
throw unsupported_op("Unsupported op '" + op->description() + throw unsupported_op("Unsupported op '" + op->description() +
......
src/ngraph/runtime/intelgpu/intelgpu_kernels.hpp
View file @ 1ec803ee
...@@ -31,6 +31,7 @@ ...@@ -31,6 +31,7 @@
#include "ngraph/op/convolution.hpp" #include "ngraph/op/convolution.hpp"
#include "ngraph/op/equal.hpp" #include "ngraph/op/equal.hpp"
#include "ngraph/op/fused/conv_fused.hpp" #include "ngraph/op/fused/conv_fused.hpp"
#include "ngraph/op/fused/group_conv.hpp"
#include "ngraph/op/greater.hpp" #include "ngraph/op/greater.hpp"
#include "ngraph/op/greater_eq.hpp" #include "ngraph/op/greater_eq.hpp"
#include "ngraph/op/less.hpp" #include "ngraph/op/less.hpp"
...@@ -126,6 +127,7 @@ private: ...@@ -126,6 +127,7 @@ private:
krnl_info build_krnl(const std::shared_ptr<op::BatchNormTrainingBackprop>& op) const; krnl_info build_krnl(const std::shared_ptr<op::BatchNormTrainingBackprop>& op) const;
krnl_info build_krnl(const std::shared_ptr<op::Broadcast>& op) const; krnl_info build_krnl(const std::shared_ptr<op::Broadcast>& op) const;
krnl_info build_krnl(const std::shared_ptr<op::Convolution>& op) const; krnl_info build_krnl(const std::shared_ptr<op::Convolution>& op) const;
krnl_info build_krnl(const std::shared_ptr<op::GroupConvolution>& op) const;
krnl_info build_krnl(const std::shared_ptr<op::ConvolutionBackpropData>& op) const; krnl_info build_krnl(const std::shared_ptr<op::ConvolutionBackpropData>& op) const;
krnl_info build_krnl(const std::shared_ptr<op::ConvolutionBackpropFilters>& op) const; krnl_info build_krnl(const std::shared_ptr<op::ConvolutionBackpropFilters>& op) const;
krnl_info build_krnl(const std::shared_ptr<op::ConvolutionBias>& op) const; krnl_info build_krnl(const std::shared_ptr<op::ConvolutionBias>& op) const;
......
src/ngraph/runtime/intelgpu/intelgpu_op_convolution.cpp
View file @ 1ec803ee
...@@ -108,7 +108,8 @@ static CustomKernels::krnl_info do_convolution_operation(const string& input_nam ...@@ -108,7 +108,8 @@ static CustomKernels::krnl_info do_convolution_operation(const string& input_nam
const string& input_order, const string& input_order,
const string& filter_order, const string& filter_order,
const string& output_order, const string& output_order,
bool reverse_filter) bool reverse_filter,
size_t group_count)
{ {
const string kernel_type_name = get_opencl_type_name(output_type); const string kernel_type_name = get_opencl_type_name(output_type);
const string entry_point_name = "convolution_" + output_name; const string entry_point_name = "convolution_" + output_name;
...@@ -140,6 +141,8 @@ static CustomKernels::krnl_info do_convolution_operation(const string& input_nam ...@@ -140,6 +141,8 @@ static CustomKernels::krnl_info do_convolution_operation(const string& input_nam
writer.block_begin(); writer.block_begin();
{ // Main function body { // Main function body
writer << "const unsigned group_size = " << input_shape.at(input_channel_axis_data) << " / "
<< group_count << " /*group_count*/;\n";
writer << "const unsigned batch = get_global_id(0); /*batch trip count: " writer << "const unsigned batch = get_global_id(0); /*batch trip count: "
<< output_shape.at(batch_axis_data) << "*/\n"; << output_shape.at(batch_axis_data) << "*/\n";
gws.push_back(output_shape.at(batch_axis_data)); gws.push_back(output_shape.at(batch_axis_data));
...@@ -174,8 +177,8 @@ static CustomKernels::krnl_info do_convolution_operation(const string& input_nam ...@@ -174,8 +177,8 @@ static CustomKernels::krnl_info do_convolution_operation(const string& input_nam
} }
writer << kernel_type_name << " result = " << acc_init << ";\n\n" writer << kernel_type_name << " result = " << acc_init << ";\n\n"
<< "// Loop over input_channel\n" << "// Loop over input_channel\n"
<< "for (uint input_channel = 0; input_channel < " << "for (uint input_channel = 0; input_channel < group_size; "
<< input_shape.at(input_channel_axis_data) << "; ++input_channel)\n"; "++input_channel)\n";
writer.block_begin(); writer.block_begin();
{ {
// Loop over filter // Loop over filter
...@@ -326,7 +329,35 @@ CustomKernels::krnl_info CustomKernels::build_krnl(const shared_ptr<op::Convolut ...@@ -326,7 +329,35 @@ CustomKernels::krnl_info CustomKernels::build_krnl(const shared_ptr<op::Convolut
"input[batch][input_channel]", "input[batch][input_channel]",
"filter[output_channel][input_channel]", "filter[output_channel][input_channel]",
"output[batch][output_channel]", "output[batch][output_channel]",
false); false,
1);
}
CustomKernels::krnl_info CustomKernels::build_krnl(const shared_ptr<op::GroupConvolution>& op) const
{
return do_convolution_operation(op->get_input_tensor_name(0),
op->get_input_shape(0),
op->get_input_tensor_name(1),
op->get_input_shape(1),
string(),
{},
string(),
{},
op->get_output_tensor_name(0),
op->get_output_shape(0),
op->get_output_element_type(0),
op->get_padding_below(),
op->get_window_movement_strides(),
op->get_window_dilation_strides(),
op->get_data_dilation_strides(),
0,
1,
1,
"input[batch][(output_channel * group_size) + input_channel]",
"filter[output_channel][input_channel]",
"output[batch][output_channel]",
false,
op->get_groups());
} }
CustomKernels::krnl_info CustomKernels::build_krnl(const shared_ptr<op::ConvolutionBias>& op) const CustomKernels::krnl_info CustomKernels::build_krnl(const shared_ptr<op::ConvolutionBias>& op) const
...@@ -352,7 +383,8 @@ CustomKernels::krnl_info CustomKernels::build_krnl(const shared_ptr<op::Convolut ...@@ -352,7 +383,8 @@ CustomKernels::krnl_info CustomKernels::build_krnl(const shared_ptr<op::Convolut
"input[batch][input_channel]", "input[batch][input_channel]",
"filter[output_channel][input_channel]", "filter[output_channel][input_channel]",
"output[batch][output_channel]", "output[batch][output_channel]",
false); false,
1);
} }
CustomKernels::krnl_info CustomKernels::krnl_info
...@@ -379,7 +411,8 @@ CustomKernels::krnl_info ...@@ -379,7 +411,8 @@ CustomKernels::krnl_info
"input[batch][input_channel]", "input[batch][input_channel]",
"filter[output_channel][input_channel]", "filter[output_channel][input_channel]",
"output[batch][output_channel]", "output[batch][output_channel]",
false); false,
1);
} }
CustomKernels::krnl_info CustomKernels::krnl_info
...@@ -406,7 +439,8 @@ CustomKernels::krnl_info ...@@ -406,7 +439,8 @@ CustomKernels::krnl_info
"input[input_channel][batch]", "input[input_channel][batch]",
"filter[input_channel][output_channel]", "filter[input_channel][output_channel]",
"output[output_channel][batch]", "output[output_channel][batch]",
false); false,
1);
} }
CustomKernels::krnl_info CustomKernels::krnl_info
...@@ -436,7 +470,8 @@ CustomKernels::krnl_info ...@@ -436,7 +470,8 @@ CustomKernels::krnl_info
"input[input_channel][batch]", "input[input_channel][batch]",
"filter[input_channel][output_channel]", "filter[input_channel][output_channel]",
"output[output_channel][batch]", "output[output_channel][batch]",
false); false,
1);
result.insert(result.end(), filter.begin(), filter.end()); result.insert(result.end(), filter.begin(), filter.end());
AxisSet reduce_axes; AxisSet reduce_axes;
...@@ -478,5 +513,6 @@ CustomKernels::krnl_info ...@@ -478,5 +513,6 @@ CustomKernels::krnl_info
"input[batch][input_channel]", "input[batch][input_channel]",
"filter[input_channel][output_channel]", "filter[input_channel][output_channel]",
"output[batch][output_channel]", "output[batch][output_channel]",
true); true,
1);
} }
src/ngraph/runtime/intelgpu/unit_test.manifest
View file @ 1ec803ee
...@@ -28,7 +28,6 @@ convert_float32_bool ...@@ -28,7 +28,6 @@ convert_float32_bool
prelu prelu
prelu_shared_slope prelu_shared_slope
prelu_negative_slope prelu_negative_slope
group_conv
elu elu
elu_negative_alpha elu_negative_alpha
space_to_depth space_to_depth
......
src/ngraph/runtime/intelgpu/visualize_tree.cpp
View file @ 1ec803ee
...@@ -17,6 +17,7 @@ ...@@ -17,6 +17,7 @@
#include <fstream> #include <fstream>
#include <map> #include <map>
#include <memory> #include <memory>
#include <type_traits>
#include "ngraph/runtime/intelgpu/visualize_tree.hpp" #include "ngraph/runtime/intelgpu/visualize_tree.hpp"
...@@ -33,6 +34,8 @@ ...@@ -33,6 +34,8 @@
#include "ngraph/op/convolution.hpp" #include "ngraph/op/convolution.hpp"
#include "ngraph/op/dequantize.hpp" #include "ngraph/op/dequantize.hpp"
#include "ngraph/op/dot.hpp" #include "ngraph/op/dot.hpp"
#include "ngraph/op/fused/conv_fused.hpp"
#include "ngraph/op/fused/group_conv.hpp"
#include "ngraph/op/get_output_element.hpp" #include "ngraph/op/get_output_element.hpp"
#include "ngraph/op/lrn.hpp" #include "ngraph/op/lrn.hpp"
#include "ngraph/op/max.hpp" #include "ngraph/op/max.hpp"
...@@ -331,20 +334,59 @@ void print_node_parameters(ostringstream& writer, const shared_ptr<Node>& node) ...@@ -331,20 +334,59 @@ void print_node_parameters(ostringstream& writer, const shared_ptr<Node>& node)
break; break;
} }
case OP_TYPEID::Convolution: case OP_TYPEID::Convolution:
{
const shared_ptr<op::Convolution> conv_op = static_pointer_cast<op::Convolution>(node);
writer << print_table_row_dims("win_stride", conv_op->get_window_movement_strides())
<< print_table_row_dims("win_dilation", conv_op->get_window_dilation_strides())
<< print_table_row_dims("data_dilation", conv_op->get_data_dilation_strides())
<< print_table_row_value(
"pad_type",
static_cast<underlying_type<op::PadType>::type>(conv_op->get_pad_type()))
<< print_table_row_dims("pad_above", conv_op->get_padding_above())
<< print_table_row_dims("pad_below", conv_op->get_padding_below());
break;
}
case OP_TYPEID::GroupConvolution:
{
const shared_ptr<op::GroupConvolution> conv_op =
static_pointer_cast<op::GroupConvolution>(node);
writer << print_table_row_dims("win_stride", conv_op->get_window_movement_strides())
<< print_table_row_dims("win_dilation", conv_op->get_window_dilation_strides())
<< print_table_row_dims("data_dilation", conv_op->get_data_dilation_strides())
<< print_table_row_value("groups_count", conv_op->get_groups())
<< print_table_row_dims("pad_above", conv_op->get_padding_above())
<< print_table_row_dims("pad_below", conv_op->get_padding_below());
break;
}
case OP_TYPEID::ConvolutionBias: case OP_TYPEID::ConvolutionBias:
{
const shared_ptr<op::ConvolutionBias> conv_op =
static_pointer_cast<op::ConvolutionBias>(node);
writer << print_table_row_dims("win_stride", conv_op->get_window_movement_strides())
<< print_table_row_dims("win_dilation", conv_op->get_window_dilation_strides())
<< print_table_row_dims("data_dilation", conv_op->get_data_dilation_strides())
<< print_table_row_value("with_relu", conv_op->with_relu())
<< print_table_row_dims("pad_above", conv_op->get_padding_above())
<< print_table_row_dims("pad_below", conv_op->get_padding_below());
break;
}
case OP_TYPEID::ConvolutionBiasAdd: case OP_TYPEID::ConvolutionBiasAdd:
{ {
const shared_ptr<op::Convolution> conv_op = static_pointer_cast<op::Convolution>(node); const shared_ptr<op::ConvolutionBiasAdd> conv_op =
static_pointer_cast<op::ConvolutionBiasAdd>(node);
writer << print_table_row_dims("win_stride", conv_op->get_window_movement_strides()) writer << print_table_row_dims("win_stride", conv_op->get_window_movement_strides())
<< print_table_row_dims("win_dilation", conv_op->get_window_dilation_strides()) << print_table_row_dims("win_dilation", conv_op->get_window_dilation_strides())
<< print_table_row_dims("data_dilation", conv_op->get_data_dilation_strides()) << print_table_row_dims("data_dilation", conv_op->get_data_dilation_strides())
<< print_table_row_value("with_relu", conv_op->with_relu())
<< print_table_row_dims("pad_above", conv_op->get_padding_above()) << print_table_row_dims("pad_above", conv_op->get_padding_above())
<< print_table_row_dims("pad_below", conv_op->get_padding_below()); << print_table_row_dims("pad_below", conv_op->get_padding_below());
break; break;
} }
case OP_TYPEID::ConvolutionBackpropFilters: case OP_TYPEID::ConvolutionBackpropFilters:
case OP_TYPEID::ConvolutionBiasBackpropFiltersBias:
{ {
const shared_ptr<op::ConvolutionBackpropFilters> conv_op_filt = const shared_ptr<op::ConvolutionBackpropFilters> conv_op_filt =
static_pointer_cast<op::ConvolutionBackpropFilters>(node); static_pointer_cast<op::ConvolutionBackpropFilters>(node);
...@@ -362,6 +404,24 @@ void print_node_parameters(ostringstream& writer, const shared_ptr<Node>& node) ...@@ -362,6 +404,24 @@ void print_node_parameters(ostringstream& writer, const shared_ptr<Node>& node)
conv_op_filt->get_padding_below_forward()); conv_op_filt->get_padding_below_forward());
break; break;
} }
case OP_TYPEID::ConvolutionBiasBackpropFiltersBias:
{
const shared_ptr<op::ConvolutionBiasBackpropFiltersBias> conv_op_filt =
static_pointer_cast<op::ConvolutionBiasBackpropFiltersBias>(node);
writer << print_table_row_dims("filters_shape", conv_op_filt->get_filters_shape())
<< print_table_row_dims("window_movement_strides_forward",
conv_op_filt->get_window_movement_strides_forward())
<< print_table_row_dims("window_dilation_strides_forward",
conv_op_filt->get_window_dilation_strides_forward())
<< print_table_row_dims("data_dilation_strides_forward",
conv_op_filt->get_data_dilation_strides_forward())
<< print_table_row_dims("pad_above_forward",
conv_op_filt->get_padding_above_forward())
<< print_table_row_dims("pad_below_forward",
conv_op_filt->get_padding_below_forward());
break;
}
case OP_TYPEID::ConvolutionBackpropData: case OP_TYPEID::ConvolutionBackpropData:
{ {
const shared_ptr<op::ConvolutionBackpropData> conv_op_data = const shared_ptr<op::ConvolutionBackpropData> conv_op_data =
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment