From 42c9d66723f377ee9524ecca3044a8fd91f7a498 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Fri, 28 Mar 2014 15:04:36 -0700 Subject: [PATCH 01/18] add LRN within map layer and dependencies (eltwise product and power) --- include/caffe/util/math_functions.hpp | 18 +++ include/caffe/vision_layers.hpp | 83 ++++++++++++++ src/caffe/layer_factory.cpp | 6 + src/caffe/layers/eltwise_product_layer.cpp | 62 +++++++++++ src/caffe/layers/eltwise_product_layer.cu | 42 +++++++ src/caffe/layers/lrn_map_layer.cpp | 123 +++++++++++++++++++++ src/caffe/layers/lrn_map_layer.cu | 42 +++++++ src/caffe/layers/power_layer.cpp | 105 ++++++++++++++++++ src/caffe/layers/power_layer.cu | 92 +++++++++++++++ src/caffe/proto/caffe.proto | 25 ++++- src/caffe/util/math_functions.cpp | 36 ++++++ src/caffe/util/math_functions.cu | 98 ++++++++++++++++ 12 files changed, 729 insertions(+), 3 deletions(-) create mode 100644 src/caffe/layers/eltwise_product_layer.cpp create mode 100644 src/caffe/layers/eltwise_product_layer.cu create mode 100644 src/caffe/layers/lrn_map_layer.cpp create mode 100644 src/caffe/layers/lrn_map_layer.cu create mode 100644 src/caffe/layers/power_layer.cpp create mode 100644 src/caffe/layers/power_layer.cu diff --git a/include/caffe/util/math_functions.hpp b/include/caffe/util/math_functions.hpp index 5ac6f8a9ac5..c501f2383a1 100644 --- a/include/caffe/util/math_functions.hpp +++ b/include/caffe/util/math_functions.hpp @@ -57,9 +57,21 @@ void caffe_gpu_axpby(const int N, const Dtype alpha, const Dtype* X, template void caffe_copy(const int N, const Dtype *X, Dtype *Y); +template +void caffe_set(const int N, const Dtype alpha, Dtype *X); + +template +void caffe_gpu_set(const int N, const Dtype alpha, Dtype *X); + template void caffe_gpu_copy(const int N, const Dtype *X, Dtype *Y); +template +void caffe_add_scalar(const int N, const Dtype alpha, Dtype *X); + +template +void caffe_gpu_add_scalar(const int N, const Dtype alpha, Dtype *X); + template void caffe_scal(const int N, const Dtype alpha, Dtype *X); @@ -84,9 +96,15 @@ void caffe_gpu_mul(const int N, const Dtype* a, const Dtype* b, Dtype* y); template void caffe_div(const int N, const Dtype* a, const Dtype* b, Dtype* y); +template +void caffe_gpu_div(const int N, const Dtype* a, const Dtype* b, Dtype* y); + template void caffe_powx(const int n, const Dtype* a, const Dtype b, Dtype* y); +template +void caffe_gpu_powx(const int n, const Dtype* a, const Dtype b, Dtype* y); + template Dtype caffe_nextafter(const Dtype b); diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index 4f6dfa70be2..1d963dd5dd2 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -75,6 +75,30 @@ class DropoutLayer : public NeuronLayer { unsigned int uint_thres_; }; +template +class PowerLayer : public NeuronLayer { + public: + explicit PowerLayer(const LayerParameter& param) + : NeuronLayer(param) {} + virtual void SetUp(const vector*>& bottom, + vector*>* top); + + protected: + virtual Dtype Forward_cpu(const vector*>& bottom, + vector*>* top); + virtual Dtype Forward_gpu(const vector*>& bottom, + vector*>* top); + virtual void Backward_cpu(const vector*>& top, + const bool propagate_down, vector*>* bottom); + virtual void Backward_gpu(const vector*>& top, + const bool propagate_down, vector*>* bottom); + + Dtype power_; + Dtype scale_; + Dtype shift_; + Dtype diff_scale_; +}; + template class ReLULayer : public NeuronLayer { public: @@ -246,6 +270,30 @@ class DataLayer : public Layer { Blob data_mean_; }; +template +class EltwiseProductLayer : public Layer { + public: + explicit EltwiseProductLayer(const LayerParameter& param) + : Layer(param) {} + virtual void SetUp(const vector*>& bottom, + vector*>* top); + + protected: + virtual Dtype Forward_cpu(const vector*>& bottom, + vector*>* top); + virtual Dtype Forward_gpu(const vector*>& bottom, + vector*>* top); + virtual void Backward_cpu(const vector*>& top, + const bool propagate_down, vector*>* bottom); + virtual void Backward_gpu(const vector*>& top, + const bool propagate_down, vector*>* bottom); + + int num_; + int channels_; + int height_; + int width_; +}; + template class EuclideanLossLayer : public Layer { public: @@ -482,6 +530,41 @@ class LRNLayer : public Layer { int width_; }; +template +class LRNMapLayer : public Layer { + public: + explicit LRNMapLayer(const LayerParameter& param) + : Layer(param) {} + virtual void SetUp(const vector*>& bottom, + vector*>* top); + + protected: + virtual Dtype Forward_cpu(const vector*>& bottom, + vector*>* top); + virtual Dtype Forward_gpu(const vector*>& bottom, + vector*>* top); + virtual void Backward_cpu(const vector*>& top, + const bool propagate_down, vector*>* bottom); + virtual void Backward_gpu(const vector*>& top, + const bool propagate_down, vector*>* bottom); + + shared_ptr > square_layer_; + Blob square_input_; + Blob square_output_; + vector*> square_bottom_vec_; + vector*> square_top_vec_; + shared_ptr > conv_layer_; + Blob conv_output_; + vector*> conv_top_vec_; + shared_ptr > power_layer_; + Blob power_output_; + vector*> power_top_vec_; + shared_ptr > product_layer_; + Blob product_data_input_; + vector*> product_bottom_vec_; + vector*> product_top_vec_; +}; + template class MultinomialLogisticLossLayer : public Layer { public: diff --git a/src/caffe/layer_factory.cpp b/src/caffe/layer_factory.cpp index f3e52a68237..91c5850b9c7 100644 --- a/src/caffe/layer_factory.cpp +++ b/src/caffe/layer_factory.cpp @@ -36,6 +36,8 @@ Layer* GetLayer(const LayerParameter& param) { return new DropoutLayer(param); case LayerParameter_LayerType_EUCLIDEAN_LOSS: return new EuclideanLossLayer(param); + case LayerParameter_LayerType_ELTWISE_PRODUCT: + return new EltwiseProductLayer(param); case LayerParameter_LayerType_FLATTEN: return new FlattenLayer(param); case LayerParameter_LayerType_HDF5_DATA: @@ -52,10 +54,14 @@ Layer* GetLayer(const LayerParameter& param) { return new InnerProductLayer(param); case LayerParameter_LayerType_LRN: return new LRNLayer(param); + case LayerParameter_LayerType_LRN_MAP: + return new LRNMapLayer(param); case LayerParameter_LayerType_MULTINOMIAL_LOGISTIC_LOSS: return new MultinomialLogisticLossLayer(param); case LayerParameter_LayerType_POOLING: return new PoolingLayer(param); + case LayerParameter_LayerType_POWER: + return new PowerLayer(param); case LayerParameter_LayerType_RELU: return new ReLULayer(param); case LayerParameter_LayerType_SIGMOID: diff --git a/src/caffe/layers/eltwise_product_layer.cpp b/src/caffe/layers/eltwise_product_layer.cpp new file mode 100644 index 00000000000..c663d10ffb8 --- /dev/null +++ b/src/caffe/layers/eltwise_product_layer.cpp @@ -0,0 +1,62 @@ +// Copyright 2013 Yangqing Jia + +#include + +#include "caffe/layer.hpp" +#include "caffe/vision_layers.hpp" +#include "caffe/util/math_functions.hpp" + +namespace caffe { + +template +void EltwiseProductLayer::SetUp(const vector*>& bottom, + vector*>* top) { + CHECK_GE(bottom.size(), 2) << + "Eltwise Product Layer takes at least 2 blobs as input."; + CHECK_EQ(top->size(), 1) << + "Eltwise Product Layer takes a single blob as output."; + num_ = bottom[0]->num(); + channels_ = bottom[0]->channels(); + height_ = bottom[0]->height(); + width_ = bottom[0]->width(); + for (int i = 1; i < bottom.size(); ++i) { + CHECK_EQ(num_, bottom[i]->num()); + CHECK_EQ(channels_, bottom[i]->channels()); + CHECK_EQ(height_, bottom[i]->height()); + CHECK_EQ(width_, bottom[i]->width()); + } + (*top)[0]->Reshape(num_, channels_, height_, width_); +} + +template +Dtype EltwiseProductLayer::Forward_cpu( + const vector*>& bottom, vector*>* top) { + const int count = (*top)[0]->count(); + Dtype* top_data = (*top)[0]->mutable_cpu_data(); + caffe_mul(count, bottom[0]->cpu_data(), bottom[1]->cpu_data(), top_data); + for (int i = 2; i < bottom.size(); ++i) { + caffe_mul(count, top_data, bottom[i]->cpu_data(), top_data); + } + return Dtype(0.); +} + +template +void EltwiseProductLayer::Backward_cpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { + if (propagate_down) { + const int count = top[0]->count(); + const Dtype* top_data = top[0]->cpu_data(); + const Dtype* top_diff = top[0]->cpu_diff(); + for (int i = 0; i < bottom->size(); ++i) { + const Dtype* bottom_data = (*bottom)[i]->cpu_data(); + Dtype* bottom_diff = (*bottom)[i]->mutable_cpu_diff(); + caffe_div(count, top_data, bottom_data, bottom_diff); + caffe_mul(count, bottom_diff, top_diff, bottom_diff); + } + } +} + +INSTANTIATE_CLASS(EltwiseProductLayer); + + +} // namespace caffe diff --git a/src/caffe/layers/eltwise_product_layer.cu b/src/caffe/layers/eltwise_product_layer.cu new file mode 100644 index 00000000000..d3227294eea --- /dev/null +++ b/src/caffe/layers/eltwise_product_layer.cu @@ -0,0 +1,42 @@ +// Copyright 2013 Yangqing Jia + +#include + +#include "caffe/layer.hpp" +#include "caffe/vision_layers.hpp" +#include "caffe/util/math_functions.hpp" + +namespace caffe { + +template +Dtype EltwiseProductLayer::Forward_gpu( + const vector*>& bottom, vector*>* top) { + const int count = (*top)[0]->count(); + Dtype* top_data = (*top)[0]->mutable_gpu_data(); + caffe_gpu_mul(count, bottom[0]->gpu_data(), bottom[1]->gpu_data(), top_data); + for (int i = 2; i < bottom.size(); ++i) { + caffe_gpu_mul(count, top_data, bottom[i]->gpu_data(), top_data); + } + return Dtype(0.); +} + +template +void EltwiseProductLayer::Backward_gpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { + if (propagate_down) { + const int count = top[0]->count(); + const Dtype* top_data = top[0]->gpu_data(); + const Dtype* top_diff = top[0]->gpu_diff(); + for (int i = 0; i < bottom->size(); ++i) { + const Dtype* bottom_data = (*bottom)[i]->gpu_data(); + Dtype* bottom_diff = (*bottom)[i]->mutable_gpu_diff(); + caffe_gpu_div(count, top_data, bottom_data, bottom_diff); + caffe_gpu_mul(count, bottom_diff, top_diff, bottom_diff); + } + } +} + +INSTANTIATE_CLASS(EltwiseProductLayer); + + +} // namespace caffe diff --git a/src/caffe/layers/lrn_map_layer.cpp b/src/caffe/layers/lrn_map_layer.cpp new file mode 100644 index 00000000000..786168c76c1 --- /dev/null +++ b/src/caffe/layers/lrn_map_layer.cpp @@ -0,0 +1,123 @@ +// Copyright 2014 BVLC and contributors. + +#include + +#include "caffe/filler.hpp" +#include "caffe/layer.hpp" +#include "caffe/vision_layers.hpp" +#include "caffe/util/math_functions.hpp" + +namespace caffe { + +template +void LRNMapLayer::SetUp(const vector*>& bottom, + vector*>* top) { + CHECK_EQ(bottom.size(), 1) << + "Local Response Normalization Layer takes a single blob as input."; + CHECK_EQ(top->size(), 1) << + "Local Response Normalization Layer takes a single blob as output."; + const int num = bottom[0]->num(); + const int channels = bottom[0]->channels(); + const int height = bottom[0]->height(); + const int width = bottom[0]->width(); + const int size_ = this->layer_param_.lrn_map_param().local_size(); + const Dtype pre_pad = (size_ - 1) / 2; + const Dtype alpha = this->layer_param_.lrn_map_param().alpha(); + const Dtype beta = this->layer_param_.lrn_map_param().beta(); + // Set up square layer to square the inputs. + square_input_.Reshape(num, channels, height, width); + square_bottom_vec_.clear(); + square_top_vec_.clear(); + square_bottom_vec_.push_back(&square_input_); + square_top_vec_.push_back(&square_output_); + LayerParameter square_param; + square_param.mutable_power_param()->set_power(Dtype(2)); + square_layer_.reset(new PowerLayer(square_param)); + square_layer_->SetUp(square_bottom_vec_, &square_top_vec_); + CHECK_EQ(square_output_.num(), num); + CHECK_EQ(square_output_.channels(), channels); + CHECK_EQ(square_output_.height(), height); + CHECK_EQ(square_output_.width(), width); + // Set up conv layer to have N filters with N groups, all of which are 1's. + // (With #filters == #groups, each filter looks at exactly 1 input channel, + // which is what we want for this layer type.) + // Output of conv layer gives us the neighborhood response. + conv_top_vec_.clear(); + conv_top_vec_.push_back(&conv_output_); + LayerParameter conv_param; + conv_param.mutable_convolution_param()->set_pad(pre_pad); + conv_param.mutable_convolution_param()->set_kernel_size(size_); + conv_param.mutable_convolution_param()->set_num_output(channels); + conv_param.mutable_convolution_param()->set_group(channels); + conv_param.mutable_convolution_param()->set_bias_term(false); + conv_layer_.reset(new ConvolutionLayer(conv_param)); + conv_layer_->SetUp(square_top_vec_, &conv_top_vec_); + CHECK_EQ(conv_output_.num(), num); + CHECK_EQ(conv_output_.channels(), channels); + CHECK_EQ(conv_output_.height(), height); + CHECK_EQ(conv_output_.width(), width); + FillerParameter one_filler_param; + one_filler_param.set_value(1); + ConstantFiller one_filler(one_filler_param); + one_filler.Fill(conv_layer_->blobs()[0].get()); + // Set up power layer to compute (1 + alpha/N^2 s)^-beta, where s is the sum + // of a squared neighborhood (as output by the conv layer). + power_top_vec_.clear(); + power_top_vec_.push_back(&power_output_); + LayerParameter power_param; + power_param.mutable_power_param()->set_power(-beta); + power_param.mutable_power_param()->set_scale(alpha / (size_ * size_)); + power_param.mutable_power_param()->set_shift(Dtype(1)); + power_layer_.reset(new PowerLayer(power_param)); + power_layer_->SetUp(conv_top_vec_, &power_top_vec_); + CHECK_EQ(power_output_.num(), num); + CHECK_EQ(power_output_.channels(), channels); + CHECK_EQ(power_output_.height(), height); + CHECK_EQ(power_output_.width(), width); + // Set up a product layer to compute outputs by multiplying inputs by scale. + product_bottom_vec_.clear(); + product_bottom_vec_.push_back(bottom[0]); + product_bottom_vec_.push_back(&power_output_); + product_top_vec_.clear(); + product_top_vec_.push_back((*top)[0]); + LayerParameter product_param; + product_layer_.reset(new EltwiseProductLayer(product_param)); + product_layer_->SetUp(product_bottom_vec_, &product_top_vec_); + CHECK_EQ((*top)[0]->num(), num); + CHECK_EQ((*top)[0]->channels(), channels); + CHECK_EQ((*top)[0]->height(), height); + CHECK_EQ((*top)[0]->width(), width); +} + +template +Dtype LRNMapLayer::Forward_cpu(const vector*>& bottom, + vector*>* top) { + const int count = bottom[0]->count(); + const Dtype* bottom_data = bottom[0]->cpu_data(); + Dtype* square_bottom_data = square_input_.mutable_cpu_data(); + caffe_copy(count, bottom_data, square_bottom_data); + square_layer_->Forward(square_bottom_vec_, &square_top_vec_); + conv_layer_->Forward(square_top_vec_, &conv_top_vec_); + power_layer_->Forward(conv_top_vec_, &power_top_vec_); + product_layer_->Forward(product_bottom_vec_, &product_top_vec_); + return Dtype(0.); +} + +template +void LRNMapLayer::Backward_cpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { + if (propagate_down) { + product_layer_->Backward(product_top_vec_, true, &product_bottom_vec_); + power_layer_->Backward(power_top_vec_, true, &conv_top_vec_); + conv_layer_->Backward(conv_top_vec_, true, &square_top_vec_); + square_layer_->Backward(square_top_vec_, true, &square_bottom_vec_); + const int count = (*bottom)[0]->count(); + const Dtype* scale_diff = square_input_.cpu_diff(); + Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff(); + caffe_axpy(count, Dtype(1), scale_diff, bottom_diff); + } +} + +INSTANTIATE_CLASS(LRNMapLayer); + +} // namespace caffe diff --git a/src/caffe/layers/lrn_map_layer.cu b/src/caffe/layers/lrn_map_layer.cu new file mode 100644 index 00000000000..0aba8e1c8d4 --- /dev/null +++ b/src/caffe/layers/lrn_map_layer.cu @@ -0,0 +1,42 @@ +// Copyright 2014 BVLC and contributors. + +#include + +#include "caffe/layer.hpp" +#include "caffe/vision_layers.hpp" +#include "caffe/util/math_functions.hpp" + +namespace caffe { + +template +Dtype LRNMapLayer::Forward_gpu(const vector*>& bottom, + vector*>* top) { + const int count = bottom[0]->count(); + const Dtype* bottom_data = bottom[0]->gpu_data(); + Dtype* square_bottom_data = square_input_.mutable_gpu_data(); + caffe_gpu_copy(count, bottom_data, square_bottom_data); + square_layer_->Forward(square_bottom_vec_, &square_top_vec_); + conv_layer_->Forward(square_top_vec_, &conv_top_vec_); + power_layer_->Forward(conv_top_vec_, &power_top_vec_); + product_layer_->Forward(product_bottom_vec_, &product_top_vec_); + return Dtype(0.); +} + +template +void LRNMapLayer::Backward_gpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { + if (propagate_down) { + product_layer_->Backward(product_top_vec_, true, &product_bottom_vec_); + power_layer_->Backward(power_top_vec_, true, &conv_top_vec_); + conv_layer_->Backward(conv_top_vec_, true, &square_top_vec_); + square_layer_->Backward(square_top_vec_, true, &square_bottom_vec_); + const int count = (*bottom)[0]->count(); + const Dtype* scale_diff = square_input_.gpu_diff(); + Dtype* bottom_diff = (*bottom)[0]->mutable_gpu_diff(); + caffe_gpu_axpy(count, Dtype(1), scale_diff, bottom_diff); + } +} + +INSTANTIATE_CLASS(LRNMapLayer); + +} // namespace caffe diff --git a/src/caffe/layers/power_layer.cpp b/src/caffe/layers/power_layer.cpp new file mode 100644 index 00000000000..3be59d36575 --- /dev/null +++ b/src/caffe/layers/power_layer.cpp @@ -0,0 +1,105 @@ +// Copyright 2014 Jeff Donahue + +#include +#include + +#include "caffe/layer.hpp" +#include "caffe/vision_layers.hpp" +#include "caffe/util/math_functions.hpp" + +using std::max; + +namespace caffe { + +template +void PowerLayer::SetUp(const vector*>& bottom, + vector*>* top) { + NeuronLayer::SetUp(bottom, top); + power_ = this->layer_param_.power_param().power(); + scale_ = this->layer_param_.power_param().scale(); + shift_ = this->layer_param_.power_param().shift(); + diff_scale_ = power_ * scale_; +} + +// Compute y = (shift + scale * x)^power +template +Dtype PowerLayer::Forward_cpu(const vector*>& bottom, + vector*>* top) { + Dtype* top_data = (*top)[0]->mutable_cpu_data(); + const int count = bottom[0]->count(); + // Special case where we can ignore the input: scale or power is 0. + if (diff_scale_ == Dtype(0)) { + Dtype value = (power_ == 0) ? Dtype(1) : pow(shift_, power_); + caffe_set(count, value, top_data); + return Dtype(0); + } + const Dtype* bottom_data = bottom[0]->cpu_data(); + caffe_copy(count, bottom_data, top_data); + if (scale_ != Dtype(1)) { + caffe_scal(count, scale_, top_data); + } + if (shift_ != Dtype(0)) { + caffe_add_scalar(count, shift_, top_data); + } + if (power_ != Dtype(1)) { + caffe_powx(count, top_data, power_, top_data); + } + return Dtype(0); +} + +template +void PowerLayer::Backward_cpu(const vector*>& top, + const bool propagate_down, + vector*>* bottom) { + if (propagate_down) { + Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff(); + const int count = (*bottom)[0]->count(); + const Dtype* top_diff = top[0]->cpu_diff(); + if (diff_scale_ == Dtype(0) || power_ == Dtype(1)) { + caffe_set(count, diff_scale_, bottom_diff); + } else { + const Dtype* bottom_data = (*bottom)[0]->cpu_data(); + // Compute dy/dx = scale * power * (shift + scale * x)^(power - 1) + // = diff_scale * y / (shift + scale * x) + if (power_ == Dtype(2)) { + // Special case for y = (shift + scale * x)^2 + // -> dy/dx = 2 * scale * (shift + scale * x) + // = diff_scale * shift + diff_scale * scale * x + caffe_cpu_axpby(count, diff_scale_ * scale_, bottom_data, + Dtype(0), bottom_diff); + if (shift_ != Dtype(0)) { + caffe_add_scalar(count, diff_scale_ * shift_, bottom_diff); + } + } else if (shift_ == Dtype(0)) { + // Special case for y = (scale * x)^power + // -> dy/dx = scale * power * (scale * x)^(power - 1) + // = scale * power * (scale * x)^power * (scale * x)^(-1) + // = power * y / x + const Dtype* top_data = top[0]->cpu_data(); + caffe_div(count, top_data, bottom_data, bottom_diff); + caffe_scal(count, power_, bottom_diff); + } else { + caffe_copy(count, bottom_data, bottom_diff); + if (scale_ != Dtype(1)) { + caffe_scal(count, scale_, bottom_diff); + } + if (shift_ != Dtype(0)) { + caffe_add_scalar(count, shift_, bottom_diff); + } + const Dtype* top_data = top[0]->cpu_data(); + caffe_div(count, top_data, bottom_diff, bottom_diff); + if (diff_scale_ != Dtype(1)) { + caffe_scal(count, diff_scale_, bottom_diff); + } + } + } + if (diff_scale_ != Dtype(0)) { + caffe_mul(count, top_diff, bottom_diff, bottom_diff); + } + } +} + +INSTANTIATE_CLASS(PowerLayer); + + +} // namespace caffe diff --git a/src/caffe/layers/power_layer.cu b/src/caffe/layers/power_layer.cu new file mode 100644 index 00000000000..9df44612cdf --- /dev/null +++ b/src/caffe/layers/power_layer.cu @@ -0,0 +1,92 @@ +// Copyright 2014 Jeff Donahue + +#include +#include + +#include "caffe/layer.hpp" +#include "caffe/vision_layers.hpp" +#include "caffe/util/math_functions.hpp" + +using std::max; + +namespace caffe { + +template +Dtype PowerLayer::Forward_gpu(const vector*>& bottom, + vector*>* top) { + Dtype* top_data = (*top)[0]->mutable_gpu_data(); + const int count = bottom[0]->count(); + // Special case where we can ignore the input: scale or power is 0. + if (diff_scale_ == Dtype(0)) { + Dtype value = (power_ == 0) ? Dtype(1) : pow(shift_, power_); + caffe_gpu_set(count, value, top_data); + return Dtype(0); + } + const Dtype* bottom_data = bottom[0]->gpu_data(); + caffe_gpu_copy(count, bottom_data, top_data); + if (scale_ != Dtype(1)) { + caffe_gpu_scal(count, scale_, top_data); + } + if (shift_ != Dtype(0)) { + caffe_gpu_add_scalar(count, shift_, top_data); + } + if (power_ != Dtype(1)) { + caffe_gpu_powx(count, top_data, power_, top_data); + } + return Dtype(0); +} + +template +void PowerLayer::Backward_gpu(const vector*>& top, + const bool propagate_down, + vector*>* bottom) { + if (propagate_down) { + Dtype* bottom_diff = (*bottom)[0]->mutable_gpu_diff(); + const int count = (*bottom)[0]->count(); + const Dtype* top_diff = top[0]->gpu_diff(); + if (diff_scale_ == Dtype(0) || power_ == Dtype(1)) { + caffe_gpu_set(count, diff_scale_, bottom_diff); + } else { + const Dtype* bottom_data = (*bottom)[0]->gpu_data(); + // Compute dy/dx = scale * power * (shift + scale * x)^(power - 1) + // = diff_scale * y / (shift + scale * x) + if (power_ == Dtype(2)) { + // Special case for y = (shift + scale * x)^2 + // -> dy/dx = 2 * scale * (shift + scale * x) + // = diff_scale * shift + diff_scale * scale * x + caffe_gpu_axpby(count, diff_scale_ * scale_, bottom_data, + Dtype(0), bottom_diff); + if (shift_ != Dtype(0)) { + caffe_gpu_add_scalar(count, diff_scale_ * shift_, bottom_diff); + } + } else if (shift_ == Dtype(0)) { + // Special case for y = (scale * x)^power + // -> dy/dx = scale * power * (scale * x)^(power - 1) + // = scale * power * (scale * x)^power * (scale * x)^(-1) + // = power * y / x + const Dtype* top_data = top[0]->gpu_data(); + caffe_gpu_div(count, top_data, bottom_data, bottom_diff); + caffe_gpu_scal(count, power_, bottom_diff); + } else { + caffe_gpu_copy(count, bottom_data, bottom_diff); + if (scale_ != Dtype(1)) { + caffe_gpu_scal(count, scale_, bottom_diff); + } + if (shift_ != Dtype(0)) { + caffe_gpu_add_scalar(count, shift_, bottom_diff); + } + const Dtype* top_data = top[0]->gpu_data(); + caffe_gpu_div(count, top_data, bottom_diff, bottom_diff); + if (diff_scale_ != Dtype(1)) { + caffe_gpu_scal(count, diff_scale_, bottom_diff); + } + } + } + caffe_gpu_mul(count, top_diff, bottom_diff, bottom_diff); + } +} + +INSTANTIATE_CLASS(PowerLayer); + + +} // namespace caffe diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 1a31109f784..bda35c1e34d 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -112,6 +112,7 @@ message LayerParameter { DATA = 5; DROPOUT = 6; EUCLIDEAN_LOSS = 7; + ELTWISE_PRODUCT = 1000; FLATTEN = 8; HDF5_DATA = 9; HDF5_OUTPUT = 10; @@ -120,8 +121,10 @@ message LayerParameter { INFOGAIN_LOSS = 13; INNER_PRODUCT = 14; LRN = 15; + LRN_MAP = 1001; MULTINOMIAL_LOGISTIC_LOSS = 16; POOLING = 17; + POWER = 1002; RELU = 18; SIGMOID = 19; SOFTMAX = 20; @@ -151,7 +154,9 @@ message LayerParameter { optional InfogainLossParameter infogain_loss_param = 16; optional InnerProductParameter inner_product_param = 17; optional LRNParameter lrn_param = 18; + optional LRNMapParameter lrn_map_param = 1000; optional PoolingParameter pooling_param = 19; + optional PowerParameter power_param = 1001; optional WindowDataParameter window_data_param = 20; // DEPRECATED: The layer parameters specified as a V0LayerParameter. @@ -263,9 +268,16 @@ message InnerProductParameter { // Message that stores parameters used by LRNLayer message LRNParameter { - optional uint32 local_size = 1 [default = 5]; // for local response norm - optional float alpha = 2 [default = 1.]; // for local response norm - optional float beta = 3 [default = 0.75]; // for local response norm + optional uint32 local_size = 1 [default = 5]; + optional float alpha = 2 [default = 1.]; + optional float beta = 3 [default = 0.75]; +} + +// Message that stores parameters used by LRNMapLayer +message LRNMapParameter { + optional uint32 local_size = 1 [default = 5]; + optional float alpha = 2 [default = 1.]; + optional float beta = 3 [default = 0.75]; } // Message that stores parameters used by PoolingLayer @@ -280,6 +292,13 @@ message PoolingParameter { optional uint32 stride = 3 [default = 1]; // The stride } +// Message that stores parameters used by PowerLayer +message PowerParameter { + optional float power = 1 [default = 1.0]; + optional float scale = 2 [default = 1.0]; + optional float shift = 3 [default = 0.0]; +} + // Message that stores parameters used by WindowDataLayer message WindowDataParameter { // Specify the data source. diff --git a/src/caffe/util/math_functions.cpp b/src/caffe/util/math_functions.cpp index c621255734b..3062dbfde74 100644 --- a/src/caffe/util/math_functions.cpp +++ b/src/caffe/util/math_functions.cpp @@ -120,6 +120,42 @@ void caffe_gpu_axpy(const int N, const double alpha, const double* X, CUBLAS_CHECK(cublasDaxpy(Caffe::cublas_handle(), N, &alpha, X, 1, Y, 1)); } +template <> +void caffe_set(const int N, const float alpha, float* Y) { + if (alpha == 0) { + memset(Y, 0, sizeof(float) * N); + return; + } + for (int i = 0; i < N; ++i) { + Y[i] = alpha; + } +} + +template <> +void caffe_set(const int N, const double alpha, double* Y) { + if (alpha == 0) { + memset(Y, 0, sizeof(double) * N); + return; + } + for (int i = 0; i < N; ++i) { + Y[i] = alpha; + } +} + +template <> +void caffe_add_scalar(const int N, const float alpha, float* Y) { + for (int i = 0; i < N; ++i) { + Y[i] += alpha; + } +} + +template <> +void caffe_add_scalar(const int N, const double alpha, double* Y) { + for (int i = 0; i < N; ++i) { + Y[i] += alpha; + } +} + template <> void caffe_copy(const int N, const float* X, float* Y) { cblas_scopy(N, X, 1, Y, 1); diff --git a/src/caffe/util/math_functions.cu b/src/caffe/util/math_functions.cu index c385890f35e..0ac10d966fd 100644 --- a/src/caffe/util/math_functions.cu +++ b/src/caffe/util/math_functions.cu @@ -13,6 +13,56 @@ namespace caffe { +template +__global__ void set_kernel(const int n, const Dtype alpha, Dtype* y) { + CUDA_KERNEL_LOOP(index, n) { + y[index] = alpha; + } +} + +template <> +void caffe_gpu_set(const int N, const float alpha, float* Y) { + if (alpha == 0) { + CUDA_CHECK(cudaMemset(Y, 0, sizeof(float) * N)); + return; + } + // NOLINT_NEXT_LINE(whitespace/operators) + set_kernel<<>>( + N, alpha, Y); +} + +template <> +void caffe_gpu_set(const int N, const double alpha, double* Y) { + if (alpha == 0) { + CUDA_CHECK(cudaMemset(Y, 0, sizeof(double) * N)); + return; + } + // NOLINT_NEXT_LINE(whitespace/operators) + set_kernel<<>>( + N, alpha, Y); +} + +template +__global__ void add_scalar_kernel(const int n, const Dtype alpha, Dtype* y) { + CUDA_KERNEL_LOOP(index, n) { + y[index] += alpha; + } +} + +template <> +void caffe_gpu_add_scalar(const int N, const float alpha, float* Y) { + // NOLINT_NEXT_LINE(whitespace/operators) + add_scalar_kernel<<>>( + N, alpha, Y); +} + +template <> +void caffe_gpu_add_scalar(const int N, const double alpha, double* Y) { + // NOLINT_NEXT_LINE(whitespace/operators) + add_scalar_kernel<<>>( + N, alpha, Y); +} + template __global__ void mul_kernel(const int n, const Dtype* a, const Dtype* b, Dtype* y) { @@ -37,6 +87,54 @@ void caffe_gpu_mul(const int N, const double* a, N, a, b, y); } +template +__global__ void div_kernel(const int n, const Dtype* a, + const Dtype* b, Dtype* y) { + CUDA_KERNEL_LOOP(index, n) { + y[index] = a[index] / b[index]; + } +} + +template <> +void caffe_gpu_div(const int N, const float* a, + const float* b, float* y) { + // NOLINT_NEXT_LINE(whitespace/operators) + div_kernel<<>>( + N, a, b, y); +} + +template <> +void caffe_gpu_div(const int N, const double* a, + const double* b, double* y) { + // NOLINT_NEXT_LINE(whitespace/operators) + div_kernel<<>>( + N, a, b, y); +} + +template +__global__ void powx_kernel(const int n, const Dtype* a, + const Dtype alpha, Dtype* y) { + CUDA_KERNEL_LOOP(index, n) { + y[index] = pow(a[index], alpha); + } +} + +template <> +void caffe_gpu_powx(const int N, const float* a, + const float alpha, float* y) { + // NOLINT_NEXT_LINE(whitespace/operators) + powx_kernel<<>>( + N, a, alpha, y); +} + +template <> +void caffe_gpu_powx(const int N, const double* a, + const double alpha, double* y) { + // NOLINT_NEXT_LINE(whitespace/operators) + powx_kernel<<>>( + N, a, alpha, y); +} + DEFINE_AND_INSTANTIATE_GPU_UNARY_FUNC(sign, y[index] = (Dtype(0) < x[index]) - (x[index] < Dtype(0))); DEFINE_AND_INSTANTIATE_GPU_UNARY_FUNC(sgnbit, y[index] = signbit(x[index])); From f3e2fe67ae87dbcc0dfa2e812050ebc9fc55c85d Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Fri, 28 Mar 2014 15:39:33 -0700 Subject: [PATCH 02/18] add unit tests for new layer types --- src/caffe/test/test_eltwise_product_layer.cpp | 122 +++++ src/caffe/test/test_lrn_map_layer.cpp | 160 +++++++ src/caffe/test/test_power_layer.cpp | 429 ++++++++++++++++++ 3 files changed, 711 insertions(+) create mode 100644 src/caffe/test/test_eltwise_product_layer.cpp create mode 100644 src/caffe/test/test_lrn_map_layer.cpp create mode 100644 src/caffe/test/test_power_layer.cpp diff --git a/src/caffe/test/test_eltwise_product_layer.cpp b/src/caffe/test/test_eltwise_product_layer.cpp new file mode 100644 index 00000000000..8766f39da86 --- /dev/null +++ b/src/caffe/test/test_eltwise_product_layer.cpp @@ -0,0 +1,122 @@ +// Copyright 2014 BVLC and contributors. + +#include + +#include "cuda_runtime.h" +#include "gtest/gtest.h" +#include "caffe/blob.hpp" +#include "caffe/common.hpp" +#include "caffe/filler.hpp" +#include "caffe/vision_layers.hpp" +#include "caffe/test/test_gradient_check_util.hpp" + +#include "caffe/test/test_caffe_main.hpp" + +namespace caffe { + +extern cudaDeviceProp CAFFE_TEST_CUDA_PROP; + +template +class EltwiseProductLayerTest : public ::testing::Test { + protected: + EltwiseProductLayerTest() + : blob_bottom_a_(new Blob(2, 3, 4, 5)), + blob_bottom_b_(new Blob(2, 3, 4, 5)), + blob_bottom_c_(new Blob(2, 3, 4, 5)), + blob_top_(new Blob()) { + // fill the values + FillerParameter filler_param; + UniformFiller filler(filler_param); + filler.Fill(this->blob_bottom_a_); + filler.Fill(this->blob_bottom_b_); + filler.Fill(this->blob_bottom_c_); + blob_bottom_vec_.push_back(blob_bottom_a_); + blob_bottom_vec_.push_back(blob_bottom_b_); + blob_bottom_vec_.push_back(blob_bottom_c_); + blob_top_vec_.push_back(blob_top_); + } + virtual ~EltwiseProductLayerTest() { + delete blob_bottom_a_; + delete blob_bottom_b_; + delete blob_bottom_c_; + delete blob_top_; + } + Blob* const blob_bottom_a_; + Blob* const blob_bottom_b_; + Blob* const blob_bottom_c_; + Blob* const blob_top_; + vector*> blob_bottom_vec_; + vector*> blob_top_vec_; +}; + +typedef ::testing::Types Dtypes; +TYPED_TEST_CASE(EltwiseProductLayerTest, Dtypes); + +TYPED_TEST(EltwiseProductLayerTest, TestSetUp) { + LayerParameter layer_param; + shared_ptr > layer( + new EltwiseProductLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + EXPECT_EQ(this->blob_top_->num(), 2); + EXPECT_EQ(this->blob_top_->channels(), 3); + EXPECT_EQ(this->blob_top_->height(), 4); + EXPECT_EQ(this->blob_top_->width(), 5); +} + +TYPED_TEST(EltwiseProductLayerTest, TestCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + shared_ptr > layer( + new EltwiseProductLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + const TypeParam* data = this->blob_top_->cpu_data(); + const int count = this->blob_top_->count(); + const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); + const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); + const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_EQ(data[i], in_data_a[i] * in_data_b[i] * in_data_c[i]); + } +} + +TYPED_TEST(EltwiseProductLayerTest, TestGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + shared_ptr > layer( + new EltwiseProductLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + const TypeParam* data = this->blob_top_->cpu_data(); + const int count = this->blob_top_->count(); + const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); + const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); + const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_EQ(data[i], in_data_a[i] * in_data_b[i] * in_data_c[i]); + } +} + +TYPED_TEST(EltwiseProductLayerTest, TestCPUGradient) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + EltwiseProductLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(EltwiseProductLayerTest, TestGPUGradient) { + if (sizeof(TypeParam) == 4 || CAFFE_TEST_CUDA_PROP.major >= 2) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + EltwiseProductLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); + } else { + LOG(ERROR) << "Skipping test due to old architecture."; + } +} + +} // namespace caffe diff --git a/src/caffe/test/test_lrn_map_layer.cpp b/src/caffe/test/test_lrn_map_layer.cpp new file mode 100644 index 00000000000..93377cc1beb --- /dev/null +++ b/src/caffe/test/test_lrn_map_layer.cpp @@ -0,0 +1,160 @@ +// Copyright 2014 BVLC and contributors. + +#include +#include + +#include "cuda_runtime.h" +#include "gtest/gtest.h" + +#include "caffe/blob.hpp" +#include "caffe/common.hpp" +#include "caffe/filler.hpp" +#include "caffe/vision_layers.hpp" +#include "caffe/test/test_gradient_check_util.hpp" + +#include "caffe/test/test_caffe_main.hpp" + +using std::min; +using std::max; + +namespace caffe { + +extern cudaDeviceProp CAFFE_TEST_CUDA_PROP; + +template +class LRNMapLayerTest : public ::testing::Test { + protected: + LRNMapLayerTest() + : blob_bottom_(new Blob()), + blob_top_(new Blob()) {} + virtual void SetUp() { + Caffe::set_random_seed(1701); + blob_bottom_->Reshape(2, 2, 7, 7); + // fill the values + FillerParameter filler_param; + GaussianFiller filler(filler_param); + filler.Fill(this->blob_bottom_); + blob_bottom_vec_.push_back(blob_bottom_); + blob_top_vec_.push_back(blob_top_); + epsilon_ = 1e-5; + } + virtual ~LRNMapLayerTest() { delete blob_bottom_; delete blob_top_; } + void ReferenceLRNMapForward(const Blob& blob_bottom, + const LayerParameter& layer_param, Blob* blob_top); + Blob* const blob_bottom_; + Blob* const blob_top_; + vector*> blob_bottom_vec_; + vector*> blob_top_vec_; + Dtype epsilon_; +}; + +template +void LRNMapLayerTest::ReferenceLRNMapForward( + const Blob& blob_bottom, const LayerParameter& layer_param, + Blob* blob_top) { + blob_top->Reshape(blob_bottom.num(), blob_bottom.channels(), + blob_bottom.height(), blob_bottom.width()); + const Dtype* bottom_data = blob_bottom.cpu_data(); + Dtype* top_data = blob_top->mutable_cpu_data(); + const Dtype alpha = layer_param.lrn_map_param().alpha(); + const Dtype beta = layer_param.lrn_map_param().beta(); + const int size = layer_param.lrn_map_param().local_size(); + for (int n = 0; n < blob_bottom.num(); ++n) { + for (int c = 0; c < blob_bottom.channels(); ++c) { + for (int h = 0; h < blob_bottom.height(); ++h) { + int h_start = h - (size - 1) / 2; + int h_end = min(h_start + size, blob_bottom.height()); + h_start = max(h_start, 0); + for (int w = 0; w < blob_bottom.width(); ++w) { + Dtype scale = 1.; + int w_start = w - (size - 1) / 2; + int w_end = min(w_start + size, blob_bottom.width()); + w_start = max(w_start, 0); + for (int nh = h_start; nh < h_end; ++nh) { + for (int nw = w_start; nw < w_end; ++nw) { + Dtype value = blob_bottom.data_at(n, c, nh, nw); + scale += value * value * alpha / (size * size); + } + } + *(top_data + blob_top->offset(n, c, h, w)) = + blob_bottom.data_at(n, c, h, w) / pow(scale, beta); + } + } + } + } +} + +typedef ::testing::Types Dtypes; +TYPED_TEST_CASE(LRNMapLayerTest, Dtypes); + +TYPED_TEST(LRNMapLayerTest, TestSetup) { + LayerParameter layer_param; + LRNMapLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + EXPECT_EQ(this->blob_top_->num(), 2); + EXPECT_EQ(this->blob_top_->channels(), 2); + EXPECT_EQ(this->blob_top_->height(), 7); + EXPECT_EQ(this->blob_top_->width(), 7); +} + +TYPED_TEST(LRNMapLayerTest, TestCPUForward) { + LayerParameter layer_param; + LRNMapLayer layer(layer_param); + Caffe::set_mode(Caffe::CPU); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + Blob top_reference; + this->ReferenceLRNMapForward(*(this->blob_bottom_), layer_param, + &top_reference); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_NEAR(this->blob_top_->cpu_data()[i], top_reference.cpu_data()[i], + this->epsilon_); + } +} + +TYPED_TEST(LRNMapLayerTest, TestGPUForward) { + LayerParameter layer_param; + LRNMapLayer layer(layer_param); + Caffe::set_mode(Caffe::GPU); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + Blob top_reference; + this->ReferenceLRNMapForward(*(this->blob_bottom_), layer_param, + &top_reference); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_NEAR(this->blob_top_->cpu_data()[i], top_reference.cpu_data()[i], + this->epsilon_); + } +} + +TYPED_TEST(LRNMapLayerTest, TestCPUGradient) { + LayerParameter layer_param; + LRNMapLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + Caffe::set_mode(Caffe::CPU); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + for (int i = 0; i < this->blob_top_->count(); ++i) { + this->blob_top_->mutable_cpu_diff()[i] = 1.; + } + layer.Backward(this->blob_top_vec_, true, &(this->blob_bottom_vec_)); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(LRNMapLayerTest, TestGPUGradient) { + LayerParameter layer_param; + LRNMapLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + Caffe::set_mode(Caffe::GPU); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + for (int i = 0; i < this->blob_top_->count(); ++i) { + this->blob_top_->mutable_cpu_diff()[i] = 1.; + } + layer.Backward(this->blob_top_vec_, true, &(this->blob_bottom_vec_)); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +} // namespace caffe diff --git a/src/caffe/test/test_power_layer.cpp b/src/caffe/test/test_power_layer.cpp new file mode 100644 index 00000000000..40734b65fbf --- /dev/null +++ b/src/caffe/test/test_power_layer.cpp @@ -0,0 +1,429 @@ +// Copyright 2014 BVLC and contributors. + +#include + +#include "cuda_runtime.h" +#include "gtest/gtest.h" + +#include "caffe/blob.hpp" +#include "caffe/common.hpp" +#include "caffe/filler.hpp" +#include "caffe/vision_layers.hpp" +#include "caffe/test/test_gradient_check_util.hpp" + +#include "caffe/test/test_caffe_main.hpp" + +namespace caffe { + +extern cudaDeviceProp CAFFE_TEST_CUDA_PROP; + +template +class PowerLayerTest : public ::testing::Test { + protected: + PowerLayerTest() + : blob_bottom_(new Blob(2, 3, 4, 5)), + blob_top_(new Blob()) { + // fill the values + FillerParameter filler_param; + GaussianFiller filler(filler_param); + filler.Fill(this->blob_bottom_); + blob_bottom_vec_.push_back(blob_bottom_); + blob_top_vec_.push_back(blob_top_); + } + virtual ~PowerLayerTest() { delete blob_bottom_; delete blob_top_; } + Blob* const blob_bottom_; + Blob* const blob_top_; + vector*> blob_bottom_vec_; + vector*> blob_top_vec_; +}; + +typedef ::testing::Types Dtypes; +TYPED_TEST_CASE(PowerLayerTest, Dtypes); + +TYPED_TEST(PowerLayerTest, TestPowerCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 0.37; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + // Now, check values + const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); + const TypeParam* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + TypeParam expected_value = pow(shift + scale * bottom_data[i], power); + if (isnan(expected_value)) { + EXPECT_TRUE(isnan(top_data[i])); + } else { + TypeParam precision = expected_value * 0.0001; + precision *= (precision < 0) ? -1 : 1; + EXPECT_NEAR(expected_value, top_data[i], precision); + } + } +} + +TYPED_TEST(PowerLayerTest, TestPowerGradientCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 0.37; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + // Avoid NaNs by forcing (shift + scale * x) >= 0 + TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); + TypeParam min_value = -shift / scale; + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + if (bottom_data[i] < min_value) { + bottom_data[i] = min_value + (min_value - bottom_data[i]); + } + } + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerGradientShiftZeroCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 0.37; + TypeParam scale = 0.83; + TypeParam shift = 0.0; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + // Flip negative values in bottom vector as x < 0 -> x^0.37 = nan + TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + bottom_data[i] *= (bottom_data[i] < 0) ? -1 : 1; + } + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerZeroCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 0.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + // Now, check values + const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); + const TypeParam* top_data = this->blob_top_->cpu_data(); + TypeParam expected_value = TypeParam(1); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_EQ(expected_value, top_data[i]); + } +} + +TYPED_TEST(PowerLayerTest, TestPowerZeroGradientCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 0.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerOneCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 1.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + // Now, check values + const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); + const TypeParam* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + TypeParam expected_value = shift + scale * bottom_data[i]; + EXPECT_NEAR(expected_value, top_data[i], 0.001); + } +} + +TYPED_TEST(PowerLayerTest, TestPowerOneGradientCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 1.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerTwoCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 2.0; + TypeParam scale = 0.34; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + // Now, check values + const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); + const TypeParam* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + TypeParam expected_value = pow(shift + scale * bottom_data[i], 2); + EXPECT_NEAR(expected_value, top_data[i], 0.001); + } +} + +TYPED_TEST(PowerLayerTest, TestPowerTwoGradientCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 2.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerTwoScaleHalfGradientCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + TypeParam power = 2.0; + TypeParam scale = 0.5; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 0.37; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + // Now, check values + const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); + const TypeParam* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + TypeParam expected_value = pow(shift + scale * bottom_data[i], power); + if (isnan(expected_value)) { + EXPECT_TRUE(isnan(top_data[i])); + } else { + TypeParam precision = expected_value * 0.0001; + precision *= (precision < 0) ? -1 : 1; + EXPECT_NEAR(expected_value, top_data[i], precision); + } + } +} + +TYPED_TEST(PowerLayerTest, TestPowerGradientGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 0.37; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + // Avoid NaNs by forcing (shift + scale * x) >= 0 + TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); + TypeParam min_value = -shift / scale; + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + if (bottom_data[i] < min_value) { + bottom_data[i] = min_value + (min_value - bottom_data[i]); + } + } + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerGradientShiftZeroGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 0.37; + TypeParam scale = 0.83; + TypeParam shift = 0.0; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + // Flip negative values in bottom vector as x < 0 -> x^0.37 = nan + TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + bottom_data[i] *= (bottom_data[i] < 0) ? -1 : 1; + } + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerZeroGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 0.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + // Now, check values + const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); + const TypeParam* top_data = this->blob_top_->cpu_data(); + TypeParam expected_value = TypeParam(1); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_EQ(expected_value, top_data[i]); + } +} + +TYPED_TEST(PowerLayerTest, TestPowerZeroGradientGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 0.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerOneGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 1.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + // Now, check values + const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); + const TypeParam* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + TypeParam expected_value = shift + scale * bottom_data[i]; + EXPECT_NEAR(expected_value, top_data[i], 0.001); + } +} + +TYPED_TEST(PowerLayerTest, TestPowerOneGradientGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 1.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerTwoGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 2.0; + TypeParam scale = 0.34; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + // Now, check values + const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); + const TypeParam* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + TypeParam expected_value = pow(shift + scale * bottom_data[i], 2); + EXPECT_NEAR(expected_value, top_data[i], 0.001); + } +} + +TYPED_TEST(PowerLayerTest, TestPowerTwoGradientGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 2.0; + TypeParam scale = 0.83; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(PowerLayerTest, TestPowerTwoScaleHalfGradientGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + TypeParam power = 2.0; + TypeParam scale = 0.5; + TypeParam shift = -2.4; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +} // namespace caffe From 6e92f47ef250b5469c34576724e479ce2d98552d Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Fri, 28 Mar 2014 16:29:21 -0700 Subject: [PATCH 03/18] use average pool instead of conv --- include/caffe/vision_layers.hpp | 10 +++++-- src/caffe/layers/lrn_map_layer.cpp | 45 ++++++++++++------------------ src/caffe/layers/lrn_map_layer.cu | 8 +++--- src/caffe/proto/caffe.proto | 1 + 4 files changed, 30 insertions(+), 34 deletions(-) diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index 1d963dd5dd2..e82cd82b9c9 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -530,6 +530,9 @@ class LRNLayer : public Layer { int width_; }; +template +class PoolingLayer; + template class LRNMapLayer : public Layer { public: @@ -553,9 +556,9 @@ class LRNMapLayer : public Layer { Blob square_output_; vector*> square_bottom_vec_; vector*> square_top_vec_; - shared_ptr > conv_layer_; - Blob conv_output_; - vector*> conv_top_vec_; + shared_ptr > pool_layer_; + Blob pool_output_; + vector*> pool_top_vec_; shared_ptr > power_layer_; Blob power_output_; vector*> power_top_vec_; @@ -604,6 +607,7 @@ class PoolingLayer : public Layer { int kernel_size_; int stride_; + int pad_; int channels_; int height_; int width_; diff --git a/src/caffe/layers/lrn_map_layer.cpp b/src/caffe/layers/lrn_map_layer.cpp index 786168c76c1..9370bff26ee 100644 --- a/src/caffe/layers/lrn_map_layer.cpp +++ b/src/caffe/layers/lrn_map_layer.cpp @@ -38,28 +38,19 @@ void LRNMapLayer::SetUp(const vector*>& bottom, CHECK_EQ(square_output_.channels(), channels); CHECK_EQ(square_output_.height(), height); CHECK_EQ(square_output_.width(), width); - // Set up conv layer to have N filters with N groups, all of which are 1's. - // (With #filters == #groups, each filter looks at exactly 1 input channel, - // which is what we want for this layer type.) - // Output of conv layer gives us the neighborhood response. - conv_top_vec_.clear(); - conv_top_vec_.push_back(&conv_output_); - LayerParameter conv_param; - conv_param.mutable_convolution_param()->set_pad(pre_pad); - conv_param.mutable_convolution_param()->set_kernel_size(size_); - conv_param.mutable_convolution_param()->set_num_output(channels); - conv_param.mutable_convolution_param()->set_group(channels); - conv_param.mutable_convolution_param()->set_bias_term(false); - conv_layer_.reset(new ConvolutionLayer(conv_param)); - conv_layer_->SetUp(square_top_vec_, &conv_top_vec_); - CHECK_EQ(conv_output_.num(), num); - CHECK_EQ(conv_output_.channels(), channels); - CHECK_EQ(conv_output_.height(), height); - CHECK_EQ(conv_output_.width(), width); - FillerParameter one_filler_param; - one_filler_param.set_value(1); - ConstantFiller one_filler(one_filler_param); - one_filler.Fill(conv_layer_->blobs()[0].get()); + // Output of pool layer gives us the neighborhood response. + pool_top_vec_.clear(); + pool_top_vec_.push_back(&pool_output_); + LayerParameter pool_param; + pool_param.mutable_pooling_param()->set_pool(PoolingParameter_PoolMethod_AVE); + pool_param.mutable_pooling_param()->set_pad(pre_pad); + pool_param.mutable_pooling_param()->set_kernel_size(size_); + pool_layer_.reset(new PoolingLayer(pool_param)); + pool_layer_->SetUp(square_top_vec_, &pool_top_vec_); + CHECK_EQ(pool_output_.num(), num); + CHECK_EQ(pool_output_.channels(), channels); + CHECK_EQ(pool_output_.height(), height); + CHECK_EQ(pool_output_.width(), width); // Set up power layer to compute (1 + alpha/N^2 s)^-beta, where s is the sum // of a squared neighborhood (as output by the conv layer). power_top_vec_.clear(); @@ -69,7 +60,7 @@ void LRNMapLayer::SetUp(const vector*>& bottom, power_param.mutable_power_param()->set_scale(alpha / (size_ * size_)); power_param.mutable_power_param()->set_shift(Dtype(1)); power_layer_.reset(new PowerLayer(power_param)); - power_layer_->SetUp(conv_top_vec_, &power_top_vec_); + power_layer_->SetUp(pool_top_vec_, &power_top_vec_); CHECK_EQ(power_output_.num(), num); CHECK_EQ(power_output_.channels(), channels); CHECK_EQ(power_output_.height(), height); @@ -97,8 +88,8 @@ Dtype LRNMapLayer::Forward_cpu(const vector*>& bottom, Dtype* square_bottom_data = square_input_.mutable_cpu_data(); caffe_copy(count, bottom_data, square_bottom_data); square_layer_->Forward(square_bottom_vec_, &square_top_vec_); - conv_layer_->Forward(square_top_vec_, &conv_top_vec_); - power_layer_->Forward(conv_top_vec_, &power_top_vec_); + pool_layer_->Forward(square_top_vec_, &pool_top_vec_); + power_layer_->Forward(pool_top_vec_, &power_top_vec_); product_layer_->Forward(product_bottom_vec_, &product_top_vec_); return Dtype(0.); } @@ -108,8 +99,8 @@ void LRNMapLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { product_layer_->Backward(product_top_vec_, true, &product_bottom_vec_); - power_layer_->Backward(power_top_vec_, true, &conv_top_vec_); - conv_layer_->Backward(conv_top_vec_, true, &square_top_vec_); + power_layer_->Backward(power_top_vec_, true, &pool_top_vec_); + pool_layer_->Backward(pool_top_vec_, true, &square_top_vec_); square_layer_->Backward(square_top_vec_, true, &square_bottom_vec_); const int count = (*bottom)[0]->count(); const Dtype* scale_diff = square_input_.cpu_diff(); diff --git a/src/caffe/layers/lrn_map_layer.cu b/src/caffe/layers/lrn_map_layer.cu index 0aba8e1c8d4..727728e0d4b 100644 --- a/src/caffe/layers/lrn_map_layer.cu +++ b/src/caffe/layers/lrn_map_layer.cu @@ -16,8 +16,8 @@ Dtype LRNMapLayer::Forward_gpu(const vector*>& bottom, Dtype* square_bottom_data = square_input_.mutable_gpu_data(); caffe_gpu_copy(count, bottom_data, square_bottom_data); square_layer_->Forward(square_bottom_vec_, &square_top_vec_); - conv_layer_->Forward(square_top_vec_, &conv_top_vec_); - power_layer_->Forward(conv_top_vec_, &power_top_vec_); + pool_layer_->Forward(square_top_vec_, &pool_top_vec_); + power_layer_->Forward(pool_top_vec_, &power_top_vec_); product_layer_->Forward(product_bottom_vec_, &product_top_vec_); return Dtype(0.); } @@ -27,8 +27,8 @@ void LRNMapLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { product_layer_->Backward(product_top_vec_, true, &product_bottom_vec_); - power_layer_->Backward(power_top_vec_, true, &conv_top_vec_); - conv_layer_->Backward(conv_top_vec_, true, &square_top_vec_); + power_layer_->Backward(power_top_vec_, true, &pool_top_vec_); + pool_layer_->Backward(pool_top_vec_, true, &square_top_vec_); square_layer_->Backward(square_top_vec_, true, &square_bottom_vec_); const int count = (*bottom)[0]->count(); const Dtype* scale_diff = square_input_.gpu_diff(); diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index bda35c1e34d..e0011483f6b 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -290,6 +290,7 @@ message PoolingParameter { optional PoolMethod pool = 1 [default = MAX]; // The pooling method optional uint32 kernel_size = 2; // The kernel size optional uint32 stride = 3 [default = 1]; // The stride + optional uint32 pad = 4 [default = 0]; // The padding size } // Message that stores parameters used by PowerLayer From bd756fe1f805b0a731707b7103d9ac0b203a1479 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Fri, 28 Mar 2014 20:18:06 -0700 Subject: [PATCH 04/18] add padding for average pooling --- src/caffe/layers/pooling_layer.cpp | 46 +++++++---- src/caffe/layers/pooling_layer.cu | 91 +++++++++++---------- src/caffe/proto/caffe.proto | 3 +- src/caffe/test/test_pooling_layer.cpp | 111 ++++++++++++++++++++++++++ 4 files changed, 193 insertions(+), 58 deletions(-) diff --git a/src/caffe/layers/pooling_layer.cpp b/src/caffe/layers/pooling_layer.cpp index a186741232f..7e880a27b69 100644 --- a/src/caffe/layers/pooling_layer.cpp +++ b/src/caffe/layers/pooling_layer.cpp @@ -20,13 +20,19 @@ void PoolingLayer::SetUp(const vector*>& bottom, CHECK_EQ(top->size(), 1) << "PoolingLayer takes a single blob as output."; kernel_size_ = this->layer_param_.pooling_param().kernel_size(); stride_ = this->layer_param_.pooling_param().stride(); + pad_ = this->layer_param_.pooling_param().pad(); + if (pad_ != 0) { + CHECK_EQ(this->layer_param_.pooling_param().pool(), + PoolingParameter_PoolMethod_AVE) + << "Padding implemented only for average pooling."; + } channels_ = bottom[0]->channels(); height_ = bottom[0]->height(); width_ = bottom[0]->width(); - pooled_height_ = static_cast( - ceil(static_cast(height_ - kernel_size_) / stride_)) + 1; - pooled_width_ = static_cast( - ceil(static_cast(width_ - kernel_size_) / stride_)) + 1; + pooled_height_ = static_cast(ceil(static_cast( + height_ + 2 * pad_ - kernel_size_) / stride_)) + 1; + pooled_width_ = static_cast(ceil(static_cast( + width_ + 2 * pad_ - kernel_size_) / stride_)) + 1; (*top)[0]->Reshape(bottom[0]->num(), channels_, pooled_height_, pooled_width_); // If stochastic pooling, we will initialize the random index part. @@ -86,18 +92,22 @@ Dtype PoolingLayer::Forward_cpu(const vector*>& bottom, for (int c = 0; c < channels_; ++c) { for (int ph = 0; ph < pooled_height_; ++ph) { for (int pw = 0; pw < pooled_width_; ++pw) { - int hstart = ph * stride_; - int wstart = pw * stride_; - int hend = min(hstart + kernel_size_, height_); - int wend = min(wstart + kernel_size_, width_); + int hstart = ph * stride_ - pad_; + int wstart = pw * stride_ - pad_; + int hend = min(hstart + kernel_size_, height_ + pad_); + int wend = min(wstart + kernel_size_, width_ + pad_); + int pool_size = (hend - hstart) * (wend - wstart); + hstart = max(hstart, 0); + wstart = max(wstart, 0); + hend = min(hend, height_); + wend = min(wend, width_); for (int h = hstart; h < hend; ++h) { for (int w = wstart; w < wend; ++w) { top_data[ph * pooled_width_ + pw] += bottom_data[h * width_ + w]; } } - top_data[ph * pooled_width_ + pw] /= - (hend - hstart) * (wend - wstart); + top_data[ph * pooled_width_ + pw] /= pool_size; } } // compute offset @@ -163,15 +173,19 @@ void PoolingLayer::Backward_cpu(const vector*>& top, for (int c = 0; c < channels_; ++c) { for (int ph = 0; ph < pooled_height_; ++ph) { for (int pw = 0; pw < pooled_width_; ++pw) { - int hstart = ph * stride_; - int wstart = pw * stride_; - int hend = min(hstart + kernel_size_, height_); - int wend = min(wstart + kernel_size_, width_); - int poolsize = (hend - hstart) * (wend - wstart); + int hstart = ph * stride_ - pad_; + int wstart = pw * stride_ - pad_; + int hend = min(hstart + kernel_size_, height_ + pad_); + int wend = min(wstart + kernel_size_, width_ + pad_); + int pool_size = (hend - hstart) * (wend - wstart); + hstart = max(hstart, 0); + wstart = max(wstart, 0); + hend = min(hend, height_); + wend = min(wend, width_); for (int h = hstart; h < hend; ++h) { for (int w = wstart; w < wend; ++w) { bottom_diff[h * width_ + w] += - top_diff[ph * pooled_width_ + pw] / poolsize; + top_diff[ph * pooled_width_ + pw] / pool_size; } } } diff --git a/src/caffe/layers/pooling_layer.cu b/src/caffe/layers/pooling_layer.cu index 7adf348be34..74150df2493 100644 --- a/src/caffe/layers/pooling_layer.cu +++ b/src/caffe/layers/pooling_layer.cu @@ -16,13 +16,13 @@ namespace caffe { template __global__ void MaxPoolForward(const int nthreads, const Dtype* bottom_data, const int num, const int channels, const int height, - const int width, const int pooled_height_, const int pooled_width, + const int width, const int pooled_height, const int pooled_width, const int kernel_size, const int stride, Dtype* top_data) { CUDA_KERNEL_LOOP(index, nthreads) { int pw = index % pooled_width; - int ph = (index / pooled_width) % pooled_height_; - int c = (index / pooled_width / pooled_height_) % channels; - int n = index / pooled_width / pooled_height_ / channels; + int ph = (index / pooled_width) % pooled_height; + int c = (index / pooled_width / pooled_height) % channels; + int n = index / pooled_width / pooled_height / channels; int hstart = ph * stride; int hend = min(hstart + kernel_size, height); int wstart = pw * stride; @@ -41,17 +41,22 @@ __global__ void MaxPoolForward(const int nthreads, const Dtype* bottom_data, template __global__ void AvePoolForward(const int nthreads, const Dtype* bottom_data, const int num, const int channels, const int height, - const int width, const int pooled_height_, const int pooled_width, - const int kernel_size, const int stride, Dtype* top_data) { + const int width, const int pooled_height, const int pooled_width, + const int kernel_size, const int stride, const int pad, Dtype* top_data) { CUDA_KERNEL_LOOP(index, nthreads) { int pw = index % pooled_width; - int ph = (index / pooled_width) % pooled_height_; - int c = (index / pooled_width / pooled_height_) % channels; - int n = index / pooled_width / pooled_height_ / channels; - int hstart = ph * stride; - int hend = min(hstart + kernel_size, height); - int wstart = pw * stride; - int wend = min(wstart + kernel_size, width); + int ph = (index / pooled_width) % pooled_height; + int c = (index / pooled_width / pooled_height) % channels; + int n = index / pooled_width / pooled_height / channels; + int hstart = ph * stride - pad; + int wstart = pw * stride - pad; + int hend = min(hstart + kernel_size, height + pad); + int wend = min(wstart + kernel_size, width + pad); + int pool_size = (hend - hstart) * (wend - wstart); + hstart = max(hstart, 0); + wstart = max(wstart, 0); + hend = min(hend, height); + wend = min(wend, width); Dtype aveval = 0; bottom_data += (n * channels + c) * height * width; for (int h = hstart; h < hend; ++h) { @@ -59,7 +64,7 @@ __global__ void AvePoolForward(const int nthreads, const Dtype* bottom_data, aveval += bottom_data[h * width + w]; } } - top_data[index] = aveval / (hend - hstart) / (wend - wstart); + top_data[index] = aveval / pool_size; } } @@ -67,13 +72,13 @@ template __global__ void StoPoolForwardTrain(const int nthreads, const Dtype* bottom_data, const int num, const int channels, const int height, - const int width, const int pooled_height_, const int pooled_width, + const int width, const int pooled_height, const int pooled_width, const int kernel_size, const int stride, float* rand_idx, Dtype* top_data) { CUDA_KERNEL_LOOP(index, nthreads) { int pw = index % pooled_width; - int ph = (index / pooled_width) % pooled_height_; - int c = (index / pooled_width / pooled_height_) % channels; - int n = index / pooled_width / pooled_height_ / channels; + int ph = (index / pooled_width) % pooled_height; + int c = (index / pooled_width / pooled_height) % channels; + int n = index / pooled_width / pooled_height / channels; int hstart = ph * stride; int hend = min(hstart + kernel_size, height); int wstart = pw * stride; @@ -107,13 +112,13 @@ template __global__ void StoPoolForwardTest(const int nthreads, const Dtype* bottom_data, const int num, const int channels, const int height, - const int width, const int pooled_height_, const int pooled_width, + const int width, const int pooled_height, const int pooled_width, const int kernel_size, const int stride, Dtype* top_data) { CUDA_KERNEL_LOOP(index, nthreads) { int pw = index % pooled_width; - int ph = (index / pooled_width) % pooled_height_; - int c = (index / pooled_width / pooled_height_) % channels; - int n = index / pooled_width / pooled_height_ / channels; + int ph = (index / pooled_width) % pooled_height; + int c = (index / pooled_width / pooled_height) % channels; + int n = index / pooled_width / pooled_height / channels; int hstart = ph * stride; int hend = min(hstart + kernel_size, height); int wstart = pw * stride; @@ -153,7 +158,7 @@ Dtype PoolingLayer::Forward_gpu(const vector*>& bottom, AvePoolForward<<>>( count, bottom_data, bottom[0]->num(), channels_, height_, width_, pooled_height_, pooled_width_, kernel_size_, stride_, - top_data); + pad_, top_data); break; case PoolingParameter_PoolMethod_STOCHASTIC: if (Caffe::phase() == Caffe::TRAIN) { @@ -186,7 +191,7 @@ template __global__ void MaxPoolBackward(const int nthreads, const Dtype* bottom_data, const Dtype* top_data, const Dtype* top_diff, const int num, const int channels, const int height, - const int width, const int pooled_height_, const int pooled_width, + const int width, const int pooled_height, const int pooled_width, const int kernel_size, const int stride, Dtype* bottom_diff) { CUDA_KERNEL_LOOP(index, nthreads) { // find out the local index @@ -196,14 +201,14 @@ __global__ void MaxPoolBackward(const int nthreads, const Dtype* bottom_data, int c = (index / width / height) % channels; int n = index / width / height / channels; int phstart = (h < kernel_size) ? 0 : (h - kernel_size) / stride + 1; - int phend = min(h / stride + 1, pooled_height_); + int phend = min(h / stride + 1, pooled_height); int pwstart = (w < kernel_size) ? 0 : (w - kernel_size) / stride + 1; int pwend = min(w / stride + 1, pooled_width); Dtype gradient = 0; Dtype bottom_datum = bottom_data[((n * channels + c) * height + h) * width + w]; - top_data += (n * channels + c) * pooled_height_ * pooled_width; - top_diff += (n * channels + c) * pooled_height_ * pooled_width; + top_data += (n * channels + c) * pooled_height * pooled_width; + top_diff += (n * channels + c) * pooled_height * pooled_width; for (int ph = phstart; ph < phend; ++ph) { for (int pw = pwstart; pw < pwend; ++pw) { gradient += top_diff[ph * pooled_width + pw] * @@ -218,27 +223,31 @@ __global__ void MaxPoolBackward(const int nthreads, const Dtype* bottom_data, template __global__ void AvePoolBackward(const int nthreads, const Dtype* top_diff, const int num, const int channels, const int height, - const int width, const int pooled_height_, const int pooled_width, - const int kernel_size, const int stride, Dtype* bottom_diff) { + const int width, const int pooled_height, const int pooled_width, + const int kernel_size, const int stride, const int pad, + Dtype* bottom_diff) { CUDA_KERNEL_LOOP(index, nthreads) { // find out the local index // find out the local offset - int w = index % width; - int h = (index / width) % height; + int w = index % width + pad; + int h = (index / width) % height + pad; int c = (index / width / height) % channels; int n = index / width / height / channels; int phstart = (h < kernel_size) ? 0 : (h - kernel_size) / stride + 1; - int phend = min(h / stride + 1, pooled_height_); + int phend = min(h / stride + 1, pooled_height); int pwstart = (w < kernel_size) ? 0 : (w - kernel_size) / stride + 1; int pwend = min(w / stride + 1, pooled_width); Dtype gradient = 0; - top_diff += (n * channels + c) * pooled_height_ * pooled_width; + top_diff += (n * channels + c) * pooled_height * pooled_width; for (int ph = phstart; ph < phend; ++ph) { for (int pw = pwstart; pw < pwend; ++pw) { // figure out the pooling size - int poolsize = (min(ph * stride + kernel_size, height) - ph * stride) * - (min(pw * stride + kernel_size, width) - pw * stride); - gradient += top_diff[ph * pooled_width + pw] / poolsize; + int hstart = ph * stride - pad; + int wstart = pw * stride - pad; + int hend = min(hstart + kernel_size, height + pad); + int wend = min(wstart + kernel_size, width + pad); + int pool_size = (hend - hstart) * (wend - wstart); + gradient += top_diff[ph * pooled_width + pw] / pool_size; } } bottom_diff[index] = gradient; @@ -250,7 +259,7 @@ template __global__ void StoPoolBackward(const int nthreads, const float* rand_idx, const Dtype* top_diff, const int num, const int channels, const int height, - const int width, const int pooled_height_, const int pooled_width, + const int width, const int pooled_height, const int pooled_width, const int kernel_size, const int stride, Dtype* bottom_diff) { CUDA_KERNEL_LOOP(index, nthreads) { // find out the local index @@ -260,12 +269,12 @@ __global__ void StoPoolBackward(const int nthreads, int c = (index / width / height) % channels; int n = index / width / height / channels; int phstart = (h < kernel_size) ? 0 : (h - kernel_size) / stride + 1; - int phend = min(h / stride + 1, pooled_height_); + int phend = min(h / stride + 1, pooled_height); int pwstart = (w < kernel_size) ? 0 : (w - kernel_size) / stride + 1; int pwend = min(w / stride + 1, pooled_width); Dtype gradient = 0; - rand_idx += (n * channels + c) * pooled_height_ * pooled_width; - top_diff += (n * channels + c) * pooled_height_ * pooled_width; + rand_idx += (n * channels + c) * pooled_height * pooled_width; + top_diff += (n * channels + c) * pooled_height * pooled_width; for (int ph = phstart; ph < phend; ++ph) { for (int pw = pwstart; pw < pwend; ++pw) { gradient += top_diff[ph * pooled_width + pw] * @@ -299,7 +308,7 @@ void PoolingLayer::Backward_gpu(const vector*>& top, AvePoolBackward<<>>( count, top_diff, top[0]->num(), channels_, height_, width_, pooled_height_, pooled_width_, kernel_size_, stride_, - bottom_diff); + pad_, bottom_diff); break; case PoolingParameter_PoolMethod_STOCHASTIC: // NOLINT_NEXT_LINE(whitespace/operators) diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index e0011483f6b..54f2743ee4f 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -290,7 +290,8 @@ message PoolingParameter { optional PoolMethod pool = 1 [default = MAX]; // The pooling method optional uint32 kernel_size = 2; // The kernel size optional uint32 stride = 3 [default = 1]; // The stride - optional uint32 pad = 4 [default = 0]; // The padding size + // The padding size -- currently implemented only for average pooling. + optional uint32 pad = 4 [default = 0]; } // Message that stores parameters used by PowerLayer diff --git a/src/caffe/test/test_pooling_layer.cpp b/src/caffe/test/test_pooling_layer.cpp index d1246a098c8..a57110491ec 100644 --- a/src/caffe/test/test_pooling_layer.cpp +++ b/src/caffe/test/test_pooling_layer.cpp @@ -56,6 +56,21 @@ TYPED_TEST(PoolingLayerTest, TestSetup) { EXPECT_EQ(this->blob_top_->width(), 2); } +TYPED_TEST(PoolingLayerTest, TestSetupPadded) { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->set_kernel_size(3); + pooling_param->set_stride(2); + pooling_param->set_pad(1); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + PoolingLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + EXPECT_EQ(this->blob_top_->num(), this->blob_bottom_->num()); + EXPECT_EQ(this->blob_top_->channels(), this->blob_bottom_->channels()); + EXPECT_EQ(this->blob_top_->height(), 4); + EXPECT_EQ(this->blob_top_->width(), 3); +} + /* TYPED_TEST(PoolingLayerTest, PrintGPUBackward) { LayerParameter layer_param; @@ -111,6 +126,72 @@ TYPED_TEST(PoolingLayerTest, TestGPUGradientMax) { } +TYPED_TEST(PoolingLayerTest, TestCPUForwardAve) { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->set_kernel_size(3); + pooling_param->set_stride(1); + pooling_param->set_pad(1); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + Caffe::set_mode(Caffe::CPU); + this->blob_bottom_->Reshape(1, 1, 3, 3); + FillerParameter filler_param; + filler_param.set_value(TypeParam(2)); + ConstantFiller filler(filler_param); + filler.Fill(this->blob_bottom_); + PoolingLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + EXPECT_EQ(this->blob_top_->num(), 1); + EXPECT_EQ(this->blob_top_->channels(), 1); + EXPECT_EQ(this->blob_top_->height(), 3); + EXPECT_EQ(this->blob_top_->width(), 3); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + TypeParam epsilon = 1e-5; + EXPECT_NEAR(this->blob_top_->cpu_data()[0], 8.0 / 9, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[1], 4.0 / 3, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[2], 8.0 / 9, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[3], 4.0 / 3, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[4], 2.0, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[5], 4.0 / 3, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[6], 8.0 / 9, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[7], 4.0 / 3, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[8], 8.0 / 9, epsilon); +} + + +TYPED_TEST(PoolingLayerTest, TestGPUForwardAve) { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->set_kernel_size(3); + pooling_param->set_stride(1); + pooling_param->set_pad(1); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + Caffe::set_mode(Caffe::GPU); + this->blob_bottom_->Reshape(1, 1, 3, 3); + FillerParameter filler_param; + filler_param.set_value(TypeParam(2)); + ConstantFiller filler(filler_param); + filler.Fill(this->blob_bottom_); + PoolingLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + EXPECT_EQ(this->blob_top_->num(), 1); + EXPECT_EQ(this->blob_top_->channels(), 1); + EXPECT_EQ(this->blob_top_->height(), 3); + EXPECT_EQ(this->blob_top_->width(), 3); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + TypeParam epsilon = 1e-5; + EXPECT_NEAR(this->blob_top_->cpu_data()[0], 8.0 / 9, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[1], 4.0 / 3, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[2], 8.0 / 9, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[3], 4.0 / 3, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[4], 2.0, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[5], 4.0 / 3, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[6], 8.0 / 9, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[7], 4.0 / 3, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[8], 8.0 / 9, epsilon); +} + + TYPED_TEST(PoolingLayerTest, TestCPUGradientAve) { LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); @@ -139,4 +220,34 @@ TYPED_TEST(PoolingLayerTest, TestGPUGradientAve) { } +TYPED_TEST(PoolingLayerTest, TestCPUGradientAvePadded) { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->set_kernel_size(3); + pooling_param->set_stride(2); + pooling_param->set_pad(2); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + Caffe::set_mode(Caffe::CPU); + PoolingLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + + +TYPED_TEST(PoolingLayerTest, TestGPUGradientAvePadded) { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->set_kernel_size(3); + pooling_param->set_stride(2); + pooling_param->set_pad(2); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + Caffe::set_mode(Caffe::GPU); + PoolingLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + + } // namespace caffe From d047befce4256fa0b8914369910806fb0ad70647 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Fri, 28 Mar 2014 20:20:47 -0700 Subject: [PATCH 05/18] bug fix: average pooling already divides by N^2 --- src/caffe/layers/lrn_map_layer.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/caffe/layers/lrn_map_layer.cpp b/src/caffe/layers/lrn_map_layer.cpp index 9370bff26ee..3ae3e5f6abc 100644 --- a/src/caffe/layers/lrn_map_layer.cpp +++ b/src/caffe/layers/lrn_map_layer.cpp @@ -57,7 +57,7 @@ void LRNMapLayer::SetUp(const vector*>& bottom, power_top_vec_.push_back(&power_output_); LayerParameter power_param; power_param.mutable_power_param()->set_power(-beta); - power_param.mutable_power_param()->set_scale(alpha / (size_ * size_)); + power_param.mutable_power_param()->set_scale(alpha); power_param.mutable_power_param()->set_shift(Dtype(1)); power_layer_.reset(new PowerLayer(power_param)); power_layer_->SetUp(pool_top_vec_, &power_top_vec_); From 6c844cc712dbd052e07dfa65f589041dbfb7acf7 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Fri, 28 Mar 2014 21:25:15 -0700 Subject: [PATCH 06/18] use split layer in LRNMapLayer --- include/caffe/vision_layers.hpp | 7 ++++--- src/caffe/layers/lrn_map_layer.cpp | 28 ++++++++++++++-------------- src/caffe/layers/lrn_map_layer.cu | 14 ++++---------- 3 files changed, 22 insertions(+), 27 deletions(-) diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index e82cd82b9c9..22c0fd06457 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -530,8 +530,8 @@ class LRNLayer : public Layer { int width_; }; -template -class PoolingLayer; +template class PoolingLayer; +template class SplitLayer; template class LRNMapLayer : public Layer { @@ -551,6 +551,8 @@ class LRNMapLayer : public Layer { virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); + shared_ptr > split_layer_; + vector*> split_top_vec_; shared_ptr > square_layer_; Blob square_input_; Blob square_output_; @@ -565,7 +567,6 @@ class LRNMapLayer : public Layer { shared_ptr > product_layer_; Blob product_data_input_; vector*> product_bottom_vec_; - vector*> product_top_vec_; }; template diff --git a/src/caffe/layers/lrn_map_layer.cpp b/src/caffe/layers/lrn_map_layer.cpp index 3ae3e5f6abc..15119adf64e 100644 --- a/src/caffe/layers/lrn_map_layer.cpp +++ b/src/caffe/layers/lrn_map_layer.cpp @@ -24,6 +24,13 @@ void LRNMapLayer::SetUp(const vector*>& bottom, const Dtype pre_pad = (size_ - 1) / 2; const Dtype alpha = this->layer_param_.lrn_map_param().alpha(); const Dtype beta = this->layer_param_.lrn_map_param().beta(); + // Set up split layer to use inputs in the numerator and denominator. + split_top_vec_.clear(); + split_top_vec_.push_back(bottom[0]); + split_top_vec_.push_back(&square_input_); + LayerParameter split_param; + split_layer_.reset(new SplitLayer(split_param)); + split_layer_->SetUp(bottom, &split_top_vec_); // Set up square layer to square the inputs. square_input_.Reshape(num, channels, height, width); square_bottom_vec_.clear(); @@ -65,15 +72,14 @@ void LRNMapLayer::SetUp(const vector*>& bottom, CHECK_EQ(power_output_.channels(), channels); CHECK_EQ(power_output_.height(), height); CHECK_EQ(power_output_.width(), width); - // Set up a product layer to compute outputs by multiplying inputs by scale. + // Set up a product layer to compute outputs by multiplying inputs by the + // demoninator computed by the power layer. product_bottom_vec_.clear(); product_bottom_vec_.push_back(bottom[0]); product_bottom_vec_.push_back(&power_output_); - product_top_vec_.clear(); - product_top_vec_.push_back((*top)[0]); LayerParameter product_param; product_layer_.reset(new EltwiseProductLayer(product_param)); - product_layer_->SetUp(product_bottom_vec_, &product_top_vec_); + product_layer_->SetUp(product_bottom_vec_, top); CHECK_EQ((*top)[0]->num(), num); CHECK_EQ((*top)[0]->channels(), channels); CHECK_EQ((*top)[0]->height(), height); @@ -83,14 +89,11 @@ void LRNMapLayer::SetUp(const vector*>& bottom, template Dtype LRNMapLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { - const int count = bottom[0]->count(); - const Dtype* bottom_data = bottom[0]->cpu_data(); - Dtype* square_bottom_data = square_input_.mutable_cpu_data(); - caffe_copy(count, bottom_data, square_bottom_data); + split_layer_->Forward(bottom, &split_top_vec_); square_layer_->Forward(square_bottom_vec_, &square_top_vec_); pool_layer_->Forward(square_top_vec_, &pool_top_vec_); power_layer_->Forward(pool_top_vec_, &power_top_vec_); - product_layer_->Forward(product_bottom_vec_, &product_top_vec_); + product_layer_->Forward(product_bottom_vec_, top); return Dtype(0.); } @@ -98,14 +101,11 @@ template void LRNMapLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { - product_layer_->Backward(product_top_vec_, true, &product_bottom_vec_); + product_layer_->Backward(top, true, &product_bottom_vec_); power_layer_->Backward(power_top_vec_, true, &pool_top_vec_); pool_layer_->Backward(pool_top_vec_, true, &square_top_vec_); square_layer_->Backward(square_top_vec_, true, &square_bottom_vec_); - const int count = (*bottom)[0]->count(); - const Dtype* scale_diff = square_input_.cpu_diff(); - Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff(); - caffe_axpy(count, Dtype(1), scale_diff, bottom_diff); + split_layer_->Backward(split_top_vec_, true, bottom); } } diff --git a/src/caffe/layers/lrn_map_layer.cu b/src/caffe/layers/lrn_map_layer.cu index 727728e0d4b..b41adfc96be 100644 --- a/src/caffe/layers/lrn_map_layer.cu +++ b/src/caffe/layers/lrn_map_layer.cu @@ -11,14 +11,11 @@ namespace caffe { template Dtype LRNMapLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { - const int count = bottom[0]->count(); - const Dtype* bottom_data = bottom[0]->gpu_data(); - Dtype* square_bottom_data = square_input_.mutable_gpu_data(); - caffe_gpu_copy(count, bottom_data, square_bottom_data); + split_layer_->Forward(bottom, &split_top_vec_); square_layer_->Forward(square_bottom_vec_, &square_top_vec_); pool_layer_->Forward(square_top_vec_, &pool_top_vec_); power_layer_->Forward(pool_top_vec_, &power_top_vec_); - product_layer_->Forward(product_bottom_vec_, &product_top_vec_); + product_layer_->Forward(product_bottom_vec_, top); return Dtype(0.); } @@ -26,14 +23,11 @@ template void LRNMapLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { - product_layer_->Backward(product_top_vec_, true, &product_bottom_vec_); + product_layer_->Backward(top, true, &product_bottom_vec_); power_layer_->Backward(power_top_vec_, true, &pool_top_vec_); pool_layer_->Backward(pool_top_vec_, true, &square_top_vec_); square_layer_->Backward(square_top_vec_, true, &square_bottom_vec_); - const int count = (*bottom)[0]->count(); - const Dtype* scale_diff = square_input_.gpu_diff(); - Dtype* bottom_diff = (*bottom)[0]->mutable_gpu_diff(); - caffe_gpu_axpy(count, Dtype(1), scale_diff, bottom_diff); + split_layer_->Backward(split_top_vec_, true, bottom); } } From 2031663f903a7aaedb8df89039eaae97ec278399 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Fri, 28 Mar 2014 22:55:12 -0700 Subject: [PATCH 07/18] use bvlc copyright --- src/caffe/layers/eltwise_product_layer.cpp | 2 +- src/caffe/layers/eltwise_product_layer.cu | 2 +- src/caffe/layers/power_layer.cpp | 2 +- src/caffe/layers/power_layer.cu | 2 +- 4 files changed, 4 insertions(+), 4 deletions(-) diff --git a/src/caffe/layers/eltwise_product_layer.cpp b/src/caffe/layers/eltwise_product_layer.cpp index c663d10ffb8..d056ab1350b 100644 --- a/src/caffe/layers/eltwise_product_layer.cpp +++ b/src/caffe/layers/eltwise_product_layer.cpp @@ -1,4 +1,4 @@ -// Copyright 2013 Yangqing Jia +// Copyright 2014 BVLC and contributors. #include diff --git a/src/caffe/layers/eltwise_product_layer.cu b/src/caffe/layers/eltwise_product_layer.cu index d3227294eea..9c66033c20f 100644 --- a/src/caffe/layers/eltwise_product_layer.cu +++ b/src/caffe/layers/eltwise_product_layer.cu @@ -1,4 +1,4 @@ -// Copyright 2013 Yangqing Jia +// Copyright 2014 BVLC and contributors. #include diff --git a/src/caffe/layers/power_layer.cpp b/src/caffe/layers/power_layer.cpp index 3be59d36575..85c84423aa2 100644 --- a/src/caffe/layers/power_layer.cpp +++ b/src/caffe/layers/power_layer.cpp @@ -1,4 +1,4 @@ -// Copyright 2014 Jeff Donahue +// Copyright 2014 BVLC and contributors. #include #include diff --git a/src/caffe/layers/power_layer.cu b/src/caffe/layers/power_layer.cu index 9df44612cdf..9a25de72d36 100644 --- a/src/caffe/layers/power_layer.cu +++ b/src/caffe/layers/power_layer.cu @@ -1,4 +1,4 @@ -// Copyright 2014 Jeff Donahue +// Copyright 2014 BVLC and contributors. #include #include From 42b832c2e8fba20d68990bbc26733b3a2ea7a8d0 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 00:51:49 -0700 Subject: [PATCH 08/18] add cifar example using LRN_MAP (just like the cuda-convnet layers-18pct architecture) instead of LRN --- .../cifar10_full_lrn_map_solver.prototxt | 28 +++ .../cifar10_full_lrn_map_solver_lr1.prototxt | 28 +++ .../cifar10_full_lrn_map_solver_lr2.prototxt | 28 +++ .../cifar10_full_lrn_map_test.prototxt | 179 ++++++++++++++++++ .../cifar10_full_lrn_map_train.prototxt | 172 +++++++++++++++++ examples/cifar10/train_full_lrn_map.sh | 17 ++ 6 files changed, 452 insertions(+) create mode 100644 examples/cifar10/cifar10_full_lrn_map_solver.prototxt create mode 100644 examples/cifar10/cifar10_full_lrn_map_solver_lr1.prototxt create mode 100644 examples/cifar10/cifar10_full_lrn_map_solver_lr2.prototxt create mode 100644 examples/cifar10/cifar10_full_lrn_map_test.prototxt create mode 100644 examples/cifar10/cifar10_full_lrn_map_train.prototxt create mode 100755 examples/cifar10/train_full_lrn_map.sh diff --git a/examples/cifar10/cifar10_full_lrn_map_solver.prototxt b/examples/cifar10/cifar10_full_lrn_map_solver.prototxt new file mode 100644 index 00000000000..492912a0ffc --- /dev/null +++ b/examples/cifar10/cifar10_full_lrn_map_solver.prototxt @@ -0,0 +1,28 @@ +# reduce learning rate after 120 epochs (60000 iters) by factor 0f 10 +# then another factor of 10 after 10 more epochs (5000 iters) + +# The training protocol buffer definition +train_net: "cifar10_full_lrn_map_train.prototxt" +# The testing protocol buffer definition +test_net: "cifar10_full_lrn_map_test.prototxt" +# test_iter specifies how many forward passes the test should carry out. +# In the case of CIFAR10, we have test batch size 100 and 100 test iterations, +# covering the full 10,000 testing images. +test_iter: 100 +# Carry out testing every 1000 training iterations. +test_interval: 1000 +# The base learning rate, momentum and the weight decay of the network. +base_lr: 0.001 +momentum: 0.9 +weight_decay: 0.004 +# The learning rate policy +lr_policy: "fixed" +# Display every 200 iterations +display: 200 +# The maximum number of iterations +max_iter: 60000 +# snapshot intermediate results +snapshot: 10000 +snapshot_prefix: "cifar10_full_lrn_map" +# solver mode: 0 for CPU and 1 for GPU +solver_mode: 1 diff --git a/examples/cifar10/cifar10_full_lrn_map_solver_lr1.prototxt b/examples/cifar10/cifar10_full_lrn_map_solver_lr1.prototxt new file mode 100644 index 00000000000..c51a4e6b9f6 --- /dev/null +++ b/examples/cifar10/cifar10_full_lrn_map_solver_lr1.prototxt @@ -0,0 +1,28 @@ +# reduce learning rate after 120 epochs (60000 iters) by factor 0f 10 +# then another factor of 10 after 10 more epochs (5000 iters) + +# The training protocol buffer definition +train_net: "cifar10_full_lrn_map_train.prototxt" +# The testing protocol buffer definition +test_net: "cifar10_full_lrn_map_test.prototxt" +# test_iter specifies how many forward passes the test should carry out. +# In the case of CIFAR10, we have test batch size 100 and 100 test iterations, +# covering the full 10,000 testing images. +test_iter: 100 +# Carry out testing every 1000 training iterations. +test_interval: 1000 +# The base learning rate, momentum and the weight decay of the network. +base_lr: 0.0001 +momentum: 0.9 +weight_decay: 0.004 +# The learning rate policy +lr_policy: "fixed" +# Display every 200 iterations +display: 200 +# The maximum number of iterations +max_iter: 65000 +# snapshot intermediate results +snapshot: 5000 +snapshot_prefix: "cifar10_full_lrn_map" +# solver mode: 0 for CPU and 1 for GPU +solver_mode: 1 diff --git a/examples/cifar10/cifar10_full_lrn_map_solver_lr2.prototxt b/examples/cifar10/cifar10_full_lrn_map_solver_lr2.prototxt new file mode 100644 index 00000000000..d15b1155e2d --- /dev/null +++ b/examples/cifar10/cifar10_full_lrn_map_solver_lr2.prototxt @@ -0,0 +1,28 @@ +# reduce learning rate after 120 epochs (60000 iters) by factor 0f 10 +# then another factor of 10 after 10 more epochs (5000 iters) + +# The training protocol buffer definition +train_net: "cifar10_full_lrn_map_train.prototxt" +# The testing protocol buffer definition +test_net: "cifar10_full_lrn_map_test.prototxt" +# test_iter specifies how many forward passes the test should carry out. +# In the case of CIFAR10, we have test batch size 100 and 100 test iterations, +# covering the full 10,000 testing images. +test_iter: 100 +# Carry out testing every 1000 training iterations. +test_interval: 1000 +# The base learning rate, momentum and the weight decay of the network. +base_lr: 0.00001 +momentum: 0.9 +weight_decay: 0.004 +# The learning rate policy +lr_policy: "fixed" +# Display every 200 iterations +display: 200 +# The maximum number of iterations +max_iter: 70000 +# snapshot intermediate results +snapshot: 5000 +snapshot_prefix: "cifar10_full_lrn_map" +# solver mode: 0 for CPU and 1 for GPU +solver_mode: 1 diff --git a/examples/cifar10/cifar10_full_lrn_map_test.prototxt b/examples/cifar10/cifar10_full_lrn_map_test.prototxt new file mode 100644 index 00000000000..a8c81eda11b --- /dev/null +++ b/examples/cifar10/cifar10_full_lrn_map_test.prototxt @@ -0,0 +1,179 @@ +name: "CIFAR10_full_test" +layers { + name: "cifar" + type: DATA + top: "data" + top: "label" + data_param { + source: "cifar10-leveldb/cifar-test-leveldb" + mean_file: "mean.binaryproto" + batch_size: 100 + } +} +layers { + name: "conv1" + type: CONVOLUTION + bottom: "data" + top: "conv1" + blobs_lr: 1 + blobs_lr: 2 + convolution_param { + num_output: 32 + pad: 2 + kernel_size: 5 + stride: 1 + weight_filler { + type: "gaussian" + std: 0.0001 + } + bias_filler { + type: "constant" + } + } +} +layers { + name: "pool1" + type: POOLING + bottom: "conv1" + top: "pool1" + pooling_param { + pool: MAX + kernel_size: 3 + stride: 2 + } +} +layers { + name: "relu1" + type: RELU + bottom: "pool1" + top: "pool1" +} +layers { + name: "norm1" + type: LRN_MAP + bottom: "pool1" + top: "norm1" + lrn_param { + local_size: 3 + alpha: 5e-05 + beta: 0.75 + } +} +layers { + name: "conv2" + type: CONVOLUTION + bottom: "norm1" + top: "conv2" + blobs_lr: 1 + blobs_lr: 2 + convolution_param { + num_output: 32 + pad: 2 + kernel_size: 5 + stride: 1 + weight_filler { + type: "gaussian" + std: 0.01 + } + bias_filler { + type: "constant" + } + } +} +layers { + name: "relu2" + type: RELU + bottom: "conv2" + top: "conv2" +} +layers { + name: "pool2" + type: POOLING + bottom: "conv2" + top: "pool2" + pooling_param { + pool: AVE + kernel_size: 3 + stride: 2 + } +} +layers { + name: "norm2" + type: LRN_MAP + bottom: "pool2" + top: "norm2" + lrn_param { + local_size: 3 + alpha: 5e-05 + beta: 0.75 + } +} +layers { + name: "conv3" + type: CONVOLUTION + bottom: "norm2" + top: "conv3" + convolution_param { + num_output: 64 + pad: 2 + kernel_size: 5 + stride: 1 + weight_filler { + type: "gaussian" + std: 0.01 + } + bias_filler { + type: "constant" + } + } +} +layers { + name: "relu3" + type: RELU + bottom: "conv3" + top: "conv3" +} +layers { + name: "pool3" + type: POOLING + bottom: "conv3" + top: "pool3" + pooling_param { + pool: AVE + kernel_size: 3 + stride: 2 + } +} +layers { + name: "ip1" + type: INNER_PRODUCT + bottom: "pool3" + top: "ip1" + blobs_lr: 1 + blobs_lr: 2 + weight_decay: 250 + weight_decay: 0 + inner_product_param { + num_output: 10 + weight_filler { + type: "gaussian" + std: 0.01 + } + bias_filler { + type: "constant" + } + } +} +layers { + name: "prob" + type: SOFTMAX + bottom: "ip1" + top: "prob" +} +layers { + name: "accuracy" + type: ACCURACY + bottom: "prob" + bottom: "label" + top: "accuracy" +} diff --git a/examples/cifar10/cifar10_full_lrn_map_train.prototxt b/examples/cifar10/cifar10_full_lrn_map_train.prototxt new file mode 100644 index 00000000000..b9c436e6b36 --- /dev/null +++ b/examples/cifar10/cifar10_full_lrn_map_train.prototxt @@ -0,0 +1,172 @@ +name: "CIFAR10_full_train" +layers { + name: "cifar" + type: DATA + top: "data" + top: "label" + data_param { + source: "cifar10-leveldb/cifar-train-leveldb" + mean_file: "mean.binaryproto" + batch_size: 100 + } +} +layers { + name: "conv1" + type: CONVOLUTION + bottom: "data" + top: "conv1" + blobs_lr: 1 + blobs_lr: 2 + convolution_param { + num_output: 32 + pad: 2 + kernel_size: 5 + stride: 1 + weight_filler { + type: "gaussian" + std: 0.0001 + } + bias_filler { + type: "constant" + } + } +} +layers { + name: "pool1" + type: POOLING + bottom: "conv1" + top: "pool1" + pooling_param { + pool: MAX + kernel_size: 3 + stride: 2 + } +} +layers { + name: "relu1" + type: RELU + bottom: "pool1" + top: "pool1" +} +layers { + name: "norm1" + type: LRN_MAP + bottom: "pool1" + top: "norm1" + lrn_param { + local_size: 3 + alpha: 5e-05 + beta: 0.75 + } +} +layers { + name: "conv2" + type: CONVOLUTION + bottom: "norm1" + top: "conv2" + blobs_lr: 1 + blobs_lr: 2 + convolution_param { + num_output: 32 + pad: 2 + kernel_size: 5 + stride: 1 + weight_filler { + type: "gaussian" + std: 0.01 + } + bias_filler { + type: "constant" + } + } +} +layers { + name: "relu2" + type: RELU + bottom: "conv2" + top: "conv2" +} +layers { + name: "pool2" + type: POOLING + bottom: "conv2" + top: "pool2" + pooling_param { + pool: AVE + kernel_size: 3 + stride: 2 + } +} +layers { + name: "norm2" + type: LRN_MAP + bottom: "pool2" + top: "norm2" + lrn_param { + local_size: 3 + alpha: 5e-05 + beta: 0.75 + } +} +layers { + name: "conv3" + type: CONVOLUTION + bottom: "norm2" + top: "conv3" + convolution_param { + num_output: 64 + pad: 2 + kernel_size: 5 + stride: 1 + weight_filler { + type: "gaussian" + std: 0.01 + } + bias_filler { + type: "constant" + } + } +} +layers { + name: "relu3" + type: RELU + bottom: "conv3" + top: "conv3" +} +layers { + name: "pool3" + type: POOLING + bottom: "conv3" + top: "pool3" + pooling_param { + pool: AVE + kernel_size: 3 + stride: 2 + } +} +layers { + name: "ip1" + type: INNER_PRODUCT + bottom: "pool3" + top: "ip1" + blobs_lr: 1 + blobs_lr: 2 + weight_decay: 250 + weight_decay: 0 + inner_product_param { + num_output: 10 + weight_filler { + type: "gaussian" + std: 0.01 + } + bias_filler { + type: "constant" + } + } +} +layers { + name: "loss" + type: SOFTMAX_LOSS + bottom: "ip1" + bottom: "label" +} diff --git a/examples/cifar10/train_full_lrn_map.sh b/examples/cifar10/train_full_lrn_map.sh new file mode 100755 index 00000000000..a8f1ad62a7e --- /dev/null +++ b/examples/cifar10/train_full_lrn_map.sh @@ -0,0 +1,17 @@ +#!/usr/bin/env sh + +TOOLS=../../build/tools + +GLOG_logtostderr=1 $TOOLS/train_net.bin \ + cifar10_full_lrn_map_solver.prototxt \ + cifar10_full_lrn_map_iter_60000.solverstate + +#reduce learning rate by factor of 10 +GLOG_logtostderr=1 $TOOLS/train_net.bin \ + cifar10_full_lrn_map_solver_lr1.prototxt \ + cifar10_full_lrn_map_iter_60000.solverstate + +#reduce learning rate by factor of 10 +GLOG_logtostderr=1 $TOOLS/train_net.bin \ + cifar10_full_lrn_map_solver_lr2.prototxt \ + cifar10_full_lrn_map_iter_65000.solverstate From 4c81ee54af139e2072affbb865bb6deef92e14c5 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 02:13:34 -0700 Subject: [PATCH 09/18] fix some param bugs --- examples/cifar10/cifar10_full_lrn_map_test.prototxt | 4 ++-- examples/cifar10/cifar10_full_lrn_map_train.prototxt | 4 ++-- examples/cifar10/train_full_lrn_map.sh | 3 +-- 3 files changed, 5 insertions(+), 6 deletions(-) diff --git a/examples/cifar10/cifar10_full_lrn_map_test.prototxt b/examples/cifar10/cifar10_full_lrn_map_test.prototxt index a8c81eda11b..62a6e102bc8 100644 --- a/examples/cifar10/cifar10_full_lrn_map_test.prototxt +++ b/examples/cifar10/cifar10_full_lrn_map_test.prototxt @@ -53,7 +53,7 @@ layers { type: LRN_MAP bottom: "pool1" top: "norm1" - lrn_param { + lrn_map_param { local_size: 3 alpha: 5e-05 beta: 0.75 @@ -102,7 +102,7 @@ layers { type: LRN_MAP bottom: "pool2" top: "norm2" - lrn_param { + lrn_map_param { local_size: 3 alpha: 5e-05 beta: 0.75 diff --git a/examples/cifar10/cifar10_full_lrn_map_train.prototxt b/examples/cifar10/cifar10_full_lrn_map_train.prototxt index b9c436e6b36..d9244dc59d5 100644 --- a/examples/cifar10/cifar10_full_lrn_map_train.prototxt +++ b/examples/cifar10/cifar10_full_lrn_map_train.prototxt @@ -53,7 +53,7 @@ layers { type: LRN_MAP bottom: "pool1" top: "norm1" - lrn_param { + lrn_map_param { local_size: 3 alpha: 5e-05 beta: 0.75 @@ -102,7 +102,7 @@ layers { type: LRN_MAP bottom: "pool2" top: "norm2" - lrn_param { + lrn_map_param { local_size: 3 alpha: 5e-05 beta: 0.75 diff --git a/examples/cifar10/train_full_lrn_map.sh b/examples/cifar10/train_full_lrn_map.sh index a8f1ad62a7e..ffd3470285e 100755 --- a/examples/cifar10/train_full_lrn_map.sh +++ b/examples/cifar10/train_full_lrn_map.sh @@ -3,8 +3,7 @@ TOOLS=../../build/tools GLOG_logtostderr=1 $TOOLS/train_net.bin \ - cifar10_full_lrn_map_solver.prototxt \ - cifar10_full_lrn_map_iter_60000.solverstate + cifar10_full_lrn_map_solver.prototxt #reduce learning rate by factor of 10 GLOG_logtostderr=1 $TOOLS/train_net.bin \ From 824c34483ec4360cd660ec9ca60561cd52da8b86 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 03:36:28 -0700 Subject: [PATCH 10/18] merge LRNMapLayer into LRNLayer with norm_region proto field --- ...far10_full_lrn_in_channel_solver.prototxt} | 6 +- ...0_full_lrn_in_channel_solver_lr1.prototxt} | 6 +- ...0_full_lrn_in_channel_solver_lr2.prototxt} | 6 +- ...cifar10_full_lrn_in_channel_test.prototxt} | 10 +- ...ifar10_full_lrn_in_channel_train.prototxt} | 10 +- examples/cifar10/train_full_lrn_in_channel.sh | 16 ++ examples/cifar10/train_full_lrn_map.sh | 16 -- include/caffe/vision_layers.hpp | 44 +++-- src/caffe/layer_factory.cpp | 2 - src/caffe/layers/lrn_layer.cpp | 130 +++++++++++++- src/caffe/layers/lrn_layer.cu | 38 +++- src/caffe/layers/lrn_map_layer.cpp | 114 ------------ src/caffe/layers/lrn_map_layer.cu | 36 ---- src/caffe/proto/caffe.proto | 18 +- src/caffe/test/test_lrn_layer.cpp | 170 +++++++++++++++--- src/caffe/test/test_lrn_map_layer.cpp | 160 ----------------- 16 files changed, 370 insertions(+), 412 deletions(-) rename examples/cifar10/{cifar10_full_lrn_map_solver.prototxt => cifar10_full_lrn_in_channel_solver.prototxt} (84%) rename examples/cifar10/{cifar10_full_lrn_map_solver_lr1.prototxt => cifar10_full_lrn_in_channel_solver_lr1.prototxt} (84%) rename examples/cifar10/{cifar10_full_lrn_map_solver_lr2.prototxt => cifar10_full_lrn_in_channel_solver_lr2.prototxt} (84%) rename examples/cifar10/{cifar10_full_lrn_map_test.prototxt => cifar10_full_lrn_in_channel_test.prototxt} (95%) rename examples/cifar10/{cifar10_full_lrn_map_train.prototxt => cifar10_full_lrn_in_channel_train.prototxt} (95%) create mode 100755 examples/cifar10/train_full_lrn_in_channel.sh delete mode 100755 examples/cifar10/train_full_lrn_map.sh delete mode 100644 src/caffe/layers/lrn_map_layer.cpp delete mode 100644 src/caffe/layers/lrn_map_layer.cu delete mode 100644 src/caffe/test/test_lrn_map_layer.cpp diff --git a/examples/cifar10/cifar10_full_lrn_map_solver.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_solver.prototxt similarity index 84% rename from examples/cifar10/cifar10_full_lrn_map_solver.prototxt rename to examples/cifar10/cifar10_full_lrn_in_channel_solver.prototxt index 492912a0ffc..a217ec1a0dd 100644 --- a/examples/cifar10/cifar10_full_lrn_map_solver.prototxt +++ b/examples/cifar10/cifar10_full_lrn_in_channel_solver.prototxt @@ -2,9 +2,9 @@ # then another factor of 10 after 10 more epochs (5000 iters) # The training protocol buffer definition -train_net: "cifar10_full_lrn_map_train.prototxt" +train_net: "cifar10_full_lrn_in_channel_train.prototxt" # The testing protocol buffer definition -test_net: "cifar10_full_lrn_map_test.prototxt" +test_net: "cifar10_full_lrn_in_channel_test.prototxt" # test_iter specifies how many forward passes the test should carry out. # In the case of CIFAR10, we have test batch size 100 and 100 test iterations, # covering the full 10,000 testing images. @@ -23,6 +23,6 @@ display: 200 max_iter: 60000 # snapshot intermediate results snapshot: 10000 -snapshot_prefix: "cifar10_full_lrn_map" +snapshot_prefix: "cifar10_full_lrn_in_channel" # solver mode: 0 for CPU and 1 for GPU solver_mode: 1 diff --git a/examples/cifar10/cifar10_full_lrn_map_solver_lr1.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr1.prototxt similarity index 84% rename from examples/cifar10/cifar10_full_lrn_map_solver_lr1.prototxt rename to examples/cifar10/cifar10_full_lrn_in_channel_solver_lr1.prototxt index c51a4e6b9f6..6fcae73e7e9 100644 --- a/examples/cifar10/cifar10_full_lrn_map_solver_lr1.prototxt +++ b/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr1.prototxt @@ -2,9 +2,9 @@ # then another factor of 10 after 10 more epochs (5000 iters) # The training protocol buffer definition -train_net: "cifar10_full_lrn_map_train.prototxt" +train_net: "cifar10_full_lrn_in_channel_train.prototxt" # The testing protocol buffer definition -test_net: "cifar10_full_lrn_map_test.prototxt" +test_net: "cifar10_full_lrn_in_channel_test.prototxt" # test_iter specifies how many forward passes the test should carry out. # In the case of CIFAR10, we have test batch size 100 and 100 test iterations, # covering the full 10,000 testing images. @@ -23,6 +23,6 @@ display: 200 max_iter: 65000 # snapshot intermediate results snapshot: 5000 -snapshot_prefix: "cifar10_full_lrn_map" +snapshot_prefix: "cifar10_full_lrn_in_channel" # solver mode: 0 for CPU and 1 for GPU solver_mode: 1 diff --git a/examples/cifar10/cifar10_full_lrn_map_solver_lr2.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr2.prototxt similarity index 84% rename from examples/cifar10/cifar10_full_lrn_map_solver_lr2.prototxt rename to examples/cifar10/cifar10_full_lrn_in_channel_solver_lr2.prototxt index d15b1155e2d..8f53fe11a67 100644 --- a/examples/cifar10/cifar10_full_lrn_map_solver_lr2.prototxt +++ b/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr2.prototxt @@ -2,9 +2,9 @@ # then another factor of 10 after 10 more epochs (5000 iters) # The training protocol buffer definition -train_net: "cifar10_full_lrn_map_train.prototxt" +train_net: "cifar10_full_lrn_in_channel_train.prototxt" # The testing protocol buffer definition -test_net: "cifar10_full_lrn_map_test.prototxt" +test_net: "cifar10_full_lrn_in_channel_test.prototxt" # test_iter specifies how many forward passes the test should carry out. # In the case of CIFAR10, we have test batch size 100 and 100 test iterations, # covering the full 10,000 testing images. @@ -23,6 +23,6 @@ display: 200 max_iter: 70000 # snapshot intermediate results snapshot: 5000 -snapshot_prefix: "cifar10_full_lrn_map" +snapshot_prefix: "cifar10_full_lrn_in_channel" # solver mode: 0 for CPU and 1 for GPU solver_mode: 1 diff --git a/examples/cifar10/cifar10_full_lrn_map_test.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_test.prototxt similarity index 95% rename from examples/cifar10/cifar10_full_lrn_map_test.prototxt rename to examples/cifar10/cifar10_full_lrn_in_channel_test.prototxt index 62a6e102bc8..0e1957a9045 100644 --- a/examples/cifar10/cifar10_full_lrn_map_test.prototxt +++ b/examples/cifar10/cifar10_full_lrn_in_channel_test.prototxt @@ -50,10 +50,11 @@ layers { } layers { name: "norm1" - type: LRN_MAP + type: LRN bottom: "pool1" top: "norm1" - lrn_map_param { + lrn_param { + norm_region: WITHIN_CHANNEL local_size: 3 alpha: 5e-05 beta: 0.75 @@ -99,10 +100,11 @@ layers { } layers { name: "norm2" - type: LRN_MAP + type: LRN bottom: "pool2" top: "norm2" - lrn_map_param { + lrn_param { + norm_region: WITHIN_CHANNEL local_size: 3 alpha: 5e-05 beta: 0.75 diff --git a/examples/cifar10/cifar10_full_lrn_map_train.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_train.prototxt similarity index 95% rename from examples/cifar10/cifar10_full_lrn_map_train.prototxt rename to examples/cifar10/cifar10_full_lrn_in_channel_train.prototxt index d9244dc59d5..25c76060991 100644 --- a/examples/cifar10/cifar10_full_lrn_map_train.prototxt +++ b/examples/cifar10/cifar10_full_lrn_in_channel_train.prototxt @@ -50,10 +50,11 @@ layers { } layers { name: "norm1" - type: LRN_MAP + type: LRN bottom: "pool1" top: "norm1" - lrn_map_param { + lrn_param { + norm_region: WITHIN_CHANNEL local_size: 3 alpha: 5e-05 beta: 0.75 @@ -99,10 +100,11 @@ layers { } layers { name: "norm2" - type: LRN_MAP + type: LRN bottom: "pool2" top: "norm2" - lrn_map_param { + lrn_param { + norm_region: WITHIN_CHANNEL local_size: 3 alpha: 5e-05 beta: 0.75 diff --git a/examples/cifar10/train_full_lrn_in_channel.sh b/examples/cifar10/train_full_lrn_in_channel.sh new file mode 100755 index 00000000000..fcc9f4d0f92 --- /dev/null +++ b/examples/cifar10/train_full_lrn_in_channel.sh @@ -0,0 +1,16 @@ +#!/usr/bin/env sh + +TOOLS=../../build/tools + +GLOG_logtostderr=1 $TOOLS/train_net.bin \ + cifar10_full_lrn_in_channel_solver.prototxt + +#reduce learning rate by factor of 10 +GLOG_logtostderr=1 $TOOLS/train_net.bin \ + cifar10_full_lrn_in_channel_solver_lr1.prototxt \ + cifar10_full_lrn_in_channel_iter_60000.solverstate + +#reduce learning rate by factor of 10 +GLOG_logtostderr=1 $TOOLS/train_net.bin \ + cifar10_full_lrn_in_channel_solver_lr2.prototxt \ + cifar10_full_lrn_in_channel_iter_65000.solverstate diff --git a/examples/cifar10/train_full_lrn_map.sh b/examples/cifar10/train_full_lrn_map.sh deleted file mode 100755 index ffd3470285e..00000000000 --- a/examples/cifar10/train_full_lrn_map.sh +++ /dev/null @@ -1,16 +0,0 @@ -#!/usr/bin/env sh - -TOOLS=../../build/tools - -GLOG_logtostderr=1 $TOOLS/train_net.bin \ - cifar10_full_lrn_map_solver.prototxt - -#reduce learning rate by factor of 10 -GLOG_logtostderr=1 $TOOLS/train_net.bin \ - cifar10_full_lrn_map_solver_lr1.prototxt \ - cifar10_full_lrn_map_iter_60000.solverstate - -#reduce learning rate by factor of 10 -GLOG_logtostderr=1 $TOOLS/train_net.bin \ - cifar10_full_lrn_map_solver_lr2.prototxt \ - cifar10_full_lrn_map_iter_65000.solverstate diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index 22c0fd06457..93f6d25296a 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -500,6 +500,10 @@ class InnerProductLayer : public Layer { shared_ptr bias_multiplier_; }; +// Forward declare PoolingLayer and SplitLayer for use in LRNLayer. +template class PoolingLayer; +template class SplitLayer; + template class LRNLayer : public Layer { public: @@ -518,8 +522,19 @@ class LRNLayer : public Layer { virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - // scale_ stores the intermediate summing results - Blob scale_; + virtual Dtype Forward_cpu_cross_channel(const vector*>& bottom, + vector*>* top); + virtual Dtype Forward_gpu_cross_channel(const vector*>& bottom, + vector*>* top); + virtual Dtype Forward_within_channel(const vector*>& bottom, + vector*>* top); + virtual void Backward_cpu_cross_channel(const vector*>& top, + const bool propagate_down, vector*>* bottom); + virtual void Backward_gpu_cross_channel(const vector*>& top, + const bool propagate_down, vector*>* bottom); + virtual void Backward_within_channel(const vector*>& top, + const bool propagate_down, vector*>* bottom); + int size_; int pre_pad_; Dtype alpha_; @@ -528,29 +543,12 @@ class LRNLayer : public Layer { int channels_; int height_; int width_; -}; -template class PoolingLayer; -template class SplitLayer; - -template -class LRNMapLayer : public Layer { - public: - explicit LRNMapLayer(const LayerParameter& param) - : Layer(param) {} - virtual void SetUp(const vector*>& bottom, - vector*>* top); - - protected: - virtual Dtype Forward_cpu(const vector*>& bottom, - vector*>* top); - virtual Dtype Forward_gpu(const vector*>& bottom, - vector*>* top); - virtual void Backward_cpu(const vector*>& top, - const bool propagate_down, vector*>* bottom); - virtual void Backward_gpu(const vector*>& top, - const bool propagate_down, vector*>* bottom); + // Fields used for normalization ACROSS_CHANNELS + // scale_ stores the intermediate summing results + Blob scale_; + // Fields used for normalization WITHIN_CHANNEL shared_ptr > split_layer_; vector*> split_top_vec_; shared_ptr > square_layer_; diff --git a/src/caffe/layer_factory.cpp b/src/caffe/layer_factory.cpp index 91c5850b9c7..d30ffeeb88f 100644 --- a/src/caffe/layer_factory.cpp +++ b/src/caffe/layer_factory.cpp @@ -54,8 +54,6 @@ Layer* GetLayer(const LayerParameter& param) { return new InnerProductLayer(param); case LayerParameter_LayerType_LRN: return new LRNLayer(param); - case LayerParameter_LayerType_LRN_MAP: - return new LRNMapLayer(param); case LayerParameter_LayerType_MULTINOMIAL_LOGISTIC_LOSS: return new MultinomialLogisticLossLayer(param); case LayerParameter_LayerType_POOLING: diff --git a/src/caffe/layers/lrn_layer.cpp b/src/caffe/layers/lrn_layer.cpp index e95bc8a61c9..d9be541940e 100644 --- a/src/caffe/layers/lrn_layer.cpp +++ b/src/caffe/layers/lrn_layer.cpp @@ -19,17 +19,103 @@ void LRNLayer::SetUp(const vector*>& bottom, channels_ = bottom[0]->channels(); height_ = bottom[0]->height(); width_ = bottom[0]->width(); - (*top)[0]->Reshape(num_, channels_, height_, width_); - scale_.Reshape(num_, channels_, height_, width_); size_ = this->layer_param_.lrn_param().local_size(); pre_pad_ = (size_ - 1) / 2; alpha_ = this->layer_param_.lrn_param().alpha(); beta_ = this->layer_param_.lrn_param().beta(); + switch (this->layer_param_.lrn_param().norm_region()) { + case LRNParameter_NormRegion_ACROSS_CHANNELS: + (*top)[0]->Reshape(num_, channels_, height_, width_); + scale_.Reshape(num_, channels_, height_, width_); + break; + case LRNParameter_NormRegion_WITHIN_CHANNEL: + { + const Dtype pre_pad = (size_ - 1) / 2; + // Set up split layer to use inputs in the num_erator and denominator. + split_top_vec_.clear(); + split_top_vec_.push_back(bottom[0]); + split_top_vec_.push_back(&square_input_); + LayerParameter split_param; + split_layer_.reset(new SplitLayer(split_param)); + split_layer_->SetUp(bottom, &split_top_vec_); + // Set up square layer to square the inputs. + square_input_.Reshape(num_, channels_, height_, width_); + square_bottom_vec_.clear(); + square_top_vec_.clear(); + square_bottom_vec_.push_back(&square_input_); + square_top_vec_.push_back(&square_output_); + LayerParameter square_param; + square_param.mutable_power_param()->set_power(Dtype(2)); + square_layer_.reset(new PowerLayer(square_param)); + square_layer_->SetUp(square_bottom_vec_, &square_top_vec_); + CHECK_EQ(square_output_.num(), num_); + CHECK_EQ(square_output_.channels(), channels_); + CHECK_EQ(square_output_.height(), height_); + CHECK_EQ(square_output_.width(), width_); + // Output of pool layer gives us the neighborhood response. + pool_top_vec_.clear(); + pool_top_vec_.push_back(&pool_output_); + LayerParameter pool_param; + pool_param.mutable_pooling_param()->set_pool( + PoolingParameter_PoolMethod_AVE); + pool_param.mutable_pooling_param()->set_pad(pre_pad); + pool_param.mutable_pooling_param()->set_kernel_size(size_); + pool_layer_.reset(new PoolingLayer(pool_param)); + pool_layer_->SetUp(square_top_vec_, &pool_top_vec_); + CHECK_EQ(pool_output_.num(), num_); + CHECK_EQ(pool_output_.channels(), channels_); + CHECK_EQ(pool_output_.height(), height_); + CHECK_EQ(pool_output_.width(), width_); + // Set up power layer to compute (1 + alpha_/N^2 s)^-beta_, where s is the + // sum of a squared neighborhood (as output by pool_layer_). + power_top_vec_.clear(); + power_top_vec_.push_back(&power_output_); + LayerParameter power_param; + power_param.mutable_power_param()->set_power(-beta_); + power_param.mutable_power_param()->set_scale(alpha_); + power_param.mutable_power_param()->set_shift(Dtype(1)); + power_layer_.reset(new PowerLayer(power_param)); + power_layer_->SetUp(pool_top_vec_, &power_top_vec_); + CHECK_EQ(power_output_.num(), num_); + CHECK_EQ(power_output_.channels(), channels_); + CHECK_EQ(power_output_.height(), height_); + CHECK_EQ(power_output_.width(), width_); + // Set up a product layer to compute outputs by multiplying inputs by the + // demoninator computed by the power layer. + product_bottom_vec_.clear(); + product_bottom_vec_.push_back(bottom[0]); + product_bottom_vec_.push_back(&power_output_); + LayerParameter product_param; + product_layer_.reset(new EltwiseProductLayer(product_param)); + product_layer_->SetUp(product_bottom_vec_, top); + CHECK_EQ((*top)[0]->num(), num_); + CHECK_EQ((*top)[0]->channels(), channels_); + CHECK_EQ((*top)[0]->height(), height_); + CHECK_EQ((*top)[0]->width(), width_); + } + break; + default: + LOG(FATAL) << "Unknown normalization region."; + } } template Dtype LRNLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { + switch (this->layer_param_.lrn_param().norm_region()) { + case LRNParameter_NormRegion_ACROSS_CHANNELS: + return Forward_cpu_cross_channel(bottom, top); + case LRNParameter_NormRegion_WITHIN_CHANNEL: + return Forward_within_channel(bottom, top); + default: + LOG(FATAL) << "Unknown normalization region."; + return Dtype(0); + } +} + +template +Dtype LRNLayer::Forward_cpu_cross_channel( + const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); Dtype* scale_data = scale_.mutable_cpu_data(); @@ -76,9 +162,36 @@ Dtype LRNLayer::Forward_cpu(const vector*>& bottom, return Dtype(0.); } +template +Dtype LRNLayer::Forward_within_channel( + const vector*>& bottom, vector*>* top) { + split_layer_->Forward(bottom, &split_top_vec_); + square_layer_->Forward(square_bottom_vec_, &square_top_vec_); + pool_layer_->Forward(square_top_vec_, &pool_top_vec_); + power_layer_->Forward(pool_top_vec_, &power_top_vec_); + product_layer_->Forward(product_bottom_vec_, top); + return Dtype(0.); +} + template void LRNLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { + switch (this->layer_param_.lrn_param().norm_region()) { + case LRNParameter_NormRegion_ACROSS_CHANNELS: + Backward_cpu_cross_channel(top, propagate_down, bottom); + break; + case LRNParameter_NormRegion_WITHIN_CHANNEL: + Backward_within_channel(top, propagate_down, bottom); + break; + default: + LOG(FATAL) << "Unknown normalization region."; + } +} + +template +void LRNLayer::Backward_cpu_cross_channel( + const vector*>& top, const bool propagate_down, + vector*>* bottom) { const Dtype* top_diff = top[0]->cpu_diff(); const Dtype* top_data = top[0]->cpu_data(); const Dtype* bottom_data = (*bottom)[0]->cpu_data(); @@ -130,6 +243,19 @@ void LRNLayer::Backward_cpu(const vector*>& top, } } +template +void LRNLayer::Backward_within_channel( + const vector*>& top, const bool propagate_down, + vector*>* bottom) { + if (propagate_down) { + product_layer_->Backward(top, true, &product_bottom_vec_); + power_layer_->Backward(power_top_vec_, true, &pool_top_vec_); + pool_layer_->Backward(pool_top_vec_, true, &square_top_vec_); + square_layer_->Backward(square_top_vec_, true, &square_bottom_vec_); + split_layer_->Backward(split_top_vec_, true, bottom); + } +} + INSTANTIATE_CLASS(LRNLayer); diff --git a/src/caffe/layers/lrn_layer.cu b/src/caffe/layers/lrn_layer.cu index b8cd99ac3df..47eca1765de 100644 --- a/src/caffe/layers/lrn_layer.cu +++ b/src/caffe/layers/lrn_layer.cu @@ -55,6 +55,20 @@ __global__ void LRNFillScale(const int nthreads, const Dtype* in, } +template +Dtype LRNLayer::Forward_gpu(const vector*>& bottom, + vector*>* top) { + switch (this->layer_param_.lrn_param().norm_region()) { + case LRNParameter_NormRegion_ACROSS_CHANNELS: + return Forward_gpu_cross_channel(bottom, top); + case LRNParameter_NormRegion_WITHIN_CHANNEL: + return Forward_within_channel(bottom, top); + default: + LOG(FATAL) << "Unknown normalization region."; + return Dtype(0); + } +} + // TODO: check if it would be faster to just put it into the previous kernel. template __global__ void LRNComputeOutput(const int nthreads, const Dtype* in, @@ -65,8 +79,8 @@ __global__ void LRNComputeOutput(const int nthreads, const Dtype* in, } template -Dtype LRNLayer::Forward_gpu(const vector*>& bottom, - vector*>* top) { +Dtype LRNLayer::Forward_gpu_cross_channel( + const vector*>& bottom, vector*>* top) { // First, compute scale const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -88,6 +102,21 @@ Dtype LRNLayer::Forward_gpu(const vector*>& bottom, } +template +void LRNLayer::Backward_gpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { + switch (this->layer_param_.lrn_param().norm_region()) { + case LRNParameter_NormRegion_ACROSS_CHANNELS: + Backward_gpu_cross_channel(top, propagate_down, bottom); + break; + case LRNParameter_NormRegion_WITHIN_CHANNEL: + Backward_within_channel(top, propagate_down, bottom); + break; + default: + LOG(FATAL) << "Unknown normalization region."; + } +} + template __global__ void LRNComputeDiff(const int nthreads, const Dtype* bottom_data, const Dtype* top_data, const Dtype* scale, const Dtype* top_diff, @@ -150,8 +179,9 @@ __global__ void LRNComputeDiff(const int nthreads, const Dtype* bottom_data, } template -void LRNLayer::Backward_gpu(const vector*>& top, - const bool propagate_down, vector*>* bottom) { +void LRNLayer::Backward_gpu_cross_channel( + const vector*>& top, const bool propagate_down, + vector*>* bottom) { int n_threads = num_ * height_ * width_; // NOLINT_NEXT_LINE(whitespace/operators) LRNComputeDiff<<>>( diff --git a/src/caffe/layers/lrn_map_layer.cpp b/src/caffe/layers/lrn_map_layer.cpp deleted file mode 100644 index 15119adf64e..00000000000 --- a/src/caffe/layers/lrn_map_layer.cpp +++ /dev/null @@ -1,114 +0,0 @@ -// Copyright 2014 BVLC and contributors. - -#include - -#include "caffe/filler.hpp" -#include "caffe/layer.hpp" -#include "caffe/vision_layers.hpp" -#include "caffe/util/math_functions.hpp" - -namespace caffe { - -template -void LRNMapLayer::SetUp(const vector*>& bottom, - vector*>* top) { - CHECK_EQ(bottom.size(), 1) << - "Local Response Normalization Layer takes a single blob as input."; - CHECK_EQ(top->size(), 1) << - "Local Response Normalization Layer takes a single blob as output."; - const int num = bottom[0]->num(); - const int channels = bottom[0]->channels(); - const int height = bottom[0]->height(); - const int width = bottom[0]->width(); - const int size_ = this->layer_param_.lrn_map_param().local_size(); - const Dtype pre_pad = (size_ - 1) / 2; - const Dtype alpha = this->layer_param_.lrn_map_param().alpha(); - const Dtype beta = this->layer_param_.lrn_map_param().beta(); - // Set up split layer to use inputs in the numerator and denominator. - split_top_vec_.clear(); - split_top_vec_.push_back(bottom[0]); - split_top_vec_.push_back(&square_input_); - LayerParameter split_param; - split_layer_.reset(new SplitLayer(split_param)); - split_layer_->SetUp(bottom, &split_top_vec_); - // Set up square layer to square the inputs. - square_input_.Reshape(num, channels, height, width); - square_bottom_vec_.clear(); - square_top_vec_.clear(); - square_bottom_vec_.push_back(&square_input_); - square_top_vec_.push_back(&square_output_); - LayerParameter square_param; - square_param.mutable_power_param()->set_power(Dtype(2)); - square_layer_.reset(new PowerLayer(square_param)); - square_layer_->SetUp(square_bottom_vec_, &square_top_vec_); - CHECK_EQ(square_output_.num(), num); - CHECK_EQ(square_output_.channels(), channels); - CHECK_EQ(square_output_.height(), height); - CHECK_EQ(square_output_.width(), width); - // Output of pool layer gives us the neighborhood response. - pool_top_vec_.clear(); - pool_top_vec_.push_back(&pool_output_); - LayerParameter pool_param; - pool_param.mutable_pooling_param()->set_pool(PoolingParameter_PoolMethod_AVE); - pool_param.mutable_pooling_param()->set_pad(pre_pad); - pool_param.mutable_pooling_param()->set_kernel_size(size_); - pool_layer_.reset(new PoolingLayer(pool_param)); - pool_layer_->SetUp(square_top_vec_, &pool_top_vec_); - CHECK_EQ(pool_output_.num(), num); - CHECK_EQ(pool_output_.channels(), channels); - CHECK_EQ(pool_output_.height(), height); - CHECK_EQ(pool_output_.width(), width); - // Set up power layer to compute (1 + alpha/N^2 s)^-beta, where s is the sum - // of a squared neighborhood (as output by the conv layer). - power_top_vec_.clear(); - power_top_vec_.push_back(&power_output_); - LayerParameter power_param; - power_param.mutable_power_param()->set_power(-beta); - power_param.mutable_power_param()->set_scale(alpha); - power_param.mutable_power_param()->set_shift(Dtype(1)); - power_layer_.reset(new PowerLayer(power_param)); - power_layer_->SetUp(pool_top_vec_, &power_top_vec_); - CHECK_EQ(power_output_.num(), num); - CHECK_EQ(power_output_.channels(), channels); - CHECK_EQ(power_output_.height(), height); - CHECK_EQ(power_output_.width(), width); - // Set up a product layer to compute outputs by multiplying inputs by the - // demoninator computed by the power layer. - product_bottom_vec_.clear(); - product_bottom_vec_.push_back(bottom[0]); - product_bottom_vec_.push_back(&power_output_); - LayerParameter product_param; - product_layer_.reset(new EltwiseProductLayer(product_param)); - product_layer_->SetUp(product_bottom_vec_, top); - CHECK_EQ((*top)[0]->num(), num); - CHECK_EQ((*top)[0]->channels(), channels); - CHECK_EQ((*top)[0]->height(), height); - CHECK_EQ((*top)[0]->width(), width); -} - -template -Dtype LRNMapLayer::Forward_cpu(const vector*>& bottom, - vector*>* top) { - split_layer_->Forward(bottom, &split_top_vec_); - square_layer_->Forward(square_bottom_vec_, &square_top_vec_); - pool_layer_->Forward(square_top_vec_, &pool_top_vec_); - power_layer_->Forward(pool_top_vec_, &power_top_vec_); - product_layer_->Forward(product_bottom_vec_, top); - return Dtype(0.); -} - -template -void LRNMapLayer::Backward_cpu(const vector*>& top, - const bool propagate_down, vector*>* bottom) { - if (propagate_down) { - product_layer_->Backward(top, true, &product_bottom_vec_); - power_layer_->Backward(power_top_vec_, true, &pool_top_vec_); - pool_layer_->Backward(pool_top_vec_, true, &square_top_vec_); - square_layer_->Backward(square_top_vec_, true, &square_bottom_vec_); - split_layer_->Backward(split_top_vec_, true, bottom); - } -} - -INSTANTIATE_CLASS(LRNMapLayer); - -} // namespace caffe diff --git a/src/caffe/layers/lrn_map_layer.cu b/src/caffe/layers/lrn_map_layer.cu deleted file mode 100644 index b41adfc96be..00000000000 --- a/src/caffe/layers/lrn_map_layer.cu +++ /dev/null @@ -1,36 +0,0 @@ -// Copyright 2014 BVLC and contributors. - -#include - -#include "caffe/layer.hpp" -#include "caffe/vision_layers.hpp" -#include "caffe/util/math_functions.hpp" - -namespace caffe { - -template -Dtype LRNMapLayer::Forward_gpu(const vector*>& bottom, - vector*>* top) { - split_layer_->Forward(bottom, &split_top_vec_); - square_layer_->Forward(square_bottom_vec_, &square_top_vec_); - pool_layer_->Forward(square_top_vec_, &pool_top_vec_); - power_layer_->Forward(pool_top_vec_, &power_top_vec_); - product_layer_->Forward(product_bottom_vec_, top); - return Dtype(0.); -} - -template -void LRNMapLayer::Backward_gpu(const vector*>& top, - const bool propagate_down, vector*>* bottom) { - if (propagate_down) { - product_layer_->Backward(top, true, &product_bottom_vec_); - power_layer_->Backward(power_top_vec_, true, &pool_top_vec_); - pool_layer_->Backward(pool_top_vec_, true, &square_top_vec_); - square_layer_->Backward(square_top_vec_, true, &square_bottom_vec_); - split_layer_->Backward(split_top_vec_, true, bottom); - } -} - -INSTANTIATE_CLASS(LRNMapLayer); - -} // namespace caffe diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 54f2743ee4f..745763cc358 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -121,10 +121,9 @@ message LayerParameter { INFOGAIN_LOSS = 13; INNER_PRODUCT = 14; LRN = 15; - LRN_MAP = 1001; MULTINOMIAL_LOGISTIC_LOSS = 16; POOLING = 17; - POWER = 1002; + POWER = 1001; RELU = 18; SIGMOID = 19; SOFTMAX = 20; @@ -154,9 +153,8 @@ message LayerParameter { optional InfogainLossParameter infogain_loss_param = 16; optional InnerProductParameter inner_product_param = 17; optional LRNParameter lrn_param = 18; - optional LRNMapParameter lrn_map_param = 1000; optional PoolingParameter pooling_param = 19; - optional PowerParameter power_param = 1001; + optional PowerParameter power_param = 1000; optional WindowDataParameter window_data_param = 20; // DEPRECATED: The layer parameters specified as a V0LayerParameter. @@ -271,13 +269,11 @@ message LRNParameter { optional uint32 local_size = 1 [default = 5]; optional float alpha = 2 [default = 1.]; optional float beta = 3 [default = 0.75]; -} - -// Message that stores parameters used by LRNMapLayer -message LRNMapParameter { - optional uint32 local_size = 1 [default = 5]; - optional float alpha = 2 [default = 1.]; - optional float beta = 3 [default = 0.75]; + enum NormRegion { + ACROSS_CHANNELS = 0; + WITHIN_CHANNEL = 1; + } + optional NormRegion norm_region = 4 [default = ACROSS_CHANNELS]; } // Message that stores parameters used by PoolingLayer diff --git a/src/caffe/test/test_lrn_layer.cpp b/src/caffe/test/test_lrn_layer.cpp index 6ad6d020c03..1923128dd71 100644 --- a/src/caffe/test/test_lrn_layer.cpp +++ b/src/caffe/test/test_lrn_layer.cpp @@ -26,7 +26,8 @@ class LRNLayerTest : public ::testing::Test { protected: LRNLayerTest() : blob_bottom_(new Blob()), - blob_top_(new Blob()) {} + blob_top_(new Blob()), + epsilon_(Dtype(1e-5)) {} virtual void SetUp() { Caffe::set_random_seed(1701); blob_bottom_->Reshape(2, 7, 3, 3); @@ -40,6 +41,8 @@ class LRNLayerTest : public ::testing::Test { virtual ~LRNLayerTest() { delete blob_bottom_; delete blob_top_; } void ReferenceLRNForward(const Blob& blob_bottom, const LayerParameter& layer_param, Blob* blob_top); + + Dtype epsilon_; Blob* const blob_bottom_; Blob* const blob_top_; vector*> blob_bottom_vec_; @@ -58,30 +61,61 @@ void LRNLayerTest::ReferenceLRNForward( Dtype alpha = lrn_param.alpha(); Dtype beta = lrn_param.beta(); int size = lrn_param.local_size(); - for (int n = 0; n < blob_bottom.num(); ++n) { - for (int c = 0; c < blob_bottom.channels(); ++c) { - for (int h = 0; h < blob_bottom.height(); ++h) { - for (int w = 0; w < blob_bottom.width(); ++w) { - int c_start = c - (size - 1) / 2; - int c_end = min(c_start + size, blob_bottom.channels()); - c_start = max(c_start, 0); - Dtype scale = 1.; - for (int i = c_start; i < c_end; ++i) { - Dtype value = blob_bottom.data_at(n, i, h, w); - scale += value * value * alpha / size; + switch (lrn_param.norm_region()) { + case LRNParameter_NormRegion_ACROSS_CHANNELS: + for (int n = 0; n < blob_bottom.num(); ++n) { + for (int c = 0; c < blob_bottom.channels(); ++c) { + for (int h = 0; h < blob_bottom.height(); ++h) { + for (int w = 0; w < blob_bottom.width(); ++w) { + int c_start = c - (size - 1) / 2; + int c_end = min(c_start + size, blob_bottom.channels()); + c_start = max(c_start, 0); + Dtype scale = 1.; + for (int i = c_start; i < c_end; ++i) { + Dtype value = blob_bottom.data_at(n, i, h, w); + scale += value * value * alpha / size; + } + *(top_data + blob_top->offset(n, c, h, w)) = + blob_bottom.data_at(n, c, h, w) / pow(scale, beta); + } + } + } + } + break; + case LRNParameter_NormRegion_WITHIN_CHANNEL: + for (int n = 0; n < blob_bottom.num(); ++n) { + for (int c = 0; c < blob_bottom.channels(); ++c) { + for (int h = 0; h < blob_bottom.height(); ++h) { + int h_start = h - (size - 1) / 2; + int h_end = min(h_start + size, blob_bottom.height()); + h_start = max(h_start, 0); + for (int w = 0; w < blob_bottom.width(); ++w) { + Dtype scale = 1.; + int w_start = w - (size - 1) / 2; + int w_end = min(w_start + size, blob_bottom.width()); + w_start = max(w_start, 0); + for (int nh = h_start; nh < h_end; ++nh) { + for (int nw = w_start; nw < w_end; ++nw) { + Dtype value = blob_bottom.data_at(n, c, nh, nw); + scale += value * value * alpha / (size * size); + } + } + *(top_data + blob_top->offset(n, c, h, w)) = + blob_bottom.data_at(n, c, h, w) / pow(scale, beta); } - *(top_data + blob_top->offset(n, c, h, w)) = - blob_bottom.data_at(n, c, h, w) / pow(scale, beta); } } } + break; + default: + LOG(FATAL) << "Unknown normalization region."; } } typedef ::testing::Types Dtypes; TYPED_TEST_CASE(LRNLayerTest, Dtypes); -TYPED_TEST(LRNLayerTest, TestSetup) { +TYPED_TEST(LRNLayerTest, TestSetupAcrossChannels) { LayerParameter layer_param; LRNLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); @@ -91,7 +125,7 @@ TYPED_TEST(LRNLayerTest, TestSetup) { EXPECT_EQ(this->blob_top_->width(), 3); } -TYPED_TEST(LRNLayerTest, TestCPUForward) { +TYPED_TEST(LRNLayerTest, TestCPUForwardAcrossChannels) { LayerParameter layer_param; LRNLayer layer(layer_param); Caffe::set_mode(Caffe::CPU); @@ -101,14 +135,12 @@ TYPED_TEST(LRNLayerTest, TestCPUForward) { this->ReferenceLRNForward(*(this->blob_bottom_), layer_param, &top_reference); for (int i = 0; i < this->blob_bottom_->count(); ++i) { - EXPECT_GE(this->blob_top_->cpu_data()[i], - top_reference.cpu_data()[i] - 1e-5); - EXPECT_LE(this->blob_top_->cpu_data()[i], - top_reference.cpu_data()[i] + 1e-5); + EXPECT_NEAR(this->blob_top_->cpu_data()[i], top_reference.cpu_data()[i], + this->epsilon_); } } -TYPED_TEST(LRNLayerTest, TestGPUForward) { +TYPED_TEST(LRNLayerTest, TestGPUForwardAcrossChannels) { LayerParameter layer_param; LRNLayer layer(layer_param); Caffe::set_mode(Caffe::GPU); @@ -118,14 +150,12 @@ TYPED_TEST(LRNLayerTest, TestGPUForward) { this->ReferenceLRNForward(*(this->blob_bottom_), layer_param, &top_reference); for (int i = 0; i < this->blob_bottom_->count(); ++i) { - EXPECT_GE(this->blob_top_->cpu_data()[i], - top_reference.cpu_data()[i] - 1e-5); - EXPECT_LE(this->blob_top_->cpu_data()[i], - top_reference.cpu_data()[i] + 1e-5); + EXPECT_NEAR(this->blob_top_->cpu_data()[i], top_reference.cpu_data()[i], + this->epsilon_); } } -TYPED_TEST(LRNLayerTest, TestCPUGradient) { +TYPED_TEST(LRNLayerTest, TestCPUGradientAcrossChannels) { LayerParameter layer_param; LRNLayer layer(layer_param); GradientChecker checker(1e-2, 1e-2); @@ -144,7 +174,7 @@ TYPED_TEST(LRNLayerTest, TestCPUGradient) { &(this->blob_top_vec_)); } -TYPED_TEST(LRNLayerTest, TestGPUGradient) { +TYPED_TEST(LRNLayerTest, TestGPUGradientAcrossChannels) { LayerParameter layer_param; LRNLayer layer(layer_param); GradientChecker checker(1e-2, 1e-2); @@ -163,4 +193,90 @@ TYPED_TEST(LRNLayerTest, TestGPUGradient) { &(this->blob_top_vec_)); } +TYPED_TEST(LRNLayerTest, TestSetupWithinChannel) { + LayerParameter layer_param; + layer_param.mutable_lrn_param()->set_norm_region( + LRNParameter_NormRegion_WITHIN_CHANNEL); + layer_param.mutable_lrn_param()->set_local_size(3); + LRNLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + EXPECT_EQ(this->blob_top_->num(), 2); + EXPECT_EQ(this->blob_top_->channels(), 7); + EXPECT_EQ(this->blob_top_->height(), 3); + EXPECT_EQ(this->blob_top_->width(), 3); +} + +TYPED_TEST(LRNLayerTest, TestCPUForwardWithinChannel) { + LayerParameter layer_param; + layer_param.mutable_lrn_param()->set_norm_region( + LRNParameter_NormRegion_WITHIN_CHANNEL); + layer_param.mutable_lrn_param()->set_local_size(3); + LRNLayer layer(layer_param); + Caffe::set_mode(Caffe::CPU); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + Blob top_reference; + this->ReferenceLRNForward(*(this->blob_bottom_), layer_param, + &top_reference); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_NEAR(this->blob_top_->cpu_data()[i], top_reference.cpu_data()[i], + this->epsilon_); + } +} + +TYPED_TEST(LRNLayerTest, TestGPUForwardWithinChannel) { + LayerParameter layer_param; + layer_param.mutable_lrn_param()->set_norm_region( + LRNParameter_NormRegion_WITHIN_CHANNEL); + layer_param.mutable_lrn_param()->set_local_size(3); + LRNLayer layer(layer_param); + Caffe::set_mode(Caffe::GPU); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + Blob top_reference; + this->ReferenceLRNForward(*(this->blob_bottom_), layer_param, + &top_reference); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_NEAR(this->blob_top_->cpu_data()[i], top_reference.cpu_data()[i], + this->epsilon_); + } +} + +TYPED_TEST(LRNLayerTest, TestCPUGradientWithinChannel) { + LayerParameter layer_param; + layer_param.mutable_lrn_param()->set_norm_region( + LRNParameter_NormRegion_WITHIN_CHANNEL); + layer_param.mutable_lrn_param()->set_local_size(3); + LRNLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + Caffe::set_mode(Caffe::CPU); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + for (int i = 0; i < this->blob_top_->count(); ++i) { + this->blob_top_->mutable_cpu_diff()[i] = 1.; + } + layer.Backward(this->blob_top_vec_, true, &(this->blob_bottom_vec_)); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(LRNLayerTest, TestGPUGradientWithinChannel) { + LayerParameter layer_param; + layer_param.mutable_lrn_param()->set_norm_region( + LRNParameter_NormRegion_WITHIN_CHANNEL); + layer_param.mutable_lrn_param()->set_local_size(3); + LRNLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + Caffe::set_mode(Caffe::GPU); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + for (int i = 0; i < this->blob_top_->count(); ++i) { + this->blob_top_->mutable_cpu_diff()[i] = 1.; + } + layer.Backward(this->blob_top_vec_, true, &(this->blob_bottom_vec_)); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + + } // namespace caffe diff --git a/src/caffe/test/test_lrn_map_layer.cpp b/src/caffe/test/test_lrn_map_layer.cpp deleted file mode 100644 index 93377cc1beb..00000000000 --- a/src/caffe/test/test_lrn_map_layer.cpp +++ /dev/null @@ -1,160 +0,0 @@ -// Copyright 2014 BVLC and contributors. - -#include -#include - -#include "cuda_runtime.h" -#include "gtest/gtest.h" - -#include "caffe/blob.hpp" -#include "caffe/common.hpp" -#include "caffe/filler.hpp" -#include "caffe/vision_layers.hpp" -#include "caffe/test/test_gradient_check_util.hpp" - -#include "caffe/test/test_caffe_main.hpp" - -using std::min; -using std::max; - -namespace caffe { - -extern cudaDeviceProp CAFFE_TEST_CUDA_PROP; - -template -class LRNMapLayerTest : public ::testing::Test { - protected: - LRNMapLayerTest() - : blob_bottom_(new Blob()), - blob_top_(new Blob()) {} - virtual void SetUp() { - Caffe::set_random_seed(1701); - blob_bottom_->Reshape(2, 2, 7, 7); - // fill the values - FillerParameter filler_param; - GaussianFiller filler(filler_param); - filler.Fill(this->blob_bottom_); - blob_bottom_vec_.push_back(blob_bottom_); - blob_top_vec_.push_back(blob_top_); - epsilon_ = 1e-5; - } - virtual ~LRNMapLayerTest() { delete blob_bottom_; delete blob_top_; } - void ReferenceLRNMapForward(const Blob& blob_bottom, - const LayerParameter& layer_param, Blob* blob_top); - Blob* const blob_bottom_; - Blob* const blob_top_; - vector*> blob_bottom_vec_; - vector*> blob_top_vec_; - Dtype epsilon_; -}; - -template -void LRNMapLayerTest::ReferenceLRNMapForward( - const Blob& blob_bottom, const LayerParameter& layer_param, - Blob* blob_top) { - blob_top->Reshape(blob_bottom.num(), blob_bottom.channels(), - blob_bottom.height(), blob_bottom.width()); - const Dtype* bottom_data = blob_bottom.cpu_data(); - Dtype* top_data = blob_top->mutable_cpu_data(); - const Dtype alpha = layer_param.lrn_map_param().alpha(); - const Dtype beta = layer_param.lrn_map_param().beta(); - const int size = layer_param.lrn_map_param().local_size(); - for (int n = 0; n < blob_bottom.num(); ++n) { - for (int c = 0; c < blob_bottom.channels(); ++c) { - for (int h = 0; h < blob_bottom.height(); ++h) { - int h_start = h - (size - 1) / 2; - int h_end = min(h_start + size, blob_bottom.height()); - h_start = max(h_start, 0); - for (int w = 0; w < blob_bottom.width(); ++w) { - Dtype scale = 1.; - int w_start = w - (size - 1) / 2; - int w_end = min(w_start + size, blob_bottom.width()); - w_start = max(w_start, 0); - for (int nh = h_start; nh < h_end; ++nh) { - for (int nw = w_start; nw < w_end; ++nw) { - Dtype value = blob_bottom.data_at(n, c, nh, nw); - scale += value * value * alpha / (size * size); - } - } - *(top_data + blob_top->offset(n, c, h, w)) = - blob_bottom.data_at(n, c, h, w) / pow(scale, beta); - } - } - } - } -} - -typedef ::testing::Types Dtypes; -TYPED_TEST_CASE(LRNMapLayerTest, Dtypes); - -TYPED_TEST(LRNMapLayerTest, TestSetup) { - LayerParameter layer_param; - LRNMapLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - EXPECT_EQ(this->blob_top_->num(), 2); - EXPECT_EQ(this->blob_top_->channels(), 2); - EXPECT_EQ(this->blob_top_->height(), 7); - EXPECT_EQ(this->blob_top_->width(), 7); -} - -TYPED_TEST(LRNMapLayerTest, TestCPUForward) { - LayerParameter layer_param; - LRNMapLayer layer(layer_param); - Caffe::set_mode(Caffe::CPU); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - Blob top_reference; - this->ReferenceLRNMapForward(*(this->blob_bottom_), layer_param, - &top_reference); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - EXPECT_NEAR(this->blob_top_->cpu_data()[i], top_reference.cpu_data()[i], - this->epsilon_); - } -} - -TYPED_TEST(LRNMapLayerTest, TestGPUForward) { - LayerParameter layer_param; - LRNMapLayer layer(layer_param); - Caffe::set_mode(Caffe::GPU); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - Blob top_reference; - this->ReferenceLRNMapForward(*(this->blob_bottom_), layer_param, - &top_reference); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - EXPECT_NEAR(this->blob_top_->cpu_data()[i], top_reference.cpu_data()[i], - this->epsilon_); - } -} - -TYPED_TEST(LRNMapLayerTest, TestCPUGradient) { - LayerParameter layer_param; - LRNMapLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2); - Caffe::set_mode(Caffe::CPU); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - for (int i = 0; i < this->blob_top_->count(); ++i) { - this->blob_top_->mutable_cpu_diff()[i] = 1.; - } - layer.Backward(this->blob_top_vec_, true, &(this->blob_bottom_vec_)); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); -} - -TYPED_TEST(LRNMapLayerTest, TestGPUGradient) { - LayerParameter layer_param; - LRNMapLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2); - Caffe::set_mode(Caffe::GPU); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - for (int i = 0; i < this->blob_top_->count(); ++i) { - this->blob_top_->mutable_cpu_diff()[i] = 1.; - } - layer.Backward(this->blob_top_vec_, true, &(this->blob_bottom_vec_)); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); -} - -} // namespace caffe From 24f7318ef2a50e59d2fc8862f4c7426392b9c8a7 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 04:26:35 -0700 Subject: [PATCH 11/18] replace old cifar full with within channel LRN (per cuda-convnet layers-18pct) -- slightly slower (5000 iters now takes 6:57; took 6:43 previously), but slightly more accurate (exactly 82% test accuracy; got 81.65% before) --- ...ifar10_full_lrn_in_channel_solver.prototxt | 28 --- ...10_full_lrn_in_channel_solver_lr1.prototxt | 28 --- ...10_full_lrn_in_channel_solver_lr2.prototxt | 28 --- .../cifar10_full_lrn_in_channel_test.prototxt | 181 ------------------ ...cifar10_full_lrn_in_channel_train.prototxt | 174 ----------------- examples/cifar10/cifar10_full_test.prototxt | 2 + examples/cifar10/cifar10_full_train.prototxt | 2 + examples/cifar10/train_full.sh | 11 +- examples/cifar10/train_full_lrn_in_channel.sh | 16 -- 9 files changed, 12 insertions(+), 458 deletions(-) delete mode 100644 examples/cifar10/cifar10_full_lrn_in_channel_solver.prototxt delete mode 100644 examples/cifar10/cifar10_full_lrn_in_channel_solver_lr1.prototxt delete mode 100644 examples/cifar10/cifar10_full_lrn_in_channel_solver_lr2.prototxt delete mode 100644 examples/cifar10/cifar10_full_lrn_in_channel_test.prototxt delete mode 100644 examples/cifar10/cifar10_full_lrn_in_channel_train.prototxt delete mode 100755 examples/cifar10/train_full_lrn_in_channel.sh diff --git a/examples/cifar10/cifar10_full_lrn_in_channel_solver.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_solver.prototxt deleted file mode 100644 index a217ec1a0dd..00000000000 --- a/examples/cifar10/cifar10_full_lrn_in_channel_solver.prototxt +++ /dev/null @@ -1,28 +0,0 @@ -# reduce learning rate after 120 epochs (60000 iters) by factor 0f 10 -# then another factor of 10 after 10 more epochs (5000 iters) - -# The training protocol buffer definition -train_net: "cifar10_full_lrn_in_channel_train.prototxt" -# The testing protocol buffer definition -test_net: "cifar10_full_lrn_in_channel_test.prototxt" -# test_iter specifies how many forward passes the test should carry out. -# In the case of CIFAR10, we have test batch size 100 and 100 test iterations, -# covering the full 10,000 testing images. -test_iter: 100 -# Carry out testing every 1000 training iterations. -test_interval: 1000 -# The base learning rate, momentum and the weight decay of the network. -base_lr: 0.001 -momentum: 0.9 -weight_decay: 0.004 -# The learning rate policy -lr_policy: "fixed" -# Display every 200 iterations -display: 200 -# The maximum number of iterations -max_iter: 60000 -# snapshot intermediate results -snapshot: 10000 -snapshot_prefix: "cifar10_full_lrn_in_channel" -# solver mode: 0 for CPU and 1 for GPU -solver_mode: 1 diff --git a/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr1.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr1.prototxt deleted file mode 100644 index 6fcae73e7e9..00000000000 --- a/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr1.prototxt +++ /dev/null @@ -1,28 +0,0 @@ -# reduce learning rate after 120 epochs (60000 iters) by factor 0f 10 -# then another factor of 10 after 10 more epochs (5000 iters) - -# The training protocol buffer definition -train_net: "cifar10_full_lrn_in_channel_train.prototxt" -# The testing protocol buffer definition -test_net: "cifar10_full_lrn_in_channel_test.prototxt" -# test_iter specifies how many forward passes the test should carry out. -# In the case of CIFAR10, we have test batch size 100 and 100 test iterations, -# covering the full 10,000 testing images. -test_iter: 100 -# Carry out testing every 1000 training iterations. -test_interval: 1000 -# The base learning rate, momentum and the weight decay of the network. -base_lr: 0.0001 -momentum: 0.9 -weight_decay: 0.004 -# The learning rate policy -lr_policy: "fixed" -# Display every 200 iterations -display: 200 -# The maximum number of iterations -max_iter: 65000 -# snapshot intermediate results -snapshot: 5000 -snapshot_prefix: "cifar10_full_lrn_in_channel" -# solver mode: 0 for CPU and 1 for GPU -solver_mode: 1 diff --git a/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr2.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr2.prototxt deleted file mode 100644 index 8f53fe11a67..00000000000 --- a/examples/cifar10/cifar10_full_lrn_in_channel_solver_lr2.prototxt +++ /dev/null @@ -1,28 +0,0 @@ -# reduce learning rate after 120 epochs (60000 iters) by factor 0f 10 -# then another factor of 10 after 10 more epochs (5000 iters) - -# The training protocol buffer definition -train_net: "cifar10_full_lrn_in_channel_train.prototxt" -# The testing protocol buffer definition -test_net: "cifar10_full_lrn_in_channel_test.prototxt" -# test_iter specifies how many forward passes the test should carry out. -# In the case of CIFAR10, we have test batch size 100 and 100 test iterations, -# covering the full 10,000 testing images. -test_iter: 100 -# Carry out testing every 1000 training iterations. -test_interval: 1000 -# The base learning rate, momentum and the weight decay of the network. -base_lr: 0.00001 -momentum: 0.9 -weight_decay: 0.004 -# The learning rate policy -lr_policy: "fixed" -# Display every 200 iterations -display: 200 -# The maximum number of iterations -max_iter: 70000 -# snapshot intermediate results -snapshot: 5000 -snapshot_prefix: "cifar10_full_lrn_in_channel" -# solver mode: 0 for CPU and 1 for GPU -solver_mode: 1 diff --git a/examples/cifar10/cifar10_full_lrn_in_channel_test.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_test.prototxt deleted file mode 100644 index 0e1957a9045..00000000000 --- a/examples/cifar10/cifar10_full_lrn_in_channel_test.prototxt +++ /dev/null @@ -1,181 +0,0 @@ -name: "CIFAR10_full_test" -layers { - name: "cifar" - type: DATA - top: "data" - top: "label" - data_param { - source: "cifar10-leveldb/cifar-test-leveldb" - mean_file: "mean.binaryproto" - batch_size: 100 - } -} -layers { - name: "conv1" - type: CONVOLUTION - bottom: "data" - top: "conv1" - blobs_lr: 1 - blobs_lr: 2 - convolution_param { - num_output: 32 - pad: 2 - kernel_size: 5 - stride: 1 - weight_filler { - type: "gaussian" - std: 0.0001 - } - bias_filler { - type: "constant" - } - } -} -layers { - name: "pool1" - type: POOLING - bottom: "conv1" - top: "pool1" - pooling_param { - pool: MAX - kernel_size: 3 - stride: 2 - } -} -layers { - name: "relu1" - type: RELU - bottom: "pool1" - top: "pool1" -} -layers { - name: "norm1" - type: LRN - bottom: "pool1" - top: "norm1" - lrn_param { - norm_region: WITHIN_CHANNEL - local_size: 3 - alpha: 5e-05 - beta: 0.75 - } -} -layers { - name: "conv2" - type: CONVOLUTION - bottom: "norm1" - top: "conv2" - blobs_lr: 1 - blobs_lr: 2 - convolution_param { - num_output: 32 - pad: 2 - kernel_size: 5 - stride: 1 - weight_filler { - type: "gaussian" - std: 0.01 - } - bias_filler { - type: "constant" - } - } -} -layers { - name: "relu2" - type: RELU - bottom: "conv2" - top: "conv2" -} -layers { - name: "pool2" - type: POOLING - bottom: "conv2" - top: "pool2" - pooling_param { - pool: AVE - kernel_size: 3 - stride: 2 - } -} -layers { - name: "norm2" - type: LRN - bottom: "pool2" - top: "norm2" - lrn_param { - norm_region: WITHIN_CHANNEL - local_size: 3 - alpha: 5e-05 - beta: 0.75 - } -} -layers { - name: "conv3" - type: CONVOLUTION - bottom: "norm2" - top: "conv3" - convolution_param { - num_output: 64 - pad: 2 - kernel_size: 5 - stride: 1 - weight_filler { - type: "gaussian" - std: 0.01 - } - bias_filler { - type: "constant" - } - } -} -layers { - name: "relu3" - type: RELU - bottom: "conv3" - top: "conv3" -} -layers { - name: "pool3" - type: POOLING - bottom: "conv3" - top: "pool3" - pooling_param { - pool: AVE - kernel_size: 3 - stride: 2 - } -} -layers { - name: "ip1" - type: INNER_PRODUCT - bottom: "pool3" - top: "ip1" - blobs_lr: 1 - blobs_lr: 2 - weight_decay: 250 - weight_decay: 0 - inner_product_param { - num_output: 10 - weight_filler { - type: "gaussian" - std: 0.01 - } - bias_filler { - type: "constant" - } - } -} -layers { - name: "prob" - type: SOFTMAX - bottom: "ip1" - top: "prob" -} -layers { - name: "accuracy" - type: ACCURACY - bottom: "prob" - bottom: "label" - top: "accuracy" -} diff --git a/examples/cifar10/cifar10_full_lrn_in_channel_train.prototxt b/examples/cifar10/cifar10_full_lrn_in_channel_train.prototxt deleted file mode 100644 index 25c76060991..00000000000 --- a/examples/cifar10/cifar10_full_lrn_in_channel_train.prototxt +++ /dev/null @@ -1,174 +0,0 @@ -name: "CIFAR10_full_train" -layers { - name: "cifar" - type: DATA - top: "data" - top: "label" - data_param { - source: "cifar10-leveldb/cifar-train-leveldb" - mean_file: "mean.binaryproto" - batch_size: 100 - } -} -layers { - name: "conv1" - type: CONVOLUTION - bottom: "data" - top: "conv1" - blobs_lr: 1 - blobs_lr: 2 - convolution_param { - num_output: 32 - pad: 2 - kernel_size: 5 - stride: 1 - weight_filler { - type: "gaussian" - std: 0.0001 - } - bias_filler { - type: "constant" - } - } -} -layers { - name: "pool1" - type: POOLING - bottom: "conv1" - top: "pool1" - pooling_param { - pool: MAX - kernel_size: 3 - stride: 2 - } -} -layers { - name: "relu1" - type: RELU - bottom: "pool1" - top: "pool1" -} -layers { - name: "norm1" - type: LRN - bottom: "pool1" - top: "norm1" - lrn_param { - norm_region: WITHIN_CHANNEL - local_size: 3 - alpha: 5e-05 - beta: 0.75 - } -} -layers { - name: "conv2" - type: CONVOLUTION - bottom: "norm1" - top: "conv2" - blobs_lr: 1 - blobs_lr: 2 - convolution_param { - num_output: 32 - pad: 2 - kernel_size: 5 - stride: 1 - weight_filler { - type: "gaussian" - std: 0.01 - } - bias_filler { - type: "constant" - } - } -} -layers { - name: "relu2" - type: RELU - bottom: "conv2" - top: "conv2" -} -layers { - name: "pool2" - type: POOLING - bottom: "conv2" - top: "pool2" - pooling_param { - pool: AVE - kernel_size: 3 - stride: 2 - } -} -layers { - name: "norm2" - type: LRN - bottom: "pool2" - top: "norm2" - lrn_param { - norm_region: WITHIN_CHANNEL - local_size: 3 - alpha: 5e-05 - beta: 0.75 - } -} -layers { - name: "conv3" - type: CONVOLUTION - bottom: "norm2" - top: "conv3" - convolution_param { - num_output: 64 - pad: 2 - kernel_size: 5 - stride: 1 - weight_filler { - type: "gaussian" - std: 0.01 - } - bias_filler { - type: "constant" - } - } -} -layers { - name: "relu3" - type: RELU - bottom: "conv3" - top: "conv3" -} -layers { - name: "pool3" - type: POOLING - bottom: "conv3" - top: "pool3" - pooling_param { - pool: AVE - kernel_size: 3 - stride: 2 - } -} -layers { - name: "ip1" - type: INNER_PRODUCT - bottom: "pool3" - top: "ip1" - blobs_lr: 1 - blobs_lr: 2 - weight_decay: 250 - weight_decay: 0 - inner_product_param { - num_output: 10 - weight_filler { - type: "gaussian" - std: 0.01 - } - bias_filler { - type: "constant" - } - } -} -layers { - name: "loss" - type: SOFTMAX_LOSS - bottom: "ip1" - bottom: "label" -} diff --git a/examples/cifar10/cifar10_full_test.prototxt b/examples/cifar10/cifar10_full_test.prototxt index ada373a55cb..0e1957a9045 100644 --- a/examples/cifar10/cifar10_full_test.prototxt +++ b/examples/cifar10/cifar10_full_test.prototxt @@ -54,6 +54,7 @@ layers { bottom: "pool1" top: "norm1" lrn_param { + norm_region: WITHIN_CHANNEL local_size: 3 alpha: 5e-05 beta: 0.75 @@ -103,6 +104,7 @@ layers { bottom: "pool2" top: "norm2" lrn_param { + norm_region: WITHIN_CHANNEL local_size: 3 alpha: 5e-05 beta: 0.75 diff --git a/examples/cifar10/cifar10_full_train.prototxt b/examples/cifar10/cifar10_full_train.prototxt index 56c9306e2dc..25c76060991 100644 --- a/examples/cifar10/cifar10_full_train.prototxt +++ b/examples/cifar10/cifar10_full_train.prototxt @@ -54,6 +54,7 @@ layers { bottom: "pool1" top: "norm1" lrn_param { + norm_region: WITHIN_CHANNEL local_size: 3 alpha: 5e-05 beta: 0.75 @@ -103,6 +104,7 @@ layers { bottom: "pool2" top: "norm2" lrn_param { + norm_region: WITHIN_CHANNEL local_size: 3 alpha: 5e-05 beta: 0.75 diff --git a/examples/cifar10/train_full.sh b/examples/cifar10/train_full.sh index 1767da6798d..4db7b9a98f1 100755 --- a/examples/cifar10/train_full.sh +++ b/examples/cifar10/train_full.sh @@ -2,10 +2,15 @@ TOOLS=../../build/tools -GLOG_logtostderr=1 $TOOLS/train_net.bin cifar10_full_solver.prototxt +GLOG_logtostderr=1 $TOOLS/train_net.bin \ + cifar10_full_solver.prototxt #reduce learning rate by factor of 10 -GLOG_logtostderr=1 $TOOLS/train_net.bin cifar10_full_solver_lr1.prototxt cifar10_full_iter_60000.solverstate +GLOG_logtostderr=1 $TOOLS/train_net.bin \ + cifar10_full_solver_lr1.prototxt \ + cifar10_full_iter_60000.solverstate #reduce learning rate by factor of 10 -GLOG_logtostderr=1 $TOOLS/train_net.bin cifar10_full_solver_lr2.prototxt cifar10_full_iter_65000.solverstate +GLOG_logtostderr=1 $TOOLS/train_net.bin \ + cifar10_full_solver_lr2.prototxt \ + cifar10_full_iter_65000.solverstate diff --git a/examples/cifar10/train_full_lrn_in_channel.sh b/examples/cifar10/train_full_lrn_in_channel.sh deleted file mode 100755 index fcc9f4d0f92..00000000000 --- a/examples/cifar10/train_full_lrn_in_channel.sh +++ /dev/null @@ -1,16 +0,0 @@ -#!/usr/bin/env sh - -TOOLS=../../build/tools - -GLOG_logtostderr=1 $TOOLS/train_net.bin \ - cifar10_full_lrn_in_channel_solver.prototxt - -#reduce learning rate by factor of 10 -GLOG_logtostderr=1 $TOOLS/train_net.bin \ - cifar10_full_lrn_in_channel_solver_lr1.prototxt \ - cifar10_full_lrn_in_channel_iter_60000.solverstate - -#reduce learning rate by factor of 10 -GLOG_logtostderr=1 $TOOLS/train_net.bin \ - cifar10_full_lrn_in_channel_solver_lr2.prototxt \ - cifar10_full_lrn_in_channel_iter_65000.solverstate From 69ac4f4366e2388afaa491fb6105570489523aaf Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 12:01:11 -0700 Subject: [PATCH 12/18] minor polishing --- src/caffe/layers/lrn_layer.cpp | 14 +++++++------- src/caffe/proto/caffe.proto | 1 + src/caffe/test/test_pooling_layer.cpp | 4 ++-- 3 files changed, 10 insertions(+), 9 deletions(-) diff --git a/src/caffe/layers/lrn_layer.cpp b/src/caffe/layers/lrn_layer.cpp index d9be541940e..98f95041fbd 100644 --- a/src/caffe/layers/lrn_layer.cpp +++ b/src/caffe/layers/lrn_layer.cpp @@ -31,14 +31,14 @@ void LRNLayer::SetUp(const vector*>& bottom, case LRNParameter_NormRegion_WITHIN_CHANNEL: { const Dtype pre_pad = (size_ - 1) / 2; - // Set up split layer to use inputs in the num_erator and denominator. + // Set up split_layer_ to use inputs in the numerator and denominator. split_top_vec_.clear(); split_top_vec_.push_back(bottom[0]); split_top_vec_.push_back(&square_input_); LayerParameter split_param; split_layer_.reset(new SplitLayer(split_param)); split_layer_->SetUp(bottom, &split_top_vec_); - // Set up square layer to square the inputs. + // Set up square_layer_ to square the inputs. square_input_.Reshape(num_, channels_, height_, width_); square_bottom_vec_.clear(); square_top_vec_.clear(); @@ -52,7 +52,7 @@ void LRNLayer::SetUp(const vector*>& bottom, CHECK_EQ(square_output_.channels(), channels_); CHECK_EQ(square_output_.height(), height_); CHECK_EQ(square_output_.width(), width_); - // Output of pool layer gives us the neighborhood response. + // Set up pool_layer_ to sum over square neighborhoods of the input. pool_top_vec_.clear(); pool_top_vec_.push_back(&pool_output_); LayerParameter pool_param; @@ -66,8 +66,8 @@ void LRNLayer::SetUp(const vector*>& bottom, CHECK_EQ(pool_output_.channels(), channels_); CHECK_EQ(pool_output_.height(), height_); CHECK_EQ(pool_output_.width(), width_); - // Set up power layer to compute (1 + alpha_/N^2 s)^-beta_, where s is the - // sum of a squared neighborhood (as output by pool_layer_). + // Set up power_layer_ to compute (1 + alpha_/N^2 s)^-beta_, where s is the + // sum of a squared neighborhood (the output of pool_layer_). power_top_vec_.clear(); power_top_vec_.push_back(&power_output_); LayerParameter power_param; @@ -80,8 +80,8 @@ void LRNLayer::SetUp(const vector*>& bottom, CHECK_EQ(power_output_.channels(), channels_); CHECK_EQ(power_output_.height(), height_); CHECK_EQ(power_output_.width(), width_); - // Set up a product layer to compute outputs by multiplying inputs by the - // demoninator computed by the power layer. + // Set up a product_layer_ to compute outputs by multiplying inputs by the + // inverse demoninator computed by the power layer. product_bottom_vec_.clear(); product_bottom_vec_.push_back(bottom[0]); product_bottom_vec_.push_back(&power_output_); diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 745763cc358..6b54bbb0faa 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -292,6 +292,7 @@ message PoolingParameter { // Message that stores parameters used by PowerLayer message PowerParameter { + // PowerLayer computes outputs y = (shift + scale * x) ^ power. optional float power = 1 [default = 1.0]; optional float scale = 2 [default = 1.0]; optional float shift = 3 [default = 0.0]; diff --git a/src/caffe/test/test_pooling_layer.cpp b/src/caffe/test/test_pooling_layer.cpp index a57110491ec..41d4841f444 100644 --- a/src/caffe/test/test_pooling_layer.cpp +++ b/src/caffe/test/test_pooling_layer.cpp @@ -151,7 +151,7 @@ TYPED_TEST(PoolingLayerTest, TestCPUForwardAve) { EXPECT_NEAR(this->blob_top_->cpu_data()[1], 4.0 / 3, epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[2], 8.0 / 9, epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[3], 4.0 / 3, epsilon); - EXPECT_NEAR(this->blob_top_->cpu_data()[4], 2.0, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[4], 2.0 , epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[5], 4.0 / 3, epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[6], 8.0 / 9, epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[7], 4.0 / 3, epsilon); @@ -184,7 +184,7 @@ TYPED_TEST(PoolingLayerTest, TestGPUForwardAve) { EXPECT_NEAR(this->blob_top_->cpu_data()[1], 4.0 / 3, epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[2], 8.0 / 9, epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[3], 4.0 / 3, epsilon); - EXPECT_NEAR(this->blob_top_->cpu_data()[4], 2.0, epsilon); + EXPECT_NEAR(this->blob_top_->cpu_data()[4], 2.0 , epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[5], 4.0 / 3, epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[6], 8.0 / 9, epsilon); EXPECT_NEAR(this->blob_top_->cpu_data()[7], 4.0 / 3, epsilon); From d61adb8212e027a164da87bb83bcc18367a8a9f9 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 12:23:03 -0700 Subject: [PATCH 13/18] remove unnecessary local variables from EltwiseProductLayer --- include/caffe/vision_layers.hpp | 5 ----- src/caffe/layers/eltwise_product_layer.cpp | 18 +++++++++--------- 2 files changed, 9 insertions(+), 14 deletions(-) diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index 93f6d25296a..bbad3193f0b 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -287,11 +287,6 @@ class EltwiseProductLayer : public Layer { const bool propagate_down, vector*>* bottom); virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - - int num_; - int channels_; - int height_; - int width_; }; template diff --git a/src/caffe/layers/eltwise_product_layer.cpp b/src/caffe/layers/eltwise_product_layer.cpp index d056ab1350b..b394450d6ae 100644 --- a/src/caffe/layers/eltwise_product_layer.cpp +++ b/src/caffe/layers/eltwise_product_layer.cpp @@ -15,17 +15,17 @@ void EltwiseProductLayer::SetUp(const vector*>& bottom, "Eltwise Product Layer takes at least 2 blobs as input."; CHECK_EQ(top->size(), 1) << "Eltwise Product Layer takes a single blob as output."; - num_ = bottom[0]->num(); - channels_ = bottom[0]->channels(); - height_ = bottom[0]->height(); - width_ = bottom[0]->width(); + const int num = bottom[0]->num(); + const int channels = bottom[0]->channels(); + const int height = bottom[0]->height(); + const int width = bottom[0]->width(); for (int i = 1; i < bottom.size(); ++i) { - CHECK_EQ(num_, bottom[i]->num()); - CHECK_EQ(channels_, bottom[i]->channels()); - CHECK_EQ(height_, bottom[i]->height()); - CHECK_EQ(width_, bottom[i]->width()); + CHECK_EQ(num, bottom[i]->num()); + CHECK_EQ(channels, bottom[i]->channels()); + CHECK_EQ(height, bottom[i]->height()); + CHECK_EQ(width, bottom[i]->width()); } - (*top)[0]->Reshape(num_, channels_, height_, width_); + (*top)[0]->Reshape(num, channels, height, width); } template From 9fb7818cf0484cd62197714e91815141097709de Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 12:31:49 -0700 Subject: [PATCH 14/18] don't recompute pre_pad --- src/caffe/layers/lrn_layer.cpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/src/caffe/layers/lrn_layer.cpp b/src/caffe/layers/lrn_layer.cpp index 98f95041fbd..d4addb5e1a5 100644 --- a/src/caffe/layers/lrn_layer.cpp +++ b/src/caffe/layers/lrn_layer.cpp @@ -30,7 +30,6 @@ void LRNLayer::SetUp(const vector*>& bottom, break; case LRNParameter_NormRegion_WITHIN_CHANNEL: { - const Dtype pre_pad = (size_ - 1) / 2; // Set up split_layer_ to use inputs in the numerator and denominator. split_top_vec_.clear(); split_top_vec_.push_back(bottom[0]); @@ -58,7 +57,7 @@ void LRNLayer::SetUp(const vector*>& bottom, LayerParameter pool_param; pool_param.mutable_pooling_param()->set_pool( PoolingParameter_PoolMethod_AVE); - pool_param.mutable_pooling_param()->set_pad(pre_pad); + pool_param.mutable_pooling_param()->set_pad(pre_pad_); pool_param.mutable_pooling_param()->set_kernel_size(size_); pool_layer_.reset(new PoolingLayer(pool_param)); pool_layer_->SetUp(square_top_vec_, &pool_top_vec_); @@ -66,8 +65,8 @@ void LRNLayer::SetUp(const vector*>& bottom, CHECK_EQ(pool_output_.channels(), channels_); CHECK_EQ(pool_output_.height(), height_); CHECK_EQ(pool_output_.width(), width_); - // Set up power_layer_ to compute (1 + alpha_/N^2 s)^-beta_, where s is the - // sum of a squared neighborhood (the output of pool_layer_). + // Set up power_layer_ to compute (1 + alpha_/N^2 s)^-beta_, where s is + // the sum of a squared neighborhood (the output of pool_layer_). power_top_vec_.clear(); power_top_vec_.push_back(&power_output_); LayerParameter power_param; From 45f8626bae88df646f9e5047ee3c28fd5a01590d Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 13:12:12 -0700 Subject: [PATCH 15/18] cleanup extra LRN method names --- include/caffe/vision_layers.hpp | 12 ++++++------ src/caffe/layers/lrn_layer.cpp | 16 ++++++++-------- src/caffe/layers/lrn_layer.cu | 12 ++++++------ 3 files changed, 20 insertions(+), 20 deletions(-) diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index bbad3193f0b..feda44d85df 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -517,17 +517,17 @@ class LRNLayer : public Layer { virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Forward_cpu_cross_channel(const vector*>& bottom, + virtual Dtype CrossChannelForward_cpu(const vector*>& bottom, vector*>* top); - virtual Dtype Forward_gpu_cross_channel(const vector*>& bottom, + virtual Dtype CrossChannelForward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Forward_within_channel(const vector*>& bottom, + virtual Dtype WithinChannelForward(const vector*>& bottom, vector*>* top); - virtual void Backward_cpu_cross_channel(const vector*>& top, + virtual void CrossChannelBackward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual void Backward_gpu_cross_channel(const vector*>& top, + virtual void CrossChannelBackward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual void Backward_within_channel(const vector*>& top, + virtual void WithinChannelBackward(const vector*>& top, const bool propagate_down, vector*>* bottom); int size_; diff --git a/src/caffe/layers/lrn_layer.cpp b/src/caffe/layers/lrn_layer.cpp index d4addb5e1a5..cfcc59c9feb 100644 --- a/src/caffe/layers/lrn_layer.cpp +++ b/src/caffe/layers/lrn_layer.cpp @@ -103,9 +103,9 @@ Dtype LRNLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { switch (this->layer_param_.lrn_param().norm_region()) { case LRNParameter_NormRegion_ACROSS_CHANNELS: - return Forward_cpu_cross_channel(bottom, top); + return CrossChannelForward_cpu(bottom, top); case LRNParameter_NormRegion_WITHIN_CHANNEL: - return Forward_within_channel(bottom, top); + return WithinChannelForward(bottom, top); default: LOG(FATAL) << "Unknown normalization region."; return Dtype(0); @@ -113,7 +113,7 @@ Dtype LRNLayer::Forward_cpu(const vector*>& bottom, } template -Dtype LRNLayer::Forward_cpu_cross_channel( +Dtype LRNLayer::CrossChannelForward_cpu( const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -162,7 +162,7 @@ Dtype LRNLayer::Forward_cpu_cross_channel( } template -Dtype LRNLayer::Forward_within_channel( +Dtype LRNLayer::WithinChannelForward( const vector*>& bottom, vector*>* top) { split_layer_->Forward(bottom, &split_top_vec_); square_layer_->Forward(square_bottom_vec_, &square_top_vec_); @@ -177,10 +177,10 @@ void LRNLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { switch (this->layer_param_.lrn_param().norm_region()) { case LRNParameter_NormRegion_ACROSS_CHANNELS: - Backward_cpu_cross_channel(top, propagate_down, bottom); + CrossChannelBackward_cpu(top, propagate_down, bottom); break; case LRNParameter_NormRegion_WITHIN_CHANNEL: - Backward_within_channel(top, propagate_down, bottom); + WithinChannelBackward(top, propagate_down, bottom); break; default: LOG(FATAL) << "Unknown normalization region."; @@ -188,7 +188,7 @@ void LRNLayer::Backward_cpu(const vector*>& top, } template -void LRNLayer::Backward_cpu_cross_channel( +void LRNLayer::CrossChannelBackward_cpu( const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->cpu_diff(); @@ -243,7 +243,7 @@ void LRNLayer::Backward_cpu_cross_channel( } template -void LRNLayer::Backward_within_channel( +void LRNLayer::WithinChannelBackward( const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { diff --git a/src/caffe/layers/lrn_layer.cu b/src/caffe/layers/lrn_layer.cu index 47eca1765de..b2097eb99cf 100644 --- a/src/caffe/layers/lrn_layer.cu +++ b/src/caffe/layers/lrn_layer.cu @@ -60,9 +60,9 @@ Dtype LRNLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { switch (this->layer_param_.lrn_param().norm_region()) { case LRNParameter_NormRegion_ACROSS_CHANNELS: - return Forward_gpu_cross_channel(bottom, top); + return CrossChannelForward_gpu(bottom, top); case LRNParameter_NormRegion_WITHIN_CHANNEL: - return Forward_within_channel(bottom, top); + return WithinChannelForward(bottom, top); default: LOG(FATAL) << "Unknown normalization region."; return Dtype(0); @@ -79,7 +79,7 @@ __global__ void LRNComputeOutput(const int nthreads, const Dtype* in, } template -Dtype LRNLayer::Forward_gpu_cross_channel( +Dtype LRNLayer::CrossChannelForward_gpu( const vector*>& bottom, vector*>* top) { // First, compute scale const Dtype* bottom_data = bottom[0]->gpu_data(); @@ -107,10 +107,10 @@ void LRNLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { switch (this->layer_param_.lrn_param().norm_region()) { case LRNParameter_NormRegion_ACROSS_CHANNELS: - Backward_gpu_cross_channel(top, propagate_down, bottom); + CrossChannelBackward_gpu(top, propagate_down, bottom); break; case LRNParameter_NormRegion_WITHIN_CHANNEL: - Backward_within_channel(top, propagate_down, bottom); + WithinChannelBackward(top, propagate_down, bottom); break; default: LOG(FATAL) << "Unknown normalization region."; @@ -179,7 +179,7 @@ __global__ void LRNComputeDiff(const int nthreads, const Dtype* bottom_data, } template -void LRNLayer::Backward_gpu_cross_channel( +void LRNLayer::CrossChannelBackward_gpu( const vector*>& top, const bool propagate_down, vector*>* bottom) { int n_threads = num_ * height_ * width_; From 86db2a9fadfea399d5ee382e5f30f33da04a5b6f Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 13:21:36 -0700 Subject: [PATCH 16/18] minor unit test cleanup --- src/caffe/test/test_eltwise_product_layer.cpp | 16 ++++++---------- src/caffe/test/test_power_layer.cpp | 10 ++++------ 2 files changed, 10 insertions(+), 16 deletions(-) diff --git a/src/caffe/test/test_eltwise_product_layer.cpp b/src/caffe/test/test_eltwise_product_layer.cpp index 8766f39da86..8255a579b15 100644 --- a/src/caffe/test/test_eltwise_product_layer.cpp +++ b/src/caffe/test/test_eltwise_product_layer.cpp @@ -107,16 +107,12 @@ TYPED_TEST(EltwiseProductLayerTest, TestCPUGradient) { } TYPED_TEST(EltwiseProductLayerTest, TestGPUGradient) { - if (sizeof(TypeParam) == 4 || CAFFE_TEST_CUDA_PROP.major >= 2) { - Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; - EltwiseProductLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); - } else { - LOG(ERROR) << "Skipping test due to old architecture."; - } + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + EltwiseProductLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); } } // namespace caffe diff --git a/src/caffe/test/test_power_layer.cpp b/src/caffe/test/test_power_layer.cpp index 40734b65fbf..9d2948f509b 100644 --- a/src/caffe/test/test_power_layer.cpp +++ b/src/caffe/test/test_power_layer.cpp @@ -60,8 +60,7 @@ TYPED_TEST(PowerLayerTest, TestPowerCPU) { if (isnan(expected_value)) { EXPECT_TRUE(isnan(top_data[i])); } else { - TypeParam precision = expected_value * 0.0001; - precision *= (precision < 0) ? -1 : 1; + TypeParam precision = abs(expected_value * 0.0001); EXPECT_NEAR(expected_value, top_data[i], precision); } } @@ -103,7 +102,7 @@ TYPED_TEST(PowerLayerTest, TestPowerGradientShiftZeroCPU) { // Flip negative values in bottom vector as x < 0 -> x^0.37 = nan TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); for (int i = 0; i < this->blob_bottom_->count(); ++i) { - bottom_data[i] *= (bottom_data[i] < 0) ? -1 : 1; + bottom_data[i] = abs(bottom_data[i]); } GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), @@ -253,8 +252,7 @@ TYPED_TEST(PowerLayerTest, TestPowerGPU) { if (isnan(expected_value)) { EXPECT_TRUE(isnan(top_data[i])); } else { - TypeParam precision = expected_value * 0.0001; - precision *= (precision < 0) ? -1 : 1; + TypeParam precision = abs(expected_value * 0.0001); EXPECT_NEAR(expected_value, top_data[i], precision); } } @@ -296,7 +294,7 @@ TYPED_TEST(PowerLayerTest, TestPowerGradientShiftZeroGPU) { // Flip negative values in bottom vector as x < 0 -> x^0.37 = nan TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); for (int i = 0; i < this->blob_bottom_->count(); ++i) { - bottom_data[i] *= (bottom_data[i] < 0) ? -1 : 1; + bottom_data[i] = abs(bottom_data[i]); } GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), From bf511f7bdfd40577917b3797bb84ccc1de8db6f8 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 29 Mar 2014 13:46:20 -0700 Subject: [PATCH 17/18] cleanup power layer test suite --- src/caffe/test/test_power_layer.cpp | 314 +++++++--------------------- 1 file changed, 70 insertions(+), 244 deletions(-) diff --git a/src/caffe/test/test_power_layer.cpp b/src/caffe/test/test_power_layer.cpp index 9d2948f509b..451c415f143 100644 --- a/src/caffe/test/test_power_layer.cpp +++ b/src/caffe/test/test_power_layer.cpp @@ -1,5 +1,6 @@ // Copyright 2014 BVLC and contributors. +#include #include #include "cuda_runtime.h" @@ -31,6 +32,55 @@ class PowerLayerTest : public ::testing::Test { blob_top_vec_.push_back(blob_top_); } virtual ~PowerLayerTest() { delete blob_bottom_; delete blob_top_; } + + void TestForward(Dtype power, Dtype scale, Dtype shift) { + LayerParameter layer_param; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + // Now, check values + const Dtype* bottom_data = this->blob_bottom_->cpu_data(); + const Dtype* top_data = this->blob_top_->cpu_data(); + const Dtype min_precision = 1e-5; + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + Dtype expected_value = pow(shift + scale * bottom_data[i], power); + if (power == Dtype(0) || power == Dtype(1) || power == Dtype(2)) { + EXPECT_FALSE(isnan(top_data[i])); + } + if (isnan(expected_value)) { + EXPECT_TRUE(isnan(top_data[i])); + } else { + Dtype precision = max(Dtype(abs(expected_value * 0.0001)), + min_precision); + EXPECT_NEAR(expected_value, top_data[i], precision); + } + } + } + + void TestBackward(Dtype power, Dtype scale, Dtype shift) { + LayerParameter layer_param; + layer_param.mutable_power_param()->set_power(power); + layer_param.mutable_power_param()->set_scale(scale); + layer_param.mutable_power_param()->set_shift(shift); + PowerLayer layer(layer_param); + if (power != Dtype(0) && power != Dtype(1) && power != Dtype(2)) { + // Avoid NaNs by forcing (shift + scale * x) >= 0 + Dtype* bottom_data = this->blob_bottom_->mutable_cpu_data(); + Dtype min_value = -shift / scale; + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + if (bottom_data[i] < min_value) { + bottom_data[i] = min_value + (min_value - bottom_data[i]); + } + } + } + GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); + } + Blob* const blob_bottom_; Blob* const blob_top_; vector*> blob_bottom_vec_; @@ -42,386 +92,162 @@ TYPED_TEST_CASE(PowerLayerTest, Dtypes); TYPED_TEST(PowerLayerTest, TestPowerCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 0.37; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - // Now, check values - const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); - const TypeParam* top_data = this->blob_top_->cpu_data(); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - TypeParam expected_value = pow(shift + scale * bottom_data[i], power); - if (isnan(expected_value)) { - EXPECT_TRUE(isnan(top_data[i])); - } else { - TypeParam precision = abs(expected_value * 0.0001); - EXPECT_NEAR(expected_value, top_data[i], precision); - } - } + TestForward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerGradientCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 0.37; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - // Avoid NaNs by forcing (shift + scale * x) >= 0 - TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); - TypeParam min_value = -shift / scale; - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - if (bottom_data[i] < min_value) { - bottom_data[i] = min_value + (min_value - bottom_data[i]); - } - } - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerGradientShiftZeroCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 0.37; TypeParam scale = 0.83; TypeParam shift = 0.0; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - // Flip negative values in bottom vector as x < 0 -> x^0.37 = nan - TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - bottom_data[i] = abs(bottom_data[i]); - } - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerZeroCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 0.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - // Now, check values - const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); - const TypeParam* top_data = this->blob_top_->cpu_data(); - TypeParam expected_value = TypeParam(1); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - EXPECT_EQ(expected_value, top_data[i]); - } + TestForward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerZeroGradientCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 0.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerOneCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 1.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - // Now, check values - const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); - const TypeParam* top_data = this->blob_top_->cpu_data(); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - TypeParam expected_value = shift + scale * bottom_data[i]; - EXPECT_NEAR(expected_value, top_data[i], 0.001); - } + TestForward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerOneGradientCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 1.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerTwoCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 2.0; TypeParam scale = 0.34; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - // Now, check values - const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); - const TypeParam* top_data = this->blob_top_->cpu_data(); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - TypeParam expected_value = pow(shift + scale * bottom_data[i], 2); - EXPECT_NEAR(expected_value, top_data[i], 0.001); - } + TestForward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerTwoGradientCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 2.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerTwoScaleHalfGradientCPU) { Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; TypeParam power = 2.0; TypeParam scale = 0.5; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 0.37; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - // Now, check values - const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); - const TypeParam* top_data = this->blob_top_->cpu_data(); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - TypeParam expected_value = pow(shift + scale * bottom_data[i], power); - if (isnan(expected_value)) { - EXPECT_TRUE(isnan(top_data[i])); - } else { - TypeParam precision = abs(expected_value * 0.0001); - EXPECT_NEAR(expected_value, top_data[i], precision); - } - } + TestForward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerGradientGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 0.37; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - // Avoid NaNs by forcing (shift + scale * x) >= 0 - TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); - TypeParam min_value = -shift / scale; - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - if (bottom_data[i] < min_value) { - bottom_data[i] = min_value + (min_value - bottom_data[i]); - } - } - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerGradientShiftZeroGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 0.37; TypeParam scale = 0.83; TypeParam shift = 0.0; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - // Flip negative values in bottom vector as x < 0 -> x^0.37 = nan - TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data(); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - bottom_data[i] = abs(bottom_data[i]); - } - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerZeroGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 0.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - // Now, check values - const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); - const TypeParam* top_data = this->blob_top_->cpu_data(); - TypeParam expected_value = TypeParam(1); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - EXPECT_EQ(expected_value, top_data[i]); - } + TestForward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerZeroGradientGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 0.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerOneGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 1.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - // Now, check values - const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); - const TypeParam* top_data = this->blob_top_->cpu_data(); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - TypeParam expected_value = shift + scale * bottom_data[i]; - EXPECT_NEAR(expected_value, top_data[i], 0.001); - } + TestForward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerOneGradientGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 1.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerTwoGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 2.0; TypeParam scale = 0.34; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - // Now, check values - const TypeParam* bottom_data = this->blob_bottom_->cpu_data(); - const TypeParam* top_data = this->blob_top_->cpu_data(); - for (int i = 0; i < this->blob_bottom_->count(); ++i) { - TypeParam expected_value = pow(shift + scale * bottom_data[i], 2); - EXPECT_NEAR(expected_value, top_data[i], 0.001); - } + TestForward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerTwoGradientGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 2.0; TypeParam scale = 0.83; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } TYPED_TEST(PowerLayerTest, TestPowerTwoScaleHalfGradientGPU) { Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; TypeParam power = 2.0; TypeParam scale = 0.5; TypeParam shift = -2.4; - layer_param.mutable_power_param()->set_power(power); - layer_param.mutable_power_param()->set_scale(scale); - layer_param.mutable_power_param()->set_shift(shift); - PowerLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2, 1701, 0., 0.01); - checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); + TestBackward(power, scale, shift); } } // namespace caffe From 404f22d08297570523548e2cf433a9524da81a9a Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Tue, 8 Apr 2014 11:46:49 -0700 Subject: [PATCH 18/18] update proto field IDs from placeholder values --- src/caffe/proto/caffe.proto | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 6b54bbb0faa..da7824c8e27 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -89,6 +89,9 @@ message SolverState { repeated BlobProto history = 3; // The history for sgd solvers } +// Update the next available ID when you add a new LayerParameter field. +// +// LayerParameter next available ID: 22 message LayerParameter { repeated string bottom = 2; // the name of the bottom blobs repeated string top = 3; // the name of the top blobs @@ -99,7 +102,7 @@ message LayerParameter { // line above the enum. Update the next available ID when you add a new // LayerType. // - // LayerType next available ID: 25 + // LayerType next available ID: 27 enum LayerType { // "NONE" layer type is 0th enum element so that we don't cause confusion // by defaulting to an existent LayerType (instead, should usually error if @@ -112,7 +115,7 @@ message LayerParameter { DATA = 5; DROPOUT = 6; EUCLIDEAN_LOSS = 7; - ELTWISE_PRODUCT = 1000; + ELTWISE_PRODUCT = 25; FLATTEN = 8; HDF5_DATA = 9; HDF5_OUTPUT = 10; @@ -123,7 +126,7 @@ message LayerParameter { LRN = 15; MULTINOMIAL_LOGISTIC_LOSS = 16; POOLING = 17; - POWER = 1001; + POWER = 26; RELU = 18; SIGMOID = 19; SOFTMAX = 20; @@ -154,7 +157,7 @@ message LayerParameter { optional InnerProductParameter inner_product_param = 17; optional LRNParameter lrn_param = 18; optional PoolingParameter pooling_param = 19; - optional PowerParameter power_param = 1000; + optional PowerParameter power_param = 21; optional WindowDataParameter window_data_param = 20; // DEPRECATED: The layer parameters specified as a V0LayerParameter.