From f8256c99fff3848efda073681ffddfb066f6904d Mon Sep 17 00:00:00 2001 From: to3i Date: Fri, 11 Jul 2014 13:19:23 +0200 Subject: [PATCH 1/5] Implemented elementwise max layer -> GPU test OK, CPU test fails --- include/caffe/vision_layers.hpp | 1 + src/caffe/layers/eltwise_layer.cpp | 65 +++++++++++ src/caffe/layers/eltwise_layer.cu | 159 +++++++++++++++++--------- src/caffe/proto/caffe.proto | 1 + src/caffe/test/test_eltwise_layer.cpp | 139 ++++++++++++++++++++++ 5 files changed, 314 insertions(+), 51 deletions(-) diff --git a/include/caffe/vision_layers.hpp b/include/caffe/vision_layers.hpp index 1016f761af5..f8be14cc3c6 100644 --- a/include/caffe/vision_layers.hpp +++ b/include/caffe/vision_layers.hpp @@ -91,6 +91,7 @@ class EltwiseLayer : public Layer { EltwiseParameter_EltwiseOp op_; vector coeffs_; + shared_ptr > max_idx_; }; /* Im2colLayer diff --git a/src/caffe/layers/eltwise_layer.cpp b/src/caffe/layers/eltwise_layer.cpp index 2c265f6678f..963b3384203 100644 --- a/src/caffe/layers/eltwise_layer.cpp +++ b/src/caffe/layers/eltwise_layer.cpp @@ -1,6 +1,7 @@ // Copyright 2014 BVLC and contributors. #include +#include #include "caffe/layer.hpp" #include "caffe/vision_layers.hpp" @@ -38,11 +39,20 @@ void EltwiseLayer::SetUp(const vector*>& bottom, coeffs_[i] = this->layer_param().eltwise_param().coeff(i); } } + // If max operation, we will initialize the vector index part. + if (this->layer_param_.eltwise_param().operation() == + EltwiseParameter_EltwiseOp_MAX && top->size() == 1) { + max_idx_.reset(new Blob(bottom[0]->num(), channels, + height, width)); + } } template Dtype EltwiseLayer::Forward_cpu( const vector*>& bottom, vector*>* top) { + int* mask = NULL; + const Dtype* bottom_data_a = NULL; + const Dtype* bottom_data_b = NULL; const int count = (*top)[0]->count(); Dtype* top_data = (*top)[0]->mutable_cpu_data(); switch (op_) { @@ -59,6 +69,50 @@ Dtype EltwiseLayer::Forward_cpu( caffe_axpy(count, coeffs_[i], bottom[i]->cpu_data(), top_data); } break; + case EltwiseParameter_EltwiseOp_MAX: + // Initialize + mask = max_idx_->mutable_cpu_data(); + caffe_set(count, -1, mask); + caffe_set(count, Dtype(-FLT_MAX), top_data); + // bottom 0 & 1 + bottom_data_a = bottom[0]->gpu_data(); + bottom_data_b = bottom[1]->gpu_data(); + for (int idx = 0; idx < count; ++idx) { + if (bottom_data_a[idx] > bottom_data_b[idx]) { + top_data[idx] = bottom_data_a[idx]; // maxval + mask[idx] = 0; // maxid + } else { + top_data[idx] = bottom_data_b[idx]; // maxval + mask[idx] = 1; // maxid + } + } + // bottom 2++ + bottom_data_a = top_data; + for (int blob_idx = 2; blob_idx < bottom.size(); ++blob_idx) { + bottom_data_b = bottom[blob_idx]->gpu_data(); + for (int idx = 0; idx < count; ++idx) { + if (bottom_data_a[idx] < bottom_data_b[idx]) { + top_data[idx] = bottom_data_b[idx]; // maxval + mask[idx] = blob_idx; // maxid + } + } + } + /* + NOT_IMPLEMENTED; + int channels = bottom[0]->channels(); + int height = bottom[0]->height(); + int width = bottom[0]->width(); + // The main loop + for (int n = 0; n < bottom[0]->num(); ++n) { + for (int c = 0; c < channels; ++c) { + for (int h = 0; h < height; ++h) { + for (int w = 0; w < width; ++w) { + } + } + } + } + */ + break; default: LOG(FATAL) << "Unknown elementwise operation."; } @@ -68,6 +122,7 @@ Dtype EltwiseLayer::Forward_cpu( template void EltwiseLayer::Backward_cpu(const vector*>& top, const vector& propagate_down, vector*>* bottom) { + const int* mask = NULL; const int count = top[0]->count(); const Dtype* top_data = top[0]->cpu_data(); const Dtype* top_diff = top[0]->cpu_diff(); @@ -87,6 +142,16 @@ void EltwiseLayer::Backward_cpu(const vector*>& top, caffe_cpu_scale(count, coeffs_[i], top_diff, bottom_diff); } break; + case EltwiseParameter_EltwiseOp_MAX: + mask = max_idx_->cpu_data(); + for (int index = 0; index < count; ++index) { + Dtype gradient = 0; + if (mask[index] == i) { + gradient += top_diff[index]; + } + bottom_diff[index] = gradient; + } + break; default: LOG(FATAL) << "Unknown elementwise operation."; } diff --git a/src/caffe/layers/eltwise_layer.cu b/src/caffe/layers/eltwise_layer.cu index 99c14feace1..0e2b6ce39f0 100644 --- a/src/caffe/layers/eltwise_layer.cu +++ b/src/caffe/layers/eltwise_layer.cu @@ -1,6 +1,7 @@ // Copyright 2014 BVLC and contributors. #include +#include #include "caffe/layer.hpp" #include "caffe/vision_layers.hpp" @@ -8,62 +9,118 @@ namespace caffe { -template -Dtype EltwiseLayer::Forward_gpu( - const vector*>& bottom, vector*>* top) { - const int count = (*top)[0]->count(); - Dtype* top_data = (*top)[0]->mutable_gpu_data(); - switch (op_) { - case EltwiseParameter_EltwiseOp_PROD: - caffe_gpu_mul(count, bottom[0]->gpu_data(), - bottom[1]->gpu_data(), top_data); - for (int i = 2; i < bottom.size(); ++i) { - caffe_gpu_mul(count, top_data, bottom[i]->gpu_data(), top_data); - } - break; - case EltwiseParameter_EltwiseOp_SUM: - caffe_gpu_set(count, Dtype(0.), top_data); - // TODO(shelhamer) does cuBLAS optimize to sum for coeff = 1? - for (int i = 0; i < bottom.size(); ++i) { - caffe_gpu_axpy(count, coeffs_[i], bottom[i]->gpu_data(), top_data); - } - break; - default: - LOG(FATAL) << "Unknown elementwise operation."; - } - return Dtype(0.); +template +__global__ void MaxForward(const int nthreads, const Dtype* bottom_data_a, + const Dtype* bottom_data_b, const int blob_idx, Dtype* top_data, + int* mask) { + CUDA_KERNEL_LOOP(index, nthreads) + { + Dtype maxval = -FLT_MAX; + int maxidx = -1; + if (bottom_data_a[index] > bottom_data_b[index]) { + // only update for very first bottom_data blob (blob_idx == 0) + if (blob_idx == 0) { + maxval = bottom_data_a[index]; + top_data[index] = maxval; + maxidx = blob_idx; + mask[index] = maxidx; + } + } else { + maxval = bottom_data_b[index]; + top_data[index] = maxval; + maxidx = blob_idx + 1; + mask[index] = maxidx; + } + } } -template +template +Dtype EltwiseLayer::Forward_gpu(const vector*>& bottom, + vector*>* top) { + int* mask = NULL; + const int count = (*top)[0]->count(); + Dtype* top_data = (*top)[0]->mutable_gpu_data(); + switch (op_) { + case EltwiseParameter_EltwiseOp_PROD: + caffe_gpu_mul(count, bottom[0]->gpu_data(), bottom[1]->gpu_data(), + top_data); + for (int i = 2; i < bottom.size(); ++i) { + caffe_gpu_mul(count, top_data, bottom[i]->gpu_data(), top_data); + } + break; + case EltwiseParameter_EltwiseOp_SUM: + caffe_gpu_set(count, Dtype(0.), top_data); + // TODO(shelhamer) does cuBLAS optimize to sum for coeff = 1? + for (int i = 0; i < bottom.size(); ++i) { + caffe_gpu_axpy(count, coeffs_[i], bottom[i]->gpu_data(), top_data); + } + break; + case EltwiseParameter_EltwiseOp_MAX: + mask = max_idx_->mutable_gpu_data(); + // NOLINT_NEXT_LINE(whitespace/operators) + MaxForward <<>>( + count, bottom[0]->gpu_data(), bottom[1]->gpu_data(), 0, top_data, mask); + for (int i = 2; i < bottom.size(); ++i) { + // NOLINT_NEXT_LINE(whitespace/operators) + MaxForward<<>>( + count, top_data, bottom[i]->gpu_data(), i-1, top_data, mask); + } + break; + default: + LOG(FATAL) << "Unknown elementwise operation."; + } + return Dtype(0.); +} + +template +__global__ void MaxBackward(const int nthreads, const Dtype* top_diff, + const int blob_idx, const int* mask, Dtype* bottom_diff) { + CUDA_KERNEL_LOOP(index, nthreads) + { + Dtype gradient = 0; + if (mask[index] == blob_idx) { + gradient += top_diff[index]; + } + bottom_diff[index] = gradient; + } +} + +template void EltwiseLayer::Backward_gpu(const vector*>& top, - const vector& propagate_down, vector*>* bottom) { - const int count = top[0]->count(); - const Dtype* top_data = top[0]->gpu_data(); - const Dtype* top_diff = top[0]->gpu_diff(); - for (int i = 0; i < bottom->size(); ++i) { - if (propagate_down[i]) { - const Dtype* bottom_data = (*bottom)[i]->gpu_data(); - Dtype* bottom_diff = (*bottom)[i]->mutable_gpu_diff(); - switch (op_) { - case EltwiseParameter_EltwiseOp_PROD: - caffe_gpu_div(count, top_data, bottom_data, bottom_diff); - caffe_gpu_mul(count, bottom_diff, top_diff, bottom_diff); - break; - case EltwiseParameter_EltwiseOp_SUM: - if (coeffs_[i] == Dtype(1.)) { - caffe_copy(count, top_diff, bottom_diff); - } else { - caffe_gpu_scale(count, coeffs_[i], top_diff, bottom_diff); - } - break; - default: - LOG(FATAL) << "Unknown elementwise operation."; - } - } - } + const vector& propagate_down, vector*>* bottom) { + const int* mask = NULL; + const int count = top[0]->count(); + const Dtype* top_data = top[0]->gpu_data(); + const Dtype* top_diff = top[0]->gpu_diff(); + for (int i = 0; i < bottom->size(); ++i) { + if (propagate_down[i]) { + const Dtype* bottom_data = (*bottom)[i]->gpu_data(); + Dtype* bottom_diff = (*bottom)[i]->mutable_gpu_diff(); + switch (op_) { + case EltwiseParameter_EltwiseOp_PROD: + caffe_gpu_div(count, top_data, bottom_data, bottom_diff); + caffe_gpu_mul(count, bottom_diff, top_diff, bottom_diff); + break; + case EltwiseParameter_EltwiseOp_SUM: + if (coeffs_[i] == Dtype(1.)) { + caffe_copy(count, top_diff, bottom_diff); + } else { + caffe_gpu_scale(count, coeffs_[i], top_diff, bottom_diff); + } + break; + case EltwiseParameter_EltwiseOp_MAX: + mask = max_idx_->gpu_data(); + // NOLINT_NEXT_LINE(whitespace/operators) + MaxBackward <<>>( + count, top_diff, i, mask, bottom_diff); + break; + default: + LOG(FATAL) << "Unknown elementwise operation."; + } + } + } } INSTANTIATE_CLASS(EltwiseLayer); - } // namespace caffe diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 70a4ab2e8e9..76cd358c20b 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -304,6 +304,7 @@ message EltwiseParameter { enum EltwiseOp { PROD = 0; SUM = 1; + MAX = 2; } optional EltwiseOp operation = 1 [default = SUM]; // element-wise operation repeated float coeff = 2; // blob-wise coefficient for SUM operation diff --git a/src/caffe/test/test_eltwise_layer.cpp b/src/caffe/test/test_eltwise_layer.cpp index 66490d2bc4e..a93f7c9913f 100644 --- a/src/caffe/test/test_eltwise_layer.cpp +++ b/src/caffe/test/test_eltwise_layer.cpp @@ -128,6 +128,7 @@ TYPED_TEST(EltwiseLayerTest, TestSumCoeff) { } } +<<<<<<< HEAD TYPED_TEST(EltwiseLayerTest, TestProdGradient) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; @@ -135,6 +136,144 @@ TYPED_TEST(EltwiseLayerTest, TestProdGradient) { eltwise_param->set_operation(EltwiseParameter_EltwiseOp_PROD); EltwiseLayer layer(layer_param); GradientChecker checker(1e-2, 1e-3); +======= +TYPED_TEST(EltwiseLayerTest, TestMaxCPU) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_MAX); + shared_ptr > layer( + new EltwiseLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + const TypeParam* data = this->blob_top_->cpu_data(); + const int count = this->blob_top_->count(); + const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); + const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); + const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_GE(data[i], in_data_a[i]); + EXPECT_GE(data[i], in_data_b[i]); + EXPECT_GE(data[i], in_data_c[i]); + } +} + +TYPED_TEST(EltwiseLayerTest, TestProdGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_PROD); + shared_ptr > layer( + new EltwiseLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + const TypeParam* data = this->blob_top_->cpu_data(); + const int count = this->blob_top_->count(); + const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); + const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); + const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_EQ(data[i], in_data_a[i] * in_data_b[i] * in_data_c[i]); + } +} + +TYPED_TEST(EltwiseLayerTest, TestSumGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); + shared_ptr > layer( + new EltwiseLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + const TypeParam* data = this->blob_top_->cpu_data(); + const int count = this->blob_top_->count(); + const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); + const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); + const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_EQ(data[i], in_data_a[i] + in_data_b[i] + in_data_c[i]); + } +} + +TYPED_TEST(EltwiseLayerTest, TestSumCoeffGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); + eltwise_param->add_coeff(1); + eltwise_param->add_coeff(-0.5); + eltwise_param->add_coeff(2); + shared_ptr > layer( + new EltwiseLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + const TypeParam* data = this->blob_top_->cpu_data(); + const int count = this->blob_top_->count(); + const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); + const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); + const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_NEAR(data[i], in_data_a[i] - 0.5*in_data_b[i] + 2*in_data_c[i], + 1e-4); + } +} + + +TYPED_TEST(EltwiseLayerTest, TestMaxGPU) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_MAX); + shared_ptr > layer( + new EltwiseLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); + layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); + const TypeParam* data = this->blob_top_->cpu_data(); + const int count = this->blob_top_->count(); + const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); + const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); + const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_GE(data[i], in_data_a[i]); + EXPECT_GE(data[i], in_data_b[i]); + EXPECT_GE(data[i], in_data_c[i]); + } +} + +TYPED_TEST(EltwiseLayerTest, TestProdCPUGradient) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_PROD); + EltwiseLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); + checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(EltwiseLayerTest, TestSumCPUGradient) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); + EltwiseLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); + checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(EltwiseLayerTest, TestSumCoeffCPUGradient) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); + eltwise_param->add_coeff(1); + eltwise_param->add_coeff(-0.5); + eltwise_param->add_coeff(2); + EltwiseLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); +>>>>>>> Implemented elementwise max layer -> GPU test OK, CPU test fails checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), &(this->blob_top_vec_)); } From 2bea03b488e1907c06814f03fbd125f50a061bbf Mon Sep 17 00:00:00 2001 From: to3i Date: Fri, 11 Jul 2014 16:50:44 +0200 Subject: [PATCH 2/5] changed gpu to cpu_data() --- src/caffe/layers/eltwise_layer.cpp | 15 +++++++------- src/caffe/test/test_eltwise_layer.cpp | 28 ++++++++++++++++++++++++++- 2 files changed, 35 insertions(+), 8 deletions(-) diff --git a/src/caffe/layers/eltwise_layer.cpp b/src/caffe/layers/eltwise_layer.cpp index 963b3384203..af00dcd0b6c 100644 --- a/src/caffe/layers/eltwise_layer.cpp +++ b/src/caffe/layers/eltwise_layer.cpp @@ -75,9 +75,10 @@ Dtype EltwiseLayer::Forward_cpu( caffe_set(count, -1, mask); caffe_set(count, Dtype(-FLT_MAX), top_data); // bottom 0 & 1 - bottom_data_a = bottom[0]->gpu_data(); - bottom_data_b = bottom[1]->gpu_data(); + bottom_data_a = bottom[0]->cpu_data(); + bottom_data_b = bottom[1]->cpu_data(); for (int idx = 0; idx < count; ++idx) { + bottom_data_a[idx]; if (bottom_data_a[idx] > bottom_data_b[idx]) { top_data[idx] = bottom_data_a[idx]; // maxval mask[idx] = 0; // maxid @@ -89,7 +90,7 @@ Dtype EltwiseLayer::Forward_cpu( // bottom 2++ bottom_data_a = top_data; for (int blob_idx = 2; blob_idx < bottom.size(); ++blob_idx) { - bottom_data_b = bottom[blob_idx]->gpu_data(); + bottom_data_b = bottom[blob_idx]->cpu_data(); for (int idx = 0; idx < count; ++idx) { if (bottom_data_a[idx] < bottom_data_b[idx]) { top_data[idx] = bottom_data_b[idx]; // maxval @@ -144,12 +145,12 @@ void EltwiseLayer::Backward_cpu(const vector*>& top, break; case EltwiseParameter_EltwiseOp_MAX: mask = max_idx_->cpu_data(); - for (int index = 0; index < count; ++index) { + for (int idx = 0; idx < count; ++idx) { Dtype gradient = 0; - if (mask[index] == i) { - gradient += top_diff[index]; + if (mask[idx] == i) { + gradient += top_diff[idx]; } - bottom_diff[index] = gradient; + bottom_diff[idx] = gradient; } break; default: diff --git a/src/caffe/test/test_eltwise_layer.cpp b/src/caffe/test/test_eltwise_layer.cpp index a93f7c9913f..fa4761c9844 100644 --- a/src/caffe/test/test_eltwise_layer.cpp +++ b/src/caffe/test/test_eltwise_layer.cpp @@ -219,7 +219,6 @@ TYPED_TEST(EltwiseLayerTest, TestSumCoeffGPU) { } } - TYPED_TEST(EltwiseLayerTest, TestMaxGPU) { Caffe::set_mode(Caffe::GPU); LayerParameter layer_param; @@ -278,8 +277,24 @@ TYPED_TEST(EltwiseLayerTest, TestSumCoeffCPUGradient) { &(this->blob_top_vec_)); } +<<<<<<< HEAD TYPED_TEST(EltwiseLayerTest, TestSumGradient) { typedef typename TypeParam::Dtype Dtype; +======= +TYPED_TEST(EltwiseLayerTest, TestMaxCPUGradient) { + Caffe::set_mode(Caffe::CPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_MAX); + EltwiseLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); + checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + +TYPED_TEST(EltwiseLayerTest, TestSumGPUGradient) { + Caffe::set_mode(Caffe::GPU); +>>>>>>> changed gpu to cpu_data() LayerParameter layer_param; EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); @@ -303,4 +318,15 @@ TYPED_TEST(EltwiseLayerTest, TestSumCoeffGradient) { &(this->blob_top_vec_)); } +TYPED_TEST(EltwiseLayerTest, TestMaxGPUGradient) { + Caffe::set_mode(Caffe::GPU); + LayerParameter layer_param; + EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); + eltwise_param->set_operation(EltwiseParameter_EltwiseOp_MAX); + EltwiseLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), + &(this->blob_top_vec_)); +} + } // namespace caffe From 56b7cadcb2a491787ecb34dee4e4668c70b7f14f Mon Sep 17 00:00:00 2001 From: to3i Date: Fri, 11 Jul 2014 16:58:46 +0200 Subject: [PATCH 3/5] remove unused statement --- src/caffe/layers/eltwise_layer.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/src/caffe/layers/eltwise_layer.cpp b/src/caffe/layers/eltwise_layer.cpp index af00dcd0b6c..f3446062f63 100644 --- a/src/caffe/layers/eltwise_layer.cpp +++ b/src/caffe/layers/eltwise_layer.cpp @@ -78,7 +78,6 @@ Dtype EltwiseLayer::Forward_cpu( bottom_data_a = bottom[0]->cpu_data(); bottom_data_b = bottom[1]->cpu_data(); for (int idx = 0; idx < count; ++idx) { - bottom_data_a[idx]; if (bottom_data_a[idx] > bottom_data_b[idx]) { top_data[idx] = bottom_data_a[idx]; // maxval mask[idx] = 0; // maxid From 7bedf36d4b59e39551c1f7533980018d82ae533c Mon Sep 17 00:00:00 2001 From: to3i Date: Mon, 14 Jul 2014 13:44:56 +0200 Subject: [PATCH 4/5] add init for bottom_diff (obsolete?) --- src/caffe/layers/eltwise_layer.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/src/caffe/layers/eltwise_layer.cpp b/src/caffe/layers/eltwise_layer.cpp index f3446062f63..474b8e49781 100644 --- a/src/caffe/layers/eltwise_layer.cpp +++ b/src/caffe/layers/eltwise_layer.cpp @@ -144,6 +144,7 @@ void EltwiseLayer::Backward_cpu(const vector*>& top, break; case EltwiseParameter_EltwiseOp_MAX: mask = max_idx_->cpu_data(); + caffe_set(count, Dtype(0), bottom_diff); for (int idx = 0; idx < count; ++idx) { Dtype gradient = 0; if (mask[idx] == i) { From 9ba949e3555e7bfcb35f32e8997f257d06d34c49 Mon Sep 17 00:00:00 2001 From: to3i Date: Mon, 14 Jul 2014 14:16:00 +0200 Subject: [PATCH 5/5] Merge changes from rebase --- src/caffe/test/test_eltwise_layer.cpp | 165 +++----------------------- 1 file changed, 16 insertions(+), 149 deletions(-) diff --git a/src/caffe/test/test_eltwise_layer.cpp b/src/caffe/test/test_eltwise_layer.cpp index fa4761c9844..abc309ca9f9 100644 --- a/src/caffe/test/test_eltwise_layer.cpp +++ b/src/caffe/test/test_eltwise_layer.cpp @@ -128,173 +128,40 @@ TYPED_TEST(EltwiseLayerTest, TestSumCoeff) { } } -<<<<<<< HEAD -TYPED_TEST(EltwiseLayerTest, TestProdGradient) { +TYPED_TEST(EltwiseLayerTest, TestMax) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); - eltwise_param->set_operation(EltwiseParameter_EltwiseOp_PROD); - EltwiseLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-3); -======= -TYPED_TEST(EltwiseLayerTest, TestMaxCPU) { - Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; - EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); eltwise_param->set_operation(EltwiseParameter_EltwiseOp_MAX); - shared_ptr > layer( - new EltwiseLayer(layer_param)); + shared_ptr > layer( + new EltwiseLayer(layer_param)); layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - const TypeParam* data = this->blob_top_->cpu_data(); + const Dtype* data = this->blob_top_->cpu_data(); const int count = this->blob_top_->count(); - const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); - const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); - const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); + const Dtype* in_data_a = this->blob_bottom_a_->cpu_data(); + const Dtype* in_data_b = this->blob_bottom_b_->cpu_data(); + const Dtype* in_data_c = this->blob_bottom_c_->cpu_data(); for (int i = 0; i < count; ++i) { - EXPECT_GE(data[i], in_data_a[i]); + EXPECT_GE(data[i], in_data_a[i]); EXPECT_GE(data[i], in_data_b[i]); EXPECT_GE(data[i], in_data_c[i]); } } -TYPED_TEST(EltwiseLayerTest, TestProdGPU) { - Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; - EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); - eltwise_param->set_operation(EltwiseParameter_EltwiseOp_PROD); - shared_ptr > layer( - new EltwiseLayer(layer_param)); - layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - const TypeParam* data = this->blob_top_->cpu_data(); - const int count = this->blob_top_->count(); - const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); - const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); - const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); - for (int i = 0; i < count; ++i) { - EXPECT_EQ(data[i], in_data_a[i] * in_data_b[i] * in_data_c[i]); - } -} - -TYPED_TEST(EltwiseLayerTest, TestSumGPU) { - Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; - EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); - eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); - shared_ptr > layer( - new EltwiseLayer(layer_param)); - layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - const TypeParam* data = this->blob_top_->cpu_data(); - const int count = this->blob_top_->count(); - const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); - const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); - const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); - for (int i = 0; i < count; ++i) { - EXPECT_EQ(data[i], in_data_a[i] + in_data_b[i] + in_data_c[i]); - } -} - -TYPED_TEST(EltwiseLayerTest, TestSumCoeffGPU) { - Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; - EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); - eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); - eltwise_param->add_coeff(1); - eltwise_param->add_coeff(-0.5); - eltwise_param->add_coeff(2); - shared_ptr > layer( - new EltwiseLayer(layer_param)); - layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - const TypeParam* data = this->blob_top_->cpu_data(); - const int count = this->blob_top_->count(); - const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); - const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); - const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); - for (int i = 0; i < count; ++i) { - EXPECT_NEAR(data[i], in_data_a[i] - 0.5*in_data_b[i] + 2*in_data_c[i], - 1e-4); - } -} - -TYPED_TEST(EltwiseLayerTest, TestMaxGPU) { - Caffe::set_mode(Caffe::GPU); - LayerParameter layer_param; - EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); - eltwise_param->set_operation(EltwiseParameter_EltwiseOp_MAX); - shared_ptr > layer( - new EltwiseLayer(layer_param)); - layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_)); - layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); - const TypeParam* data = this->blob_top_->cpu_data(); - const int count = this->blob_top_->count(); - const TypeParam* in_data_a = this->blob_bottom_a_->cpu_data(); - const TypeParam* in_data_b = this->blob_bottom_b_->cpu_data(); - const TypeParam* in_data_c = this->blob_bottom_c_->cpu_data(); - for (int i = 0; i < count; ++i) { - EXPECT_GE(data[i], in_data_a[i]); - EXPECT_GE(data[i], in_data_b[i]); - EXPECT_GE(data[i], in_data_c[i]); - } -} - -TYPED_TEST(EltwiseLayerTest, TestProdCPUGradient) { - Caffe::set_mode(Caffe::CPU); +TYPED_TEST(EltwiseLayerTest, TestProdGradient) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); eltwise_param->set_operation(EltwiseParameter_EltwiseOp_PROD); - EltwiseLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-3); - checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); -} - -TYPED_TEST(EltwiseLayerTest, TestSumCPUGradient) { - Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; - EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); - eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); - EltwiseLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-3); - checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); -} - -TYPED_TEST(EltwiseLayerTest, TestSumCoeffCPUGradient) { - Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; - EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); - eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); - eltwise_param->add_coeff(1); - eltwise_param->add_coeff(-0.5); - eltwise_param->add_coeff(2); - EltwiseLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-3); ->>>>>>> Implemented elementwise max layer -> GPU test OK, CPU test fails + EltwiseLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), &(this->blob_top_vec_)); } -<<<<<<< HEAD TYPED_TEST(EltwiseLayerTest, TestSumGradient) { typedef typename TypeParam::Dtype Dtype; -======= -TYPED_TEST(EltwiseLayerTest, TestMaxCPUGradient) { - Caffe::set_mode(Caffe::CPU); - LayerParameter layer_param; - EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); - eltwise_param->set_operation(EltwiseParameter_EltwiseOp_MAX); - EltwiseLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-3); - checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), - &(this->blob_top_vec_)); -} - -TYPED_TEST(EltwiseLayerTest, TestSumGPUGradient) { - Caffe::set_mode(Caffe::GPU); ->>>>>>> changed gpu to cpu_data() LayerParameter layer_param; EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); eltwise_param->set_operation(EltwiseParameter_EltwiseOp_SUM); @@ -318,13 +185,13 @@ TYPED_TEST(EltwiseLayerTest, TestSumCoeffGradient) { &(this->blob_top_vec_)); } -TYPED_TEST(EltwiseLayerTest, TestMaxGPUGradient) { - Caffe::set_mode(Caffe::GPU); +TYPED_TEST(EltwiseLayerTest, TestMaxGradient) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; EltwiseParameter* eltwise_param = layer_param.mutable_eltwise_param(); eltwise_param->set_operation(EltwiseParameter_EltwiseOp_MAX); - EltwiseLayer layer(layer_param); - GradientChecker checker(1e-2, 1e-2); + EltwiseLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); checker.CheckGradientEltwise(&layer, &(this->blob_bottom_vec_), &(this->blob_top_vec_)); }