From 53130b33e13d33b94f4a40fcca979c6ed70948ca Mon Sep 17 00:00:00 2001 From: aravindh_m Date: Thu, 20 Feb 2014 23:07:57 +0000 Subject: [PATCH 1/4] Bugs fixed in the euclidean loss layer. --- include/caffe/vision_layers.hpp | 7 ++- src/caffe/layers/loss_layer.cu | 43 +++++++++++-- src/caffe/test/test_euclidean_loss_layer.cpp | 66 +++++++++++++++++++- src/caffe/util/math_functions.cu | 33 ++++++++++ 4 files changed, 139 insertions(+), 10 deletions(-) diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index 4db2556de62..53804fed6e5 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -467,7 +467,7 @@ template class EuclideanLossLayer : public Layer { public: explicit EuclideanLossLayer(const LayerParameter& param) - : Layer(param), difference_() {} + : Layer(param), difference_(), scale_(param.scale()) {} virtual void SetUp(const vector*>& bottom, vector*>* top); @@ -475,14 +475,15 @@ class EuclideanLossLayer : public Layer { // 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; } + vector*>* top); virtual void Forward_gpu(const vector*>& bottom, - vector*>* top) { return; } + 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); Blob difference_; + Dtype scale_; }; diff --git a/src/caffe/layers/loss_layer.cu b/src/caffe/layers/loss_layer.cu index ac05ba41b84..4ade6d7cda3 100644 --- a/src/caffe/layers/loss_layer.cu +++ b/src/caffe/layers/loss_layer.cu @@ -97,14 +97,45 @@ template void EuclideanLossLayer::SetUp( const vector*>& bottom, vector*>* top) { CHECK_EQ(bottom.size(), 2) << "Loss Layer takes two blobs as input."; - CHECK_EQ(top->size(), 0) << "Loss Layer takes no as output."; + CHECK_LE(top->size(), 1) << "Loss Layer takes atmost one output."; CHECK_EQ(bottom[0]->num(), bottom[1]->num()) - << "The data and label should have the same number."; + << "The data1 and data2 should have the same number."; CHECK_EQ(bottom[0]->channels(), bottom[1]->channels()); CHECK_EQ(bottom[0]->height(), bottom[1]->height()); CHECK_EQ(bottom[0]->width(), bottom[1]->width()); difference_.Reshape(bottom[0]->num(), bottom[0]->channels(), bottom[0]->height(), bottom[0]->width()); + if(top->size() == 1) { + (*top)[0]->Reshape(1, 1, 1, 1); + } +} + + +template +void EuclideanLossLayer::Forward_cpu(const vector*>& bottom, + vector*>* top) { + if(top->size() == 1) { + 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()); + ((Dtype*)(*top)[0]->mutable_cpu_data())[0] = scale_ * loss / num; + } +} + +template +void EuclideanLossLayer::Forward_gpu(const vector*>& bottom, + vector*>* top) { + if(top->size() == 1) { + int count = bottom[0]->count(); + int num = bottom[0]->num(); + caffe_gpu_sub(count, bottom[0]->gpu_data(), bottom[1]->gpu_data(), + difference_.mutable_gpu_data()); + Dtype loss; + caffe_gpu_dot(count, difference_.gpu_data(), difference_.gpu_data(), &loss); + ((Dtype*)(*top)[0]->mutable_cpu_data())[0] = scale_ * loss / num; + } } template @@ -115,11 +146,13 @@ Dtype EuclideanLossLayer::Backward_cpu(const vector*>& top, 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); + count, difference_.cpu_data(), difference_.cpu_data()) / num / Dtype(2.); // Compute the gradient - caffe_axpby(count, Dtype(1) / num, difference_.cpu_data(), Dtype(0), + caffe_axpby(count, scale_ / num, difference_.cpu_data(), Dtype(0.), (*bottom)[0]->mutable_cpu_diff()); - return loss; + caffe_axpby(count, -scale_ / num, difference_.cpu_data(), Dtype(0.), + (*bottom)[1]->mutable_cpu_diff()); + return scale_*loss*Dtype(2.); } template diff --git a/src/caffe/test/test_euclidean_loss_layer.cpp b/src/caffe/test/test_euclidean_loss_layer.cpp index 82ea682fcc0..04877d57ef3 100644 --- a/src/caffe/test/test_euclidean_loss_layer.cpp +++ b/src/caffe/test/test_euclidean_loss_layer.cpp @@ -22,8 +22,9 @@ template class EuclideanLossLayerTest : public ::testing::Test { protected: EuclideanLossLayerTest() - : blob_bottom_data_(new Blob(10, 5, 1, 1)), - blob_bottom_label_(new Blob(10, 5, 1, 1)) { + : blob_bottom_data_(new Blob(10, 5, 1, 10)), + blob_bottom_label_(new Blob(10, 5, 1, 10)), + blob_top_(new Blob()) { // fill the values FillerParameter filler_param; GaussianFiller filler(filler_param); @@ -31,13 +32,16 @@ class EuclideanLossLayerTest : public ::testing::Test { blob_bottom_vec_.push_back(blob_bottom_data_); filler.Fill(this->blob_bottom_label_); blob_bottom_vec_.push_back(blob_bottom_label_); + blob_top_vec_.push_back(blob_top_); } virtual ~EuclideanLossLayerTest() { delete blob_bottom_data_; delete blob_bottom_label_; + delete blob_top_; } Blob* const blob_bottom_data_; Blob* const blob_bottom_label_; + Blob* blob_top_; vector*> blob_bottom_vec_; vector*> blob_top_vec_; }; @@ -45,6 +49,64 @@ class EuclideanLossLayerTest : public ::testing::Test { typedef ::testing::Types Dtypes; TYPED_TEST_CASE(EuclideanLossLayerTest, Dtypes); +TYPED_TEST(EuclideanLossLayerTest, TestSetUp) { + LayerParameter layer_param; + shared_ptr > layer( + new EuclideanLossLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + EXPECT_EQ(this->blob_top_->num(), 1); + EXPECT_EQ(this->blob_top_->height(), 1); + EXPECT_EQ(this->blob_top_->width(), 1); + EXPECT_EQ(this->blob_top_->channels(), 1); +} + + +TYPED_TEST(EuclideanLossLayerTest, TestCPU) { + LayerParameter layer_param; + Caffe::set_mode(Caffe::CPU); + shared_ptr > layer( + new EuclideanLossLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + TypeParam sum = 0; + for (int n = 0; n < 10; ++n) { + for (int c = 0; c < 5; ++c) { + for (int h = 0; h < 1; ++h) { + for (int w = 0; w < 10; ++w) { + sum += pow((this->blob_bottom_vec_[0]->data_at(n,c,h,w) - + this->blob_bottom_vec_[1]->data_at(n,c,h,w)), 2); + } + } + } + } + sum = sum /10; + EXPECT_LE(this->blob_top_vec_[0]->data_at(0, 0, 0, 0) - 1e-4, sum); + EXPECT_GE(this->blob_top_vec_[0]->data_at(0, 0, 0, 0) + 1e-4, sum); +} + +TYPED_TEST(EuclideanLossLayerTest, TestGPU) { + LayerParameter layer_param; + Caffe::set_mode(Caffe::GPU); + shared_ptr > layer( + new EuclideanLossLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + TypeParam sum = 0; + for (int n = 0; n < 10; ++n) { + for (int c = 0; c < 5; ++c) { + for (int h = 0; h < 1; ++h) { + for (int w = 0; w < 10; ++w) { + sum += pow((this->blob_bottom_vec_[0]->data_at(n,c,h,w) - + this->blob_bottom_vec_[1]->data_at(n,c,h,w)), 2); + } + } + } + } + sum = sum / 10; + EXPECT_LE(this->blob_top_vec_[0]->data_at(0, 0, 0, 0) - 1e-4, sum); + EXPECT_GE(this->blob_top_vec_[0]->data_at(0, 0, 0, 0) + 1e-4, sum); +} + TYPED_TEST(EuclideanLossLayerTest, TestGradientCPU) { LayerParameter layer_param; Caffe::set_mode(Caffe::CPU); diff --git a/src/caffe/util/math_functions.cu b/src/caffe/util/math_functions.cu index e9305810e81..1c655079cab 100644 --- a/src/caffe/util/math_functions.cu +++ b/src/caffe/util/math_functions.cu @@ -32,5 +32,38 @@ void caffe_gpu_mul(const int N, const double* a, N, a, b, y); } +/* grid stride kernel */ +template +__global__ void sub_kernel(const int n, const Dtype* a, + const Dtype* b, Dtype* y) { + for(int i = threadIdx.x + blockIdx.x * blockDim.x; + i < n; + i += blockDim.x + gridDim.x) + { + y[i] = a[i] - b[i]; + } +} + +template <> +void caffe_gpu_sub(const int N, const float* a, + const float* b, float* y) { + int deviceid; + cudaGetDevice(&deviceid); + int numSMs; + cudaDeviceGetAttribute(&numSMs, cudaDevAttrMultiProcessorCount, deviceid); + sub_kernel<<>>( + N, a, b, y); +} + +template <> +void caffe_gpu_sub(const int N, const double* a, + const double* b, double* y) { + int deviceid; + cudaGetDevice(&deviceid); + int numSMs; + cudaDeviceGetAttribute(&numSMs, cudaDevAttrMultiProcessorCount, deviceid); + sub_kernel<<>>( + N, a, b, y); +} } // namespace caffe From 94dbb0a3d536b96755a383f113589a52c60a5be0 Mon Sep 17 00:00:00 2001 From: aravindh_m Date: Thu, 20 Feb 2014 23:46:55 +0000 Subject: [PATCH 2/4] Forgot to add the definition for caffe_gpu_sub() --- include/caffe/util/math_functions.hpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/include/caffe/util/math_functions.hpp b/include/caffe/util/math_functions.hpp index e9e2db8f274..ed2aad3afe4 100644 --- a/include/caffe/util/math_functions.hpp +++ b/include/caffe/util/math_functions.hpp @@ -78,6 +78,9 @@ void caffe_mul(const int N, const Dtype* a, const Dtype* b, Dtype* y); template void caffe_gpu_mul(const int N, const Dtype* a, const Dtype* b, Dtype* y); +template +void caffe_gpu_sub(int N, const Dtype* a, const Dtype* b, Dtype* y); + template void caffe_div(const int N, const Dtype* a, const Dtype* b, Dtype* y); From f4bae0b8a5b5f505d279c5ee82157c72c467bdc3 Mon Sep 17 00:00:00 2001 From: aravindh_m Date: Tue, 25 Feb 2014 23:32:12 +0000 Subject: [PATCH 3/4] Changed loss function to 1/2 ||x-y||^2 and updated the test cases to test the same. Changed testing macros to EXPECT_NEAR. --- src/caffe/layers/loss_layer.cu | 8 ++++---- src/caffe/test/test_euclidean_loss_layer.cpp | 10 ++++------ 2 files changed, 8 insertions(+), 10 deletions(-) diff --git a/src/caffe/layers/loss_layer.cu b/src/caffe/layers/loss_layer.cu index 4ade6d7cda3..68cc2c61406 100644 --- a/src/caffe/layers/loss_layer.cu +++ b/src/caffe/layers/loss_layer.cu @@ -120,7 +120,7 @@ void EuclideanLossLayer::Forward_cpu(const vector*>& bottom, 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()); - ((Dtype*)(*top)[0]->mutable_cpu_data())[0] = scale_ * loss / num; + ((Dtype*)(*top)[0]->mutable_cpu_data())[0] = scale_ * loss / (2*num); } } @@ -134,7 +134,7 @@ void EuclideanLossLayer::Forward_gpu(const vector*>& bottom, difference_.mutable_gpu_data()); Dtype loss; caffe_gpu_dot(count, difference_.gpu_data(), difference_.gpu_data(), &loss); - ((Dtype*)(*top)[0]->mutable_cpu_data())[0] = scale_ * loss / num; + ((Dtype*)(*top)[0]->mutable_cpu_data())[0] = scale_ * loss / (2*num); } } @@ -146,13 +146,13 @@ Dtype EuclideanLossLayer::Backward_cpu(const vector*>& top, 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.); + count, difference_.cpu_data(), difference_.cpu_data()) / num; // Compute the gradient caffe_axpby(count, scale_ / num, difference_.cpu_data(), Dtype(0.), (*bottom)[0]->mutable_cpu_diff()); caffe_axpby(count, -scale_ / num, difference_.cpu_data(), Dtype(0.), (*bottom)[1]->mutable_cpu_diff()); - return scale_*loss*Dtype(2.); + return scale_*loss/Dtype(2.); } template diff --git a/src/caffe/test/test_euclidean_loss_layer.cpp b/src/caffe/test/test_euclidean_loss_layer.cpp index 04877d57ef3..b25a5d479ce 100644 --- a/src/caffe/test/test_euclidean_loss_layer.cpp +++ b/src/caffe/test/test_euclidean_loss_layer.cpp @@ -79,9 +79,8 @@ TYPED_TEST(EuclideanLossLayerTest, TestCPU) { } } } - sum = sum /10; - EXPECT_LE(this->blob_top_vec_[0]->data_at(0, 0, 0, 0) - 1e-4, sum); - EXPECT_GE(this->blob_top_vec_[0]->data_at(0, 0, 0, 0) + 1e-4, sum); + sum = sum /(10*2); + EXPECT_NEAR(this->blob_top_vec_[0]->data_at(0, 0, 0, 0), sum); } TYPED_TEST(EuclideanLossLayerTest, TestGPU) { @@ -102,9 +101,8 @@ TYPED_TEST(EuclideanLossLayerTest, TestGPU) { } } } - sum = sum / 10; - EXPECT_LE(this->blob_top_vec_[0]->data_at(0, 0, 0, 0) - 1e-4, sum); - EXPECT_GE(this->blob_top_vec_[0]->data_at(0, 0, 0, 0) + 1e-4, sum); + sum = sum / (10*2); + EXPECT_NEAR(this->blob_top_vec_[0]->data_at(0, 0, 0, 0), sum); } TYPED_TEST(EuclideanLossLayerTest, TestGradientCPU) { From 88f8597b05087e65a207cea2c04c4fff5e58e28f Mon Sep 17 00:00:00 2001 From: aravindh_m Date: Wed, 26 Feb 2014 00:17:04 +0000 Subject: [PATCH 4/4] Fixed the bug with EXPECT_NEAR argument. --- src/caffe/test/test_euclidean_loss_layer.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/caffe/test/test_euclidean_loss_layer.cpp b/src/caffe/test/test_euclidean_loss_layer.cpp index b25a5d479ce..b540e5b4623 100644 --- a/src/caffe/test/test_euclidean_loss_layer.cpp +++ b/src/caffe/test/test_euclidean_loss_layer.cpp @@ -80,7 +80,7 @@ TYPED_TEST(EuclideanLossLayerTest, TestCPU) { } } sum = sum /(10*2); - EXPECT_NEAR(this->blob_top_vec_[0]->data_at(0, 0, 0, 0), sum); + EXPECT_NEAR(this->blob_top_vec_[0]->data_at(0, 0, 0, 0), sum, 1e-4); } TYPED_TEST(EuclideanLossLayerTest, TestGPU) { @@ -102,7 +102,7 @@ TYPED_TEST(EuclideanLossLayerTest, TestGPU) { } } sum = sum / (10*2); - EXPECT_NEAR(this->blob_top_vec_[0]->data_at(0, 0, 0, 0), sum); + EXPECT_NEAR(this->blob_top_vec_[0]->data_at(0, 0, 0, 0), sum, 1e-4); } TYPED_TEST(EuclideanLossLayerTest, TestGradientCPU) {