Skip to content

Conv1D and Pool1D for CUDA backend #19058

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
alalek merged 9 commits into from
Jan 21, 2021
Merged

Conv1D and Pool1D for CUDA backend #19058

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
a26e240
CUDA-independent changes
sl-sergei Dec 8, 2020
82bbd02
Add Conv1D and Pool1D for CUDA backend
sl-sergei Dec 8, 2020
03a51d8
CUDA-independent changes
sl-sergei Dec 8, 2020
5836880
Fix typo
sl-sergei Dec 10, 2020
b55e382
fix comment
sl-sergei Jan 12, 2021
d791541
Update fix
sl-sergei Jan 12, 2021
57809b5
make changes more correct for pooling layer
sl-sergei Jan 12, 2021
b8de02a
Minor fixes for review
sl-sergei Jan 20, 2021
a642bc0
Split skip blocks
sl-sergei Jan 21, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
25 changes: 21 additions & 4 deletions modules/dnn/src/cuda/max_unpooling.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,7 +31,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {

namespace raw {
template <class T, std::size_t Order,
typename std::enable_if<Order == 2 || Order == 3, bool>::type = true> /* Order has been hardcoded; see code */
typename std::enable_if<Order == 1 || Order == 2 || Order == 3, bool>::type = true> /* Order has been hardcoded; see code */
__global__ void max_pooling_with_indices(
Span<T> output, Span<T> indices, View<T> input, size_type channels,
array<size_type, Order> out_spatial_dims, array<size_type, Order> in_spatial_dims,
Expand Down Expand Up @@ -72,7 +72,22 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
in_spatial_size *= in_spatial_dims[i];

const auto outer_offset = (n * channels + c) * in_spatial_size;
if (Order == 2) {
if (Order == 1) {
array<index_type, Order> idx;
for (idx[0] = start[0]; idx[0] != end[0]; idx[0]++) {
index_type offset = 0;
index_type stride = 1;
for (int i = Order - 1; i >= 0; i--) {
offset += stride * idx[i];
stride *= in_spatial_dims[i];
}

if (input[outer_offset + offset] > max_value) {
max_idx = offset;
max_value = input[outer_offset + offset];
}
}
} else if (Order == 2) {
array<index_type, Order> idx;
for (idx[0] = start[0]; idx[0] != end[0]; idx[0]++) {
for (idx[1] = start[1]; idx[1] != end[1]; idx[1]++) {
Expand Down Expand Up @@ -206,15 +221,17 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
out_spatial_dims[i] = output.get_axis_size(2 + i);
}

/* only max_pooling2d and max_pooling3d are supported */
CV_Assert(2 <= order && order <= 3);
CV_Assert(1 <= order && order <= 3);
std::size_t channels = input.get_axis_size(1);
if (order == 3) {
launch_max_pooling_kernel<T, 3>(stream, output, indices, input, channels,
out_spatial_dims, in_spatial_dims, window_size, strides, padding_left);
} else if (order == 2) {
launch_max_pooling_kernel<T, 2>(stream, output, indices, input, channels,
out_spatial_dims, in_spatial_dims, window_size, strides, padding_left);
} else if (order == 1) {
launch_max_pooling_kernel<T, 1>(stream, output, indices, input, channels,
out_spatial_dims, in_spatial_dims, window_size, strides, padding_left);
}
}

Expand Down
2 changes: 1 addition & 1 deletion modules/dnn/src/cuda4dnn/primitives/convolution.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -103,7 +103,7 @@ namespace cv { namespace dnn { namespace cuda4dnn {

const auto groups = config.groups;

CV_Assert (1 < convolution_order && convolution_order <= 3);
CV_Assert (1 <= convolution_order && convolution_order <= 3);

const auto rank = input_shape.size();
const auto output_feature_maps = output_shape[1];
Expand Down
5 changes: 2 additions & 3 deletions modules/dnn/src/cuda4dnn/primitives/max_unpooling.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,13 +50,12 @@ namespace cv { namespace dnn { namespace cuda4dnn {
window_size = config.window_size;

const auto pooling_order = window_size.size();
CV_Assert(pooling_order >= 1);

strides = config.strides;
CV_Assert(pooling_order == strides.size());

if (pooling_order != 2 && pooling_order != 3)
CV_Error(Error::StsNotImplemented, "Only 2D/3D max-pooling are supported.");
if (pooling_order < 1 || pooling_order > 3)
CV_Error(Error::StsNotImplemented, "Only 1D/2D/3D max-pooling are supported.");

padding_left.resize(pooling_order);
if (config.padMode == MaxPoolingConfiguration::PaddingMode::MANUAL)
Expand Down
22 changes: 20 additions & 2 deletions modules/dnn/src/layers/convolution_layer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -125,6 +125,9 @@ class BaseConvolutionLayerImpl : public ConvolutionLayer
{
kernel_size.assign(1, kernel_size[0]);
strides.assign(1, strides[0]);
dilations.assign(1, dilations[0]);
pads_begin.assign(1, pads_begin[0]);
pads_end.assign(1, pads_end[0]);
}
CV_Assert(weightShape.dims() == kernel_size.size() + 2);
for (int i = 0; i < kernel_size.size(); i++) {
Expand Down Expand Up @@ -311,8 +314,8 @@ class ConvolutionLayerImpl CV_FINAL : public BaseConvolutionLayerImpl
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA)
{
/* only convolution 2d and 3d supported */
if (ksize == 2 || ksize == 3)
/* only 1d, 2d and 3d convolutions supported */
if (ksize > 0 && ksize <= 3)
return true;

return false;
Expand Down Expand Up @@ -2001,6 +2004,21 @@ class ConvolutionLayerImpl CV_FINAL : public BaseConvolutionLayerImpl
const auto groups = input_feature_maps / input_feature_maps_per_group;

ConvolutionConfiguration config;

if (input_shape.size() == 3)
{
// Conv1D
// We add an extra dim for input and output tensors, because CuDNN doesn't support convolution with 3D tensors
input_shape.insert(std::end(input_shape) - 1, 1);
output_shape.insert(std::end(output_shape) - 1, 1);

// Do the similar thing for the other parameters
pads_begin.insert(std::begin(pads_begin), 0);
pads_end.insert(std::begin(pads_end), 0);
strides.insert(std::begin(strides), 1);
dilations.insert(std::begin(dilations), 1);
kernel_size.insert(std::begin(kernel_size), 1);
}
config.kernel_size.assign(std::begin(kernel_size), std::end(kernel_size));
config.dilations.assign(std::begin(dilations), std::end(dilations));
config.strides.assign(std::begin(strides), std::end(strides));
Expand Down
24 changes: 18 additions & 6 deletions modules/dnn/src/layers/pooling_layer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -178,14 +178,13 @@ class PoolingLayerImpl CV_FINAL : public PoolingLayer

if (inputs[0].dims == 3)
{
//Pool1D
kernel_size.erase(kernel_size.begin() + 1);
strides.erase(strides.begin() + 1);
pads_begin.erase(pads_begin.begin() + 1);
pads_end.erase(pads_end.begin() + 1);
// Pool1D
kernel_size.assign(1, kernel_size[0]);
strides.assign(1, strides[0]);
pads_begin.assign(1, pads_begin[0]);
pads_end.assign(1, pads_end[0]);
}


#ifdef HAVE_OPENCL
poolOp.release();
#endif
Expand Down Expand Up @@ -392,6 +391,19 @@ class PoolingLayerImpl CV_FINAL : public PoolingLayer
return make_cuda_node<cuda4dnn::MaxPoolingOp>(preferableTarget, std::move(context->stream), config);
}

if (input_shape.size() == 3)
{
// Pool1D
// We add an extra dim for input tensor, because CuDNN support pooling only with 2 and 3 spatial dimensions
input_shape.insert(std::end(input_shape) - 1, 1);

// Do the similar thing for the other parameters
pads_begin.insert(std::begin(pads_begin), 0);
pads_end.insert(std::begin(pads_end), 0);
strides.insert(std::begin(strides), 1);
kernel_size.insert(std::begin(kernel_size), 1);
}

PoolingConfiguration config;
if (type == MAX)
{
Expand Down
13 changes: 12 additions & 1 deletion modules/dnn/test/test_onnx_importer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -122,7 +122,8 @@ TEST_P(Test_ONNX_layers, Convolution_variable_weight)

if (backend == DNN_BACKEND_CUDA)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA); // not supported

if (backend == DNN_BACKEND_VKCOM)
applyTestTag(CV_TEST_TAG_DNN_SKIP_VULKAN); // not supported
String basename = "conv_variable_w";
Net net = readNetFromONNX(_tf("models/" + basename + ".onnx"));
ASSERT_FALSE(net.empty());
Expand Down Expand Up @@ -152,6 +153,8 @@ TEST_P(Test_ONNX_layers, Convolution_variable_weight_bias)

if (backend == DNN_BACKEND_CUDA)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA); // not supported
if (backend == DNN_BACKEND_VKCOM)
applyTestTag(CV_TEST_TAG_DNN_SKIP_VULKAN); // not supported

String basename = "conv_variable_wb";
Net net = readNetFromONNX(_tf("models/" + basename + ".onnx"));
Expand Down Expand Up @@ -710,6 +713,10 @@ TEST_P(Test_ONNX_layers, Conv1d_bias)

TEST_P(Test_ONNX_layers, Conv1d_variable_weight)
{
if (backend == DNN_BACKEND_CUDA)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA); // not supported
if (backend == DNN_BACKEND_VKCOM)
applyTestTag(CV_TEST_TAG_DNN_SKIP_VULKAN); // not supported
String basename = "conv1d_variable_w";
Net net = readNetFromONNX(_tf("models/" + basename + ".onnx"));
ASSERT_FALSE(net.empty());
Expand All @@ -730,6 +737,10 @@ TEST_P(Test_ONNX_layers, Conv1d_variable_weight)

TEST_P(Test_ONNX_layers, Conv1d_variable_weight_bias)
{
if (backend == DNN_BACKEND_CUDA)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA); // not supported
if (backend == DNN_BACKEND_VKCOM)
applyTestTag(CV_TEST_TAG_DNN_SKIP_VULKAN); // not supported
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
{
if (target == DNN_TARGET_MYRIAD) applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NGRAPH);
Expand Down