From 748aaff59c6169cc0cb443380b68e4153a47db00 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Thu, 13 Mar 2014 17:50:15 -0700 Subject: [PATCH 01/11] change specification of forward/backward function and fix layer definitions appropriately --- include/caffe/layer.hpp | 34 ++--- include/caffe/net.hpp | 13 +- include/caffe/vision_layers.hpp | 173 +++++++++++------------ src/caffe/layers/bnll_layer.cpp | 6 +- src/caffe/layers/bnll_layer.cu | 6 +- src/caffe/layers/conv_layer.cpp | 6 +- src/caffe/layers/conv_layer.cu | 6 +- src/caffe/layers/data_layer.cpp | 8 +- src/caffe/layers/data_layer.cu | 8 +- src/caffe/layers/dropout_layer.cpp | 6 +- src/caffe/layers/dropout_layer.cu | 6 +- src/caffe/layers/flatten_layer.cpp | 6 +- src/caffe/layers/flatten_layer.cu | 6 +- src/caffe/layers/hdf5_data_layer.cpp | 9 +- src/caffe/layers/hdf5_data_layer.cu | 6 +- src/caffe/layers/im2col_layer.cpp | 6 +- src/caffe/layers/im2col_layer.cu | 6 +- src/caffe/layers/inner_product_layer.cpp | 6 +- src/caffe/layers/inner_product_layer.cu | 6 +- src/caffe/layers/loss_layer.cpp | 72 +++++++--- src/caffe/layers/lrn_layer.cpp | 7 +- src/caffe/layers/lrn_layer.cu | 6 +- src/caffe/layers/pooling_layer.cpp | 8 +- src/caffe/layers/pooling_layer.cu | 8 +- src/caffe/layers/relu_layer.cpp | 6 +- src/caffe/layers/relu_layer.cu | 6 +- src/caffe/layers/sigmoid_layer.cpp | 6 +- src/caffe/layers/sigmoid_layer.cu | 6 +- src/caffe/layers/softmax_layer.cpp | 6 +- src/caffe/layers/softmax_layer.cu | 6 +- src/caffe/layers/softmax_loss_layer.cpp | 22 +-- src/caffe/layers/softmax_loss_layer.cu | 7 +- src/caffe/layers/split_layer.cpp | 6 +- src/caffe/layers/split_layer.cu | 6 +- src/caffe/layers/tanh_layer.cpp | 6 +- src/caffe/layers/tanh_layer.cu | 6 +- src/caffe/net.cpp | 31 ++-- 37 files changed, 295 insertions(+), 249 deletions(-) diff --git a/include/caffe/layer.hpp b/include/caffe/layer.hpp index a0cb487e50d..ad36c827734 100644 --- a/include/caffe/layer.hpp +++ b/include/caffe/layer.hpp @@ -34,12 +34,13 @@ class Layer { virtual void SetUp(const vector*>& bottom, vector*>* top) = 0; + // Forward and backward wrappers. You should implement the cpu and // gpu specific implementations instead, and should not change these // functions. - inline void Forward(const vector*>& bottom, + inline Dtype Forward(const vector*>& bottom, vector*>* top); - inline Dtype Backward(const vector*>& top, + inline void Backward(const vector*>& top, const bool propagate_down, vector*>* bottom); @@ -60,26 +61,26 @@ class Layer { vector > > blobs_; // Forward functions - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top) = 0; // If no gpu code is provided, we will simply use cpu code. - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top) { // LOG(WARNING) << "Using CPU code as backup."; - Forward_cpu(bottom, top); + return Forward_cpu(bottom, top); } // Backward functions: the backward function will compute the gradients for // any parameters and also for the bottom blobs if propagate_down is true. // It will return the loss produced from this layer. - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) = 0; - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { // LOG(WARNING) << "Using CPU code as backup."; - return Backward_cpu(top, propagate_down, bottom); + Backward_cpu(top, propagate_down, bottom); } DISABLE_COPY_AND_ASSIGN(Layer); @@ -89,29 +90,30 @@ class Layer { // gpu specific implementations instead, and should not change these // functions. template -inline void Layer::Forward(const vector*>& bottom, +inline Dtype Layer::Forward(const vector*>& bottom, vector*>* top) { switch (Caffe::mode()) { case Caffe::CPU: - Forward_cpu(bottom, top); - break; + return Forward_cpu(bottom, top); case Caffe::GPU: - Forward_gpu(bottom, top); - break; + return Forward_gpu(bottom, top); default: LOG(FATAL) << "Unknown caffe mode."; + return Dtype(0); } } template -inline Dtype Layer::Backward(const vector*>& top, +inline void Layer::Backward(const vector*>& top, const bool propagate_down, vector*>* bottom) { switch (Caffe::mode()) { case Caffe::CPU: - return Backward_cpu(top, propagate_down, bottom); + Backward_cpu(top, propagate_down, bottom); + break; case Caffe::GPU: - return Backward_gpu(top, propagate_down, bottom); + Backward_gpu(top, propagate_down, bottom); + break; default: LOG(FATAL) << "Unknown caffe mode."; } diff --git a/include/caffe/net.hpp b/include/caffe/net.hpp index b5a57b3c5a4..d6c892fc3e6 100644 --- a/include/caffe/net.hpp +++ b/include/caffe/net.hpp @@ -31,21 +31,26 @@ class Net { // Run forward with the input blobs already fed separately. You can get the // input blobs using input_blobs(). + const vector*>& ForwardPrefilled(Dtype* loss); const vector*>& ForwardPrefilled(); // Run forward using a set of bottom blobs, and return the result. + const vector*>& Forward(const vector* > & bottom, + Dtype* loss); const vector*>& Forward(const vector* > & bottom); // Run forward using a serialized BlobProtoVector and return the result // as a serialized BlobProtoVector - string Forward(const string& input_blob_protos); + string Forward(const string& input_blob_protos, Dtype* loss); // The network backward should take no input and output, since it solely // computes the gradient w.r.t the parameters, and the data has already // been provided during the forward pass. - Dtype Backward(); + void Backward(); Dtype ForwardBackward(const vector* > & bottom) { - Forward(bottom); - return Backward(); + Dtype loss; + Forward(bottom, &loss); + Backward(); + return loss; } // Updates the network weights based on the diff values computed. diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index 622556396c1..c1729ddc899 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -37,14 +37,14 @@ class ReLULayer : public NeuronLayer { : NeuronLayer(param) {} protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); }; @@ -55,14 +55,14 @@ class TanHLayer : public NeuronLayer { : NeuronLayer(param) {} protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); }; @@ -73,14 +73,14 @@ class SigmoidLayer : public NeuronLayer { : NeuronLayer(param) {} protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); }; @@ -92,14 +92,14 @@ class BNLLLayer : public NeuronLayer { : NeuronLayer(param) {} protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); }; @@ -113,14 +113,14 @@ class DropoutLayer : public NeuronLayer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); shared_ptr rand_vec_; float threshold_; @@ -138,13 +138,13 @@ class SplitLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); int count_; }; @@ -159,13 +159,13 @@ class FlattenLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); int count_; }; @@ -180,14 +180,14 @@ class InnerProductLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); int M_; int K_; @@ -233,13 +233,13 @@ class LRNLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); // scale_ stores the intermediate summing results Blob scale_; @@ -263,13 +263,13 @@ class Im2colLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); int KSIZE_; int STRIDE_; @@ -288,13 +288,13 @@ class PoolingLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); int KSIZE_; int STRIDE_; @@ -316,13 +316,13 @@ class ConvolutionLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); Blob col_bob_; @@ -387,14 +387,14 @@ class DataLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, - const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, - const bool propagate_down, vector*>* bottom); + virtual void Backward_cpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { return; } + virtual void Backward_gpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { return; } shared_ptr db_; shared_ptr iter_; @@ -457,13 +457,13 @@ class HDF5DataLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); virtual void load_hdf5_file_data(const char* filename); @@ -486,13 +486,13 @@ class SoftmaxLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); // sum_multiplier is just used to carry out sum using blas @@ -513,13 +513,13 @@ class MultinomialLogisticLossLayer : public Layer { protected: // The loss layer will do nothing during forward - all computation are // carried out in the backward pass. - virtual void Forward_cpu(const vector*>& bottom, - vector*>* top) { return; } - virtual void Forward_gpu(const vector*>& bottom, - vector*>* top) { return; } - virtual Dtype Backward_cpu(const vector*>& top, + 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 Dtype Backward_gpu(const vector*>& top, + // virtual void Backward_gpu(const vector*>& top, // const bool propagate_down, vector*>* bottom); }; @@ -534,13 +534,13 @@ class InfogainLossLayer : public Layer { protected: // The loss layer will do nothing during forward - all computation are // carried out in the backward pass. - virtual void Forward_cpu(const vector*>& bottom, - vector*>* top) { return; } - virtual void Forward_gpu(const vector*>& bottom, - vector*>* top) { return; } - virtual Dtype Backward_cpu(const vector*>& top, + 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 Dtype Backward_gpu(const vector*>& top, + // virtual void Backward_gpu(const vector*>& top, // const bool propagate_down, vector*>* bottom); Blob infogain_; @@ -561,13 +561,13 @@ class SoftmaxWithLossLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); shared_ptr > softmax_layer_; @@ -590,13 +590,13 @@ class EuclideanLossLayer : public Layer { protected: // The loss layer will do nothing during forward - all computation are // carried out in the backward pass. - virtual void Forward_cpu(const vector*>& bottom, - vector*>* top) { return; } - virtual void Forward_gpu(const vector*>& bottom, - vector*>* top) { return; } - virtual Dtype Backward_cpu(const vector*>& top, + 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 Dtype Backward_gpu(const vector*>& top, + // virtual void Backward_gpu(const vector*>& top, // const bool propagate_down, vector*>* bottom); Blob difference_; }; @@ -611,13 +611,12 @@ class AccuracyLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); // The accuracy layer should not be used to compute backward operations. - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { NOT_IMPLEMENTED; - return Dtype(0.); } }; diff --git a/src/caffe/layers/bnll_layer.cpp b/src/caffe/layers/bnll_layer.cpp index b769a35212a..e7a4fba2d67 100644 --- a/src/caffe/layers/bnll_layer.cpp +++ b/src/caffe/layers/bnll_layer.cpp @@ -13,7 +13,7 @@ namespace caffe { const float kBNLL_THRESHOLD = 50.; template -void BNLLLayer::Forward_cpu(const vector*>& bottom, +Dtype BNLLLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -23,10 +23,11 @@ void BNLLLayer::Forward_cpu(const vector*>& bottom, bottom_data[i] + log(1. + exp(-bottom_data[i])) : log(1. + exp(bottom_data[i])); } + return Dtype(0); } template -Dtype BNLLLayer::Backward_cpu(const vector*>& top, +void BNLLLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { @@ -40,7 +41,6 @@ Dtype BNLLLayer::Backward_cpu(const vector*>& top, bottom_diff[i] = top_diff[i] * expval / (expval + 1.); } } - return Dtype(0); } diff --git a/src/caffe/layers/bnll_layer.cu b/src/caffe/layers/bnll_layer.cu index 1fd200894c3..7252c0222c0 100644 --- a/src/caffe/layers/bnll_layer.cu +++ b/src/caffe/layers/bnll_layer.cu @@ -22,7 +22,7 @@ __global__ void BNLLForward(const int n, const Dtype* in, Dtype* out) { } template -void BNLLLayer::Forward_gpu(const vector*>& bottom, +Dtype BNLLLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -31,6 +31,7 @@ void BNLLLayer::Forward_gpu(const vector*>& bottom, BNLLForward<<>>( count, bottom_data, top_data); CUDA_POST_KERNEL_CHECK; + return Dtype(0); } template @@ -43,7 +44,7 @@ __global__ void BNLLBackward(const int n, const Dtype* in_diff, } template -Dtype BNLLLayer::Backward_gpu(const vector*>& top, +void BNLLLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { @@ -56,7 +57,6 @@ Dtype BNLLLayer::Backward_gpu(const vector*>& top, count, top_diff, bottom_data, bottom_diff); CUDA_POST_KERNEL_CHECK; } - return Dtype(0); } INSTANTIATE_CLASS(BNLLLayer); diff --git a/src/caffe/layers/conv_layer.cpp b/src/caffe/layers/conv_layer.cpp index 64a652a8e1d..cb1bca6579c 100644 --- a/src/caffe/layers/conv_layer.cpp +++ b/src/caffe/layers/conv_layer.cpp @@ -77,7 +77,7 @@ void ConvolutionLayer::SetUp(const vector*>& bottom, template -void ConvolutionLayer::Forward_cpu(const vector*>& bottom, +Dtype ConvolutionLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -104,10 +104,11 @@ void ConvolutionLayer::Forward_cpu(const vector*>& bottom, (Dtype)1., top_data + (*top)[0]->offset(n)); } } + return Dtype(0.); } template -Dtype ConvolutionLayer::Backward_cpu(const vector*>& top, +void ConvolutionLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->cpu_diff(); const Dtype* weight = this->blobs_[0]->cpu_data(); @@ -159,7 +160,6 @@ Dtype ConvolutionLayer::Backward_cpu(const vector*>& top, bottom_diff + (*bottom)[0]->offset(n)); } } - return Dtype(0.); } INSTANTIATE_CLASS(ConvolutionLayer); diff --git a/src/caffe/layers/conv_layer.cu b/src/caffe/layers/conv_layer.cu index a7f56faa97b..f8f605584d1 100644 --- a/src/caffe/layers/conv_layer.cu +++ b/src/caffe/layers/conv_layer.cu @@ -11,7 +11,7 @@ namespace caffe { template -void ConvolutionLayer::Forward_gpu(const vector*>& bottom, +Dtype ConvolutionLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -38,10 +38,11 @@ void ConvolutionLayer::Forward_gpu(const vector*>& bottom, (Dtype)1., top_data + (*top)[0]->offset(n)); } } + return Dtype(0.); } template -Dtype ConvolutionLayer::Backward_gpu(const vector*>& top, +void ConvolutionLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->gpu_diff(); const Dtype* weight = this->blobs_[0]->gpu_data(); @@ -95,7 +96,6 @@ Dtype ConvolutionLayer::Backward_gpu(const vector*>& top, bottom_diff + (*bottom)[0]->offset(n)); } } - return Dtype(0.); } diff --git a/src/caffe/layers/data_layer.cpp b/src/caffe/layers/data_layer.cpp index cc03cdbf0b7..f2ff7ff1d93 100644 --- a/src/caffe/layers/data_layer.cpp +++ b/src/caffe/layers/data_layer.cpp @@ -213,7 +213,7 @@ void DataLayer::SetUp(const vector*>& bottom, } template -void DataLayer::Forward_cpu(const vector*>& bottom, +Dtype DataLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { // First, join the thread CHECK(!pthread_join(thread_, NULL)) << "Pthread joining failed."; @@ -225,12 +225,6 @@ void DataLayer::Forward_cpu(const vector*>& bottom, // Start a new prefetch thread CHECK(!pthread_create(&thread_, NULL, DataLayerPrefetch, reinterpret_cast(this))) << "Pthread execution failed."; -} - -// The backward operations are dummy - they do not carry any computation. -template -Dtype DataLayer::Backward_cpu(const vector*>& top, - const bool propagate_down, vector*>* bottom) { return Dtype(0.); } diff --git a/src/caffe/layers/data_layer.cu b/src/caffe/layers/data_layer.cu index 946f30f3b7f..57a375ea205 100644 --- a/src/caffe/layers/data_layer.cu +++ b/src/caffe/layers/data_layer.cu @@ -16,7 +16,7 @@ using std::string; namespace caffe { template -void DataLayer::Forward_gpu(const vector*>& bottom, +Dtype DataLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { // First, join the thread CHECK(!pthread_join(thread_, NULL)) << "Pthread joining failed."; @@ -30,12 +30,6 @@ void DataLayer::Forward_gpu(const vector*>& bottom, // Start a new prefetch thread CHECK(!pthread_create(&thread_, NULL, DataLayerPrefetch, reinterpret_cast(this))) << "Pthread execution failed."; -} - -// The backward operations are dummy - they do not carry any computation. -template -Dtype DataLayer::Backward_gpu(const vector*>& top, - const bool propagate_down, vector*>* bottom) { return Dtype(0.); } diff --git a/src/caffe/layers/dropout_layer.cpp b/src/caffe/layers/dropout_layer.cpp index f480853cdf3..6cd6ffa8e6a 100644 --- a/src/caffe/layers/dropout_layer.cpp +++ b/src/caffe/layers/dropout_layer.cpp @@ -23,7 +23,7 @@ void DropoutLayer::SetUp(const vector*>& bottom, } template -void DropoutLayer::Forward_cpu(const vector*>& bottom, +Dtype DropoutLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -39,10 +39,11 @@ void DropoutLayer::Forward_cpu(const vector*>& bottom, } else { memcpy(top_data, bottom_data, bottom[0]->count() * sizeof(Dtype)); } + return Dtype(0); } template -Dtype DropoutLayer::Backward_cpu(const vector*>& top, +void DropoutLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { CHECK(Caffe::phase() == Caffe::TRAIN); @@ -55,7 +56,6 @@ Dtype DropoutLayer::Backward_cpu(const vector*>& top, bottom_diff[i] = top_diff[i] * mask[i] * scale_; } } - return Dtype(0); } diff --git a/src/caffe/layers/dropout_layer.cu b/src/caffe/layers/dropout_layer.cu index 0b38ae2a576..dc1f3cf8740 100644 --- a/src/caffe/layers/dropout_layer.cu +++ b/src/caffe/layers/dropout_layer.cu @@ -24,7 +24,7 @@ __global__ void DropoutForward(const int n, const Dtype* in, } template -void DropoutLayer::Forward_gpu(const vector*>& bottom, +Dtype DropoutLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -42,6 +42,7 @@ void DropoutLayer::Forward_gpu(const vector*>& bottom, CUDA_CHECK(cudaMemcpy(top_data, bottom_data, count * sizeof(Dtype), cudaMemcpyDeviceToDevice)); } + return Dtype(0); } template @@ -54,7 +55,7 @@ __global__ void DropoutBackward(const int n, const Dtype* in_diff, } template -Dtype DropoutLayer::Backward_gpu(const vector*>& top, +void DropoutLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { CHECK(Caffe::phase() == Caffe::TRAIN); @@ -68,7 +69,6 @@ Dtype DropoutLayer::Backward_gpu(const vector*>& top, count, top_diff, mask, uint_thres_, scale_, bottom_diff); CUDA_POST_KERNEL_CHECK; } - return Dtype(0); } INSTANTIATE_CLASS(DropoutLayer); diff --git a/src/caffe/layers/flatten_layer.cpp b/src/caffe/layers/flatten_layer.cpp index 9e17a8200c1..d8d5c4b6053 100644 --- a/src/caffe/layers/flatten_layer.cpp +++ b/src/caffe/layers/flatten_layer.cpp @@ -22,20 +22,20 @@ void FlattenLayer::SetUp(const vector*>& bottom, } template -void FlattenLayer::Forward_cpu(const vector*>& bottom, +Dtype FlattenLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); caffe_copy(count_, bottom_data, top_data); + return Dtype(0.); } template -Dtype FlattenLayer::Backward_cpu(const vector*>& top, +void FlattenLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->cpu_diff(); Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff(); caffe_copy(count_, top_diff, bottom_diff); - return Dtype(0.); } INSTANTIATE_CLASS(FlattenLayer); diff --git a/src/caffe/layers/flatten_layer.cu b/src/caffe/layers/flatten_layer.cu index 571e22e2417..fa1e6aa3141 100644 --- a/src/caffe/layers/flatten_layer.cu +++ b/src/caffe/layers/flatten_layer.cu @@ -9,20 +9,20 @@ namespace caffe { template -void FlattenLayer::Forward_gpu(const vector*>& bottom, +Dtype FlattenLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); caffe_gpu_copy(count_, bottom_data, top_data); + return Dtype(0.); } template -Dtype FlattenLayer::Backward_gpu(const vector*>& top, +void FlattenLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->gpu_diff(); Dtype* bottom_diff = (*bottom)[0]->mutable_gpu_diff(); caffe_gpu_copy(count_, top_diff, bottom_diff); - return Dtype(0.); } INSTANTIATE_CLASS(FlattenLayer); diff --git a/src/caffe/layers/hdf5_data_layer.cpp b/src/caffe/layers/hdf5_data_layer.cpp index e5b17fedb20..3f87dbc512e 100644 --- a/src/caffe/layers/hdf5_data_layer.cpp +++ b/src/caffe/layers/hdf5_data_layer.cpp @@ -89,7 +89,7 @@ void HDF5DataLayer::SetUp(const vector*>& bottom, } template -void HDF5DataLayer::Forward_cpu(const vector*>& bottom, +Dtype HDF5DataLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const int batchsize = this->layer_param_.batchsize(); const int data_count = (*top)[0]->count() / (*top)[0]->num(); @@ -118,14 +118,13 @@ void HDF5DataLayer::Forward_cpu(const vector*>& bottom, &label_blob_.cpu_data()[current_row_ * label_data_count], sizeof(Dtype) * label_data_count); } + return Dtype(0.); } // The backward operations are dummy - they do not carry any computation. template -Dtype HDF5DataLayer::Backward_cpu(const vector*>& top, - const bool propagate_down, vector*>* bottom) { - return Dtype(0.); -} +void HDF5DataLayer::Backward_cpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { } INSTANTIATE_CLASS(HDF5DataLayer); diff --git a/src/caffe/layers/hdf5_data_layer.cu b/src/caffe/layers/hdf5_data_layer.cu index bed7f35a156..261d404d551 100644 --- a/src/caffe/layers/hdf5_data_layer.cu +++ b/src/caffe/layers/hdf5_data_layer.cu @@ -20,7 +20,7 @@ using std::string; namespace caffe { template -void HDF5DataLayer::Forward_gpu(const vector*>& bottom, +Dtype HDF5DataLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const int batchsize = this->layer_param_.batchsize(); const int data_count = (*top)[0]->count() / (*top)[0]->num(); @@ -53,12 +53,12 @@ void HDF5DataLayer::Forward_gpu(const vector*>& bottom, sizeof(Dtype) * label_data_count, cudaMemcpyHostToDevice)); } + return Dtype(0.); } template -Dtype HDF5DataLayer::Backward_gpu(const vector*>& top, +void HDF5DataLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { - return Dtype(0.); } INSTANTIATE_CLASS(HDF5DataLayer); diff --git a/src/caffe/layers/im2col_layer.cpp b/src/caffe/layers/im2col_layer.cpp index e711713b895..a01bfb7c21c 100644 --- a/src/caffe/layers/im2col_layer.cpp +++ b/src/caffe/layers/im2col_layer.cpp @@ -26,7 +26,7 @@ void Im2colLayer::SetUp(const vector*>& bottom, } template -void Im2colLayer::Forward_cpu(const vector*>& bottom, +Dtype Im2colLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -34,10 +34,11 @@ void Im2colLayer::Forward_cpu(const vector*>& bottom, im2col_cpu(bottom_data + bottom[0]->offset(n), CHANNELS_, HEIGHT_, WIDTH_, KSIZE_, PAD_, STRIDE_, top_data + (*top)[0]->offset(n)); } + return Dtype(0.); } template -Dtype Im2colLayer::Backward_cpu(const vector*>& top, +void Im2colLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->cpu_diff(); Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff(); @@ -45,7 +46,6 @@ Dtype Im2colLayer::Backward_cpu(const vector*>& top, col2im_cpu(top_diff + top[0]->offset(n), CHANNELS_, HEIGHT_, WIDTH_, KSIZE_, PAD_, STRIDE_, bottom_diff + (*bottom)[0]->offset(n)); } - return Dtype(0.); } INSTANTIATE_CLASS(Im2colLayer); diff --git a/src/caffe/layers/im2col_layer.cu b/src/caffe/layers/im2col_layer.cu index 2d949b12296..64731cc53d8 100644 --- a/src/caffe/layers/im2col_layer.cu +++ b/src/caffe/layers/im2col_layer.cu @@ -10,7 +10,7 @@ namespace caffe { template -void Im2colLayer::Forward_gpu(const vector*>& bottom, +Dtype Im2colLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -18,10 +18,11 @@ void Im2colLayer::Forward_gpu(const vector*>& bottom, im2col_gpu(bottom_data + bottom[0]->offset(n), CHANNELS_, HEIGHT_, WIDTH_, KSIZE_, PAD_, STRIDE_, top_data + (*top)[0]->offset(n)); } + return Dtype(0.); } template -Dtype Im2colLayer::Backward_gpu(const vector*>& top, +void Im2colLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->gpu_diff(); Dtype* bottom_diff = (*bottom)[0]->mutable_gpu_diff(); @@ -29,7 +30,6 @@ Dtype Im2colLayer::Backward_gpu(const vector*>& top, col2im_gpu(top_diff + top[0]->offset(n), CHANNELS_, HEIGHT_, WIDTH_, KSIZE_, PAD_, STRIDE_, bottom_diff + (*bottom)[0]->offset(n)); } - return Dtype(0.); } diff --git a/src/caffe/layers/inner_product_layer.cpp b/src/caffe/layers/inner_product_layer.cpp index 6987a787ed3..92723ef3b6c 100644 --- a/src/caffe/layers/inner_product_layer.cpp +++ b/src/caffe/layers/inner_product_layer.cpp @@ -61,7 +61,7 @@ void InnerProductLayer::SetUp(const vector*>& bottom, } template -void InnerProductLayer::Forward_cpu(const vector*>& bottom, +Dtype InnerProductLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -73,10 +73,11 @@ void InnerProductLayer::Forward_cpu(const vector*>& bottom, reinterpret_cast(bias_multiplier_->cpu_data()), this->blobs_[1]->cpu_data(), (Dtype)1., top_data); } + return Dtype(0); } template -Dtype InnerProductLayer::Backward_cpu(const vector*>& top, +void InnerProductLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->cpu_diff(); @@ -96,7 +97,6 @@ Dtype InnerProductLayer::Backward_cpu(const vector*>& top, top_diff, this->blobs_[0]->cpu_data(), (Dtype)0., (*bottom)[0]->mutable_cpu_diff()); } - return Dtype(0); } INSTANTIATE_CLASS(InnerProductLayer); diff --git a/src/caffe/layers/inner_product_layer.cu b/src/caffe/layers/inner_product_layer.cu index c7c3e2a99fd..178b488bc60 100644 --- a/src/caffe/layers/inner_product_layer.cu +++ b/src/caffe/layers/inner_product_layer.cu @@ -16,7 +16,7 @@ namespace caffe { template -void InnerProductLayer::Forward_gpu(const vector*>& bottom, +Dtype InnerProductLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -28,10 +28,11 @@ void InnerProductLayer::Forward_gpu(const vector*>& bottom, reinterpret_cast(bias_multiplier_->gpu_data()), this->blobs_[1]->gpu_data(), (Dtype)1., top_data); } + return Dtype(0); } template -Dtype InnerProductLayer::Backward_gpu(const vector*>& top, +void InnerProductLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->gpu_diff(); @@ -51,7 +52,6 @@ Dtype InnerProductLayer::Backward_gpu(const vector*>& top, top_diff, this->blobs_[0]->gpu_data(), (Dtype)0., (*bottom)[0]->mutable_gpu_diff()); } - return Dtype(0); } INSTANTIATE_CLASS(InnerProductLayer); diff --git a/src/caffe/layers/loss_layer.cpp b/src/caffe/layers/loss_layer.cpp index 1c4303d9bd4..3c0f15fb3b3 100644 --- a/src/caffe/layers/loss_layer.cpp +++ b/src/caffe/layers/loss_layer.cpp @@ -28,9 +28,24 @@ void MultinomialLogisticLossLayer::SetUp( CHECK_EQ(bottom[1]->width(), 1); } +template +Dtype MultinomialLogisticLossLayer::Forward_cpu( + const vector*>& bottom, vector*>* top) { + const Dtype* bottom_data = bottom[0]->cpu_data(); + const Dtype* bottom_label = bottom[1]->cpu_data(); + int num = bottom[0]->num(); + int dim = bottom[0]->count() / bottom[0]->num(); + Dtype loss = 0; + for (int i = 0; i < num; ++i) { + int label = static_cast(bottom_label[i]); + Dtype prob = max(bottom_data[i * dim + label], Dtype(kLOG_THRESHOLD)); + loss -= log(prob); + } + return loss / num; +} template -Dtype MultinomialLogisticLossLayer::Backward_cpu( +void MultinomialLogisticLossLayer::Backward_cpu( const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* bottom_data = (*bottom)[0]->cpu_data(); @@ -39,18 +54,13 @@ Dtype MultinomialLogisticLossLayer::Backward_cpu( int num = (*bottom)[0]->num(); int dim = (*bottom)[0]->count() / (*bottom)[0]->num(); memset(bottom_diff, 0, sizeof(Dtype) * (*bottom)[0]->count()); - Dtype loss = 0; for (int i = 0; i < num; ++i) { int label = static_cast(bottom_label[i]); Dtype prob = max(bottom_data[i * dim + label], Dtype(kLOG_THRESHOLD)); - loss -= log(prob); - bottom_diff[i * dim + label] = - 1. / prob / num; + bottom_diff[i * dim + label] = -1. / prob / num; } - return loss / num; } -// TODO: implement the GPU version for multinomial loss - template void InfogainLossLayer::SetUp( @@ -72,7 +82,27 @@ void InfogainLossLayer::SetUp( template -Dtype InfogainLossLayer::Backward_cpu(const vector*>& top, +Dtype InfogainLossLayer::Forward_cpu(const vector*>& bottom, + vector*>* top) { + const Dtype* bottom_data = bottom[0]->cpu_data(); + const Dtype* bottom_label = bottom[1]->cpu_data(); + const Dtype* infogain_mat = infogain_.cpu_data(); + int num = bottom[0]->num(); + int dim = bottom[0]->count() / bottom[0]->num(); + CHECK_EQ(infogain_.height(), dim); + Dtype loss = 0; + for (int i = 0; i < num; ++i) { + int label = static_cast(bottom_label[i]); + for (int j = 0; j < dim; ++j) { + Dtype prob = max(bottom_data[i * dim + j], Dtype(kLOG_THRESHOLD)); + loss -= infogain_mat[label * dim + j] * log(prob); + } + } + return loss / num; +} + +template +void InfogainLossLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* bottom_data = (*bottom)[0]->cpu_data(); @@ -82,16 +112,13 @@ Dtype InfogainLossLayer::Backward_cpu(const vector*>& top, int num = (*bottom)[0]->num(); int dim = (*bottom)[0]->count() / (*bottom)[0]->num(); CHECK_EQ(infogain_.height(), dim); - Dtype loss = 0; for (int i = 0; i < num; ++i) { int label = static_cast(bottom_label[i]); for (int j = 0; j < dim; ++j) { Dtype prob = max(bottom_data[i * dim + j], Dtype(kLOG_THRESHOLD)); - loss -= infogain_mat[label * dim + j] * log(prob); bottom_diff[i * dim + j] = - infogain_mat[label * dim + j] / prob / num; } } - return loss / num; } @@ -110,18 +137,25 @@ void EuclideanLossLayer::SetUp( } template -Dtype EuclideanLossLayer::Backward_cpu(const vector*>& top, - const bool propagate_down, vector*>* bottom) { - int count = (*bottom)[0]->count(); - int num = (*bottom)[0]->num(); - caffe_sub(count, (*bottom)[0]->cpu_data(), (*bottom)[1]->cpu_data(), +Dtype EuclideanLossLayer::Forward_cpu(const vector*>& bottom, + vector*>* top) { + int count = bottom[0]->count(); + int num = bottom[0]->num(); + caffe_sub(count, bottom[0]->cpu_data(), bottom[1]->cpu_data(), difference_.mutable_cpu_data()); Dtype loss = caffe_cpu_dot( count, difference_.cpu_data(), difference_.cpu_data()) / num / Dtype(2); + return loss; +} + +template +void EuclideanLossLayer::Backward_cpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { + int count = (*bottom)[0]->count(); + int num = (*bottom)[0]->num(); // Compute the gradient caffe_axpby(count, Dtype(1) / num, difference_.cpu_data(), Dtype(0), (*bottom)[0]->mutable_cpu_diff()); - return loss; } template @@ -138,7 +172,7 @@ void AccuracyLayer::SetUp( } template -void AccuracyLayer::Forward_cpu(const vector*>& bottom, +Dtype AccuracyLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { Dtype accuracy = 0; Dtype logprob = 0; @@ -166,6 +200,8 @@ void AccuracyLayer::Forward_cpu(const vector*>& bottom, // LOG(INFO) << "Accuracy: " << accuracy; (*top)[0]->mutable_cpu_data()[0] = accuracy / num; (*top)[0]->mutable_cpu_data()[1] = logprob / num; + // Accuracy layer should not be used as a loss function. + return Dtype(0); } INSTANTIATE_CLASS(MultinomialLogisticLossLayer); diff --git a/src/caffe/layers/lrn_layer.cpp b/src/caffe/layers/lrn_layer.cpp index 36dbe41ea8c..698debab6a6 100644 --- a/src/caffe/layers/lrn_layer.cpp +++ b/src/caffe/layers/lrn_layer.cpp @@ -28,7 +28,7 @@ void LRNLayer::SetUp(const vector*>& bottom, } template -void LRNLayer::Forward_cpu(const vector*>& bottom, +Dtype LRNLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -72,10 +72,12 @@ void LRNLayer::Forward_cpu(const vector*>& bottom, // In the end, compute output caffe_powx(scale_.count(), scale_data, -beta_, top_data); caffe_mul(scale_.count(), top_data, bottom_data, top_data); + + return Dtype(0.); } template -Dtype LRNLayer::Backward_cpu(const vector*>& top, +void LRNLayer::Backward_cpu(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(); @@ -126,7 +128,6 @@ Dtype LRNLayer::Backward_cpu(const vector*>& top, padded_ratio_data + padded_ratio.offset(0, c), accum_ratio_data); } } - return Dtype(0.); } INSTANTIATE_CLASS(LRNLayer); diff --git a/src/caffe/layers/lrn_layer.cu b/src/caffe/layers/lrn_layer.cu index 028aa8fa47e..1dcd0c087c0 100644 --- a/src/caffe/layers/lrn_layer.cu +++ b/src/caffe/layers/lrn_layer.cu @@ -65,7 +65,7 @@ __global__ void LRNComputeOutput(const int nthreads, const Dtype* in, } template -void LRNLayer::Forward_gpu(const vector*>& bottom, +Dtype LRNLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { // First, compute scale const Dtype* bottom_data = bottom[0]->gpu_data(); @@ -84,6 +84,7 @@ void LRNLayer::Forward_gpu(const vector*>& bottom, LRNComputeOutput<<>>( n_threads, bottom_data, scale_data, -beta_, top_data); CUDA_POST_KERNEL_CHECK; + return Dtype(0.); } @@ -149,7 +150,7 @@ __global__ void LRNComputeDiff(const int nthreads, const Dtype* bottom_data, } template -Dtype LRNLayer::Backward_gpu(const vector*>& top, +void LRNLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { int n_threads = num_ * height_ * width_; // NOLINT_NEXT_LINE(whitespace/operators) @@ -158,7 +159,6 @@ Dtype LRNLayer::Backward_gpu(const vector*>& top, scale_.gpu_data(), top[0]->gpu_diff(), num_, channels_, height_, width_, size_, -beta_, Dtype(2. * alpha_ * beta_ / size_), (*bottom)[0]->mutable_gpu_diff()); - return Dtype(0.); } diff --git a/src/caffe/layers/pooling_layer.cpp b/src/caffe/layers/pooling_layer.cpp index ce30e842c58..3fd421cd640 100644 --- a/src/caffe/layers/pooling_layer.cpp +++ b/src/caffe/layers/pooling_layer.cpp @@ -39,7 +39,7 @@ void PoolingLayer::SetUp(const vector*>& bottom, // TODO(Yangqing): Is there a faster way to do pooling in the channel-first // case? template -void PoolingLayer::Forward_cpu(const vector*>& bottom, +Dtype PoolingLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -111,13 +111,14 @@ void PoolingLayer::Forward_cpu(const vector*>& bottom, default: LOG(FATAL) << "Unknown pooling method."; } + return Dtype(0.); } template -Dtype PoolingLayer::Backward_cpu(const vector*>& top, +void PoolingLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (!propagate_down) { - return Dtype(0.); + return; } const Dtype* top_diff = top[0]->cpu_diff(); const Dtype* top_data = top[0]->cpu_data(); @@ -188,7 +189,6 @@ Dtype PoolingLayer::Backward_cpu(const vector*>& top, default: LOG(FATAL) << "Unknown pooling method."; } - return Dtype(0.); } diff --git a/src/caffe/layers/pooling_layer.cu b/src/caffe/layers/pooling_layer.cu index 357a392976d..63b4d0dbad7 100644 --- a/src/caffe/layers/pooling_layer.cu +++ b/src/caffe/layers/pooling_layer.cu @@ -135,7 +135,7 @@ __global__ void StoPoolForwardTest(const int nthreads, template -void PoolingLayer::Forward_gpu(const vector*>& bottom, +Dtype PoolingLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -179,6 +179,7 @@ void PoolingLayer::Forward_gpu(const vector*>& bottom, LOG(FATAL) << "Unknown pooling method."; } CUDA_POST_KERNEL_CHECK; + return Dtype(0.); } template @@ -277,10 +278,10 @@ __global__ void StoPoolBackward(const int nthreads, template -Dtype PoolingLayer::Backward_gpu(const vector*>& top, +void PoolingLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (!propagate_down) { - return Dtype(0.); + return; } const Dtype* top_diff = top[0]->gpu_diff(); Dtype* bottom_diff = (*bottom)[0]->mutable_gpu_diff(); @@ -311,7 +312,6 @@ Dtype PoolingLayer::Backward_gpu(const vector*>& top, LOG(FATAL) << "Unknown pooling method."; } CUDA_POST_KERNEL_CHECK; - return Dtype(0.); } diff --git a/src/caffe/layers/relu_layer.cpp b/src/caffe/layers/relu_layer.cpp index 27ae94b7cb0..18c675c98c7 100644 --- a/src/caffe/layers/relu_layer.cpp +++ b/src/caffe/layers/relu_layer.cpp @@ -11,7 +11,7 @@ using std::max; namespace caffe { template -void ReLULayer::Forward_cpu(const vector*>& bottom, +Dtype ReLULayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -19,10 +19,11 @@ void ReLULayer::Forward_cpu(const vector*>& bottom, for (int i = 0; i < count; ++i) { top_data[i] = max(bottom_data[i], Dtype(0)); } + return Dtype(0); } template -Dtype ReLULayer::Backward_cpu(const vector*>& top, +void ReLULayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { @@ -34,7 +35,6 @@ Dtype ReLULayer::Backward_cpu(const vector*>& top, bottom_diff[i] = top_diff[i] * (bottom_data[i] > 0); } } - return Dtype(0); } diff --git a/src/caffe/layers/relu_layer.cu b/src/caffe/layers/relu_layer.cu index 20a5a45e2f4..27f5da5cc89 100644 --- a/src/caffe/layers/relu_layer.cu +++ b/src/caffe/layers/relu_layer.cu @@ -18,7 +18,7 @@ __global__ void ReLUForward(const int n, const Dtype* in, Dtype* out) { } template -void ReLULayer::Forward_gpu(const vector*>& bottom, +Dtype ReLULayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -32,6 +32,7 @@ void ReLULayer::Forward_gpu(const vector*>& bottom, // << " top_data: " << (unsigned long)top_data // << " blocks: " << CAFFE_GET_BLOCKS(count) // << " threads: " << CAFFE_CUDA_NUM_THREADS; + return Dtype(0); } template @@ -43,7 +44,7 @@ __global__ void ReLUBackward(const int n, const Dtype* in_diff, } template -Dtype ReLULayer::Backward_gpu(const vector*>& top, +void ReLULayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { @@ -56,7 +57,6 @@ Dtype ReLULayer::Backward_gpu(const vector*>& top, count, top_diff, bottom_data, bottom_diff); CUDA_POST_KERNEL_CHECK; } - return Dtype(0); } INSTANTIATE_CLASS(ReLULayer); diff --git a/src/caffe/layers/sigmoid_layer.cpp b/src/caffe/layers/sigmoid_layer.cpp index ba6ec84e717..44897954677 100644 --- a/src/caffe/layers/sigmoid_layer.cpp +++ b/src/caffe/layers/sigmoid_layer.cpp @@ -15,7 +15,7 @@ inline Dtype sigmoid(Dtype x) { } template -void SigmoidLayer::Forward_cpu(const vector*>& bottom, +Dtype SigmoidLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -23,10 +23,11 @@ void SigmoidLayer::Forward_cpu(const vector*>& bottom, for (int i = 0; i < count; ++i) { top_data[i] = sigmoid(bottom_data[i]); } + return Dtype(0); } template -Dtype SigmoidLayer::Backward_cpu(const vector*>& top, +void SigmoidLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { @@ -39,7 +40,6 @@ Dtype SigmoidLayer::Backward_cpu(const vector*>& top, bottom_diff[i] = top_diff[i] * sigmoid_x * (1. - sigmoid_x); } } - return Dtype(0); } INSTANTIATE_CLASS(SigmoidLayer); diff --git a/src/caffe/layers/sigmoid_layer.cu b/src/caffe/layers/sigmoid_layer.cu index ba311f814a3..3dbdc397bee 100644 --- a/src/caffe/layers/sigmoid_layer.cu +++ b/src/caffe/layers/sigmoid_layer.cu @@ -24,7 +24,7 @@ __global__ void SigmoidForward(const int n, const Dtype* in, Dtype* out) { } template -void SigmoidLayer::Forward_gpu(const vector*>& bottom, +Dtype SigmoidLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -38,6 +38,7 @@ void SigmoidLayer::Forward_gpu(const vector*>& bottom, // << " top_data: " << (unsigned long)top_data // << " blocks: " << CAFFE_GET_BLOCKS(count) // << " threads: " << CAFFE_CUDA_NUM_THREADS; + return Dtype(0); } template @@ -50,7 +51,7 @@ __global__ void SigmoidBackward(const int n, const Dtype* in_diff, } template -Dtype SigmoidLayer::Backward_gpu(const vector*>& top, +void SigmoidLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { @@ -63,7 +64,6 @@ Dtype SigmoidLayer::Backward_gpu(const vector*>& top, count, top_diff, bottom_data, bottom_diff); CUDA_POST_KERNEL_CHECK; } - return Dtype(0); } INSTANTIATE_CLASS(SigmoidLayer); diff --git a/src/caffe/layers/softmax_layer.cpp b/src/caffe/layers/softmax_layer.cpp index 69e95ff6385..0d2e4572c76 100644 --- a/src/caffe/layers/softmax_layer.cpp +++ b/src/caffe/layers/softmax_layer.cpp @@ -28,7 +28,7 @@ void SoftmaxLayer::SetUp(const vector*>& bottom, } template -void SoftmaxLayer::Forward_cpu(const vector*>& bottom, +Dtype SoftmaxLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -56,10 +56,11 @@ void SoftmaxLayer::Forward_cpu(const vector*>& bottom, for (int i = 0; i < num; ++i) { caffe_scal(dim, Dtype(1.) / scale_data[i], top_data + i * dim); } + return Dtype(0); } template -Dtype SoftmaxLayer::Backward_cpu(const vector*>& top, +void SoftmaxLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->cpu_diff(); @@ -79,7 +80,6 @@ Dtype SoftmaxLayer::Backward_cpu(const vector*>& top, scale_data, sum_multiplier_.cpu_data(), 1., bottom_diff); // elementwise multiplication caffe_mul(top[0]->count(), bottom_diff, top_data, bottom_diff); - return Dtype(0); } diff --git a/src/caffe/layers/softmax_layer.cu b/src/caffe/layers/softmax_layer.cu index 2e41a1794df..5efa4909263 100644 --- a/src/caffe/layers/softmax_layer.cu +++ b/src/caffe/layers/softmax_layer.cu @@ -43,7 +43,7 @@ __global__ void kernel_exp(const int num, const Dtype* data, Dtype* out) { } template -void SoftmaxLayer::Forward_gpu(const vector*>& bottom, +Dtype SoftmaxLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -73,11 +73,12 @@ void SoftmaxLayer::Forward_gpu(const vector*>& bottom, kernel_softmax_div<<>>( num, dim, scale_data, top_data); + return Dtype(0); } // TODO(Yangqing): implement the GPU version of softmax. template -Dtype SoftmaxLayer::Backward_gpu(const vector*>& top, +void SoftmaxLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->gpu_diff(); const Dtype* top_data = top[0]->gpu_data(); @@ -103,7 +104,6 @@ Dtype SoftmaxLayer::Backward_gpu(const vector*>& top, scale_.gpu_data(), sum_multiplier_.gpu_data(), 1., bottom_diff); // elementwise multiplication caffe_gpu_mul(top[0]->count(), bottom_diff, top_data, bottom_diff); - return Dtype(0); } INSTANTIATE_CLASS(SoftmaxLayer); diff --git a/src/caffe/layers/softmax_loss_layer.cpp b/src/caffe/layers/softmax_loss_layer.cpp index 6fdaea5a1dd..4238cf6eac3 100644 --- a/src/caffe/layers/softmax_loss_layer.cpp +++ b/src/caffe/layers/softmax_loss_layer.cpp @@ -24,33 +24,39 @@ void SoftmaxWithLossLayer::SetUp(const vector*>& bottom, } template -void SoftmaxWithLossLayer::Forward_cpu( +Dtype SoftmaxWithLossLayer::Forward_cpu( const vector*>& bottom, vector*>* top) { // The forward pass computes the softmax prob values. softmax_bottom_vec_[0] = bottom[0]; softmax_layer_->Forward(softmax_bottom_vec_, &softmax_top_vec_); + const Dtype* prob_data = prob_.cpu_data(); + const Dtype* label = bottom[1]->cpu_data(); + int num = prob_.num(); + int dim = prob_.count() / num; + Dtype loss = 0; + for (int i = 0; i < num; ++i) { + loss += -log(max(prob_data[i * dim + static_cast(label[i])], + Dtype(FLT_MIN))); + } + return loss / num; } template -Dtype SoftmaxWithLossLayer::Backward_cpu(const vector*>& top, +void SoftmaxWithLossLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { - // First, compute the diff + // Compute the diff Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff(); const Dtype* prob_data = prob_.cpu_data(); - memcpy(bottom_diff, prob_data, sizeof(Dtype) * prob_.count()); const Dtype* label = (*bottom)[1]->cpu_data(); + memcpy(bottom_diff, prob_data, sizeof(Dtype) * prob_.count()); int num = prob_.num(); int dim = prob_.count() / num; - Dtype loss = 0; for (int i = 0; i < num; ++i) { bottom_diff[i * dim + static_cast(label[i])] -= 1; - loss += -log(max(prob_data[i * dim + static_cast(label[i])], - Dtype(FLT_MIN))); } // Scale down gradient caffe_scal(prob_.count(), Dtype(1) / num, bottom_diff); - return loss / num; } diff --git a/src/caffe/layers/softmax_loss_layer.cu b/src/caffe/layers/softmax_loss_layer.cu index 100393caa3d..5039524180f 100644 --- a/src/caffe/layers/softmax_loss_layer.cu +++ b/src/caffe/layers/softmax_loss_layer.cu @@ -13,18 +13,19 @@ using std::max; namespace caffe { template -void SoftmaxWithLossLayer::Forward_gpu( +Dtype SoftmaxWithLossLayer::Forward_gpu( const vector*>& bottom, vector*>* top) { // The forward pass computes the softmax prob values. softmax_bottom_vec_[0] = bottom[0]; softmax_layer_->Forward(softmax_bottom_vec_, &softmax_top_vec_); + return Dtype(0); } template -Dtype SoftmaxWithLossLayer::Backward_gpu(const vector*>& top, +void SoftmaxWithLossLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { // TODO(Yangqing): implement the GPU version of softmax. - return Backward_cpu(top, propagate_down, bottom); + Backward_cpu(top, propagate_down, bottom); } INSTANTIATE_CLASS(SoftmaxWithLossLayer); diff --git a/src/caffe/layers/split_layer.cpp b/src/caffe/layers/split_layer.cpp index f9fc461a11f..a8a240f74a6 100644 --- a/src/caffe/layers/split_layer.cpp +++ b/src/caffe/layers/split_layer.cpp @@ -28,7 +28,7 @@ void SplitLayer::SetUp(const vector*>& bottom, } template -void SplitLayer::Forward_cpu(const vector*>& bottom, +Dtype SplitLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); for (int i = 0; i < top->size(); ++i) { @@ -38,10 +38,11 @@ void SplitLayer::Forward_cpu(const vector*>& bottom, Dtype* top_data = (*top)[i]->mutable_cpu_data(); caffe_copy(count_, bottom_data, top_data); } + return Dtype(0.); } template -Dtype SplitLayer::Backward_cpu(const vector*>& top, +void SplitLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { const Dtype* top_diff = top[0]->cpu_diff(); @@ -58,7 +59,6 @@ Dtype SplitLayer::Backward_cpu(const vector*>& top, caffe_axpy(count_, Dtype(1.), top_diff, bottom_diff); } } - return Dtype(0.); } diff --git a/src/caffe/layers/split_layer.cu b/src/caffe/layers/split_layer.cu index 5f25a460a6a..deccf990a27 100644 --- a/src/caffe/layers/split_layer.cu +++ b/src/caffe/layers/split_layer.cu @@ -9,7 +9,7 @@ namespace caffe { template -void SplitLayer::Forward_gpu(const vector*>& bottom, +Dtype SplitLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); for (int i = 0; i < top->size(); ++i) { @@ -19,10 +19,11 @@ void SplitLayer::Forward_gpu(const vector*>& bottom, Dtype* top_data = (*top)[i]->mutable_gpu_data(); caffe_gpu_copy(count_, bottom_data, top_data); } + return Dtype(0.); } template -Dtype SplitLayer::Backward_gpu(const vector*>& top, +void SplitLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { const Dtype* top_diff = top[0]->gpu_diff(); @@ -39,7 +40,6 @@ Dtype SplitLayer::Backward_gpu(const vector*>& top, caffe_gpu_axpy(count_, Dtype(1.), top_diff, bottom_diff); } } - return Dtype(0.); } diff --git a/src/caffe/layers/tanh_layer.cpp b/src/caffe/layers/tanh_layer.cpp index d6f99560082..c26579234bc 100644 --- a/src/caffe/layers/tanh_layer.cpp +++ b/src/caffe/layers/tanh_layer.cpp @@ -11,7 +11,7 @@ namespace caffe { template -void TanHLayer::Forward_cpu(const vector*>& bottom, +Dtype TanHLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); @@ -21,10 +21,11 @@ void TanHLayer::Forward_cpu(const vector*>& bottom, exp2x = exp(2*bottom_data[i]); top_data[i] = (exp2x - Dtype(1))/(exp2x + Dtype(1)); } + return Dtype(0); } template -Dtype TanHLayer::Backward_cpu(const vector*>& top, +void TanHLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { @@ -40,7 +41,6 @@ Dtype TanHLayer::Backward_cpu(const vector*>& top, bottom_diff[i] = top_diff[i] * (1 - tanhx*tanhx); } } - return Dtype(0); } INSTANTIATE_CLASS(TanHLayer); diff --git a/src/caffe/layers/tanh_layer.cu b/src/caffe/layers/tanh_layer.cu index c1f8a29cc5c..899b841b069 100644 --- a/src/caffe/layers/tanh_layer.cu +++ b/src/caffe/layers/tanh_layer.cu @@ -19,7 +19,7 @@ __global__ void TanHForward(const int n, const Dtype* in, Dtype* out) { } template -void TanHLayer::Forward_gpu(const vector*>& bottom, +Dtype TanHLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -33,6 +33,7 @@ void TanHLayer::Forward_gpu(const vector*>& bottom, // << " top_data: " << (unsigned long)top_data // << " blocks: " << CAFFE_GET_BLOCKS(count) // << " threads: " << CAFFE_CUDA_NUM_THREADS; + return Dtype(0); } template @@ -46,7 +47,7 @@ __global__ void TanHBackward(const int n, const Dtype* in_diff, } template -Dtype TanHLayer::Backward_gpu(const vector*>& top, +void TanHLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { @@ -59,7 +60,6 @@ Dtype TanHLayer::Backward_gpu(const vector*>& top, count, top_diff, bottom_data, bottom_diff); CUDA_POST_KERNEL_CHECK; } - return Dtype(0); } INSTANTIATE_CLASS(TanHLayer); diff --git a/src/caffe/net.cpp b/src/caffe/net.cpp index 1837b0768ae..397ee02b890 100644 --- a/src/caffe/net.cpp +++ b/src/caffe/net.cpp @@ -208,9 +208,16 @@ void Net::GetLearningRateAndWeightDecay() { template const vector*>& Net::ForwardPrefilled() { + Dtype ignored_loss; + return ForwardPrefilled(&ignored_loss); +} + +template +const vector*>& Net::ForwardPrefilled(Dtype* loss) { + *loss = Dtype(0.); for (int i = 0; i < layers_.size(); ++i) { // LOG(ERROR) << "Forwarding " << layer_names_[i]; - layers_[i]->Forward(bottom_vecs_[i], &top_vecs_[i]); + *loss += layers_[i]->Forward(bottom_vecs_[i], &top_vecs_[i]); } return net_output_blobs_; } @@ -218,16 +225,22 @@ const vector*>& Net::ForwardPrefilled() { template const vector*>& Net::Forward( const vector*> & bottom) { + Dtype ignored_loss; + return Forward(bottom, &ignored_loss); +} + +template +const vector*>& Net::Forward( + const vector*> & bottom, Dtype* loss) { // Copy bottom to internal bottom for (int i = 0; i < bottom.size(); ++i) { net_input_blobs_[i]->CopyFrom(*bottom[i]); } - return ForwardPrefilled(); + return ForwardPrefilled(loss); } - template -string Net::Forward(const string& input_blob_protos) { +string Net::Forward(const string& input_blob_protos, Dtype* loss) { BlobProtoVector blob_proto_vec; if (net_input_blobs_.size()) { blob_proto_vec.ParseFromString(input_blob_protos); @@ -237,7 +250,7 @@ string Net::Forward(const string& input_blob_protos) { net_input_blobs_[i]->FromProto(blob_proto_vec.blobs(i)); } } - ForwardPrefilled(); + ForwardPrefilled(loss); blob_proto_vec.Clear(); for (int i = 0; i < net_output_blobs_.size(); ++i) { net_output_blobs_[i]->ToProto(blob_proto_vec.add_blobs()); @@ -249,16 +262,12 @@ string Net::Forward(const string& input_blob_protos) { template -Dtype Net::Backward() { - Dtype loss = 0; +void Net::Backward() { for (int i = layers_.size() - 1; i >= 0; --i) { if (layer_need_backward_[i]) { - Dtype layer_loss = layers_[i]->Backward( - top_vecs_[i], true, &bottom_vecs_[i]); - loss += layer_loss; + layers_[i]->Backward(top_vecs_[i], true, &bottom_vecs_[i]); } } - return loss; } template From aee5f54661f3f4678db101cef76aa5ca9b2930bc Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Thu, 13 Mar 2014 17:54:10 -0700 Subject: [PATCH 02/11] fix net_speed_benchmark so 'make all' works --- tools/net_speed_benchmark.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/tools/net_speed_benchmark.cpp b/tools/net_speed_benchmark.cpp index 96d40a2eb37..43f7b493671 100644 --- a/tools/net_speed_benchmark.cpp +++ b/tools/net_speed_benchmark.cpp @@ -58,9 +58,11 @@ int main(int argc, char** argv) { LOG(ERROR) << "Performing Forward"; // Note that for the speed benchmark, we will assume that the network does // not take any input blobs. - caffe_net.Forward(vector*>()); + float initial_loss; + caffe_net.Forward(vector*>(), &initial_loss); + LOG(ERROR) << "Initial loss: " << initial_loss; LOG(ERROR) << "Performing Backward"; - LOG(ERROR) << "Initial loss: " << caffe_net.Backward(); + caffe_net.Backward(); const vector > >& layers = caffe_net.layers(); vector*> >& bottom_vecs = caffe_net.bottom_vecs(); From 305e7314a1dfa78304f0591a820d8b8f71b6b5e5 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Thu, 13 Mar 2014 18:01:15 -0700 Subject: [PATCH 03/11] make tests compile and pass --- src/caffe/test/test_gradient_check_util.hpp | 21 +++++++++++---------- 1 file changed, 11 insertions(+), 10 deletions(-) diff --git a/src/caffe/test/test_gradient_check_util.hpp b/src/caffe/test/test_gradient_check_util.hpp index 895e9965a9a..f8ee04b1975 100644 --- a/src/caffe/test/test_gradient_check_util.hpp +++ b/src/caffe/test/test_gradient_check_util.hpp @@ -92,23 +92,24 @@ void GradientChecker::CheckGradientSingle(Layer* layer, for (int feat_id = 0; feat_id < current_blob->count(); ++feat_id) { // First, obtain the original data Caffe::set_random_seed(seed_); - layer->Forward(*bottom, top); - Dtype computed_objective = GetObjAndGradient(top, top_id, top_data_id); - // Get any additional loss from the layer - computed_objective += layer->Backward(*top, true, bottom); + // Get any loss from the layer + Dtype computed_objective = layer->Forward(*bottom, top); + // Get additional loss from the objective + computed_objective += GetObjAndGradient(top, top_id, top_data_id); + layer->Backward(*top, true, bottom); Dtype computed_gradient = current_blob->cpu_diff()[feat_id]; // compute score by adding stepsize current_blob->mutable_cpu_data()[feat_id] += stepsize_; Caffe::set_random_seed(seed_); - layer->Forward(*bottom, top); - Dtype positive_objective = GetObjAndGradient(top, top_id, top_data_id); - positive_objective += layer->Backward(*top, true, bottom); + Dtype positive_objective = layer->Forward(*bottom, top); + positive_objective += GetObjAndGradient(top, top_id, top_data_id); + layer->Backward(*top, true, bottom); // compute score by subtracting stepsize current_blob->mutable_cpu_data()[feat_id] -= stepsize_ * 2; Caffe::set_random_seed(seed_); - layer->Forward(*bottom, top); - Dtype negative_objective = GetObjAndGradient(top, top_id, top_data_id); - negative_objective += layer->Backward(*top, true, bottom); + Dtype negative_objective = layer->Forward(*bottom, top); + negative_objective += GetObjAndGradient(top, top_id, top_data_id); + layer->Backward(*top, true, bottom); // Recover stepsize current_blob->mutable_cpu_data()[feat_id] += stepsize_; Dtype estimated_gradient = (positive_objective - negative_objective) / From 5e982537a337c91f846371dd0ddbae3ee6235d6f Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Thu, 13 Mar 2014 18:08:34 -0700 Subject: [PATCH 04/11] test_gradient_check_util: blobid -> blob_id --- src/caffe/test/test_gradient_check_util.hpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/src/caffe/test/test_gradient_check_util.hpp b/src/caffe/test/test_gradient_check_util.hpp index f8ee04b1975..15a03be124a 100644 --- a/src/caffe/test/test_gradient_check_util.hpp +++ b/src/caffe/test/test_gradient_check_util.hpp @@ -84,10 +84,10 @@ void GradientChecker::CheckGradientSingle(Layer* layer, } // go through the bottom and parameter blobs // LOG(ERROR) << "Checking " << blobs_to_check.size() << " blobs."; - for (int blobid = 0; blobid < blobs_to_check.size(); ++blobid) { - Blob* current_blob = blobs_to_check[blobid]; - // LOG(ERROR) << "Blob " << blobid << ": checking " << current_blob->count() - // << " parameters."; + for (int blob_id = 0; blob_id < blobs_to_check.size(); ++blob_id) { + Blob* current_blob = blobs_to_check[blob_id]; + // LOG(ERROR) << "Blob " << blob_id << ": checking " + // << current_blob->count() << " parameters."; // go through the values for (int feat_id = 0; feat_id < current_blob->count(); ++feat_id) { // First, obtain the original data @@ -124,7 +124,7 @@ void GradientChecker::CheckGradientSingle(Layer* layer, max(fabs(computed_gradient), fabs(estimated_gradient)), 1.); EXPECT_NEAR(computed_gradient, estimated_gradient, threshold_ * scale) << "debug: (top_id, top_data_id, blob_id, feat_id)=" - << top_id << "," << top_data_id << "," << blobid << "," << feat_id; + << top_id << "," << top_data_id << "," << blob_id << "," << feat_id; } // LOG(ERROR) << "Feature: " << current_blob->cpu_data()[feat_id]; // LOG(ERROR) << "computed gradient: " << computed_gradient From d54833e7d55b814d4e4ca33c4a47cb112e71fa8b Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Thu, 13 Mar 2014 18:09:52 -0700 Subject: [PATCH 05/11] gradient checker optimization with forward pass loss: only need to run backward pass to compute analytic gradient (the thing being checked) now --- src/caffe/test/test_gradient_check_util.hpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/src/caffe/test/test_gradient_check_util.hpp b/src/caffe/test/test_gradient_check_util.hpp index 15a03be124a..6e895241f44 100644 --- a/src/caffe/test/test_gradient_check_util.hpp +++ b/src/caffe/test/test_gradient_check_util.hpp @@ -103,13 +103,11 @@ void GradientChecker::CheckGradientSingle(Layer* layer, Caffe::set_random_seed(seed_); Dtype positive_objective = layer->Forward(*bottom, top); positive_objective += GetObjAndGradient(top, top_id, top_data_id); - layer->Backward(*top, true, bottom); // compute score by subtracting stepsize current_blob->mutable_cpu_data()[feat_id] -= stepsize_ * 2; Caffe::set_random_seed(seed_); Dtype negative_objective = layer->Forward(*bottom, top); negative_objective += GetObjAndGradient(top, top_id, top_data_id); - layer->Backward(*top, true, bottom); // Recover stepsize current_blob->mutable_cpu_data()[feat_id] += stepsize_; Dtype estimated_gradient = (positive_objective - negative_objective) / From 74ed9e0148cdd1b02a524a501ef04c37562345fe Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Thu, 13 Mar 2014 18:16:21 -0700 Subject: [PATCH 06/11] revert unnecessary reordering of lines in softmaxwithlosslayer backward --- src/caffe/layers/softmax_loss_layer.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/caffe/layers/softmax_loss_layer.cpp b/src/caffe/layers/softmax_loss_layer.cpp index 4238cf6eac3..f9bd82e217a 100644 --- a/src/caffe/layers/softmax_loss_layer.cpp +++ b/src/caffe/layers/softmax_loss_layer.cpp @@ -48,8 +48,8 @@ void SoftmaxWithLossLayer::Backward_cpu(const vector*>& top, // Compute the diff Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff(); const Dtype* prob_data = prob_.cpu_data(); - const Dtype* label = (*bottom)[1]->cpu_data(); memcpy(bottom_diff, prob_data, sizeof(Dtype) * prob_.count()); + const Dtype* label = (*bottom)[1]->cpu_data(); int num = prob_.num(); int dim = prob_.count() / num; for (int i = 0; i < num; ++i) { From 8a3f0c225b623de97dffdf1f722b134464c03cfa Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Thu, 13 Mar 2014 18:27:46 -0700 Subject: [PATCH 07/11] remove accidentally added empty line --- include/caffe/layer.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/include/caffe/layer.hpp b/include/caffe/layer.hpp index ad36c827734..c3a88d50120 100644 --- a/include/caffe/layer.hpp +++ b/include/caffe/layer.hpp @@ -34,7 +34,6 @@ class Layer { virtual void SetUp(const vector*>& bottom, vector*>* top) = 0; - // Forward and backward wrappers. You should implement the cpu and // gpu specific implementations instead, and should not change these // functions. From ed23b6890612d9f21b0edb93ee0b9d0bca1b48dc Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Thu, 13 Mar 2014 22:44:41 -0700 Subject: [PATCH 08/11] fix softmax loss layer bug; all tests pass --- src/caffe/layers/softmax_loss_layer.cu | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/src/caffe/layers/softmax_loss_layer.cu b/src/caffe/layers/softmax_loss_layer.cu index 5039524180f..ab7ee6ee3bb 100644 --- a/src/caffe/layers/softmax_loss_layer.cu +++ b/src/caffe/layers/softmax_loss_layer.cu @@ -16,9 +16,7 @@ template Dtype SoftmaxWithLossLayer::Forward_gpu( const vector*>& bottom, vector*>* top) { // The forward pass computes the softmax prob values. - softmax_bottom_vec_[0] = bottom[0]; - softmax_layer_->Forward(softmax_bottom_vec_, &softmax_top_vec_); - return Dtype(0); + return Forward_cpu(bottom, top); } template From 44fbb82f477039bc3038b53fccec29a48a9c4a0c Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Fri, 14 Mar 2014 14:52:51 -0700 Subject: [PATCH 09/11] loss in forward pass for concat layer (thought i'd rebased to latest dev but apparently not) --- include/caffe/vision_layers.hpp | 8 ++++---- src/caffe/layers/concat_layer.cpp | 6 +++--- src/caffe/layers/concat_layer.cu | 6 +++--- 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index c1729ddc899..3a3bdfd68d8 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -352,13 +352,13 @@ class ConcatLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); Blob col_bob_; diff --git a/src/caffe/layers/concat_layer.cpp b/src/caffe/layers/concat_layer.cpp index dc949c14010..e65451061b0 100644 --- a/src/caffe/layers/concat_layer.cpp +++ b/src/caffe/layers/concat_layer.cpp @@ -42,7 +42,7 @@ void ConcatLayer::SetUp(const vector*>& bottom, } template -void ConcatLayer::Forward_cpu(const vector*>& bottom, +Dtype ConcatLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { Dtype* top_data = (*top)[0]->mutable_cpu_data(); if (concat_dim_== 0) { @@ -69,10 +69,11 @@ void ConcatLayer::Forward_cpu(const vector*>& bottom, LOG(FATAL) << "concat_dim along dim" << concat_dim_ << " not implemented yet"; } + return Dtype(0.); } template -Dtype ConcatLayer::Backward_cpu(const vector*>& top, +void ConcatLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->cpu_diff(); if (concat_dim_ == 0) { @@ -100,7 +101,6 @@ Dtype ConcatLayer::Backward_cpu(const vector*>& top, LOG(FATAL) << "concat_dim along dim" << concat_dim_ << " not implemented yet"; } - return Dtype(0.); } INSTANTIATE_CLASS(ConcatLayer); diff --git a/src/caffe/layers/concat_layer.cu b/src/caffe/layers/concat_layer.cu index 616a5e61683..8a20cea64cf 100644 --- a/src/caffe/layers/concat_layer.cu +++ b/src/caffe/layers/concat_layer.cu @@ -9,7 +9,7 @@ namespace caffe { template -void ConcatLayer::Forward_gpu(const vector*>& bottom, +Dtype ConcatLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { Dtype* top_data = (*top)[0]->mutable_gpu_data(); if (concat_dim_ == 0) { @@ -36,10 +36,11 @@ void ConcatLayer::Forward_gpu(const vector*>& bottom, LOG(FATAL) << "concat_dim along dim" << concat_dim_ << " not implemented yet"; } + return Dtype(0.); } template -Dtype ConcatLayer::Backward_gpu(const vector*>& top, +void ConcatLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->gpu_diff(); if (concat_dim_ == 0) { @@ -67,7 +68,6 @@ Dtype ConcatLayer::Backward_gpu(const vector*>& top, LOG(FATAL) << "concat_dim along dim" << concat_dim_ << " not implemented yet"; } - return Dtype(0.); } INSTANTIATE_CLASS(ConcatLayer); From 0551d93831ef3a293efae0ab474f459d09779aa8 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 15 Mar 2014 12:22:53 -0700 Subject: [PATCH 10/11] null pointer defaults for forward loss outputs --- include/caffe/net.hpp | 8 +++----- src/caffe/net.cpp | 22 +++++++--------------- 2 files changed, 10 insertions(+), 20 deletions(-) diff --git a/include/caffe/net.hpp b/include/caffe/net.hpp index d6c892fc3e6..a30491fbe6b 100644 --- a/include/caffe/net.hpp +++ b/include/caffe/net.hpp @@ -31,15 +31,13 @@ class Net { // Run forward with the input blobs already fed separately. You can get the // input blobs using input_blobs(). - const vector*>& ForwardPrefilled(Dtype* loss); - const vector*>& ForwardPrefilled(); + const vector*>& ForwardPrefilled(Dtype* loss = NULL); // Run forward using a set of bottom blobs, and return the result. const vector*>& Forward(const vector* > & bottom, - Dtype* loss); - const vector*>& Forward(const vector* > & bottom); + Dtype* loss = NULL); // Run forward using a serialized BlobProtoVector and return the result // as a serialized BlobProtoVector - string Forward(const string& input_blob_protos, Dtype* loss); + string Forward(const string& input_blob_protos, Dtype* loss = NULL); // The network backward should take no input and output, since it solely // computes the gradient w.r.t the parameters, and the data has already diff --git a/src/caffe/net.cpp b/src/caffe/net.cpp index 397ee02b890..f3429b222e2 100644 --- a/src/caffe/net.cpp +++ b/src/caffe/net.cpp @@ -206,29 +206,21 @@ void Net::GetLearningRateAndWeightDecay() { } } -template -const vector*>& Net::ForwardPrefilled() { - Dtype ignored_loss; - return ForwardPrefilled(&ignored_loss); -} - template const vector*>& Net::ForwardPrefilled(Dtype* loss) { - *loss = Dtype(0.); + if (loss != NULL) { + *loss = Dtype(0.); + } for (int i = 0; i < layers_.size(); ++i) { // LOG(ERROR) << "Forwarding " << layer_names_[i]; - *loss += layers_[i]->Forward(bottom_vecs_[i], &top_vecs_[i]); + Dtype layer_loss = layers_[i]->Forward(bottom_vecs_[i], &top_vecs_[i]); + if (loss != NULL) { + *loss += layer_loss; + } } return net_output_blobs_; } -template -const vector*>& Net::Forward( - const vector*> & bottom) { - Dtype ignored_loss; - return Forward(bottom, &ignored_loss); -} - template const vector*>& Net::Forward( const vector*> & bottom, Dtype* loss) { From a6ae5be95e216053574549857b77b4cb55748b78 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Tue, 18 Mar 2014 18:55:56 -0700 Subject: [PATCH 11/11] post rebase fixes: images layer and padding layer compute loss in forward --- include/caffe/vision_layers.hpp | 20 ++++++++++---------- src/caffe/layers/images_layer.cpp | 17 +++-------------- src/caffe/layers/padding_layer.cpp | 6 +++--- src/caffe/layers/padding_layer.cu | 6 +++--- 4 files changed, 19 insertions(+), 30 deletions(-) diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index 3a3bdfd68d8..9c0850e5f34 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -206,13 +206,13 @@ class PaddingLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, + virtual void Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, + virtual void Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom); unsigned int PAD_; int NUM_; @@ -425,14 +425,14 @@ class ImagesLayer : public Layer { vector*>* top); protected: - virtual void Forward_cpu(const vector*>& bottom, + virtual Dtype Forward_cpu(const vector*>& bottom, vector*>* top); - virtual void Forward_gpu(const vector*>& bottom, + virtual Dtype Forward_gpu(const vector*>& bottom, vector*>* top); - virtual Dtype Backward_cpu(const vector*>& top, - const bool propagate_down, vector*>* bottom); - virtual Dtype Backward_gpu(const vector*>& top, - const bool propagate_down, vector*>* bottom); + virtual void Backward_cpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { return; } + virtual void Backward_gpu(const vector*>& top, + const bool propagate_down, vector*>* bottom) { return; } vector > lines_; int lines_id_; diff --git a/src/caffe/layers/images_layer.cpp b/src/caffe/layers/images_layer.cpp index e750e01b266..6208a9e7fa6 100644 --- a/src/caffe/layers/images_layer.cpp +++ b/src/caffe/layers/images_layer.cpp @@ -233,7 +233,7 @@ void ImagesLayer::SetUp(const vector*>& bottom, } template -void ImagesLayer::Forward_cpu(const vector*>& bottom, +Dtype ImagesLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { // First, join the thread CHECK(!pthread_join(thread_, NULL)) << "Pthread joining failed."; @@ -245,10 +245,11 @@ void ImagesLayer::Forward_cpu(const vector*>& bottom, // Start a new prefetch thread CHECK(!pthread_create(&thread_, NULL, ImagesLayerPrefetch, reinterpret_cast(this))) << "Pthread execution failed."; + return Dtype(0.); } template -void ImagesLayer::Forward_gpu(const vector*>& bottom, +Dtype ImagesLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { // First, join the thread CHECK(!pthread_join(thread_, NULL)) << "Pthread joining failed."; @@ -262,18 +263,6 @@ void ImagesLayer::Forward_gpu(const vector*>& bottom, // Start a new prefetch thread CHECK(!pthread_create(&thread_, NULL, ImagesLayerPrefetch, reinterpret_cast(this))) << "Pthread execution failed."; -} - -// The backward operations are dummy - they do not carry any computation. -template -Dtype ImagesLayer::Backward_cpu(const vector*>& top, - const bool propagate_down, vector*>* bottom) { - return Dtype(0.); -} - -template -Dtype ImagesLayer::Backward_gpu(const vector*>& top, - const bool propagate_down, vector*>* bottom) { return Dtype(0.); } diff --git a/src/caffe/layers/padding_layer.cpp b/src/caffe/layers/padding_layer.cpp index 4cb67df0dcf..658cc6ab16c 100644 --- a/src/caffe/layers/padding_layer.cpp +++ b/src/caffe/layers/padding_layer.cpp @@ -29,7 +29,7 @@ void PaddingLayer::SetUp(const vector*>& bottom, } template -void PaddingLayer::Forward_cpu(const vector*>& bottom, +Dtype PaddingLayer::Forward_cpu(const vector*>& bottom, vector*>* top) { Dtype* top_data = (*top)[0]->mutable_cpu_data(); const Dtype* bottom_data = bottom[0]->cpu_data(); @@ -47,10 +47,11 @@ void PaddingLayer::Forward_cpu(const vector*>& bottom, } } } + return Dtype(0.); } template -Dtype PaddingLayer::Backward_cpu(const vector*>& top, +void PaddingLayer::Backward_cpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { const Dtype* top_diff = top[0]->cpu_diff(); Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff(); @@ -66,7 +67,6 @@ Dtype PaddingLayer::Backward_cpu(const vector*>& top, } } } - return Dtype(0.); } INSTANTIATE_CLASS(PaddingLayer); diff --git a/src/caffe/layers/padding_layer.cu b/src/caffe/layers/padding_layer.cu index 7ec28a9e30f..d476df501fd 100644 --- a/src/caffe/layers/padding_layer.cu +++ b/src/caffe/layers/padding_layer.cu @@ -27,7 +27,7 @@ __global__ void PaddingForward(const int count, const Dtype* in, Dtype* out, } template -void PaddingLayer::Forward_gpu(const vector*>& bottom, +Dtype PaddingLayer::Forward_gpu(const vector*>& bottom, vector*>* top) { const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = (*top)[0]->mutable_gpu_data(); @@ -39,6 +39,7 @@ void PaddingLayer::Forward_gpu(const vector*>& bottom, count, bottom_data, top_data, NUM_, CHANNEL_, HEIGHT_IN_, WIDTH_IN_, PAD_); CUDA_POST_KERNEL_CHECK; + return Dtype(0); } template @@ -61,7 +62,7 @@ __global__ void PaddingBackward(const int count, const Dtype* in, Dtype* out, } template -Dtype PaddingLayer::Backward_gpu(const vector*>& top, +void PaddingLayer::Backward_gpu(const vector*>& top, const bool propagate_down, vector*>* bottom) { if (propagate_down) { @@ -74,7 +75,6 @@ Dtype PaddingLayer::Backward_gpu(const vector*>& top, PAD_); CUDA_POST_KERNEL_CHECK; } - return Dtype(0); } INSTANTIATE_CLASS(PaddingLayer);