From: Sergio <sguada@gmail.com>
Date: Tue, 27 May 2014 17:57:26 +0000 (-0700)
Subject: Added ForwardGPU to ThresholdLayer and to the tests
X-Git-Tag: submit/tizen/20180823.020014~653^2~131^2~6
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=868dd61449a85dfed2f2807b9178584098473166;p=platform%2Fupstream%2Fcaffeonacl.git

Added ForwardGPU to ThresholdLayer and to the tests
---

diff --git a/include/caffe/neuron_layers.hpp b/include/caffe/neuron_layers.hpp
index 85d5a160..180ce7a2 100644
--- a/include/caffe/neuron_layers.hpp
+++ b/include/caffe/neuron_layers.hpp
@@ -209,7 +209,8 @@ class TanHLayer : public NeuronLayer<Dtype> {
 
   y = 1 if x > threshold
   y = 0 if x <= threshold
-  y' = don't have
+  
+  y' = don't differenciable
 */
 template <typename Dtype>
 class ThresholdLayer : public NeuronLayer<Dtype> {
@@ -222,11 +223,13 @@ class ThresholdLayer : public NeuronLayer<Dtype> {
  protected:
   virtual Dtype Forward_cpu(const vector<Blob<Dtype>*>& bottom,
       vector<Blob<Dtype>*>* top);
+  virtual Dtype Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+      vector<Blob<Dtype>*>* top);
   virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
       const bool propagate_down, vector<Blob<Dtype>*>* bottom) {
     NOT_IMPLEMENTED;
   }
-  
+
   Dtype threshold_;
 };
 
diff --git a/src/caffe/layers/threshold_layer.cu b/src/caffe/layers/threshold_layer.cu
new file mode 100644
index 00000000..04975076
--- /dev/null
+++ b/src/caffe/layers/threshold_layer.cu
@@ -0,0 +1,38 @@
+// Copyright 2014 BVLC and contributors.
+
+#include <algorithm>
+#include <vector>
+
+#include "caffe/layer.hpp"
+#include "caffe/vision_layers.hpp"
+
+using std::max;
+
+namespace caffe {
+
+template <typename Dtype>
+__global__ void ThresholdForward(const int n, const, threshold, const Dtype* in, Dtype* out) {
+  CUDA_KERNEL_LOOP(index, n) {
+    out[index] = in[index] > threshold ? 1 : 0;
+  }
+}
+
+template <typename Dtype>
+Dtype ThresholdLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+    vector<Blob<Dtype>*>* top) {
+  const Dtype* bottom_data = bottom[0]->gpu_data();
+  Dtype* top_data = (*top)[0]->mutable_gpu_data();
+  const int count = bottom[0]->count();
+  // NOLINT_NEXT_LINE(whitespace/operators)
+  ReLUForward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
+      count, threshold_, bottom_data, top_data);
+  CUDA_POST_KERNEL_CHECK;
+
+  return Dtype(0);
+}
+
+
+INSTANTIATE_CLASS(ReLULayer);
+
+
+}  // namespace caffe
diff --git a/src/caffe/test/test_threshold_layer.cpp b/src/caffe/test/test_threshold_layer.cpp
index 4f87d0d0..b8f9d969 100644
--- a/src/caffe/test/test_threshold_layer.cpp
+++ b/src/caffe/test/test_threshold_layer.cpp
@@ -99,4 +99,52 @@ TYPED_TEST(ThresholdLayerTest, TestCPU2) {
   }
 }
 
+TYPED_TEST(ThresholdLayerTest, TestGPU) {
+  LayerParameter layer_param;
+  Caffe::set_mode(Caffe::GPU);
+  ThresholdLayer<TypeParam> 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 TypeParam* bottom_data = this->blob_bottom_->cpu_data();
+  const TypeParam* top_data = this->blob_top_->cpu_data();
+  const TypeParam threshold_ = layer_param.threshold_param().threshold();
+  for (int i = 0; i < this->blob_bottom_->count(); ++i) {
+    EXPECT_GE(top_data[i], 0.);
+    EXPECT_LE(top_data[i], 1.);
+    if (top_data[i] == 0) {
+  EXPECT_LE(bottom_data[i],threshold_);
+    }
+    if (top_data[i] == 1) {
+  EXPECT_GT(bottom_data[i],threshold_);
+    }
+  }
+}
+
+TYPED_TEST(ThresholdLayerTest, TestGPU2) {
+  LayerParameter layer_param;
+  Caffe::set_mode(Caffe::GPU);
+  ThresholdParameter* threshold_param =
+    layer_param.mutable_threshold_param();
+  threshold_param->set_threshold(0.5);
+  ThresholdLayer<TypeParam> 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 TypeParam* bottom_data = this->blob_bottom_->cpu_data();
+  const TypeParam* top_data = this->blob_top_->cpu_data();
+  const TypeParam threshold_ = layer_param.threshold_param().threshold();
+  EXPECT_FLOAT_EQ(threshold_,0.5);
+  for (int i = 0; i < this->blob_bottom_->count(); ++i) {
+    EXPECT_GE(top_data[i], 0.);
+    EXPECT_LE(top_data[i], 1.);
+    if (top_data[i] == 0) {
+  EXPECT_LE(bottom_data[i],threshold_);
+    }
+    if (top_data[i] == 1) {
+  EXPECT_TRUE(bottom_data[i] > threshold_); 
+    }
+  }
+}
+
 }  // namespace caffe