corrected backward pass with stride != 1

Added OpenMP to Grouped Convolution layer
This commit is contained in:
Shubhaam Agrawal
2022-08-17 13:12:13 +08:00
parent 2173817a73
commit 36da227d26
3 changed files with 230 additions and 125 deletions
@@ -376,6 +376,31 @@ void ConvolutionType<
arma::Cube<typename MatType::elem_type> rotatedFilters(weight.n_cols,
weight.n_rows, weight.n_slices);
// To perform the backward pass, we need to dilate all the mappedError.
arma::Cube<typename MatType::elem_type> 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)
@@ -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<typename MatType::elem_type> 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<typename MatType::elem_type> 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);
}
}
}
+120 -1
View File
@@ -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<arma::colvec>(0, 35, 36)),
target(arma::mat("1"))
{
model = new FFN<NegativeLogLikelihood, RandomInitialization>();
model->ResetData(input, target);
model->Add<Convolution>(1, 3, 3, 2, 2, std::tuple<size_t, size_t>(0, 0),
std::tuple<size_t, size_t>(0, 0), "same");
model->Add<LogSoftMax>();
model->InputDimensions() = std::vector<size_t>({ 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<NegativeLogLikelihood, RandomInitialization>* 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<size_t>({ 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));
}