Constrain Prelu negative slope

include/caffe/layer.hpp

@@ -315,7 +315,22 @@ class Layer {
     }
     param_propagate_down_[param_id] = value;
   }
-
+  /**
+   * @brief Called on each layer after net's parameters have been updated
+   *        using the solver.
+   */
+  void PostUpdateProcessing() {
+    switch (Caffe::mode()) {
+    case Caffe::CPU:
+      PostUpdateProcessing_cpu();
+      break;
+    case Caffe::GPU:
+      PostUpdateProcessing_gpu();
+      break;
+    default:
+      LOG(FATAL) << "Unknown caffe mode.";
+    }
+  }
 
  protected:
   /** The protobuf that stores the layer parameters */
@@ -363,6 +378,20 @@ class Layer {
     Backward_cpu(top, propagate_down, bottom);
   }
 
+  /**
+   * @brief Perform any processing required after the solver has updated
+   *        network parameters. Called only when Caffe mode is CPU.
+   */
+  virtual void PostUpdateProcessing_cpu() { /* Default behavior: no action.*/ }
+  /**
+   * @brief Perform any processing required after the solver has updated
+   *        network parameters. Called only when Caffe mode is GPU.
+   */
+  virtual void PostUpdateProcessing_gpu() {
+    // Call cpu code as a backup.
+    PostUpdateProcessing_cpu();
+  }
+
   /**
    * Called by the parent Layer's SetUp to check that the number of bottom
    * and top Blobs provided as input match the expected numbers specified by

include/caffe/neuron_layers.hpp

@@ -795,7 +795,17 @@ class PReLULayer : public NeuronLayer<Dtype> {
   virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
       const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom);
 
+  /*********
+   * @brief Perform the post-update processing to constrain the negative slopes
+   *        of the Prelu.
+   *********/
+  virtual void PostUpdateProcessing_cpu();
+  virtual void PostUpdateProcessing_gpu();
+
   bool channel_shared_;
+  bool constrain_neg_slope_;
+  Dtype min_neg_slope_;
+  Dtype max_neg_slope_;
   Blob<Dtype> multiplier_;  // dot multiplier for backward computation of params
   Blob<Dtype> backward_buff_;  // temporary buffer for backward computation
   Blob<Dtype> bottom_memory_;  // memory for in-place computation

include/caffe/util/device_alternate.hpp

@@ -29,6 +29,10 @@ void classname<Dtype>::funcname##_##gpu(const vector<Blob<Dtype>*>& top, \
     const vector<bool>& propagate_down, \
     const vector<Blob<Dtype>*>& bottom) { NO_GPU; } \
 
+#define STUB_GPU_POSTUPDATEPROCESSING(classname) \
+template <typename Dtype> \
+void classname<Dtype>::PostUpdateProcessing_gpu() { NO_GPU; }
+
 #else  // Normal GPU + CPU Caffe.
 
 #include <cublas_v2.h>

src/caffe/layers/prelu_layer.cpp

@@ -13,6 +13,10 @@ void PReLULayer<Dtype>::LayerSetUp(const vector<Blob<Dtype>*>& bottom,
       << "Number of axes of bottom blob must be >=2.";
   PReLUParameter prelu_param = this->layer_param().prelu_param();
   int channels = bottom[0]->channels();
+  constrain_neg_slope_ = prelu_param.constrain_neg_slope();
+  min_neg_slope_ = prelu_param.min_neg_slope();
+  max_neg_slope_ = prelu_param.max_neg_slope();
+
   channel_shared_ = prelu_param.channel_shared();
   if (this->blobs_.size() > 0) {
     LOG(INFO) << "Skipping parameter initialization";
@@ -128,9 +132,29 @@ void PReLULayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
   }
 }
 
+template<typename Dtype>
+void PReLULayer<Dtype>::PostUpdateProcessing_cpu() {
+  if (!this->constrain_neg_slope_) {
+    return;
+  }
+
+  // Constrain the slopes to be between the limits.
+  Dtype* slopes = this->blobs_[0]->mutable_cpu_data();
+  int slope_count = this->blobs_[0]->count();
+  for (int i = 0; i < slope_count; ++i) {
+    Dtype slope = *slopes;
+    if (slope < this->min_neg_slope_) {
+      *slopes = this->min_neg_slope_;
+    } else if (slope > this->max_neg_slope_) {
+      *slopes = this->max_neg_slope_;
+    }
+    slopes++;
+  }
+}
 
 #ifdef CPU_ONLY
 STUB_GPU(PReLULayer);
+STUB_GPU_POSTUPDATEPROCESSING(PReLULayer);
 #endif
 
 INSTANTIATE_CLASS(PReLULayer);

src/caffe/layers/prelu_layer.cu

@@ -5,6 +5,20 @@
 
 namespace caffe {
 
+// CUDA kernel for constraining negative slope.
+template <typename Dtype>
+__global__ void ConstrainNegSlope(Dtype* slopes, int slope_count,
+                                         Dtype min_slope, Dtype max_slope) {
+  CUDA_KERNEL_LOOP(index, slope_count) {
+    Dtype slope = slopes[index];
+    if (slope < min_slope) {
+      slopes[index] = min_slope;
+    } else if (slope > max_slope) {
+      slopes[index] = max_slope;
+    }
+  }
+}
+
 // CUDA kernele for forward
 template <typename Dtype>
 __global__ void PReLUForward(const int n, const int channels, const int dim,
@@ -120,8 +134,22 @@ void PReLULayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
   }
 }
 
+template<typename Dtype>
+void PReLULayer<Dtype>::PostUpdateProcessing_gpu() {
+  if (!this->constrain_neg_slope_) {
+    return;
+  }
 
-INSTANTIATE_LAYER_GPU_FUNCS(PReLULayer);
+  Dtype* slopes = this->blobs_[0]->mutable_gpu_data();
+  int slope_count = this->blobs_[0]->count();
+  // NOLINT_NEXT_LINE(whitespace/operators)
+  ConstrainNegSlope<Dtype><<<CAFFE_GET_BLOCKS(slope_count),
+      CAFFE_CUDA_NUM_THREADS>>>(
+      slopes, slope_count, this->min_neg_slope_, this->max_neg_slope_);
+}
 
+INSTANTIATE_LAYER_GPU_FUNCS(PReLULayer);
+template void PReLULayer<float>::PostUpdateProcessing_gpu();
+template void PReLULayer<double>::PostUpdateProcessing_gpu();
 
 }  // namespace caffe

src/caffe/net.cpp

@@ -984,6 +984,9 @@ void Net<Dtype>::Update() {
   for (int i = 0; i < learnable_params_.size(); ++i) {
     learnable_params_[i]->Update();
   }
+  for (int i = 0; i < layers_.size(); ++i) {
+    layers_[i]->PostUpdateProcessing();
+  }
 }
 
 template <typename Dtype>

src/caffe/proto/caffe.proto

@@ -1231,4 +1231,11 @@ message PReLUParameter {
   optional FillerParameter filler = 1;
   // Whether or not slope paramters are shared across channels.
   optional bool channel_shared = 2 [default = false];
+  // Whether or not to constrain the negative slope between min_neg_slope and
+  // max_neg_slope.
+  optional bool constrain_neg_slope = 3 [default = true];
+  // limits on the value that the negative slope is allowed to acquire.
+  // Enforced only if constrain_neg_slope = true.
+  optional float min_neg_slope = 4 [default = 0.0];
+  optional float max_neg_slope = 5 [default = 1.0];
 }

src/caffe/test/test_neuron_layer.cpp

@@ -709,6 +709,45 @@ TYPED_TEST(NeuronLayerTest, TestPReLUInPlace) {
   }
 }
 
+TYPED_TEST(NeuronLayerTest, TestPReLU_ConstrainNegSlope) {
+  typedef typename TypeParam::Dtype Dtype;
+  LayerParameter layer_param;
+  layer_param.mutable_prelu_param()->set_constrain_neg_slope(true);
+  PReLULayer<Dtype> layer(layer_param);
+  layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_);
+  FillerParameter filler_param;
+  filler_param.set_std(10.0f);
+  GaussianFiller<Dtype> filler(filler_param);
+  filler.Fill(layer.blobs()[0].get());
+
+  int less_than_zero = 0;
+  int greater_than_one = 0;
+  for (int i = 0; i < layer.blobs()[0]->count(); ++i) {
+    Dtype slope = layer.blobs()[0]->cpu_data()[i];
+    less_than_zero += (slope < 0.0f) ? 1 : 0;
+    greater_than_one += (slope > 1.0f) ? 1 : 0;
+  }
+
+  // Expect to have some values outside the range 0.0 to 1.0.
+  EXPECT_GT(less_than_zero, 0);
+  EXPECT_GT(greater_than_one, 0);
+
+  // Run the logic that constrains the negative slope.
+  layer.PostUpdateProcessing();
+
+  less_than_zero = 0;
+  greater_than_one = 0;
+  for (int i = 0; i < layer.blobs()[0]->count(); ++i) {
+    Dtype slope = layer.blobs()[0]->cpu_data()[i];
+    less_than_zero += (slope < 0.0f) ? 1 : 0;
+    greater_than_one += (slope > 1.0f) ? 1 : 0;
+  }
+
+  // Expect to have no values outside the range 0.0 to 1.0.
+  EXPECT_EQ(less_than_zero, 0);
+  EXPECT_EQ(greater_than_one, 0);
+}
+
 #ifdef USE_CUDNN
 template <typename Dtype>
 class CuDNNNeuronLayerTest : public GPUDeviceTest<Dtype> {