summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorPrzemysław Dolata <snowball91b@gmail.com>2018-08-20 07:55:36 +0200
committerWook Song <wook16.song@samsung.com>2020-01-23 22:50:49 +0900
commit985206a68a0a859f98790102e07222cc1ceb116f (patch)
tree3f2b810f46e329b02e23efb770f2a75fab728f8a
parentfe062ef6983f8ebefd8f2d6c140fd5c548dca59c (diff)
parentb6ad8b657fb689ebc061d800feacfaf3ab1185c5 (diff)
downloadcaffe-985206a68a0a859f98790102e07222cc1ceb116f.tar.gz
caffe-985206a68a0a859f98790102e07222cc1ceb116f.tar.bz2
caffe-985206a68a0a859f98790102e07222cc1ceb116f.zip
Merge pull request #6320 from Noiredd/clip
Clip layer - resurrection
-rw-r--r--docs/tutorial/layers.md1
-rw-r--r--docs/tutorial/layers/clip.md20
-rw-r--r--include/caffe/layers/clip_layer.hpp75
-rw-r--r--src/caffe/layer_factory.cpp1
-rw-r--r--src/caffe/layers/clip_layer.cpp51
-rw-r--r--src/caffe/layers/clip_layer.cu67
-rw-r--r--src/caffe/proto/caffe.proto9
-rw-r--r--src/caffe/test/test_neuron_layer.cpp61
8 files changed, 284 insertions, 1 deletions
diff --git a/docs/tutorial/layers.md b/docs/tutorial/layers.md
index 78a46f3a..5036d4fd 100644
--- a/docs/tutorial/layers.md
+++ b/docs/tutorial/layers.md
@@ -93,6 +93,7 @@ Layers:
* [Log](layers/log.html) - f(x) = log(x).
* [BNLL](layers/bnll.html) - f(x) = log(1 + exp(x)).
* [Threshold](layers/threshold.html) - performs step function at user defined threshold.
+* [Clip](layers/clip.html) - clips a blob between a fixed minimum and maximum value.
* [Bias](layers/bias.html) - adds a bias to a blob that can either be learned or fixed.
* [Scale](layers/scale.html) - scales a blob by an amount that can either be learned or fixed.
diff --git a/docs/tutorial/layers/clip.md b/docs/tutorial/layers/clip.md
new file mode 100644
index 00000000..d6a20f5f
--- /dev/null
+++ b/docs/tutorial/layers/clip.md
@@ -0,0 +1,20 @@
+---
+title: Clip Layer
+---
+
+# Clip Layer
+
+* Layer type: `Clip`
+* [Doxygen Documentation](http://caffe.berkeleyvision.org/doxygen/classcaffe_1_1ClipLayer.html)
+* Header: [`./include/caffe/layers/clip_layer.hpp`](https://github.com/BVLC/caffe/blob/master/include/caffe/layers/clip_layer.hpp)
+* CPU implementation: [`./src/caffe/layers/clip_layer.cpp`](https://github.com/BVLC/caffe/blob/master/src/caffe/layers/clip_layer.cpp)
+* CUDA GPU implementation: [`./src/caffe/layers/clip_layer.cu`](https://github.com/BVLC/caffe/blob/master/src/caffe/layers/clip_layer.cu)
+
+## Parameters
+
+* Parameters (`ClipParameter clip_param`)
+* From [`./src/caffe/proto/caffe.proto`](https://github.com/BVLC/caffe/blob/master/src/caffe/proto/caffe.proto):
+
+{% highlight Protobuf %}
+{% include proto/ClipParameter.txt %}
+{% endhighlight %}
diff --git a/include/caffe/layers/clip_layer.hpp b/include/caffe/layers/clip_layer.hpp
new file mode 100644
index 00000000..2788193e
--- /dev/null
+++ b/include/caffe/layers/clip_layer.hpp
@@ -0,0 +1,75 @@
+#ifndef CAFFE_CLIP_LAYER_HPP_
+#define CAFFE_CLIP_LAYER_HPP_
+
+#include <vector>
+
+#include "caffe/blob.hpp"
+#include "caffe/layer.hpp"
+#include "caffe/proto/caffe.pb.h"
+
+#include "caffe/layers/neuron_layer.hpp"
+
+namespace caffe {
+
+/**
+ * @brief Clip: @f$ y = \max(min, \min(max, x)) @f$.
+ */
+template <typename Dtype>
+class ClipLayer : public NeuronLayer<Dtype> {
+ public:
+ /**
+ * @param param provides ClipParameter clip_param,
+ * with ClipLayer options:
+ * - min
+ * - max
+ */
+ explicit ClipLayer(const LayerParameter& param)
+ : NeuronLayer<Dtype>(param) {}
+
+ virtual inline const char* type() const { return "Clip"; }
+
+ protected:
+ /**
+ * @param bottom input Blob vector (length 1)
+ * -# @f$ (N \times C \times H \times W) @f$
+ * the inputs @f$ x @f$
+ * @param top output Blob vector (length 1)
+ * -# @f$ (N \times C \times H \times W) @f$
+ * the computed outputs @f$
+ * y = \max(min, \min(max, x))
+ * @f$
+ */
+ virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top);
+ virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top);
+
+ /**
+ * @brief Computes the error gradient w.r.t. the clipped inputs.
+ *
+ * @param top output Blob vector (length 1), providing the error gradient with
+ * respect to the outputs
+ * -# @f$ (N \times C \times H \times W) @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 H \times W) @f$
+ * the inputs @f$ x @f$; Backward fills their diff with
+ * gradients @f$
+ * \frac{\partial E}{\partial x} = \left\{
+ * \begin{array}{lr}
+ * 0 & \mathrm{if} \; x < min \vee x > max \\
+ * \frac{\partial E}{\partial y} & \mathrm{if} \; x \ge min \wedge x \le max
+ * \end{array} \right.
+ * @f$
+ */
+ virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
+ const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom);
+ virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
+ const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom);
+};
+
+} // namespace caffe
+
+#endif // CAFFE_CLIP_LAYER_HPP_
diff --git a/src/caffe/layer_factory.cpp b/src/caffe/layer_factory.cpp
index 9f9026b1..d9984431 100644
--- a/src/caffe/layer_factory.cpp
+++ b/src/caffe/layer_factory.cpp
@@ -7,6 +7,7 @@
#include "caffe/layer.hpp"
#include "caffe/layer_factory.hpp"
+#include "caffe/layers/clip_layer.hpp"
#include "caffe/layers/conv_layer.hpp"
#include "caffe/layers/deconv_layer.hpp"
#include "caffe/layers/lrn_layer.hpp"
diff --git a/src/caffe/layers/clip_layer.cpp b/src/caffe/layers/clip_layer.cpp
new file mode 100644
index 00000000..9d9a5967
--- /dev/null
+++ b/src/caffe/layers/clip_layer.cpp
@@ -0,0 +1,51 @@
+#include <algorithm>
+#include <vector>
+
+#include "caffe/layers/clip_layer.hpp"
+
+namespace caffe {
+
+template <typename Dtype>
+void ClipLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top) {
+ const Dtype* bottom_data = bottom[0]->cpu_data();
+ Dtype* top_data = top[0]->mutable_cpu_data();
+ const int count = bottom[0]->count();
+
+ Dtype min = this->layer_param_.clip_param().min();
+ Dtype max = this->layer_param_.clip_param().max();
+
+ for (int i = 0; i < count; ++i) {
+ top_data[i] = std::max(min, std::min(bottom_data[i], max));
+ }
+}
+
+template <typename Dtype>
+void ClipLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
+ const vector<bool>& propagate_down,
+ const vector<Blob<Dtype>*>& bottom) {
+ if (propagate_down[0]) {
+ const Dtype* bottom_data = bottom[0]->cpu_data();
+ const Dtype* top_diff = top[0]->cpu_diff();
+ Dtype* bottom_diff = bottom[0]->mutable_cpu_diff();
+ const int count = bottom[0]->count();
+
+ Dtype min = this->layer_param_.clip_param().min();
+ Dtype max = this->layer_param_.clip_param().max();
+
+ for (int i = 0; i < count; ++i) {
+ bottom_diff[i] = top_diff[i] * (
+ bottom_data[i] >= min && bottom_data[i] <= max);
+ }
+ }
+}
+
+
+#ifdef CPU_ONLY
+STUB_GPU(ClipLayer);
+#endif
+
+INSTANTIATE_CLASS(ClipLayer);
+REGISTER_LAYER_CLASS(Clip);
+
+} // namespace caffe
diff --git a/src/caffe/layers/clip_layer.cu b/src/caffe/layers/clip_layer.cu
new file mode 100644
index 00000000..56f3be32
--- /dev/null
+++ b/src/caffe/layers/clip_layer.cu
@@ -0,0 +1,67 @@
+#include <vector>
+
+#include "caffe/layers/clip_layer.hpp"
+#include "caffe/util/math_functions.hpp"
+
+namespace caffe {
+
+__global__ void ClipForward(const int n, const float* in, float* out,
+ float p_min, float p_max) {
+ CUDA_KERNEL_LOOP(index, n) {
+ out[index] = fmaxf(p_min, fminf(in[index], p_max));
+ }
+}
+
+__global__ void ClipForward(const int n, const double* in, double* out,
+ double p_min, double p_max) {
+ CUDA_KERNEL_LOOP(index, n) {
+ out[index] = fmax(p_min, fmin(in[index], p_max));
+ }
+}
+
+template <typename Dtype>
+void ClipLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+ const vector<Blob<Dtype>*>& top) {
+ const Dtype* bottom_data = bottom[0]->gpu_data();
+ Dtype* top_data = top[0]->mutable_gpu_data();
+ const int count = bottom[0]->count();
+ Dtype p_min = this->layer_param_.clip_param().min();
+ Dtype p_max = this->layer_param_.clip_param().max();
+ // NOLINT_NEXT_LINE(whitespace/operators)
+ ClipForward<<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
+ count, bottom_data, top_data, p_min, p_max);
+ CUDA_POST_KERNEL_CHECK;
+}
+
+template <typename Dtype>
+__global__ void ClipBackward(const int n, const Dtype* in_diff,
+ const Dtype* in_data, Dtype* out_diff, Dtype p_min, Dtype p_max) {
+ CUDA_KERNEL_LOOP(index, n) {
+ out_diff[index] = in_diff[index] * (
+ in_data[index] >= p_min && in_data[index] <= p_max);
+ }
+}
+
+template <typename Dtype>
+void ClipLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
+ const vector<bool>& propagate_down,
+ const vector<Blob<Dtype>*>& bottom) {
+ if (propagate_down[0]) {
+ const Dtype* bottom_data = bottom[0]->gpu_data();
+ const Dtype* top_diff = top[0]->gpu_diff();
+ Dtype* bottom_diff = bottom[0]->mutable_gpu_diff();
+ const int count = bottom[0]->count();
+ Dtype p_min = this->layer_param_.clip_param().min();
+ Dtype p_max = this->layer_param_.clip_param().max();
+ // NOLINT_NEXT_LINE(whitespace/operators)
+ ClipBackward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
+ count, top_diff, bottom_data, bottom_diff, p_min, p_max);
+ CUDA_POST_KERNEL_CHECK;
+ }
+}
+
+
+INSTANTIATE_LAYER_GPU_FUNCS(ClipLayer);
+
+
+} // namespace caffe
diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto
index f784aa96..5c235c6f 100644
--- a/src/caffe/proto/caffe.proto
+++ b/src/caffe/proto/caffe.proto
@@ -322,7 +322,7 @@ message ParamSpec {
// NOTE
// Update the next available ID when you add a new LayerParameter field.
//
-// LayerParameter next available layer-specific ID: 148 (last added: swish_param)
+// LayerParameter next available layer-specific ID: 149 (last added: clip_param)
message LayerParameter {
optional string name = 1; // the layer name
optional string type = 2; // the layer type
@@ -378,6 +378,7 @@ message LayerParameter {
optional ArgMaxParameter argmax_param = 103;
optional BatchNormParameter batch_norm_param = 139;
optional BiasParameter bias_param = 141;
+ optional ClipParameter clip_param = 148;
optional ConcatParameter concat_param = 104;
optional ContrastiveLossParameter contrastive_loss_param = 105;
optional ConvolutionParameter convolution_param = 106;
@@ -505,6 +506,12 @@ message ArgMaxParameter {
optional int32 axis = 3;
}
+// Message that stores parameters used by ClipLayer
+message ClipParameter {
+ required float min = 1;
+ required float max = 2;
+}
+
message ConcatParameter {
// The axis along which to concatenate -- may be negative to index from the
// end (e.g., -1 for the last axis). Other axes must have the
diff --git a/src/caffe/test/test_neuron_layer.cpp b/src/caffe/test/test_neuron_layer.cpp
index 83d80fcd..d1ecc37b 100644
--- a/src/caffe/test/test_neuron_layer.cpp
+++ b/src/caffe/test/test_neuron_layer.cpp
@@ -10,6 +10,7 @@
#include "caffe/layers/absval_layer.hpp"
#include "caffe/layers/bnll_layer.hpp"
+#include "caffe/layers/clip_layer.hpp"
#include "caffe/layers/dropout_layer.hpp"
#include "caffe/layers/elu_layer.hpp"
#include "caffe/layers/exp_layer.hpp"
@@ -206,6 +207,66 @@ TYPED_TEST(NeuronLayerTest, TestAbsGradient) {
this->blob_top_vec_);
}
+TYPED_TEST(NeuronLayerTest, TestClip) {
+ typedef typename TypeParam::Dtype Dtype;
+ LayerParameter layer_param;
+ CHECK(google::protobuf::TextFormat::ParseFromString(
+ "clip_param { min: -1, max: 2 }", &layer_param));
+ ClipLayer<Dtype> layer(layer_param);
+ layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_);
+ layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_);
+ // Now, check values
+ const Dtype* bottom_data = this->blob_bottom_->cpu_data();
+ const Dtype* top_data = this->blob_top_->cpu_data();
+ for (int i = 0; i < this->blob_bottom_->count(); ++i) {
+ EXPECT_GE(top_data[i], -1);
+ EXPECT_LE(top_data[i], 2);
+ EXPECT_TRUE(bottom_data[i] > -1 || top_data[i] == -1);
+ EXPECT_TRUE(bottom_data[i] < 2 || top_data[i] == 2);
+ EXPECT_TRUE(!(bottom_data[i] >= -1 && bottom_data[i] <= 2)
+ || top_data[i] == bottom_data[i]);
+ }
+}
+
+TYPED_TEST(NeuronLayerTest, TestClipGradient) {
+ typedef typename TypeParam::Dtype Dtype;
+ LayerParameter layer_param;
+ CHECK(google::protobuf::TextFormat::ParseFromString(
+ "clip_param { min: -1, max: 2 }", &layer_param));
+ ClipLayer<Dtype> layer(layer_param);
+ // Unfortunately, it might happen that an input value lands exactly within
+ // the discontinuity region of the Clip function. In this case the numeric
+ // gradient is likely to differ significantly (i.e. by a value larger than
+ // checker tolerance) from the computed gradient. To handle such cases, we
+ // eliminate such values from the input blob before the gradient check.
+ const Dtype epsilon = 1e-2;
+ const Dtype min_range_start = layer_param.clip_param().min() - epsilon;
+ const Dtype min_range_end = layer_param.clip_param().min() + epsilon;
+ const Dtype max_range_start = layer_param.clip_param().max() - epsilon;
+ const Dtype max_range_end = layer_param.clip_param().max() + epsilon;
+ // The input blob is owned by the NeuronLayerTest object, so we begin with
+ // creating a temporary blob and copying the input data there.
+ Blob<Dtype> temp_bottom;
+ temp_bottom.ReshapeLike(*this->blob_bottom_);
+ const Dtype* bottom_data = this->blob_bottom_->cpu_data();
+ Dtype* temp_data_mutable = temp_bottom.mutable_cpu_data();
+ for (int i = 0; i < this->blob_bottom_->count(); ++i) {
+ if (bottom_data[i] >= min_range_start &&
+ bottom_data[i] <= min_range_end) {
+ temp_data_mutable[i] = bottom_data[i] - epsilon;
+ } else if (bottom_data[i] >= max_range_start &&
+ bottom_data[i] <= max_range_end) {
+ temp_data_mutable[i] = bottom_data[i] + epsilon;
+ } else {
+ temp_data_mutable[i] = bottom_data[i];
+ }
+ }
+ vector<Blob<Dtype>*> temp_bottom_vec;
+ temp_bottom_vec.push_back(&temp_bottom);
+ GradientChecker<Dtype> checker(epsilon, 1e-3);
+ checker.CheckGradientEltwise(&layer, temp_bottom_vec, this->blob_top_vec_);
+}
+
TYPED_TEST(NeuronLayerTest, TestReLU) {
typedef typename TypeParam::Dtype Dtype;
LayerParameter layer_param;