protected:
virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
- virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
- virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
- virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
-};
-
-/* CaffeSoftmaxLayer
-*/
-template <typename Dtype>
-class CaffeSoftmaxLayer : public SoftmaxLayer<Dtype> {
- public:
- explicit CaffeSoftmaxLayer(const LayerParameter& param)
- : SoftmaxLayer<Dtype>(param) {}
- virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
-
- protected:
- virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top);
virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top);
public:
explicit SigmoidCrossEntropyLossLayer(const LayerParameter& param)
: LossLayer<Dtype>(param),
- sigmoid_layer_(new CaffeSigmoidLayer<Dtype>(param)),
+ sigmoid_layer_(new SigmoidLayer<Dtype>(param)),
sigmoid_output_(new Blob<Dtype>()) {}
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top);
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
/// The internal SigmoidLayer used to map predictions to probabilities.
- shared_ptr<CaffeSigmoidLayer<Dtype> > sigmoid_layer_;
+ shared_ptr<SigmoidLayer<Dtype> > sigmoid_layer_;
/// sigmoid_output stores the output of the SigmoidLayer.
shared_ptr<Blob<Dtype> > sigmoid_output_;
/// bottom vector holder to call the underlying SigmoidLayer::Forward
vector<Blob<Dtype>*> sigmoid_top_vec_;
};
-// Forward declare CaffeSoftmaxLayer for use in SoftmaxWithLossLayer.
-template <typename Dtype> class CaffeSoftmaxLayer;
+// Forward declare SoftmaxLayer for use in SoftmaxWithLossLayer.
+template <typename Dtype> class SoftmaxLayer;
/**
* @brief Computes the multinomial logistic loss for a one-of-many
public:
explicit SoftmaxWithLossLayer(const LayerParameter& param)
: LossLayer<Dtype>(param),
- softmax_layer_(new CaffeSoftmaxLayer<Dtype>(param)) {}
+ softmax_layer_(new SoftmaxLayer<Dtype>(param)) {}
virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top);
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
/// The internal SoftmaxLayer used to map predictions to a distribution.
- shared_ptr<CaffeSoftmaxLayer<Dtype> > softmax_layer_;
+ shared_ptr<SoftmaxLayer<Dtype> > softmax_layer_;
/// prob stores the output probability predictions from the SoftmaxLayer.
Blob<Dtype> prob_;
/// bottom vector holder used in call to the underlying SoftmaxLayer::Forward
* the computed outputs are @f$ y = \max(0, x) + \nu \min(0, x) @f$.
*/
virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
+ vector<Blob<Dtype>*>* top);
virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
+ vector<Blob<Dtype>*>* top);
/**
* @brief Computes the error gradient w.r.t. the ReLU inputs.
* @f$.
*/
virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
- virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
-};
-
-/**
- * @brief standard Caffe implementation of ReLULayer.
- */
-template <typename Dtype>
-class CaffeReLULayer : public ReLULayer<Dtype> {
- public:
- explicit CaffeReLULayer(const LayerParameter& param)
- : ReLULayer<Dtype>(param) {}
-
- protected:
- virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
- virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
-
- virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
* @f$
*/
virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
+ vector<Blob<Dtype>*>* top);
virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
+ vector<Blob<Dtype>*>* top);
/**
* @brief Computes the error gradient w.r.t. the sigmoid inputs.
* @f$ if propagate_down[0]
*/
virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
- virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
-};
-
-/**
- * @brief standard Caffe implementation of SigmoidLayer.
- */
-template <typename Dtype>
-class CaffeSigmoidLayer : public SigmoidLayer<Dtype> {
- public:
- explicit CaffeSigmoidLayer(const LayerParameter& param)
- : SigmoidLayer<Dtype>(param) {}
-
- protected:
- virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
- virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
-
- virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
* @f$
*/
virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
+ vector<Blob<Dtype>*>* top);
virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
+ vector<Blob<Dtype>*>* top);
/**
* @brief Computes the error gradient w.r.t. the sigmoid inputs.
* @f$ if propagate_down[0]
*/
virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
- virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
-};
-
-/**
- * @brief standard Caffe implementation of TanHLayer.
- */
-template <typename Dtype>
-class CaffeTanHLayer : public TanHLayer<Dtype> {
- public:
- explicit CaffeTanHLayer(const LayerParameter& param)
- : TanHLayer<Dtype>(param) {}
-
- protected:
- virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
- virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
-
- virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
protected:
virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
+ vector<Blob<Dtype>*>* top);
virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
+ vector<Blob<Dtype>*>* top);
virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
+ const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
+ const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
int kernel_h_, kernel_w_;
int stride_h_, stride_w_;
int num_output_;
int height_out_, width_out_;
bool bias_term_;
-};
-
-/* CaffeConvolutionLayer
-*/
-template <typename Dtype>
-class CaffeConvolutionLayer : public ConvolutionLayer<Dtype> {
- public:
- explicit CaffeConvolutionLayer(const LayerParameter& param)
- : ConvolutionLayer<Dtype>(param) {}
- virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
-
- protected:
- virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
- virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
- virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
- virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
-
+ // For the Caffe matrix multiplication convolution.
int M_, K_, N_;
Blob<Dtype> col_buffer_;
Blob<Dtype> bias_multiplier_;
virtual inline LayerParameter_LayerType type() const {
return LayerParameter_LayerType_POOLING;
}
-
- protected:
- virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
- virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) = 0;
- virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
- virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) = 0;
-
- int kernel_h_, kernel_w_;
- int stride_h_, stride_w_;
- int pad_h_, pad_w_;
- int channels_;
- int height_, width_;
- int pooled_height_, pooled_width_;
-};
-
-/* CaffePoolingLayer
-*/
-template <typename Dtype>
-class CaffePoolingLayer : public PoolingLayer<Dtype> {
- public:
- explicit CaffePoolingLayer(const LayerParameter& param)
- : PoolingLayer<Dtype>(param) {}
- virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top);
-
virtual inline int ExactNumBottomBlobs() const { return 1; }
virtual inline int MinTopBlobs() const { return 1; }
// MAX POOL layers can output an extra top blob for the mask;
virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom);
+ int kernel_h_, kernel_w_;
+ int stride_h_, stride_w_;
+ int pad_h_, pad_w_;
+ int channels_;
+ int height_, width_;
+ int pooled_height_, pooled_width_;
Blob<Dtype> rand_idx_;
Blob<int> max_idx_;
};
engine = ConvolutionParameter_Engine_CAFFE;
}
if (engine == ConvolutionParameter_Engine_CAFFE) {
- return new CaffeConvolutionLayer<Dtype>(param);
+ return new ConvolutionLayer<Dtype>(param);
} else {
LOG(FATAL) << "Layer " << name << " has unknown engine.";
}
engine = PoolingParameter_Engine_CAFFE;
}
if (engine == PoolingParameter_Engine_CAFFE) {
- return new CaffePoolingLayer<Dtype>(param);
+ return new PoolingLayer<Dtype>(param);
} else {
LOG(FATAL) << "Layer " << name << " has unknown engine.";
}
engine = ReLUParameter_Engine_CAFFE;
}
if (engine == ReLUParameter_Engine_CAFFE) {
- return new CaffeReLULayer<Dtype>(param);
+ return new ReLULayer<Dtype>(param);
} else {
LOG(FATAL) << "Layer " << name << " has unknown engine.";
}
engine = SigmoidParameter_Engine_CAFFE;
}
if (engine == SigmoidParameter_Engine_CAFFE) {
- return new CaffeSigmoidLayer<Dtype>(param);
+ return new SigmoidLayer<Dtype>(param);
} else {
LOG(FATAL) << "Layer " << name << " has unknown engine.";
}
engine = TanHParameter_Engine_CAFFE;
}
if (engine == TanHParameter_Engine_CAFFE) {
- return new CaffeTanHLayer<Dtype>(param);
+ return new TanHLayer<Dtype>(param);
} else {
LOG(FATAL) << "Layer " << name << " has unknown engine.";
}
engine = SoftmaxParameter_Engine_CAFFE;
}
if (engine == SoftmaxParameter_Engine_CAFFE) {
- return new CaffeSoftmaxLayer<Dtype>(param);
+ return new SoftmaxLayer<Dtype>(param);
} else {
LOG(FATAL) << "Layer " << name << " has unknown engine.";
}
+++ /dev/null
-#include <vector>
-
-#include "caffe/filler.hpp"
-#include "caffe/layer.hpp"
-#include "caffe/util/im2col.hpp"
-#include "caffe/util/math_functions.hpp"
-#include "caffe/vision_layers.hpp"
-
-namespace caffe {
-
-template <typename Dtype>
-void CaffeConvolutionLayer<Dtype>::LayerSetUp(const vector<Blob<Dtype>*>&
- bottom, vector<Blob<Dtype>*>* top) {
- ConvolutionLayer<Dtype>::LayerSetUp(bottom, top);
- // Figure out the dimensions for individual gemms.
- M_ = this->num_output_ / this->group_;
- K_ = this->channels_ * this->kernel_h_ * this->kernel_w_ / this->group_;
- N_ = this->height_out_ * this->width_out_;
- // The im2col result buffer would only hold one image at a time to avoid
- // overly large memory usage.
- col_buffer_.Reshape(1, this->channels_ * this->kernel_h_ * this->kernel_w_,
- this->height_out_, this->width_out_);
- // Set up the all ones "bias multiplier" for adding bias using blas
- if (this->bias_term_) {
- bias_multiplier_.Reshape(1, 1, 1, N_);
- caffe_set(N_, Dtype(1), bias_multiplier_.mutable_cpu_data());
- }
-}
-
-template <typename Dtype>
-void CaffeConvolutionLayer<Dtype>::Forward_cpu(
- const vector<Blob<Dtype>*>& bottom, vector<Blob<Dtype>*>* top) {
- for (int i = 0; i < bottom.size(); ++i) {
- const Dtype* bottom_data = bottom[i]->cpu_data();
- Dtype* top_data = (*top)[i]->mutable_cpu_data();
- Dtype* col_data = col_buffer_.mutable_cpu_data();
- const Dtype* weight = this->blobs_[0]->cpu_data();
- int weight_offset = M_ * K_;
- int col_offset = K_ * N_;
- int top_offset = M_ * N_;
- for (int n = 0; n < this->num_; ++n) {
- // First, im2col
- im2col_cpu(bottom_data + bottom[i]->offset(n), this->channels_,
- this->height_, this->width_, this->kernel_h_, this->kernel_w_,
- this->pad_h_, this->pad_w_, this->stride_h_, this->stride_w_,
- col_data);
- // Second, innerproduct with groups
- for (int g = 0; g < this->group_; ++g) {
- caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, M_, N_, K_,
- (Dtype)1., weight + weight_offset * g, col_data + col_offset * g,
- (Dtype)0., top_data + (*top)[i]->offset(n) + top_offset * g);
- }
- // third, add bias
- if (this->bias_term_) {
- caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, this->num_output_,
- N_, 1, (Dtype)1., this->blobs_[1]->cpu_data(),
- bias_multiplier_.cpu_data(),
- (Dtype)1., top_data + (*top)[i]->offset(n));
- }
- }
- }
-}
-
-template <typename Dtype>
-void CaffeConvolutionLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) {
- const Dtype* weight = NULL;
- Dtype* weight_diff = NULL;
- if (this->param_propagate_down_[0]) {
- weight = this->blobs_[0]->cpu_data();
- weight_diff = this->blobs_[0]->mutable_cpu_diff();
- caffe_set(this->blobs_[0]->count(), Dtype(0), weight_diff);
- }
- Dtype* bias_diff = NULL;
- if (this->bias_term_ && this->param_propagate_down_[1]) {
- bias_diff = this->blobs_[1]->mutable_cpu_diff();
- caffe_set(this->blobs_[1]->count(), Dtype(0), bias_diff);
- }
- const int weight_offset = M_ * K_;
- const int col_offset = K_ * N_;
- const int top_offset = M_ * N_;
- for (int i = 0; i < top.size(); ++i) {
- const Dtype* top_diff = NULL;
- // Bias gradient, if necessary.
- if (this->bias_term_ && this->param_propagate_down_[1]) {
- top_diff = top[i]->cpu_diff();
- for (int n = 0; n < this->num_; ++n) {
- caffe_cpu_gemv<Dtype>(CblasNoTrans, this->num_output_, N_,
- 1., top_diff + top[0]->offset(n),
- bias_multiplier_.cpu_data(), 1.,
- bias_diff);
- }
- }
- if (this->param_propagate_down_[0] || propagate_down[i]) {
- if (!top_diff) {
- top_diff = top[i]->cpu_diff();
- }
- Dtype* col_data = col_buffer_.mutable_cpu_data();
- Dtype* col_diff = col_buffer_.mutable_cpu_diff();
- const Dtype* bottom_data = (*bottom)[i]->cpu_data();
- Dtype* bottom_diff = (*bottom)[i]->mutable_cpu_diff();
- for (int n = 0; n < this->num_; ++n) {
- // Since we saved memory in the forward pass by not storing all col
- // data, we will need to recompute them.
- im2col_cpu(bottom_data + (*bottom)[i]->offset(n), this->channels_,
- this->height_, this->width_, this->kernel_h_, this->kernel_w_,
- this->pad_h_, this->pad_w_, this->stride_h_, this->stride_w_,
- col_data);
- // gradient w.r.t. weight. Note that we will accumulate diffs.
- if (this->param_propagate_down_[0]) {
- for (int g = 0; g < this->group_; ++g) {
- caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasTrans, M_, K_, N_,
- (Dtype)1., top_diff + top[i]->offset(n) + top_offset * g,
- col_data + col_offset * g, (Dtype)1.,
- weight_diff + weight_offset * g);
- }
- }
- // gradient w.r.t. bottom data, if necessary
- if (propagate_down[i]) {
- for (int g = 0; g < this->group_; ++g) {
- caffe_cpu_gemm<Dtype>(CblasTrans, CblasNoTrans, K_, N_, M_,
- (Dtype)1., weight + weight_offset * g,
- top_diff + top[i]->offset(n) + top_offset * g,
- (Dtype)0., col_diff + col_offset * g);
- }
- // col2im back to the data
- col2im_cpu(col_diff, this->channels_, this->height_, this->width_,
- this->kernel_h_, this->kernel_w_, this->pad_h_, this->pad_w_,
- this->stride_h_, this->stride_w_, bottom_diff +
- (*bottom)[i]->offset(n));
- }
- }
- }
- }
-}
-
-#ifdef CPU_ONLY
-STUB_GPU(CaffeConvolutionLayer);
-#endif
-
-INSTANTIATE_CLASS(CaffeConvolutionLayer);
-
-} // namespace caffe
-
+++ /dev/null
-#include <algorithm>
-#include <cfloat>
-#include <vector>
-
-#include "caffe/common.hpp"
-#include "caffe/layer.hpp"
-#include "caffe/syncedmem.hpp"
-#include "caffe/util/math_functions.hpp"
-#include "caffe/vision_layers.hpp"
-
-namespace caffe {
-
-using std::min;
-using std::max;
-
-template <typename Dtype>
-void CaffePoolingLayer<Dtype>::LayerSetUp(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) {
- PoolingLayer<Dtype>::LayerSetUp(bottom, top);
- PoolingParameter pool_param = this->layer_param_.pooling_param();
- // If max pooling, we will initialize the vector index part.
- if (this->layer_param_.pooling_param().pool() ==
- PoolingParameter_PoolMethod_MAX && top->size() == 1) {
- max_idx_.Reshape(bottom[0]->num(), this->channels_, this->pooled_height_,
- this->pooled_width_);
- }
- // If stochastic pooling, we will initialize the random index part.
- if (this->layer_param_.pooling_param().pool() ==
- PoolingParameter_PoolMethod_STOCHASTIC) {
- rand_idx_.Reshape(bottom[0]->num(), this->channels_, this->pooled_height_,
- this->pooled_width_);
- }
-}
-
-// TODO(Yangqing): Is there a faster way to do pooling in the channel-first
-// case?
-template <typename Dtype>
-void CaffePoolingLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) {
- const Dtype* bottom_data = bottom[0]->cpu_data();
- Dtype* top_data = (*top)[0]->mutable_cpu_data();
- const int top_count = (*top)[0]->count();
- // We'll output the mask to top[1] if it's of size >1.
- const bool use_top_mask = top->size() > 1;
- int* mask = NULL; // suppress warnings about uninitalized variables
- Dtype* top_mask = NULL;
- // Different pooling methods. We explicitly do the switch outside the for
- // loop to save time, although this results in more code.
- switch (this->layer_param_.pooling_param().pool()) {
- case PoolingParameter_PoolMethod_MAX:
- // Initialize
- if (use_top_mask) {
- top_mask = (*top)[1]->mutable_cpu_data();
- caffe_set(top_count, Dtype(-1), top_mask);
- } else {
- mask = max_idx_.mutable_cpu_data();
- caffe_set(top_count, -1, mask);
- }
- caffe_set(top_count, Dtype(-FLT_MAX), top_data);
- // The main loop
- for (int n = 0; n < bottom[0]->num(); ++n) {
- for (int c = 0; c < this->channels_; ++c) {
- for (int ph = 0; ph < this->pooled_height_; ++ph) {
- for (int pw = 0; pw < this->pooled_width_; ++pw) {
- int hstart = ph * this->stride_h_ - this->pad_h_;
- int wstart = pw * this->stride_w_ - this->pad_w_;
- int hend = min(hstart + this->kernel_h_, this->height_);
- int wend = min(wstart + this->kernel_w_, this->width_);
- hstart = max(hstart, 0);
- wstart = max(wstart, 0);
- const int pool_index = ph * this->pooled_width_ + pw;
- for (int h = hstart; h < hend; ++h) {
- for (int w = wstart; w < wend; ++w) {
- const int index = h * this->width_ + w;
- if (bottom_data[index] > top_data[pool_index]) {
- top_data[pool_index] = bottom_data[index];
- if (use_top_mask) {
- top_mask[pool_index] = static_cast<Dtype>(index);
- } else {
- mask[pool_index] = index;
- }
- }
- }
- }
- }
- }
- // compute offset
- bottom_data += bottom[0]->offset(0, 1);
- top_data += (*top)[0]->offset(0, 1);
- if (use_top_mask) {
- top_mask += (*top)[0]->offset(0, 1);
- } else {
- mask += (*top)[0]->offset(0, 1);
- }
- }
- }
- break;
- case PoolingParameter_PoolMethod_AVE:
- for (int i = 0; i < top_count; ++i) {
- top_data[i] = 0;
- }
- // The main loop
- for (int n = 0; n < bottom[0]->num(); ++n) {
- for (int c = 0; c < this->channels_; ++c) {
- for (int ph = 0; ph < this->pooled_height_; ++ph) {
- for (int pw = 0; pw < this->pooled_width_; ++pw) {
- int hstart = ph * this->stride_h_ - this->pad_h_;
- int wstart = pw * this->stride_w_ - this->pad_w_;
- int hend = min(hstart + this->kernel_h_,
- this->height_ + this->pad_h_);
- int wend = min(wstart + this->kernel_w_,
- this->width_ + this->pad_w_);
- int pool_size = (hend - hstart) * (wend - wstart);
- hstart = max(hstart, 0);
- wstart = max(wstart, 0);
- hend = min(hend, this->height_);
- wend = min(wend, this->width_);
- for (int h = hstart; h < hend; ++h) {
- for (int w = wstart; w < wend; ++w) {
- top_data[ph * this->pooled_width_ + pw] +=
- bottom_data[h * this->width_ + w];
- }
- }
- top_data[ph * this->pooled_width_ + pw] /= pool_size;
- }
- }
- // compute offset
- bottom_data += bottom[0]->offset(0, 1);
- top_data += (*top)[0]->offset(0, 1);
- }
- }
- break;
- case PoolingParameter_PoolMethod_STOCHASTIC:
- NOT_IMPLEMENTED;
- break;
- default:
- LOG(FATAL) << "Unknown pooling method.";
- }
-}
-
-template <typename Dtype>
-void CaffePoolingLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) {
- if (!propagate_down[0]) {
- return;
- }
- const Dtype* top_diff = top[0]->cpu_diff();
- Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff();
- // Different pooling methods. We explicitly do the switch outside the for
- // loop to save time, although this results in more codes.
- caffe_set((*bottom)[0]->count(), Dtype(0), bottom_diff);
- // We'll output the mask to top[1] if it's of size >1.
- const bool use_top_mask = top.size() > 1;
- const int* mask = NULL; // suppress warnings about uninitialized variables
- const Dtype* top_mask = NULL;
- switch (this->layer_param_.pooling_param().pool()) {
- case PoolingParameter_PoolMethod_MAX:
- // The main loop
- if (use_top_mask) {
- top_mask = top[1]->cpu_data();
- } else {
- mask = max_idx_.cpu_data();
- }
- for (int n = 0; n < top[0]->num(); ++n) {
- for (int c = 0; c < this->channels_; ++c) {
- for (int ph = 0; ph < this->pooled_height_; ++ph) {
- for (int pw = 0; pw < this->pooled_width_; ++pw) {
- const int index = ph * this->pooled_width_ + pw;
- const int bottom_index =
- use_top_mask ? top_mask[index] : mask[index];
- bottom_diff[bottom_index] += top_diff[index];
- }
- }
- bottom_diff += (*bottom)[0]->offset(0, 1);
- top_diff += top[0]->offset(0, 1);
- if (use_top_mask) {
- top_mask += top[0]->offset(0, 1);
- } else {
- mask += top[0]->offset(0, 1);
- }
- }
- }
- break;
- case PoolingParameter_PoolMethod_AVE:
- // The main loop
- for (int n = 0; n < top[0]->num(); ++n) {
- for (int c = 0; c < this->channels_; ++c) {
- for (int ph = 0; ph < this->pooled_height_; ++ph) {
- for (int pw = 0; pw < this->pooled_width_; ++pw) {
- int hstart = ph * this->stride_h_ - this->pad_h_;
- int wstart = pw * this->stride_w_ - this->pad_w_;
- int hend = min(hstart + this->kernel_h_,
- this->height_ + this->pad_h_);
- int wend = min(wstart + this->kernel_w_,
- this->width_ + this->pad_w_);
- int pool_size = (hend - hstart) * (wend - wstart);
- hstart = max(hstart, 0);
- wstart = max(wstart, 0);
- hend = min(hend, this->height_);
- wend = min(wend, this->width_);
- for (int h = hstart; h < hend; ++h) {
- for (int w = wstart; w < wend; ++w) {
- bottom_diff[h * this->width_ + w] +=
- top_diff[ph * this->pooled_width_ + pw] / pool_size;
- }
- }
- }
- }
- // offset
- bottom_diff += (*bottom)[0]->offset(0, 1);
- top_diff += top[0]->offset(0, 1);
- }
- }
- break;
- case PoolingParameter_PoolMethod_STOCHASTIC:
- NOT_IMPLEMENTED;
- break;
- default:
- LOG(FATAL) << "Unknown pooling method.";
- }
-}
-
-#ifdef CPU_ONLY
-STUB_GPU(CaffePoolingLayer);
-#endif
-
-INSTANTIATE_CLASS(CaffePoolingLayer);
-
-} // namespace caffe
-
+++ /dev/null
-//
-#include <algorithm>
-#include <vector>
-
-#include "caffe/layer.hpp"
-#include "caffe/util/math_functions.hpp"
-#include "caffe/vision_layers.hpp"
-
-namespace caffe {
-
-template <typename Dtype>
-void CaffeSoftmaxLayer<Dtype>::LayerSetUp(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) {
- SoftmaxLayer<Dtype>::LayerSetUp(bottom, top);
- sum_multiplier_.Reshape(1, bottom[0]->channels(), 1, 1);
- Dtype* multiplier_data = sum_multiplier_.mutable_cpu_data();
- for (int i = 0; i < sum_multiplier_.count(); ++i) {
- multiplier_data[i] = 1.;
- }
- scale_.Reshape(bottom[0]->num(), 1, bottom[0]->height(), bottom[0]->width());
-}
-
-template <typename Dtype>
-void CaffeSoftmaxLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
- vector<Blob<Dtype>*>* top) {
- const Dtype* bottom_data = bottom[0]->cpu_data();
- Dtype* top_data = (*top)[0]->mutable_cpu_data();
- Dtype* scale_data = scale_.mutable_cpu_data();
- int num = bottom[0]->num();
- int channels = bottom[0]->channels();
- int dim = bottom[0]->count() / bottom[0]->num();
- int spatial_dim = bottom[0]->height() * bottom[0]->width();
- caffe_copy(bottom[0]->count(), bottom_data, top_data);
- // We need to subtract the max to avoid numerical issues, compute the exp,
- // and then normalize.
- for (int i = 0; i < num; ++i) {
- // initialize scale_data to the first plane
- caffe_copy(spatial_dim, bottom_data + i * dim, scale_data);
- for (int j = 0; j < channels; j++) {
- for (int k = 0; k < spatial_dim; k++) {
- scale_data[k] = std::max(scale_data[k],
- bottom_data[i * dim + j * spatial_dim + k]);
- }
- }
- // subtraction
- caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, channels, spatial_dim,
- 1, -1., sum_multiplier_.cpu_data(), scale_data, 1., top_data + i * dim);
- // exponentiation
- caffe_exp<Dtype>(dim, top_data + i * dim, top_data + i * dim);
- // sum after exp
- caffe_cpu_gemv<Dtype>(CblasTrans, channels, spatial_dim, 1.,
- top_data + i * dim, sum_multiplier_.cpu_data(), 0., scale_data);
- // division
- for (int j = 0; j < channels; j++) {
- caffe_div(spatial_dim, top_data + (*top)[0]->offset(i, j), scale_data,
- top_data + (*top)[0]->offset(i, j));
- }
- }
-}
-
-template <typename Dtype>
-void CaffeSoftmaxLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
- const vector<bool>& propagate_down,
- vector<Blob<Dtype>*>* bottom) {
- const Dtype* top_diff = top[0]->cpu_diff();
- const Dtype* top_data = top[0]->cpu_data();
- Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff();
- Dtype* scale_data = scale_.mutable_cpu_data();
- int num = top[0]->num();
- int channels = top[0]->channels();
- int dim = top[0]->count() / top[0]->num();
- int spatial_dim = top[0]->height() * top[0]->width();
- caffe_copy(top[0]->count(), top_diff, bottom_diff);
- for (int i = 0; i < num; ++i) {
- // compute dot(top_diff, top_data) and subtract them from the bottom diff
- for (int k = 0; k < spatial_dim; ++k) {
- scale_data[k] = caffe_cpu_strided_dot<Dtype>(channels,
- bottom_diff + i * dim + k, spatial_dim,
- top_data + i * dim + k, spatial_dim);
- }
- // subtraction
- caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, channels, spatial_dim, 1,
- -1., sum_multiplier_.cpu_data(), scale_data, 1., bottom_diff + i * dim);
- }
- // elementwise multiplication
- caffe_mul(top[0]->count(), bottom_diff, top_data, bottom_diff);
-}
-
-
-#ifdef CPU_ONLY
-STUB_GPU(CaffeSoftmaxLayer);
-#endif
-
-INSTANTIATE_CLASS(CaffeSoftmaxLayer);
-
-
-} // namespace caffe
stride_h_ = conv_param.stride_h();
stride_w_ = conv_param.stride_w();
}
- group_ = conv_param.group();
+ group_ = this->layer_param_.convolution_param().group();
num_ = bottom[0]->num();
channels_ = bottom[0]->channels();
height_ = bottom[0]->height();
CHECK_EQ(width_, bottom[bottom_id]->width())
<< "Inputs must have same width.";
}
- num_output_ = conv_param.num_output();
+ num_output_ = this->layer_param_.convolution_param().num_output();
CHECK_GT(num_output_, 0);
CHECK_EQ(channels_ % group_, 0);
- // Calculate output dimensions.
- height_out_ = (height_ + 2 * pad_h_ - kernel_h_) / stride_h_ + 1;
- width_out_ = (width_ + 2 * pad_w_ - kernel_w_) / stride_w_ + 1;
+ // The im2col result buffer would only hold one image at a time to avoid
+ // overly large memory usage.
+ int height_out =
+ (height_ + 2 * pad_h_ - kernel_h_) / stride_h_ + 1;
+ int width_out = (width_ + 2 * pad_w_ - kernel_w_) / stride_w_ + 1;
+ col_buffer_.Reshape(
+ 1, channels_ * kernel_h_ * kernel_w_, height_out, width_out);
// Set the parameters
CHECK_EQ(num_output_ % group_, 0)
<< "Number of output should be multiples of group.";
- bias_term_ = conv_param.bias_term();
+ bias_term_ = this->layer_param_.convolution_param().bias_term();
+ // Figure out the dimensions for individual gemms.
+ M_ = num_output_ / group_;
+ K_ = channels_ * kernel_h_ * kernel_w_ / group_;
+ N_ = height_out * width_out;
for (int top_id = 0; top_id < top->size(); ++top_id) {
- (*top)[top_id]->Reshape(num_, num_output_, height_out_, width_out_);
+ (*top)[top_id]->Reshape(num_, num_output_, height_out, width_out);
}
// Check if we need to set up the weights
if (this->blobs_.size() > 0) {
num_output_, channels_ / group_, kernel_h_, kernel_w_));
// fill the weights
shared_ptr<Filler<Dtype> > weight_filler(GetFiller<Dtype>(
- conv_param.weight_filler()));
+ this->layer_param_.convolution_param().weight_filler()));
weight_filler->Fill(this->blobs_[0].get());
// If necessary, initialize and fill the bias term
if (bias_term_) {
this->blobs_[1].reset(new Blob<Dtype>(1, 1, 1, num_output_));
shared_ptr<Filler<Dtype> > bias_filler(GetFiller<Dtype>(
- conv_param.bias_filler()));
+ this->layer_param_.convolution_param().bias_filler()));
bias_filler->Fill(this->blobs_[1].get());
}
}
+ // Set up the all ones "bias multiplier" for adding bias using blas
+ if (bias_term_) {
+ bias_multiplier_.Reshape(1, 1, 1, N_);
+ caffe_set(N_, Dtype(1), bias_multiplier_.mutable_cpu_data());
+ }
this->param_propagate_down_.resize(this->blobs_.size(), true);
}
+
+template <typename Dtype>
+void ConvolutionLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+ vector<Blob<Dtype>*>* top) {
+ for (int i = 0; i < bottom.size(); ++i) {
+ const Dtype* bottom_data = bottom[i]->cpu_data();
+ Dtype* top_data = (*top)[i]->mutable_cpu_data();
+ Dtype* col_data = col_buffer_.mutable_cpu_data();
+ const Dtype* weight = this->blobs_[0]->cpu_data();
+ int weight_offset = M_ * K_;
+ int col_offset = K_ * N_;
+ int top_offset = M_ * N_;
+ for (int n = 0; n < num_; ++n) {
+ // First, im2col
+ im2col_cpu(bottom_data + bottom[i]->offset(n), channels_, height_,
+ width_, kernel_h_, kernel_w_, pad_h_, pad_w_, stride_h_, stride_w_,
+ col_data);
+ // Second, innerproduct with groups
+ for (int g = 0; g < group_; ++g) {
+ caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, M_, N_, K_,
+ (Dtype)1., weight + weight_offset * g, col_data + col_offset * g,
+ (Dtype)0., top_data + (*top)[i]->offset(n) + top_offset * g);
+ }
+ // third, add bias
+ if (bias_term_) {
+ caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, num_output_,
+ N_, 1, (Dtype)1., this->blobs_[1]->cpu_data(),
+ bias_multiplier_.cpu_data(),
+ (Dtype)1., top_data + (*top)[i]->offset(n));
+ }
+ }
+ }
+}
+
+template <typename Dtype>
+void ConvolutionLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
+ const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) {
+ const Dtype* weight = NULL;
+ Dtype* weight_diff = NULL;
+ if (this->param_propagate_down_[0]) {
+ weight = this->blobs_[0]->cpu_data();
+ weight_diff = this->blobs_[0]->mutable_cpu_diff();
+ caffe_set(this->blobs_[0]->count(), Dtype(0), weight_diff);
+ }
+ Dtype* bias_diff = NULL;
+ if (bias_term_ && this->param_propagate_down_[1]) {
+ bias_diff = this->blobs_[1]->mutable_cpu_diff();
+ caffe_set(this->blobs_[1]->count(), Dtype(0), bias_diff);
+ }
+ const int weight_offset = M_ * K_;
+ const int col_offset = K_ * N_;
+ const int top_offset = M_ * N_;
+ for (int i = 0; i < top.size(); ++i) {
+ const Dtype* top_diff = NULL;
+ // Bias gradient, if necessary.
+ if (bias_term_ && this->param_propagate_down_[1]) {
+ top_diff = top[i]->cpu_diff();
+ for (int n = 0; n < num_; ++n) {
+ caffe_cpu_gemv<Dtype>(CblasNoTrans, num_output_, N_,
+ 1., top_diff + top[0]->offset(n),
+ bias_multiplier_.cpu_data(), 1.,
+ bias_diff);
+ }
+ }
+ if (this->param_propagate_down_[0] || propagate_down[i]) {
+ if (!top_diff) {
+ top_diff = top[i]->cpu_diff();
+ }
+ Dtype* col_data = col_buffer_.mutable_cpu_data();
+ Dtype* col_diff = col_buffer_.mutable_cpu_diff();
+ const Dtype* bottom_data = (*bottom)[i]->cpu_data();
+ Dtype* bottom_diff = (*bottom)[i]->mutable_cpu_diff();
+ for (int n = 0; n < num_; ++n) {
+ // Since we saved memory in the forward pass by not storing all col
+ // data, we will need to recompute them.
+ im2col_cpu(bottom_data + (*bottom)[i]->offset(n), channels_, height_,
+ width_, kernel_h_, kernel_w_, pad_h_, pad_w_,
+ stride_h_, stride_w_, col_data);
+ // gradient w.r.t. weight. Note that we will accumulate diffs.
+ if (this->param_propagate_down_[0]) {
+ for (int g = 0; g < group_; ++g) {
+ caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasTrans, M_, K_, N_,
+ (Dtype)1., top_diff + top[i]->offset(n) + top_offset * g,
+ col_data + col_offset * g, (Dtype)1.,
+ weight_diff + weight_offset * g);
+ }
+ }
+ // gradient w.r.t. bottom data, if necessary
+ if (propagate_down[i]) {
+ for (int g = 0; g < group_; ++g) {
+ caffe_cpu_gemm<Dtype>(CblasTrans, CblasNoTrans, K_, N_, M_,
+ (Dtype)1., weight + weight_offset * g,
+ top_diff + top[i]->offset(n) + top_offset * g,
+ (Dtype)0., col_diff + col_offset * g);
+ }
+ // col2im back to the data
+ col2im_cpu(col_diff, channels_, height_, width_,
+ kernel_h_, kernel_w_, pad_h_, pad_w_,
+ stride_h_, stride_w_, bottom_diff + (*bottom)[i]->offset(n));
+ }
+ }
+ }
+ }
+}
+
+#ifdef CPU_ONLY
+STUB_GPU(ConvolutionLayer);
+#endif
+
INSTANTIATE_CLASS(ConvolutionLayer);
} // namespace caffe
namespace caffe {
template <typename Dtype>
-void CaffeConvolutionLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>&
- bottom, vector<Blob<Dtype>*>* top) {
+void ConvolutionLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+ vector<Blob<Dtype>*>* top) {
for (int i = 0; i < bottom.size(); ++i) {
const Dtype* bottom_data = bottom[i]->gpu_data();
Dtype* top_data = (*top)[i]->mutable_gpu_data();
int weight_offset = M_ * K_;
int col_offset = K_ * N_;
int top_offset = M_ * N_;
- for (int n = 0; n < this->num_; ++n) {
+ for (int n = 0; n < num_; ++n) {
// First, im2col
- im2col_gpu(bottom_data + bottom[i]->offset(n), this->channels_,
- this->height_, this->width_, this->kernel_h_, this->kernel_w_,
- this->pad_h_, this->pad_w_, this->stride_h_, this->stride_w_,
+ im2col_gpu(bottom_data + bottom[i]->offset(n), channels_, height_,
+ width_, kernel_h_, kernel_w_, pad_h_, pad_w_, stride_h_, stride_w_,
col_data);
// Second, innerproduct with groups
- for (int g = 0; g < this->group_; ++g) {
+ for (int g = 0; g < group_; ++g) {
caffe_gpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, M_, N_, K_,
(Dtype)1., weight + weight_offset * g, col_data + col_offset * g,
(Dtype)0., top_data + (*top)[i]->offset(n) + top_offset * g);
}
// third, add bias
- if (this->bias_term_) {
- caffe_gpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, this->num_output_,
+ if (bias_term_) {
+ caffe_gpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, num_output_,
N_, 1, (Dtype)1., this->blobs_[1]->gpu_data(),
bias_multiplier_.gpu_data(),
(Dtype)1., top_data + (*top)[i]->offset(n));
}
template <typename Dtype>
-void CaffeConvolutionLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
+void ConvolutionLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) {
const Dtype* weight = NULL;
Dtype* weight_diff = NULL;
caffe_gpu_set(this->blobs_[0]->count(), Dtype(0), weight_diff);
}
Dtype* bias_diff = NULL;
- if (this->bias_term_ && this->param_propagate_down_[1]) {
+ if (bias_term_ && this->param_propagate_down_[1]) {
bias_diff = this->blobs_[1]->mutable_gpu_diff();
caffe_gpu_set(this->blobs_[1]->count(), Dtype(0), bias_diff);
}
for (int i = 0; i < top.size(); ++i) {
const Dtype* top_diff = NULL;
// Bias gradient, if necessary.
- if (this->bias_term_ && this->param_propagate_down_[1]) {
+ if (bias_term_ && this->param_propagate_down_[1]) {
top_diff = top[i]->gpu_diff();
- for (int n = 0; n < this->num_; ++n) {
- caffe_gpu_gemv<Dtype>(CblasNoTrans, this->num_output_, N_,
+ for (int n = 0; n < num_; ++n) {
+ caffe_gpu_gemv<Dtype>(CblasNoTrans, num_output_, N_,
1., top_diff + top[0]->offset(n),
bias_multiplier_.gpu_data(), 1.,
bias_diff);
Dtype* col_diff = col_buffer_.mutable_gpu_diff();
const Dtype* bottom_data = (*bottom)[i]->gpu_data();
Dtype* bottom_diff = (*bottom)[i]->mutable_gpu_diff();
- for (int n = 0; n < this->num_; ++n) {
+ for (int n = 0; n < num_; ++n) {
// Since we saved memory in the forward pass by not storing all col
// data, we will need to recompute them.
- im2col_gpu(bottom_data + (*bottom)[i]->offset(n), this->channels_,
- this->height_, this->width_, this->kernel_h_, this->kernel_w_,
- this->pad_h_, this->pad_w_, this->stride_h_, this->stride_w_,
- col_data);
+ im2col_gpu(bottom_data + (*bottom)[i]->offset(n), channels_, height_,
+ width_, kernel_h_, kernel_w_, pad_h_, pad_w_,
+ stride_h_, stride_w_, col_data);
// gradient w.r.t. weight. Note that we will accumulate diffs.
if (this->param_propagate_down_[0]) {
- for (int g = 0; g < this->group_; ++g) {
+ for (int g = 0; g < group_; ++g) {
caffe_gpu_gemm<Dtype>(CblasNoTrans, CblasTrans, M_, K_, N_,
(Dtype)1., top_diff + top[i]->offset(n) + top_offset * g,
col_data + col_offset * g, (Dtype)1.,
}
// gradient w.r.t. bottom data, if necessary
if (propagate_down[i]) {
- for (int g = 0; g < this->group_; ++g) {
+ for (int g = 0; g < group_; ++g) {
caffe_gpu_gemm<Dtype>(CblasTrans, CblasNoTrans, K_, N_, M_,
(Dtype)1., weight + weight_offset * g,
top_diff + top[i]->offset(n) + top_offset * g,
(Dtype)0., col_diff + col_offset * g);
}
// col2im back to the data
- col2im_gpu(col_diff, this->channels_, this->height_, this->width_,
- this->kernel_h_, this->kernel_w_, this->pad_h_, this->pad_w_,
- this->stride_h_, this->stride_w_, bottom_diff +
- (*bottom)[i]->offset(n));
+ col2im_gpu(col_diff, channels_, height_, width_,
+ kernel_h_, kernel_w_, pad_h_, pad_w_, stride_h_, stride_w_,
+ bottom_diff + (*bottom)[i]->offset(n));
}
}
}
}
-INSTANTIATE_CLASS(CaffeConvolutionLayer);
+INSTANTIATE_CLASS(ConvolutionLayer);
} // namespace caffe
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 CaffePoolingLayer<Dtype>(pool_param));
+ pool_layer_.reset(new PoolingLayer<Dtype>(pool_param));
pool_layer_->SetUp(square_top_vec_, &pool_top_vec_);
CHECK_EQ(pool_output_.num(), num_);
CHECK_EQ(pool_output_.channels(), channels_);
stride_w_ = pool_param.stride_w();
}
if (pad_h_ != 0 || pad_w_ != 0) {
- CHECK(pool_param.pool()
+ CHECK(this->layer_param_.pooling_param().pool()
== PoolingParameter_PoolMethod_AVE
- || pool_param.pool()
+ || this->layer_param_.pooling_param().pool()
== PoolingParameter_PoolMethod_MAX)
<< "Padding implemented only for average and max pooling.";
CHECK_LT(pad_h_, kernel_h_);
if (top->size() > 1) {
(*top)[1]->ReshapeLike(*(*top)[0]);
}
+ // If max pooling, we will initialize the vector index part.
+ if (this->layer_param_.pooling_param().pool() ==
+ PoolingParameter_PoolMethod_MAX && top->size() == 1) {
+ max_idx_.Reshape(bottom[0]->num(), channels_, pooled_height_,
+ pooled_width_);
+ }
+ // If stochastic pooling, we will initialize the random index part.
+ if (this->layer_param_.pooling_param().pool() ==
+ PoolingParameter_PoolMethod_STOCHASTIC) {
+ rand_idx_.Reshape(bottom[0]->num(), channels_, pooled_height_,
+ pooled_width_);
+ }
+}
+
+// TODO(Yangqing): Is there a faster way to do pooling in the channel-first
+// case?
+template <typename Dtype>
+void PoolingLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+ vector<Blob<Dtype>*>* top) {
+ const Dtype* bottom_data = bottom[0]->cpu_data();
+ Dtype* top_data = (*top)[0]->mutable_cpu_data();
+ const int top_count = (*top)[0]->count();
+ // We'll output the mask to top[1] if it's of size >1.
+ const bool use_top_mask = top->size() > 1;
+ int* mask = NULL; // suppress warnings about uninitalized variables
+ Dtype* top_mask = NULL;
+ // Different pooling methods. We explicitly do the switch outside the for
+ // loop to save time, although this results in more code.
+ switch (this->layer_param_.pooling_param().pool()) {
+ case PoolingParameter_PoolMethod_MAX:
+ // Initialize
+ if (use_top_mask) {
+ top_mask = (*top)[1]->mutable_cpu_data();
+ caffe_set(top_count, Dtype(-1), top_mask);
+ } else {
+ mask = max_idx_.mutable_cpu_data();
+ caffe_set(top_count, -1, mask);
+ }
+ caffe_set(top_count, Dtype(-FLT_MAX), top_data);
+ // The main loop
+ for (int n = 0; n < bottom[0]->num(); ++n) {
+ 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_h_ - pad_h_;
+ int wstart = pw * stride_w_ - pad_w_;
+ int hend = min(hstart + kernel_h_, height_);
+ int wend = min(wstart + kernel_w_, width_);
+ hstart = max(hstart, 0);
+ wstart = max(wstart, 0);
+ const int pool_index = ph * pooled_width_ + pw;
+ for (int h = hstart; h < hend; ++h) {
+ for (int w = wstart; w < wend; ++w) {
+ const int index = h * width_ + w;
+ if (bottom_data[index] > top_data[pool_index]) {
+ top_data[pool_index] = bottom_data[index];
+ if (use_top_mask) {
+ top_mask[pool_index] = static_cast<Dtype>(index);
+ } else {
+ mask[pool_index] = index;
+ }
+ }
+ }
+ }
+ }
+ }
+ // compute offset
+ bottom_data += bottom[0]->offset(0, 1);
+ top_data += (*top)[0]->offset(0, 1);
+ if (use_top_mask) {
+ top_mask += (*top)[0]->offset(0, 1);
+ } else {
+ mask += (*top)[0]->offset(0, 1);
+ }
+ }
+ }
+ break;
+ case PoolingParameter_PoolMethod_AVE:
+ for (int i = 0; i < top_count; ++i) {
+ top_data[i] = 0;
+ }
+ // The main loop
+ for (int n = 0; n < bottom[0]->num(); ++n) {
+ 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_h_ - pad_h_;
+ int wstart = pw * stride_w_ - pad_w_;
+ int hend = min(hstart + kernel_h_, height_ + pad_h_);
+ int wend = min(wstart + kernel_w_, width_ + pad_w_);
+ 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] /= pool_size;
+ }
+ }
+ // compute offset
+ bottom_data += bottom[0]->offset(0, 1);
+ top_data += (*top)[0]->offset(0, 1);
+ }
+ }
+ break;
+ case PoolingParameter_PoolMethod_STOCHASTIC:
+ NOT_IMPLEMENTED;
+ break;
+ default:
+ LOG(FATAL) << "Unknown pooling method.";
+ }
}
+template <typename Dtype>
+void PoolingLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
+ const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) {
+ if (!propagate_down[0]) {
+ return;
+ }
+ const Dtype* top_diff = top[0]->cpu_diff();
+ Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff();
+ // Different pooling methods. We explicitly do the switch outside the for
+ // loop to save time, although this results in more codes.
+ caffe_set((*bottom)[0]->count(), Dtype(0), bottom_diff);
+ // We'll output the mask to top[1] if it's of size >1.
+ const bool use_top_mask = top.size() > 1;
+ const int* mask = NULL; // suppress warnings about uninitialized variables
+ const Dtype* top_mask = NULL;
+ switch (this->layer_param_.pooling_param().pool()) {
+ case PoolingParameter_PoolMethod_MAX:
+ // The main loop
+ if (use_top_mask) {
+ top_mask = top[1]->cpu_data();
+ } else {
+ mask = max_idx_.cpu_data();
+ }
+ for (int n = 0; n < top[0]->num(); ++n) {
+ for (int c = 0; c < channels_; ++c) {
+ for (int ph = 0; ph < pooled_height_; ++ph) {
+ for (int pw = 0; pw < pooled_width_; ++pw) {
+ const int index = ph * pooled_width_ + pw;
+ const int bottom_index =
+ use_top_mask ? top_mask[index] : mask[index];
+ bottom_diff[bottom_index] += top_diff[index];
+ }
+ }
+ bottom_diff += (*bottom)[0]->offset(0, 1);
+ top_diff += top[0]->offset(0, 1);
+ if (use_top_mask) {
+ top_mask += top[0]->offset(0, 1);
+ } else {
+ mask += top[0]->offset(0, 1);
+ }
+ }
+ }
+ break;
+ case PoolingParameter_PoolMethod_AVE:
+ // The main loop
+ for (int n = 0; n < top[0]->num(); ++n) {
+ 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_h_ - pad_h_;
+ int wstart = pw * stride_w_ - pad_w_;
+ int hend = min(hstart + kernel_h_, height_ + pad_h_);
+ int wend = min(wstart + kernel_w_, width_ + pad_w_);
+ 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] / pool_size;
+ }
+ }
+ }
+ }
+ // offset
+ bottom_diff += (*bottom)[0]->offset(0, 1);
+ top_diff += top[0]->offset(0, 1);
+ }
+ }
+ break;
+ case PoolingParameter_PoolMethod_STOCHASTIC:
+ NOT_IMPLEMENTED;
+ break;
+ default:
+ LOG(FATAL) << "Unknown pooling method.";
+ }
+}
+
+
+#ifdef CPU_ONLY
+STUB_GPU(PoolingLayer);
+#endif
+
INSTANTIATE_CLASS(PoolingLayer);
+
} // namespace caffe
template <typename Dtype>
-void CaffePoolingLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+void PoolingLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) {
const Dtype* bottom_data = bottom[0]->gpu_data();
Dtype* top_data = (*top)[0]->mutable_gpu_data();
}
// NOLINT_NEXT_LINE(whitespace/operators)
MaxPoolForward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
- count, bottom_data, bottom[0]->num(), this->channels_, this->height_,
- this->width_, this->pooled_height_, this->pooled_width_,
- this->kernel_h_, this->kernel_w_, this->stride_h_, this->stride_w_,
- this->pad_h_, this->pad_w_, top_data, mask, top_mask);
+ count, bottom_data, bottom[0]->num(), channels_,
+ height_, width_, pooled_height_, pooled_width_, kernel_h_,
+ kernel_w_, stride_h_, stride_w_, pad_h_, pad_w_, top_data,
+ mask, top_mask);
break;
case PoolingParameter_PoolMethod_AVE:
// NOLINT_NEXT_LINE(whitespace/operators)
AvePoolForward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
- count, bottom_data, bottom[0]->num(), this->channels_, this->height_,
- this->width_, this->pooled_height_, this->pooled_width_,
- this->kernel_h_, this->kernel_w_, this->stride_h_, this->stride_w_,
- this->pad_h_, this->pad_w_, top_data);
+ count, bottom_data, bottom[0]->num(), channels_,
+ height_, width_, pooled_height_, pooled_width_, kernel_h_,
+ kernel_w_, stride_h_, stride_w_, pad_h_, pad_w_, top_data);
break;
case PoolingParameter_PoolMethod_STOCHASTIC:
if (Caffe::phase() == Caffe::TRAIN) {
// NOLINT_NEXT_LINE(whitespace/operators)
StoPoolForwardTrain<Dtype><<<CAFFE_GET_BLOCKS(count),
CAFFE_CUDA_NUM_THREADS>>>(
- count, bottom_data, bottom[0]->num(), this->channels_, this->height_,
- this->width_, this->pooled_height_, this->pooled_width_,
- this->kernel_h_, this->kernel_w_, this->stride_h_, this->stride_w_,
+ count, bottom_data, bottom[0]->num(), channels_,
+ height_, width_, pooled_height_, pooled_width_, kernel_h_,
+ kernel_w_, stride_h_, stride_w_,
rand_idx_.mutable_gpu_data(), top_data);
} else {
// NOLINT_NEXT_LINE(whitespace/operators)
StoPoolForwardTest<Dtype><<<CAFFE_GET_BLOCKS(count),
CAFFE_CUDA_NUM_THREADS>>>(
- count, bottom_data, bottom[0]->num(), this->channels_, this->height_,
- this->width_, this->pooled_height_, this->pooled_width_,
- this->kernel_h_, this->kernel_w_, this->stride_h_, this->stride_w_,
- top_data);
+ count, bottom_data, bottom[0]->num(), channels_,
+ height_, width_, pooled_height_, pooled_width_, kernel_h_,
+ kernel_w_, stride_h_, stride_w_, top_data);
}
break;
default:
template <typename Dtype>
-void CaffePoolingLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
+void PoolingLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) {
if (!propagate_down[0]) {
return;
}
// NOLINT_NEXT_LINE(whitespace/operators)
MaxPoolBackward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
- count, top_diff, mask, top_mask, top[0]->num(), this->channels_,
- this->height_, this->width_, this->pooled_height_, this->pooled_width_,
- this->kernel_h_, this->kernel_w_, this->stride_h_, this->stride_w_,
- this->pad_h_, this->pad_w_, bottom_diff);
+ count, top_diff, mask, top_mask, top[0]->num(), channels_,
+ height_, width_, pooled_height_, pooled_width_,
+ kernel_h_, kernel_w_, stride_h_, stride_w_, pad_h_, pad_w_,
+ bottom_diff);
break;
case PoolingParameter_PoolMethod_AVE:
// NOLINT_NEXT_LINE(whitespace/operators)
AvePoolBackward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
- count, top_diff, top[0]->num(), this->channels_, this->height_,
- this->width_, this->pooled_height_, this->pooled_width_,
- this->kernel_h_, this->kernel_w_, this->stride_h_, this->stride_w_,
- this->pad_h_, this->pad_w_, bottom_diff);
+ count, top_diff, top[0]->num(), channels_,
+ height_, width_, pooled_height_, pooled_width_, kernel_h_,
+ kernel_w_, stride_h_, stride_w_, pad_h_, pad_w_, bottom_diff);
break;
case PoolingParameter_PoolMethod_STOCHASTIC:
// NOLINT_NEXT_LINE(whitespace/operators)
StoPoolBackward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
- count, rand_idx_.gpu_data(), top_diff, top[0]->num(), this->channels_,
- this->height_, this->width_, this->pooled_height_, this->pooled_width_,
- this->kernel_h_, this->kernel_w_, this->stride_h_, this->stride_w_,
+ count, rand_idx_.gpu_data(), top_diff,
+ top[0]->num(), channels_, height_, width_, pooled_height_,
+ pooled_width_, kernel_h_, kernel_w_, stride_h_, stride_w_,
bottom_diff);
break;
default:
}
-INSTANTIATE_CLASS(CaffePoolingLayer);
+INSTANTIATE_CLASS(PoolingLayer);
} // namespace caffe
namespace caffe {
template <typename Dtype>
-void CaffeReLULayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+void ReLULayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) {
const Dtype* bottom_data = bottom[0]->cpu_data();
Dtype* top_data = (*top)[0]->mutable_cpu_data();
}
template <typename Dtype>
-void CaffeReLULayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
+void ReLULayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down,
vector<Blob<Dtype>*>* bottom) {
if (propagate_down[0]) {
}
}
+
#ifdef CPU_ONLY
-STUB_GPU(CaffeReLULayer);
+STUB_GPU(ReLULayer);
#endif
-INSTANTIATE_CLASS(CaffeReLULayer);
+INSTANTIATE_CLASS(ReLULayer);
+
} // namespace caffe
}
template <typename Dtype>
-void CaffeReLULayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+void ReLULayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) {
const Dtype* bottom_data = bottom[0]->gpu_data();
Dtype* top_data = (*top)[0]->mutable_gpu_data();
}
template <typename Dtype>
-void CaffeReLULayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
+void ReLULayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down,
vector<Blob<Dtype>*>* bottom) {
if (propagate_down[0]) {
}
}
-INSTANTIATE_CLASS(CaffeReLULayer);
+
+INSTANTIATE_CLASS(ReLULayer);
+
} // namespace caffe
}
template <typename Dtype>
-void CaffeSigmoidLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+void SigmoidLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) {
const Dtype* bottom_data = bottom[0]->cpu_data();
Dtype* top_data = (*top)[0]->mutable_cpu_data();
}
template <typename Dtype>
-void CaffeSigmoidLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
+void SigmoidLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down,
vector<Blob<Dtype>*>* bottom) {
if (propagate_down[0]) {
}
#ifdef CPU_ONLY
-STUB_GPU(CaffeSigmoidLayer);
+STUB_GPU(SigmoidLayer);
#endif
-INSTANTIATE_CLASS(CaffeSigmoidLayer);
+INSTANTIATE_CLASS(SigmoidLayer);
} // namespace caffe
}
template <typename Dtype>
-void CaffeSigmoidLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+void SigmoidLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) {
const Dtype* bottom_data = bottom[0]->gpu_data();
Dtype* top_data = (*top)[0]->mutable_gpu_data();
}
template <typename Dtype>
-void CaffeSigmoidLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
+void SigmoidLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down,
vector<Blob<Dtype>*>* bottom) {
if (propagate_down[0]) {
}
}
-INSTANTIATE_CLASS(CaffeSigmoidLayer);
+INSTANTIATE_CLASS(SigmoidLayer);
} // namespace caffe
vector<Blob<Dtype>*>* top) {
(*top)[0]->Reshape(bottom[0]->num(), bottom[0]->channels(),
bottom[0]->height(), bottom[0]->width());
+ sum_multiplier_.Reshape(1, bottom[0]->channels(), 1, 1);
+ Dtype* multiplier_data = sum_multiplier_.mutable_cpu_data();
+ for (int i = 0; i < sum_multiplier_.count(); ++i) {
+ multiplier_data[i] = 1.;
+ }
+ scale_.Reshape(bottom[0]->num(), 1, bottom[0]->height(), bottom[0]->width());
}
+template <typename Dtype>
+void SoftmaxLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+ vector<Blob<Dtype>*>* top) {
+ const Dtype* bottom_data = bottom[0]->cpu_data();
+ Dtype* top_data = (*top)[0]->mutable_cpu_data();
+ Dtype* scale_data = scale_.mutable_cpu_data();
+ int num = bottom[0]->num();
+ int channels = bottom[0]->channels();
+ int dim = bottom[0]->count() / bottom[0]->num();
+ int spatial_dim = bottom[0]->height() * bottom[0]->width();
+ caffe_copy(bottom[0]->count(), bottom_data, top_data);
+ // We need to subtract the max to avoid numerical issues, compute the exp,
+ // and then normalize.
+ for (int i = 0; i < num; ++i) {
+ // initialize scale_data to the first plane
+ caffe_copy(spatial_dim, bottom_data + i * dim, scale_data);
+ for (int j = 0; j < channels; j++) {
+ for (int k = 0; k < spatial_dim; k++) {
+ scale_data[k] = std::max(scale_data[k],
+ bottom_data[i * dim + j * spatial_dim + k]);
+ }
+ }
+ // subtraction
+ caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, channels, spatial_dim,
+ 1, -1., sum_multiplier_.cpu_data(), scale_data, 1., top_data + i * dim);
+ // exponentiation
+ caffe_exp<Dtype>(dim, top_data + i * dim, top_data + i * dim);
+ // sum after exp
+ caffe_cpu_gemv<Dtype>(CblasTrans, channels, spatial_dim, 1.,
+ top_data + i * dim, sum_multiplier_.cpu_data(), 0., scale_data);
+ // division
+ for (int j = 0; j < channels; j++) {
+ caffe_div(spatial_dim, top_data + (*top)[0]->offset(i, j), scale_data,
+ top_data + (*top)[0]->offset(i, j));
+ }
+ }
+}
+
+template <typename Dtype>
+void SoftmaxLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
+ const vector<bool>& propagate_down,
+ vector<Blob<Dtype>*>* bottom) {
+ const Dtype* top_diff = top[0]->cpu_diff();
+ const Dtype* top_data = top[0]->cpu_data();
+ Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff();
+ Dtype* scale_data = scale_.mutable_cpu_data();
+ int num = top[0]->num();
+ int channels = top[0]->channels();
+ int dim = top[0]->count() / top[0]->num();
+ int spatial_dim = top[0]->height() * top[0]->width();
+ caffe_copy(top[0]->count(), top_diff, bottom_diff);
+ for (int i = 0; i < num; ++i) {
+ // compute dot(top_diff, top_data) and subtract them from the bottom diff
+ for (int k = 0; k < spatial_dim; ++k) {
+ scale_data[k] = caffe_cpu_strided_dot<Dtype>(channels,
+ bottom_diff + i * dim + k, spatial_dim,
+ top_data + i * dim + k, spatial_dim);
+ }
+ // subtraction
+ caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, channels, spatial_dim, 1,
+ -1., sum_multiplier_.cpu_data(), scale_data, 1., bottom_diff + i * dim);
+ }
+ // elementwise multiplication
+ caffe_mul(top[0]->count(), bottom_diff, top_data, bottom_diff);
+}
+
+
+#ifdef CPU_ONLY
+STUB_GPU(SoftmaxLayer);
+#endif
+
INSTANTIATE_CLASS(SoftmaxLayer);
+
} // namespace caffe
}
template <typename Dtype>
-void CaffeSoftmaxLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+void SoftmaxLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) {
const Dtype* bottom_data = bottom[0]->gpu_data();
Dtype* top_data = (*top)[0]->mutable_gpu_data();
}
template <typename Dtype>
-void CaffeSoftmaxLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
+void SoftmaxLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) {
const Dtype* top_diff = top[0]->gpu_diff();
const Dtype* top_data = top[0]->gpu_data();
caffe_gpu_mul<Dtype>(top[0]->count(), bottom_diff, top_data, bottom_diff);
}
-INSTANTIATE_CLASS(CaffeSoftmaxLayer);
+INSTANTIATE_CLASS(SoftmaxLayer);
+
} // namespace caffe
+// TanH neuron activation function layer.
+// Adapted from ReLU layer code written by Yangqing Jia
+
#include <algorithm>
#include <vector>
namespace caffe {
template <typename Dtype>
-void CaffeTanHLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+void TanHLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) {
const Dtype* bottom_data = bottom[0]->cpu_data();
Dtype* top_data = (*top)[0]->mutable_cpu_data();
}
template <typename Dtype>
-void CaffeTanHLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
+void TanHLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down,
vector<Blob<Dtype>*>* bottom) {
if (propagate_down[0]) {
}
#ifdef CPU_ONLY
-STUB_GPU(CaffeTanHLayer);
+STUB_GPU(TanHLayer);
#endif
-INSTANTIATE_CLASS(CaffeTanHLayer);
+INSTANTIATE_CLASS(TanHLayer);
} // namespace caffe
+// TanH neuron activation function layer.
+// Adapted from ReLU layer code written by Yangqing Jia
+
#include <algorithm>
#include <vector>
}
template <typename Dtype>
-void CaffeTanHLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+void TanHLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) {
const Dtype* bottom_data = bottom[0]->gpu_data();
Dtype* top_data = (*top)[0]->mutable_gpu_data();
}
template <typename Dtype>
-void CaffeTanHLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
+void TanHLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down,
vector<Blob<Dtype>*>* bottom) {
if (propagate_down[0]) {
}
}
-INSTANTIATE_CLASS(CaffeTanHLayer);
+INSTANTIATE_CLASS(TanHLayer);
+
} // namespace caffe
this->blob_bottom_vec_.push_back(this->blob_bottom_2_);
this->blob_top_vec_.push_back(this->blob_top_2_);
shared_ptr<Layer<Dtype> > layer(
- new CaffeConvolutionLayer<Dtype>(layer_param));
+ new ConvolutionLayer<Dtype>(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(), 4);
// setting group should not change the shape
convolution_param->set_num_output(3);
convolution_param->set_group(3);
- layer.reset(new CaffeConvolutionLayer<Dtype>(layer_param));
+ layer.reset(new ConvolutionLayer<Dtype>(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);
convolution_param->mutable_bias_filler()->set_type("constant");
convolution_param->mutable_bias_filler()->set_value(0.1);
shared_ptr<Layer<Dtype> > layer(
- new CaffeConvolutionLayer<Dtype>(layer_param));
+ new ConvolutionLayer<Dtype>(layer_param));
layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_));
// After the convolution, the output should all have output values 27.1
convolution_param->mutable_bias_filler()->set_type("constant");
convolution_param->mutable_bias_filler()->set_value(0.1);
shared_ptr<Layer<Dtype> > layer(
- new CaffeConvolutionLayer<Dtype>(layer_param));
+ new ConvolutionLayer<Dtype>(layer_param));
layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_));
// After the convolution, the output should all have output values 9.1
convolution_param->set_num_output(1);
convolution_param->set_bias_term(false);
shared_ptr<Layer<Dtype> > layer(
- new CaffeConvolutionLayer<Dtype>(layer_param));
+ new ConvolutionLayer<Dtype>(layer_param));
layer->blobs().resize(1);
layer->blobs()[0].reset(new Blob<Dtype>(1, 3, 3, 3));
Dtype* weights = layer->blobs()[0]->mutable_cpu_data();
convolution_param->set_stride_w(1);
convolution_param->set_num_output(1);
convolution_param->set_bias_term(false);
- layer.reset(new CaffeConvolutionLayer<Dtype>(layer_param));
+ layer.reset(new ConvolutionLayer<Dtype>(layer_param));
layer->blobs().resize(1);
layer->blobs()[0].reset(new Blob<Dtype>(1, 3, 3, 1));
Dtype* weights_1 = layer->blobs()[0]->mutable_cpu_data();
convolution_param->set_stride_w(2);
convolution_param->set_num_output(1);
convolution_param->set_bias_term(false);
- layer.reset(new CaffeConvolutionLayer<Dtype>(layer_param));
+ layer.reset(new ConvolutionLayer<Dtype>(layer_param));
layer->blobs().resize(1);
layer->blobs()[0].reset(new Blob<Dtype>(1, 3, 1, 3));
Dtype* weights_2 = layer->blobs()[0]->mutable_cpu_data();
convolution_param->set_num_output(2);
convolution_param->mutable_weight_filler()->set_type("gaussian");
convolution_param->mutable_bias_filler()->set_type("gaussian");
- CaffeConvolutionLayer<Dtype> layer(layer_param);
+ ConvolutionLayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-2, 1e-3);
checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
convolution_param->set_group(3);
convolution_param->mutable_weight_filler()->set_type("gaussian");
convolution_param->mutable_bias_filler()->set_type("gaussian");
- CaffeConvolutionLayer<Dtype> layer(layer_param);
+ ConvolutionLayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-2, 1e-3);
checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
PoolingParameter* pooling_param = layer_param.mutable_pooling_param();
pooling_param->set_kernel_size(3);
pooling_param->set_stride(2);
- CaffePoolingLayer<Dtype> max_layer(layer_param);
+ PoolingLayer<Dtype> max_layer(layer_param);
max_layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
DropoutLayer<Dtype> dropout_layer(layer_param);
dropout_layer.SetUp(this->blob_top_vec_, &(this->blob_top_vec_));
PoolingParameter* pooling_param = layer_param.mutable_pooling_param();
pooling_param->set_kernel_size(3);
pooling_param->set_stride(2);
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> layer(layer_param);
layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_));
const Dtype* top_data = this->blob_top_->cpu_data();
PoolingParameter* pooling_param = layer_param.mutable_pooling_param();
pooling_param->set_kernel_size(3);
pooling_param->set_stride(2);
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> layer(layer_param);
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) {
TYPED_TEST(NeuronLayerTest, TestReLU) {
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
- CaffeReLULayer<Dtype> layer(layer_param);
+ ReLULayer<Dtype> 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
TYPED_TEST(NeuronLayerTest, TestReLUGradient) {
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
- CaffeReLULayer<Dtype> layer(layer_param);
+ ReLULayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-2, 1e-3, 1701, 0., 0.01);
checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
layer_param.ParseFromString("relu_param{negative_slope:0.01}");
- CaffeReLULayer<Dtype> layer(layer_param);
+ ReLULayer<Dtype> 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
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
layer_param.ParseFromString("relu_param{negative_slope:0.01}");
- CaffeReLULayer<Dtype> layer(layer_param);
+ ReLULayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-2, 1e-3, 1701, 0., 0.01);
checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
TYPED_TEST(NeuronLayerTest, TestSigmoid) {
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
- CaffeSigmoidLayer<Dtype> layer(layer_param);
+ SigmoidLayer<Dtype> 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
TYPED_TEST(NeuronLayerTest, TestSigmoidGradient) {
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
- CaffeSigmoidLayer<Dtype> layer(layer_param);
+ SigmoidLayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-2, 1e-3, 1701, 0., 0.01);
checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
TYPED_TEST(NeuronLayerTest, TestTanH) {
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
- CaffeTanHLayer<Dtype> layer(layer_param);
+ TanHLayer<Dtype> layer(layer_param);
layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_));
// Test exact values
TYPED_TEST(NeuronLayerTest, TestTanHGradient) {
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
- CaffeTanHLayer<Dtype> layer(layer_param);
+ TanHLayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-2, 1e-3);
checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
blob_bottom_->mutable_cpu_data()[i + 13] = 2;
blob_bottom_->mutable_cpu_data()[i + 14] = 3;
}
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> layer(layer_param);
layer.SetUp(blob_bottom_vec_, &blob_top_vec_);
EXPECT_EQ(blob_top_->num(), num);
EXPECT_EQ(blob_top_->channels(), channels);
blob_bottom_->mutable_cpu_data()[i + 34] = 18;
blob_bottom_->mutable_cpu_data()[i + 35] = 11;
}
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> layer(layer_param);
layer.SetUp(blob_bottom_vec_, &blob_top_vec_);
EXPECT_EQ(blob_top_->num(), num);
EXPECT_EQ(blob_top_->channels(), channels);
blob_bottom_->mutable_cpu_data()[i + 34] = 18;
blob_bottom_->mutable_cpu_data()[i + 35] = 11;
}
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> layer(layer_param);
layer.SetUp(blob_bottom_vec_, &blob_top_vec_);
EXPECT_EQ(blob_top_->num(), num);
EXPECT_EQ(blob_top_->channels(), channels);
PoolingParameter* pooling_param = layer_param.mutable_pooling_param();
pooling_param->set_kernel_size(3);
pooling_param->set_stride(2);
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> 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());
pooling_param->set_stride(2);
pooling_param->set_pad(1);
pooling_param->set_pool(PoolingParameter_PoolMethod_AVE);
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> 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());
pooling_param->set_stride(2);
pooling_param->set_pad(1);
pooling_param->set_pool(PoolingParameter_PoolMethod_MAX);
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-4, 1e-2);
checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
this->blob_bottom_->mutable_cpu_data()[6] = 4;
this->blob_bottom_->mutable_cpu_data()[7] = 2;
this->blob_bottom_->mutable_cpu_data()[8] = 1;
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> 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);
pooling_param->set_stride(2);
pooling_param->set_pool(PoolingParameter_PoolMethod_MAX);
this->blob_top_vec_.push_back(this->blob_top_mask_);
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-4, 1e-2);
checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
filler_param.set_value(Dtype(2));
ConstantFiller<Dtype> filler(filler_param);
filler.Fill(this->blob_bottom_);
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> 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);
pooling_param->set_kernel_w(kernel_w);
pooling_param->set_stride(2);
pooling_param->set_pool(PoolingParameter_PoolMethod_AVE);
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-2, 1e-2);
checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
pooling_param->set_stride(2);
pooling_param->set_pad(2);
pooling_param->set_pool(PoolingParameter_PoolMethod_AVE);
- CaffePoolingLayer<Dtype> layer(layer_param);
+ PoolingLayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-2, 1e-2);
checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
TYPED_TEST(SoftmaxLayerTest, TestForward) {
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
- CaffeSoftmaxLayer<Dtype> layer(layer_param);
+ SoftmaxLayer<Dtype> layer(layer_param);
layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_));
// Test sum
TYPED_TEST(SoftmaxLayerTest, TestGradient) {
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;
- CaffeSoftmaxLayer<Dtype> layer(layer_param);
+ SoftmaxLayer<Dtype> layer(layer_param);
GradientChecker<Dtype> checker(1e-2, 1e-3);
checker.CheckGradientExhaustive(&layer, &(this->blob_bottom_vec_),
&(this->blob_top_vec_));
PoolingParameter* pooling_param = layer_param.mutable_pooling_param();
pooling_param->set_kernel_size(3);
pooling_param->set_stride(2);
- CaffePoolingLayer<TypeParam> layer(layer_param);
+ PoolingLayer<TypeParam> 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());
pooling_param->set_kernel_size(3);
pooling_param->set_stride(2);
pooling_param->set_pool(PoolingParameter_PoolMethod_STOCHASTIC);
- CaffePoolingLayer<TypeParam> layer(layer_param);
+ PoolingLayer<TypeParam> layer(layer_param);
layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_));
pooling_param->set_kernel_size(3);
pooling_param->set_stride(2);
pooling_param->set_pool(PoolingParameter_PoolMethod_STOCHASTIC);
- CaffePoolingLayer<TypeParam> layer(layer_param);
+ PoolingLayer<TypeParam> layer(layer_param);
layer.SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_));
pooling_param->set_kernel_size(3);
pooling_param->set_stride(2);
pooling_param->set_pool(PoolingParameter_PoolMethod_STOCHASTIC);
- CaffePoolingLayer<TypeParam> layer(layer_param);
+ PoolingLayer<TypeParam> layer(layer_param);
GradientChecker<TypeParam> checker(1e-4, 1e-2);
// it is too expensive to call curand multiple times, so we don't do an
// exhaustive gradient check.
projects / platform / upstream / caffe.git / commitdiff