From 36da227d2640e915e8fc92cd4ea0bb13ff1927cf Mon Sep 17 00:00:00 2001 From: Shubhaam Agrawal Date: Wed, 17 Aug 2022 13:12:13 +0800 Subject: [PATCH] corrected backward pass with stride != 1 Added OpenMP to Grouped Convolution layer --- .../methods/ann/layer/convolution_impl.hpp | 84 +++++----- .../ann/layer/grouped_convolution_impl.hpp | 150 +++++++++--------- src/mlpack/tests/ann/layer/convolution.cpp | 121 +++++++++++++- 3 files changed, 230 insertions(+), 125 deletions(-) diff --git a/src/mlpack/methods/ann/layer/convolution_impl.hpp b/src/mlpack/methods/ann/layer/convolution_impl.hpp index 50af5559fd..67a5476b79 100644 --- a/src/mlpack/methods/ann/layer/convolution_impl.hpp +++ b/src/mlpack/methods/ann/layer/convolution_impl.hpp @@ -376,6 +376,31 @@ void ConvolutionType< arma::Cube rotatedFilters(weight.n_cols, weight.n_rows, weight.n_slices); + // To perform the backward pass, we need to dilate all the mappedError. + arma::Cube dilatedMappedError; + if (strideHeight == 1 && strideWidth == 1) + { + dilatedMappedError = mappedError; + } + else + { + dilatedMappedError.zeros(mappedError.n_rows * strideWidth - + (strideWidth - 1), mappedError.n_cols * strideHeight - + (strideHeight - 1), mappedError.n_slices); + #pragma omp parallel for collapse(3) + for (size_t i = 0; i < mappedError.n_slices; ++i) + { + for (size_t j = 0; j < mappedError.n_cols; ++j) + { + for (size_t k = 0; k < mappedError.n_rows; ++k) + { + dilatedMappedError(k * strideWidth, j * strideHeight, i) + = mappedError(k, j, i); + } + } + } + } + #pragma omp parallel for for (size_t map = 0; map < (size_t) (maps * inMaps); ++map) { @@ -397,11 +422,11 @@ void ConvolutionType< for (size_t outMap = 0; outMap < maps; ++outMap) { BackwardConvolutionRule::Convolution( - mappedError.slice(outMap + fullOutputOffset), + dilatedMappedError.slice(outMap + fullOutputOffset), rotatedFilters.slice((outMap * inMaps) + inMap), output, - strideHeight, - strideWidth, + 1, + 1, 1, 1, outMap > 0); @@ -409,36 +434,17 @@ void ConvolutionType< // If the stride width or height is greater than 1, then we have to // insert columns and rows into the convolution output. MatType& curGTemp = gTemp.slice(inMap + fullInputOffset); - if (strideWidth == 1 && strideHeight == 1) + if (usingPadding) { - if (usingPadding) - { - curGTemp = output.submat( - padWLeft, - padHTop, - padWLeft + gTemp.n_rows - 1, - padHTop + gTemp.n_cols - 1); - } - else - { - curGTemp = output; - } + curGTemp = output.submat( + padWLeft, + padHTop, + padWLeft + gTemp.n_rows - 1, + padHTop + gTemp.n_cols - 1); } else { - // We must iterate over each element of the output and manually - // re-insert the stride. - size_t col = padWLeft; - for (size_t i = 0; i < output.n_cols; ++i) - { - size_t row = padHTop; - for (size_t j = 0; j < output.n_rows; ++j) - { - curGTemp(row, col) = output(j, i); - row += strideHeight; - } - col += strideWidth; - } + curGTemp = output; } } } @@ -500,26 +506,12 @@ void ConvolutionType< inputTemp.slice(inMap + fullInputOffset), curError, output, + 1, + 1, strideWidth, strideHeight); - // TODO: understand this conditional. Is it needed? - if (gradientTemp.n_rows < output.n_rows || - gradientTemp.n_cols < output.n_cols) - { - gradientTemp.slice((outMap * inMaps) + inMap) += output.submat(0, 0, - gradientTemp.n_rows - 1, gradientTemp.n_cols - 1); - } - else if (gradientTemp.n_rows > output.n_rows || - gradientTemp.n_cols > output.n_cols) - { - gradientTemp.slice((outMap * inMaps) + inMap).submat(0, 0, output.n_rows - 1, - output.n_cols - 1) += output; - } - else - { - gradientTemp.slice((outMap * inMaps) + inMap) += output; - } + gradientTemp.slice((outMap * inMaps) + inMap) += output; } if (useBias) diff --git a/src/mlpack/methods/ann/layer/grouped_convolution_impl.hpp b/src/mlpack/methods/ann/layer/grouped_convolution_impl.hpp index 3f9ed5075f..792e5589f4 100644 --- a/src/mlpack/methods/ann/layer/grouped_convolution_impl.hpp +++ b/src/mlpack/methods/ann/layer/grouped_convolution_impl.hpp @@ -334,30 +334,30 @@ void GroupedConvolutionType< const size_t fullInputOffset = offset * inMaps; const size_t fullOutputOffset = offset * maps; + #pragma omp parallel for collapse(2) for (size_t group = 0; group < groups; group++) { // Iterate over output maps. - for (size_t outMap = group * outGroupSize; - outMap < ((group + 1) * outGroupSize); ++outMap) + for (size_t outMap = 0; outMap < outGroupSize; ++outMap) { + MatType& convOutput = outputTemp.slice(group * outGroupSize + outMap + fullOutputOffset); // Iterate over input maps (we will apply the filter and sum). for (size_t inMap = 0; inMap < inGroupSize; ++inMap) { - MatType convOutput; - ForwardConvolutionRule::Convolution( inputTemp.slice((group * inGroupSize) + inMap + fullInputOffset), - weight.slice((outMap * inGroupSize) + inMap), + weight.slice(((group * outGroupSize + outMap) * inGroupSize) + inMap), convOutput, strideWidth, - strideHeight); - - outputTemp.slice(outMap + fullOutputOffset) += convOutput; + strideHeight, + 1, + 1, + true); } // Make sure to add the bias. if (useBias) - outputTemp.slice(outMap + fullOutputOffset) += bias(outMap); + convOutput += bias(group * outGroupSize + outMap); } } } @@ -391,11 +391,38 @@ void GroupedConvolutionType< // To perform the backward pass, we need to rotate all the filters. arma::Cube rotatedFilters(weight.n_cols, weight.n_rows, weight.n_slices); + + #pragma omp parallel for for (size_t map = 0; map < ((maps * inMaps) / groups); ++map) { Rotate180(weight.slice(map), rotatedFilters.slice(map)); } + // To perform the backward pass, we need to dilate all the mappedError. + arma::Cube dilatedMappedError; + if (strideHeight == 1 && strideWidth == 1) + { + dilatedMappedError = mappedError; + } + else + { + dilatedMappedError.zeros(mappedError.n_rows * strideWidth - + (strideWidth - 1), mappedError.n_cols * strideHeight - + (strideHeight - 1), mappedError.n_slices); + #pragma omp parallel for collapse(3) + for (size_t i = 0; i < mappedError.n_slices; ++i) + { + for (size_t j = 0; j < mappedError.n_cols; ++j) + { + for (size_t k = 0; k < mappedError.n_rows; ++k) + { + dilatedMappedError(k * strideWidth, j * strideHeight, i) + = mappedError(k, j, i); + } + } + } + } + size_t inGroupSize = inMaps / groups; size_t outGroupSize = maps / groups; @@ -405,60 +432,42 @@ void GroupedConvolutionType< const size_t fullInputOffset = offset * inMaps; const size_t fullOutputOffset = offset * maps; + #pragma omp parallel for collapse(2) for (size_t group = 0; group < groups; group++) { - // Iterate over input maps. + // Iterate over input maps. for (size_t inMap = 0; inMap < inGroupSize; ++inMap) { + MatType output; // Iterate over output maps. for (size_t outMap = group * outGroupSize; outMap < ((group + 1) * outGroupSize); ++outMap) { - MatType output; - BackwardConvolutionRule::Convolution( - mappedError.slice(outMap + fullOutputOffset), + dilatedMappedError.slice(outMap + fullOutputOffset), rotatedFilters.slice((outMap * inGroupSize) + inMap), output, - strideHeight, - strideWidth); - - // If the stride width or height is greater than 1, then we have to - // insert columns and rows into the convolution output. - if (strideWidth == 1 && strideHeight == 1) - { - if (usingPadding) - { - gTemp.slice((group * inGroupSize) + inMap + fullInputOffset) += - output.submat( - padWLeft, - padHTop, - padWLeft + gTemp.n_rows - 1, - padHTop + gTemp.n_cols - 1); - } - else - { - gTemp.slice((group * inGroupSize) + inMap + fullInputOffset) += - output; - } - } - else - { - // We must iterate over each element of the output and manually - // re-insert the stride. - size_t col = padWLeft; - for (size_t i = 0; i < output.n_cols; ++i) - { - size_t row = padHTop; - for (size_t j = 0; j < output.n_rows; ++j) - { - gTemp(row, col, (group * inGroupSize) + inMap + - fullInputOffset) += output(j, i); - row += strideHeight; - } - col += strideWidth; - } - } + 1, + 1, + 1, + 1, + outMap > group * outGroupSize); + } + // If the stride width or height is greater than 1, then we have to + // insert columns and rows into the convolution output. + MatType& curGTemp = gTemp.slice((group * inGroupSize) + inMap + + fullInputOffset); + if (usingPadding) + { + curGTemp = output.submat( + padWLeft, + padHTop, + padWLeft + gTemp.n_rows - 1, + padHTop + gTemp.n_cols - 1); + } + else + { + curGTemp = output; } } } @@ -513,49 +522,34 @@ void GroupedConvolutionType< const size_t fullInputOffset = offset * inMaps; const size_t fullOutputOffset = offset * maps; + #pragma omp parallel for collapse(2) for (size_t group = 0; group < groups; group++) { // Iterate over output maps. - for (size_t outMap = group * outGroupSize; - outMap < ((group + 1) * outGroupSize); ++outMap) + for (size_t outMap = 0; outMap < outGroupSize; ++outMap) { // Iterate over input maps (we will apply the filter and sum). + MatType& curError = mappedError.slice(group * outGroupSize + outMap + + fullOutputOffset); for (size_t inMap = 0; inMap < inGroupSize; ++inMap) { MatType output; GradientConvolutionRule::Convolution( inputTemp.slice((group * inGroupSize) + inMap + fullInputOffset), - mappedError.slice(outMap + fullOutputOffset), + curError, output, + 1, + 1, strideWidth, strideHeight); - // TODO: understand this conditional. Is it needed? - if (gradientTemp.n_rows < output.n_rows || - gradientTemp.n_cols < output.n_cols) - { - gradientTemp.slice((outMap * inGroupSize) + inMap) += - output.submat( - 0, - 0, - gradientTemp.n_rows - 1, - gradientTemp.n_cols - 1); - } - else if (gradientTemp.n_rows > output.n_rows || - gradientTemp.n_cols > output.n_cols) - { - gradientTemp.slice((outMap * inGroupSize) + inMap).submat(0, 0, - output.n_rows - 1, output.n_cols - 1) += output; - } - else - { - gradientTemp.slice((outMap * inGroupSize) + inMap) += output; - } + gradientTemp.slice(((group * outGroupSize + outMap) * + inGroupSize) + inMap) += output; } if (useBias) - gradient[weight.n_elem + outMap] += arma::accu(mappedError.slice( - outMap + fullOutputOffset)); + gradient[weight.n_elem + group * outGroupSize + outMap] += + arma::accu(curError); } } } diff --git a/src/mlpack/tests/ann/layer/convolution.cpp b/src/mlpack/tests/ann/layer/convolution.cpp index 7c6be35bfc..9dd17a82bf 100644 --- a/src/mlpack/tests/ann/layer/convolution.cpp +++ b/src/mlpack/tests/ann/layer/convolution.cpp @@ -160,7 +160,48 @@ TEST_CASE("GradientConvolutionLayerTest", "[ANNLayerTest]") arma::mat input, target; } function; - REQUIRE(CheckGradient(function) < 1e3); + REQUIRE(CheckGradient(function) < 1e-1); +} + +/** + * Convolution layer numerical gradient test with stride = 2. + */ +TEST_CASE("GradientConvolutionLayerWithStrideTest", "[ANNLayerTest]") +{ + struct GradientFunction + { + GradientFunction() : + input(arma::linspace(0, 35, 36)), + target(arma::mat("1")) + { + model = new FFN(); + model->ResetData(input, target); + model->Add(1, 3, 3, 2, 2, std::tuple(0, 0), + std::tuple(0, 0), "same"); + model->Add(); + + model->InputDimensions() = std::vector({ 6, 6 }); + } + + ~GradientFunction() + { + delete model; + } + + double Gradient(arma::mat& gradient) const + { + double error = model->Evaluate(model->Parameters(), 0, 1); + model->Gradient(model->Parameters(), 0, gradient, 1); + return error; + } + + arma::mat& Parameters() { return model->Parameters(); } + + FFN* model; + arma::mat input, target; + } function; + + REQUIRE(CheckGradient(function) < 1e-1); } TEST_CASE("ConvolutionLayerTestCase", "[ANNLayerTest]") @@ -327,3 +368,81 @@ TEST_CASE("AdvancedConvolutionLayerTest", "[ANNLayerTest]") layer.Backward(input, output, delta); REQUIRE(arma::accu(delta) == Approx(-1.9237523079).epsilon(1e-5)); } + +/** + * Advanced test for the Convolution layer with stride = 2. + */ +TEST_CASE("AdvancedConvolutionLayerWithStrideTest", "[ANNLayerTest]") +{ + arma::mat input, output; + + // The input test matrix is of the form 3 x 2 x 2 x 2 where + // number of images are 3 and number of feature maps are 2. + input = { { 1, 446, 42 }, + { 2, 16, 63 }, + { 1, 446, 42 }, + { 2, 16, 63 }, + { 3, 13, 63 }, + { 4, 21, 21 }, + { 3, 13, 63 }, + { 4, 21, 21 }, + { 1, 446, 42 }, + { 2, 16, 63 }, + { 1, 446, 42 }, + { 2, 16, 63 }, + { 3, 13, 63 }, + { 4, 21, 21 }, + { 3, 13, 63 }, + { 4, 21, 21 }, + { 1, 13, 11 }, + { 32, 45, 42 }, + { 1, 13, 11 }, + { 32, 45, 42 }, + { 22, 16 , 63 }, + { 32, 13 , 42 }, + { 22, 16 , 63 }, + { 32, 13 , 42 }, + { 1, 13, 11 }, + { 32, 45, 42 }, + { 1, 13, 11 }, + { 32, 45, 42 }, + { 22, 16 , 63 }, + { 32, 13 , 42 }, + { 22, 16 , 63 }, + { 32, 13 , 42 } }; + + Convolution layer(2, 2, 2, 2, 2, 0, 0); + layer.InputDimensions() = std::vector({ 4, 4, 2 }); + layer.ComputeOutputDimensions(); + arma::mat layerWeights(layer.WeightSize(), 1); + layerWeights(0) = 0.34526727; + layerWeights(1) = 0.10398731; + layerWeights(2) = -0.23198915; + layerWeights(3) = 0.05350551; + layerWeights(4) = -0.2239646; + layerWeights(5) = 0.30852968; + layerWeights(6) = -0.2635072; + layerWeights(7) = 0.01724506; + layerWeights(8) = -0.20932047; + layerWeights(9) = 0.2990749; + layerWeights(10) = -0.2981235; + layerWeights(11) = -0.14024211; + layerWeights(12) = -0.09744886; + layerWeights(13) = 0.16249102; + layerWeights(14) = 0.2692932; + layerWeights(15) = -0.12563613; + layerWeights(16) = -0.1114468053; + layerWeights(17) = -0.3029643595; + layer.SetWeights(layerWeights.memptr()); + output.set_size(layer.OutputSize(), 3); + + layer.Forward(input, output); + + // Value calculated using torch.nn.Conv2d(). + REQUIRE(arma::accu(output) == Approx(364.7379150391).epsilon(1e-5)); + + arma::mat delta; + delta.set_size(32, 3); + layer.Backward(input, output, delta); + REQUIRE(arma::accu(delta) == Approx(115.3515701294).epsilon(1e-5)); +}