From cf3177e573cddb6e9999fed888fecf416c8db2e4 Mon Sep 17 00:00:00 2001 From: Dmytro Mishkin Date: Thu, 14 Jan 2016 17:10:11 +0200 Subject: [PATCH 1/2] Rebased and cleaned-up Channelwise Affine for batch norm --- .../caffe/layers/channelwise_affine_layer.hpp | 103 ++++++++++ src/caffe/layers/channelwise_affine_layer.cpp | 189 ++++++++++++++++++ src/caffe/layers/channelwise_affine_layer.cu | 144 +++++++++++++ src/caffe/proto/caffe.proto | 14 +- .../test/test_channelwise_affine_layer.cpp | 105 ++++++++++ src/caffe/test/test_neuron_layer.cpp | 1 + 6 files changed, 555 insertions(+), 1 deletion(-) create mode 100644 include/caffe/layers/channelwise_affine_layer.hpp create mode 100644 src/caffe/layers/channelwise_affine_layer.cpp create mode 100644 src/caffe/layers/channelwise_affine_layer.cu create mode 100644 src/caffe/test/test_channelwise_affine_layer.cpp diff --git a/include/caffe/layers/channelwise_affine_layer.hpp b/include/caffe/layers/channelwise_affine_layer.hpp new file mode 100644 index 00000000000..6d8ac98b6ed --- /dev/null +++ b/include/caffe/layers/channelwise_affine_layer.hpp @@ -0,0 +1,103 @@ +#ifndef CAFFE_CHANNELWISE_AFFINE_LAYER_HPP_ +#define CAFFE_CHANNELWISE_AFFINE_LAYER_HPP_ + +#include +#include "caffe/blob.hpp" +#include "caffe/layer.hpp" +#include "caffe/layers/neuron_layer.hpp" +#include "caffe/proto/caffe.pb.h" + +namespace caffe { + /** + * @brief Affine non-linearity function @f$ + * y = ax+b + * @f$, could be used after batch normalization layer + * + */ +template +class ChannelwiseAffineLayer : public NeuronLayer { + public: + /** + * @param param provides ChannelwiseAffineParameter ChannelwiseAffine_param, + * with ChannelwiseAffineLayer options: + * - slope_filler (\b optional, FillerParameter, + * default {'type': constant 'value':1.0001}). + * - bias_filler (\b optional, FillerParameter, + * default {'type': constant 'value':0.0001}). + * - channel_shared (\b optional, default false). + * slopes and biases are shared across channels. + */ + explicit ChannelwiseAffineLayer(const LayerParameter& param) + : NeuronLayer(param) {} + virtual void LayerSetUp(const vector*>& bottom, + const vector*>& top); + virtual void Reshape(const vector*>& bottom, + const vector*>& top); + virtual inline const char* type() const { return "ChannelwiseAffine"; } + + protected: + /** + * @param bottom input Blob vector (length 1) + * -# @f$ (N \times C \times ...) @f$ + * the inputs @f$ x @f$ + * @param top output Blob vector (length 1) + * -# @f$ (N \times C \times ...) @f$ + * the computed outputs for each channel @f$i@f$ @f$ + * y_i = a_i x_i + b_i + * @f$. + */ + virtual void Forward_cpu(const vector*>& bottom, + const vector*>& top); + virtual void Forward_gpu(const vector*>& bottom, + const vector*>& top); + /** + * @brief Computes the error gradient w.r.t. the ChannelwiseAffine inputs. + * + * @param top output Blob vector (length 1), providing the error gradient with + * respect to the outputs + * -# @f$ (N \times C \times ...) @f$ + * containing error gradients @f$ \frac{\partial E}{\partial y} @f$ + * with respect to computed outputs @f$ y @f$ + * @param propagate_down see Layer::Backward. + * @param bottom input Blob vector (length 1) + * -# @f$ (N \times C \times ...) @f$ + * the inputs @f$ x @f$; For each channel @f$i@f$, backward fills their + * diff with gradients @f$ + * \frac{\partial E}{\partial x_i} = \left\{ + * \begin{array}{lr} + * a_i \frac{\partial E}{\partial y_i} + * \end{array} \right. + * @f$. + * If param_propagate_down_[0] is true, it fills the diff with gradients + * @f$ + * \frac{\partial E}{\partial a_i} = \left\{ + * \begin{array}{lr} + * \sum_{x_i} x_i \frac{\partial E}{\partial y_i} + * \end{array} \right. + * @f$. + * If param_propagate_down_[1] is true, it fills the diff with gradients + * @f$ + * \frac{\partial E}{\partial b_i} = \left\{ + * \begin{array}{lr} + * frac{\partial E}{\partial y_i} + * \end{array} \right. + * @f$. + */ + virtual void Backward_cpu(const vector*>& top, + const vector& propagate_down, + const vector*>& bottom); + virtual void Backward_gpu(const vector*>& top, + const vector& propagate_down, + const vector*>& bottom); + bool channel_shared_; + Blob multiplier_; + // dot multiplier for backward computation of params + Blob bias_multiplier_; + Blob backward_buff_; + // temporary buffer for backward computation + Blob bottom_memory_; + // memory for in-place computation +}; +} // namespace caffe + +#endif // CAFFE_CHANNELWISE_AFFINE_LAYER_HPP_ diff --git a/src/caffe/layers/channelwise_affine_layer.cpp b/src/caffe/layers/channelwise_affine_layer.cpp new file mode 100644 index 00000000000..e9f31fb10e3 --- /dev/null +++ b/src/caffe/layers/channelwise_affine_layer.cpp @@ -0,0 +1,189 @@ +#include +#include + +#include "caffe/filler.hpp" +#include "caffe/layer.hpp" +#include "caffe/layers/channelwise_affine_layer.hpp" + +namespace caffe { + +template +void ChannelwiseAffineLayer::LayerSetUp( + const vector*>& bottom, + const vector*>& top) { + CHECK_GE(bottom[0]->num_axes(), 2) + << "Number of axes of bottom blob must be >=2."; + ChannelwiseAffineParameter channelwise_affine_param = + this->layer_param().channelwise_affine_param(); + int channels = bottom[0]->channels(); + channel_shared_ = channelwise_affine_param.channel_shared(); + if (this->blobs_.size() > 0) { + LOG(INFO) << "Skipping parameter initialization"; + } else { + this->blobs_.resize(2); + if (channel_shared_) { + this->blobs_[0].reset(new Blob(vector(0))); + this->blobs_[1].reset(new Blob(vector(0))); + + } else { + this->blobs_[0].reset(new Blob(vector(1, channels))); + this->blobs_[1].reset(new Blob(vector(1, channels))); + } + shared_ptr > filler; + if (channelwise_affine_param.has_slope_filler()) { + filler.reset(GetFiller(channelwise_affine_param.slope_filler())); + } else { + FillerParameter filler_param; + filler_param.set_type("constant"); + filler_param.set_value(1.0001); + filler.reset(GetFiller(filler_param)); + } + filler->Fill(this->blobs_[0].get()); + + if (channelwise_affine_param.has_bias_filler()) { + filler.reset(GetFiller(channelwise_affine_param.bias_filler())); + } else { + FillerParameter filler_param; + filler_param.set_type("constant"); + filler_param.set_value(0.0001); + filler.reset(GetFiller(filler_param)); + } + filler->Fill(this->blobs_[1].get()); + } + if (channel_shared_) { + CHECK_EQ(this->blobs_[0]->count(), 1) + << "Slope size is inconsistent with prototxt config"; + } else { + CHECK_EQ(this->blobs_[0]->count(), channels) + << "Slope size is inconsistent with prototxt config"; + } + + // Propagate gradients to the parameters (as directed by backward pass). + this->param_propagate_down_.resize(this->blobs_.size(), true); + multiplier_.Reshape(vector(1, bottom[0]->count(1))); + bias_multiplier_.Reshape(vector(1, bottom[0]->count(1))); + backward_buff_.Reshape(vector(1, bottom[0]->count(1))); + caffe_set(multiplier_.count(), Dtype(1.0), + multiplier_.mutable_cpu_data()); + caffe_set(bias_multiplier_.count(), Dtype(1.0), + bias_multiplier_.mutable_cpu_data()); +} + +template +void ChannelwiseAffineLayer::Reshape( + const vector*>& bottom, + const vector*>& top) { + CHECK_GE(bottom[0]->num_axes(), 2) + << "Number of axes of bottom blob must be >=2."; + top[0]->ReshapeLike(*bottom[0]); + if (bottom[0] == top[0]) { + // For in-place computation + bottom_memory_.ReshapeLike(*bottom[0]); + } + int height = 1; + int width = 1; + if (bottom[0]->num_axes() > 2) { + height = bottom[0]->shape(2); + width = bottom[0]->shape(3); + } + vector bias_multiplier_shape(1, height * width); + bias_multiplier_.Reshape(bias_multiplier_shape); + caffe_set(bias_multiplier_.count(), Dtype(1), + bias_multiplier_.mutable_cpu_data()); +} + +template +void ChannelwiseAffineLayer::Forward_cpu( + const vector*>& bottom, + const vector*>& top) { + const Dtype* bottom_data = bottom[0]->cpu_data(); + Dtype* top_data = top[0]->mutable_cpu_data(); + const int count = bottom[0]->count(); + const int dim = bottom[0]->count(2); + const int channels = bottom[0]->channels(); + const Dtype* slope_data = this->blobs_[0]->cpu_data(); + const Dtype* bias_data = this->blobs_[1]->cpu_data(); + // For in-place computation + if (bottom[0] == top[0]) { + caffe_copy(count, bottom_data, bottom_memory_.mutable_cpu_data()); + } + // if channel_shared, channel index in the following computation becomes + // always zero. + const int div_factor = channel_shared_ ? channels : 1; + for (int i = 0; i < count; ++i) { + int c = (i / dim) % channels / div_factor; + top_data[i] = bottom_data[i] * slope_data[c] + bias_data[c]; + } +} + +template +void ChannelwiseAffineLayer::Backward_cpu( + const vector*>& top, + const vector& propagate_down, + const vector*>& bottom) { + const Dtype* bottom_data = bottom[0]->cpu_data(); + const Dtype* slope_data = this->blobs_[0]->cpu_data(); + + const Dtype* top_diff = top[0]->cpu_diff(); + const int count = bottom[0]->count(); + const int dim = bottom[0]->count(2); + const int channels = bottom[0]->shape(1); + const int num = bottom[0]->shape(0); + int height = 1; + int width = 1; + if (bottom[0]->num_axes() > 2) { + height = bottom[0]->shape(2); + width = bottom[0]->shape(3); + } + + // For in-place computation + if (top[0] == bottom[0]) { + bottom_data = bottom_memory_.cpu_data(); + } + + // if channel_shared, channel index in the following computation becomes + // always zero. + const int div_factor = channel_shared_ ? channels : 1; + + // Propagte to param + // Since to write bottom diff will affect top diff if top and bottom blobs + // are identical (in-place computaion), we first compute param backward to + // keep top_diff unchanged. + + if (this->param_propagate_down_[1]) { + Dtype* bias_diff = this->blobs_[1]->mutable_cpu_diff(); + caffe_set(this->blobs_[1]->count(), Dtype(0), bias_diff); + for (int n = 0; n < num; ++n) { + caffe_cpu_gemv(CblasNoTrans, channels, height * width, 1., + top_diff + top[0]->offset(n), + bias_multiplier_.cpu_data(), 1., bias_diff); + } + } + if (this->param_propagate_down_[0]) { + Dtype* slope_diff = this->blobs_[0]->mutable_cpu_diff(); + caffe_set(this->blobs_[0]->count(), Dtype(0), slope_diff); + for (int i = 0; i < count; ++i) { + int c = (i / dim) % channels / div_factor; + slope_diff[c] += top_diff[i] * bottom_data[i]; + } + } + + // Propagate to bottom + if (propagate_down[0]) { + Dtype* bottom_diff = bottom[0]->mutable_cpu_diff(); + for (int i = 0; i < count; ++i) { + int c = (i / dim) % channels / div_factor; + bottom_diff[i] = slope_data[c] * top_diff[i]; + } + } +} + + +#ifdef CPU_ONLY +STUB_GPU(ChannelwiseAffineLayer); +#endif + +INSTANTIATE_CLASS(ChannelwiseAffineLayer); +REGISTER_LAYER_CLASS(ChannelwiseAffine); + +} // namespace caffe diff --git a/src/caffe/layers/channelwise_affine_layer.cu b/src/caffe/layers/channelwise_affine_layer.cu new file mode 100644 index 00000000000..2066b26560b --- /dev/null +++ b/src/caffe/layers/channelwise_affine_layer.cu @@ -0,0 +1,144 @@ +#include +#include + +#include "caffe/layer.hpp" +#include "caffe/layers/channelwise_affine_layer.hpp" + +namespace caffe { + +// CUDA kernel for forward +template +__global__ void ChannelwiseAffineForward(const int n, const int channels, + const int dim, const Dtype* in, Dtype* out, const Dtype* slope_data, + const Dtype* bias_data, const int div_factor) { + CUDA_KERNEL_LOOP(index, n) { + int c = (index / dim) % channels / div_factor; + out[index] = in[index] * slope_data[c] + bias_data[c]; + } +} + +// CUDA kernel for bottom backward +template +__global__ void ChannelwiseAffineBackward(const int n, + const int channels, const int dim, const Dtype* in_diff, + Dtype* out_diff, const Dtype* slope_data, const int div_factor) { + CUDA_KERNEL_LOOP(index, n) { + int c = (index / dim) % channels / div_factor; + out_diff[index] = slope_data[c] * in_diff[index]; + } +} + +// CUDA kernel for element-wise parameter backward +template +__global__ void ChannelwiseAffineParamSlopeBackward(const int n, + const int rows, const int rowPitch, const Dtype* in_diff, + const Dtype* in_data, Dtype* out_diff) { + CUDA_KERNEL_LOOP(index, n) { + out_diff[index] = in_diff[index] * in_data[index]; + for ( int k = 1; k < rows; k++ ) { + out_diff[index] += in_diff[index + k*rowPitch] + * in_data[index + k*rowPitch]; + } + } +} + +template +void ChannelwiseAffineLayer::Forward_gpu( + const vector*>& bottom, + const vector*>& top) { + const Dtype* bottom_data = bottom[0]->gpu_data(); + Dtype* top_data = top[0]->mutable_gpu_data(); + const int count = bottom[0]->count(); + const int dim = bottom[0]->count(2); + const int channels = bottom[0]->channels(); + const Dtype* slope_data = this->blobs_[0]->gpu_data(); + const Dtype* bias_data = this->blobs_[1]->gpu_data(); + const int div_factor = channel_shared_ ? channels : 1; + + // For in-place computation + if (top[0] == bottom[0]) { + caffe_copy(count, bottom_data, bottom_memory_.mutable_gpu_data()); + } + // NOLINT_NEXT_LINE(whitespace/operators) + ChannelwiseAffineForward<<>>( + count, channels, dim, bottom_data, top_data, + slope_data, bias_data, div_factor); + CUDA_POST_KERNEL_CHECK; +} + +template +void ChannelwiseAffineLayer::Backward_gpu( + const vector*>& top, + const vector& propagate_down, + const vector*>& bottom) { + const Dtype* bottom_data = bottom[0]->gpu_data(); + const Dtype* top_diff = top[0]->gpu_diff(); + const int count = bottom[0]->count(); + const int num = bottom[0]->shape(0); + const int dim = bottom[0]->count(2); + const int channels = bottom[0]->shape(1); + int height = 1; + int width = 1; + if (bottom[0]->num_axes() > 2) { + height = bottom[0]->shape(2); + width = bottom[0]->shape(3); + } + + // For in-place computation + if (top[0] == bottom[0]) { + bottom_data = bottom_memory_.gpu_data(); + } + // Propagate to param + // Since to write bottom diff will affect top diff if top and bottom blobs + // are identical (in-place computaion), we first compute param backward to + // keep top_diff unchanged. + if (this->param_propagate_down_[1]) { + Dtype* bias_diff = this->blobs_[1]->mutable_gpu_diff(); + caffe_gpu_set(this->blobs_[1]->count(), Dtype(0.0), bias_diff); + // Gradient with respect to bias + for (int n = 0; n < num; ++n) { + caffe_gpu_gemv( + CblasNoTrans, channels, height * width, (Dtype)1., + top_diff + top[0]->offset(n), bias_multiplier_.gpu_data(), + (Dtype)1., bias_diff); + } + } + if (this->param_propagate_down_[0]) { + Dtype* slope_diff = this->blobs_[0]->mutable_gpu_diff(); + int cdim = channels * dim; + // compute element-wise diff + // NOLINT_NEXT_LINE(whitespace/operators) + ChannelwiseAffineParamSlopeBackward<<>>( + cdim, num, top[0]->offset(1), top_diff , + bottom_data, + backward_buff_.mutable_gpu_diff()); + CUDA_POST_KERNEL_CHECK; + if (channel_shared_) { + Dtype d = 0; + caffe_gpu_dot(cdim, backward_buff_.gpu_diff(), + multiplier_.gpu_data(), &d); + caffe_gpu_add_scalar(this->blobs_[0]->count(), Dtype(d), slope_diff); + } else { + caffe_gpu_gemv(CblasNoTrans, channels, dim, Dtype(1.), + backward_buff_.gpu_diff(), multiplier_.gpu_data(), Dtype(1.), + slope_diff); + } + } + // Propagate to bottom + if (propagate_down[0]) { + Dtype* bottom_diff = bottom[0]->mutable_gpu_diff(); + const Dtype* slope_data = this->blobs_[0]->gpu_data(); + int div_factor = channel_shared_ ? channels : 1; + // NOLINT_NEXT_LINE(whitespace/operators) + ChannelwiseAffineBackward<<>>( + count, channels, dim, top_diff, bottom_diff, slope_data, div_factor); + CUDA_POST_KERNEL_CHECK; + } +} + +INSTANTIATE_LAYER_GPU_FUNCS(ChannelwiseAffineLayer); + +} // namespace caffe diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 019aa614373..77270637137 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -306,7 +306,7 @@ message ParamSpec { // NOTE // Update the next available ID when you add a new LayerParameter field. // -// LayerParameter next available layer-specific ID: 140 (last added: batch_norm_param) +// LayerParameter next available layer-specific ID: 141 (last added: channelwise_affine_param) message LayerParameter { optional string name = 1; // the layer name optional string type = 2; // the layer type @@ -356,6 +356,7 @@ message LayerParameter { optional AccuracyParameter accuracy_param = 102; optional ArgMaxParameter argmax_param = 103; optional BatchNormParameter batch_norm_param = 139; + optional ChannelwiseAffineParameter channelwise_affine_param = 140; optional ConcatParameter concat_param = 104; optional ContrastiveLossParameter contrastive_loss_param = 105; optional ConvolutionParameter convolution_param = 106; @@ -497,6 +498,17 @@ message BatchNormParameter { optional float eps = 3 [default = 1e-5]; } +message ChannelwiseAffineParameter { + + // Initial value of a_i. Default is a_i=1.0 for all i. + optional FillerParameter slope_filler = 1; + + optional FillerParameter bias_filler = 2; + + // Whether or not slope paramters are shared across channels. + optional bool channel_shared = 3 [default = false]; +} + message ContrastiveLossParameter { // margin for dissimilar pair optional float margin = 1 [default = 1.0]; diff --git a/src/caffe/test/test_channelwise_affine_layer.cpp b/src/caffe/test/test_channelwise_affine_layer.cpp new file mode 100644 index 00000000000..a3e2544f77a --- /dev/null +++ b/src/caffe/test/test_channelwise_affine_layer.cpp @@ -0,0 +1,105 @@ +#include + +#include "gtest/gtest.h" + +#include "caffe/blob.hpp" +#include "caffe/common.hpp" +#include "caffe/filler.hpp" +#include "caffe/layers/channelwise_affine_layer.hpp" + +#include "caffe/test/test_caffe_main.hpp" +#include "caffe/test/test_gradient_check_util.hpp" + +namespace caffe { + +template +class ChannelwiseAffineLayerTest : public MultiDeviceTest { + typedef typename TypeParam::Dtype Dtype; + + protected: + ChannelwiseAffineLayerTest() + : blob_bottom_(new Blob(2, 3, 4, 5)), + blob_top_(new Blob()) { + Caffe::set_random_seed(1701); + // fill the values + FillerParameter filler_param; + GaussianFiller filler(filler_param); + filler.Fill(this->blob_bottom_); + blob_bottom_vec_.push_back(blob_bottom_); + blob_top_vec_.push_back(blob_top_); + } + virtual ~ChannelwiseAffineLayerTest() { + delete blob_bottom_; delete blob_top_; } + Blob* const blob_bottom_; + Blob* const blob_top_; + vector*> blob_bottom_vec_; + vector*> blob_top_vec_; + + void TestChannelwiseAffine(ChannelwiseAffineLayer *layer) { + layer->Forward(this->blob_bottom_vec_, this->blob_top_vec_); + // Now, check values + const Dtype* bottom_data = this->blob_bottom_->cpu_data(); + const Dtype* top_data = this->blob_top_->cpu_data(); + const Dtype* slope_data = layer->blobs()[0]->cpu_data(); + const Dtype* bias_data = layer->blobs()[1]->cpu_data(); + const Dtype kDelta = 2e-5; + int hw = this->blob_bottom_->height() * this->blob_bottom_->width(); + int channels = this->blob_bottom_->channels(); + bool channel_shared = + layer->layer_param().channelwise_affine_param().channel_shared(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + int c = channel_shared ? 0 : (i / hw) % channels; + EXPECT_NEAR(top_data[i], + bottom_data[i]* slope_data[c] + bias_data[c], kDelta); + } + } +}; +TYPED_TEST_CASE(ChannelwiseAffineLayerTest, TestDtypesAndDevices); + + +TYPED_TEST(ChannelwiseAffineLayerTest, TestChannelwiseAffineForward) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + ChannelwiseAffineLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + FillerParameter filler_param; + GaussianFiller filler(filler_param); + filler.Fill(layer.blobs()[0].get()); + filler.Fill(layer.blobs()[1].get()); + this->TestChannelwiseAffine(&layer); +} + +TYPED_TEST(ChannelwiseAffineLayerTest, + TestChannelwiseAffineForwardChannelShared) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + layer_param.mutable_channelwise_affine_param()->set_channel_shared(true); + ChannelwiseAffineLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + this->TestChannelwiseAffine(&layer); +} + +TYPED_TEST(ChannelwiseAffineLayerTest, TestChannelwiseAffineGradient) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + layer_param.mutable_channelwise_affine_param()->set_channel_shared(false); + ChannelwiseAffineLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + GradientChecker checker(1e-2, 1e-3, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + +TYPED_TEST(ChannelwiseAffineLayerTest, + TestChannelwiseAffineGradientChannelShared) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + layer_param.mutable_channelwise_affine_param()->set_channel_shared(true); + ChannelwiseAffineLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + GradientChecker checker(1e-2, 1e-3, 1701, 0., 0.01); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + +} // namespace caffe diff --git a/src/caffe/test/test_neuron_layer.cpp b/src/caffe/test/test_neuron_layer.cpp index 21441b4121e..c61c2b6d2d7 100644 --- a/src/caffe/test/test_neuron_layer.cpp +++ b/src/caffe/test/test_neuron_layer.cpp @@ -179,6 +179,7 @@ class NeuronLayerTest : public MultiDeviceTest { } }; + TYPED_TEST_CASE(NeuronLayerTest, TestDtypesAndDevices); TYPED_TEST(NeuronLayerTest, TestAbsVal) { From b146ee86c5279c41d0919b190347ea46fef3d71d Mon Sep 17 00:00:00 2001 From: Dmytro Mishkin Date: Thu, 14 Jan 2016 17:19:30 +0200 Subject: [PATCH 2/2] Erase empty line --- src/caffe/test/test_neuron_layer.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/src/caffe/test/test_neuron_layer.cpp b/src/caffe/test/test_neuron_layer.cpp index c61c2b6d2d7..21441b4121e 100644 --- a/src/caffe/test/test_neuron_layer.cpp +++ b/src/caffe/test/test_neuron_layer.cpp @@ -179,7 +179,6 @@ class NeuronLayerTest : public MultiDeviceTest { } }; - TYPED_TEST_CASE(NeuronLayerTest, TestDtypesAndDevices); TYPED_TEST(NeuronLayerTest, TestAbsVal) {