From a1d97e0d6de696108762e46eb3719d2f69f688d5 Mon Sep 17 00:00:00 2001 From: Jeff Donahue Date: Sat, 24 May 2014 15:29:35 -0700 Subject: [PATCH] lint and make compilable (using static_cast's found a couple bugs at compile time) --- src/caffe/layers/pooling_layer.cpp | 20 +++++++++------- src/caffe/layers/pooling_layer.cu | 15 +++++------- src/caffe/test/test_maxpool_dropout_layers.cpp | 32 ++++++++++++++------------ 3 files changed, 35 insertions(+), 32 deletions(-) diff --git a/src/caffe/layers/pooling_layer.cpp b/src/caffe/layers/pooling_layer.cpp index 97506cc..5c740de 100644 --- a/src/caffe/layers/pooling_layer.cpp +++ b/src/caffe/layers/pooling_layer.cpp @@ -38,7 +38,8 @@ void PoolingLayer::SetUp(const vector*>& bottom, (*top)[0]->Reshape(bottom[0]->num(), channels_, pooled_height_, pooled_width_); // If max pooling, we will initialize the vector index part. - if (this->layer_param_.pooling_param().pool() == PoolingParameter_PoolMethod_MAX) { + if (this->layer_param_.pooling_param().pool() == + PoolingParameter_PoolMethod_MAX) { max_idx_.reset(new SyncedMemory((*top)[0]->count() * sizeof(int))); } // If stochastic pooling, we will initialize the random index part. @@ -63,7 +64,7 @@ Dtype PoolingLayer::Forward_cpu(const vector*>& bottom, switch (this->layer_param_.pooling_param().pool()) { case PoolingParameter_PoolMethod_MAX: // Initialize - mask = (int*)max_idx_->mutable_cpu_data(); + mask = static_cast(max_idx_->mutable_cpu_data()); for (int i = 0; i < top_count; ++i) { top_data[i] = -FLT_MAX; mask[i] = -1; @@ -77,11 +78,13 @@ Dtype PoolingLayer::Forward_cpu(const vector*>& bottom, int wstart = pw * stride_; int hend = min(hstart + kernel_size_, height_); int wend = min(wstart + kernel_size_, width_); + const int pool_index = ph * pooled_width_ + pw; for (int h = hstart; h < hend; ++h) { for (int w = wstart; w < wend; ++w) { - if (bottom_data[h * width_ + w] > top_data[ph * pooled_width_ + pw]) { - top_data[ph * pooled_width_ + pw] = bottom_data[h * width_ + w]; - mask[ph * pooled_width_ + pw] = h * width_ + w; + const int index = h * width_ + w; + if (bottom_data[index] > top_data[pool_index]) { + top_data[pool_index] = bottom_data[index]; + mask[pool_index] = index; } } } @@ -147,16 +150,17 @@ void PoolingLayer::Backward_cpu(const vector*>& top, // Different pooling methods. We explicitly do the switch outside the for // loop to save time, although this results in more codes. memset(bottom_diff, 0, (*bottom)[0]->count() * sizeof(Dtype)); - int* mask; + const int* mask; switch (this->layer_param_.pooling_param().pool()) { case PoolingParameter_PoolMethod_MAX: // The main loop - mask = (int*)max_idx_->cpu_data(); + mask = static_cast(max_idx_->cpu_data()); for (int n = 0; n < top[0]->num(); ++n) { for (int c = 0; c < channels_; ++c) { for (int ph = 0; ph < pooled_height_; ++ph) { for (int pw = 0; pw < pooled_width_; ++pw) { - bottom_diff[mask[ph * pooled_width_ + pw]]+=top_diff[ph * pooled_width_ + pw]; + const int index = ph * pooled_width_ + pw; + bottom_diff[mask[index]] += top_diff[index]; } } bottom_diff += (*bottom)[0]->offset(0, 1); diff --git a/src/caffe/layers/pooling_layer.cu b/src/caffe/layers/pooling_layer.cu index 3bdb562..60c52e3 100644 --- a/src/caffe/layers/pooling_layer.cu +++ b/src/caffe/layers/pooling_layer.cu @@ -12,7 +12,7 @@ using std::max; using std::min; namespace caffe { - + template __global__ void MaxPoolForward(const int nthreads, const Dtype* bottom_data, const int num, const int channels, const int height, @@ -36,7 +36,6 @@ __global__ void MaxPoolForward(const int nthreads, const Dtype* bottom_data, maxidx = h * width + w; maxval = bottom_data[maxidx]; } - } } top_data[index] = maxval; @@ -154,8 +153,8 @@ Dtype PoolingLayer::Forward_gpu(const vector*>& bottom, int* mask; switch (this->layer_param_.pooling_param().pool()) { case PoolingParameter_PoolMethod_MAX: + mask = static_cast(max_idx_->mutable_gpu_data()); // NOLINT_NEXT_LINE(whitespace/operators) - mask = (int*)max_idx_->mutable_gpu_data(); MaxPoolForward<<>>( count, bottom_data, bottom[0]->num(), channels_, height_, width_, pooled_height_, pooled_width_, kernel_size_, stride_, @@ -215,7 +214,6 @@ __global__ void MaxPoolBackward(const int nthreads, const Dtype* top_diff, Dtype gradient = 0; top_diff += (n * channels + c) * pooled_height * pooled_width; mask += (n * channels + c) * pooled_height * pooled_width; - //bottom_diff[index] += top_diff[mask[index]]; for (int ph = phstart; ph < phend; ++ph) { for (int pw = pwstart; pw < pwend; ++pw) { if (mask[ph * pooled_width + pw] == h * width + w) @@ -223,7 +221,7 @@ __global__ void MaxPoolBackward(const int nthreads, const Dtype* top_diff, } } bottom_diff[index] = gradient; - } + } } /* @@ -354,14 +352,13 @@ void PoolingLayer::Backward_gpu(const vector*>& top, int count = (*bottom)[0]->count(); CUDA_CHECK(cudaMemset(bottom_diff, 0, sizeof(Dtype) * count)); int* mask; - switch (this->layer_param_.pooling_param().pool()) { case PoolingParameter_PoolMethod_MAX: - mask = (int*)max_idx_->gpu_data(); + mask = static_cast(max_idx_->mutable_gpu_data()); // Since we have the mask we only need count top_diff - count = top[0]->count(); + count = top[0]->count(); + caffe_gpu_set(count, Dtype(0.), bottom_diff); // NOLINT_NEXT_LINE(whitespace/operators) - caffe_gpu_set(count,Dtype(0.),bottom_diff); MaxPoolBackward<<>>( count, top_diff, top[0]->num(), channels_, height_, width_, pooled_height_, pooled_width_, diff --git a/src/caffe/test/test_maxpool_dropout_layers.cpp b/src/caffe/test/test_maxpool_dropout_layers.cpp index 9d30263..3862e12 100644 --- a/src/caffe/test/test_maxpool_dropout_layers.cpp +++ b/src/caffe/test/test_maxpool_dropout_layers.cpp @@ -1,7 +1,9 @@ -// Copyright 2014 Sergio Guadarrama +// Copyright 2014 BVLC and contributors. #include -#include +#include + +#include "cuda_runtime.h" #include "gtest/gtest.h" #include "caffe/blob.hpp" @@ -19,7 +21,7 @@ class MaxPoolingDropoutTest : public ::testing::Test { protected: MaxPoolingDropoutTest() : blob_bottom_(new Blob()), - blob_top_(new Blob()) {}; + blob_top_(new Blob()) {} virtual void SetUp() { Caffe::set_random_seed(1703); blob_bottom_->Reshape(2, 3, 6, 5); @@ -30,7 +32,7 @@ class MaxPoolingDropoutTest : public ::testing::Test { filler.Fill(this->blob_bottom_); blob_bottom_vec_.push_back(blob_bottom_); blob_top_vec_.push_back(blob_top_); - }; + } virtual ~MaxPoolingDropoutTest() { delete blob_bottom_; delete blob_top_; } Blob* const blob_bottom_; Blob* const blob_top_; @@ -69,7 +71,7 @@ TYPED_TEST(MaxPoolingDropoutTest, CPUForward) { const TypeParam* top_data = this->blob_top_->cpu_data(); TypeParam sum = 0.; for (int i = 0; i < this->blob_top_->count(); ++i) { - sum += top_data[i]; + sum += top_data[i]; } EXPECT_EQ(sum, this->blob_top_->count()); // Dropout in-place @@ -80,7 +82,7 @@ TYPED_TEST(MaxPoolingDropoutTest, CPUForward) { TypeParam scale = 1. / (1. - layer_param.dropout_param().dropout_ratio()); top_data = this->blob_top_->cpu_data(); for (int i = 0; i < this->blob_top_->count(); ++i) { - sum += top_data[i]; + sum += top_data[i]; } EXPECT_GE(sum, 0); EXPECT_LE(sum, this->blob_top_->count()*scale); @@ -98,7 +100,7 @@ TYPED_TEST(MaxPoolingDropoutTest, GPUForward) { const TypeParam* top_data = this->blob_top_->cpu_data(); TypeParam sum = 0.; for (int i = 0; i < this->blob_top_->count(); ++i) { - sum += top_data[i]; + sum += top_data[i]; } EXPECT_EQ(sum, this->blob_top_->count()); @@ -109,7 +111,7 @@ TYPED_TEST(MaxPoolingDropoutTest, GPUForward) { TypeParam scale = 1. / (1. - layer_param.dropout_param().dropout_ratio()); top_data = this->blob_top_->cpu_data(); for (int i = 0; i < this->blob_top_->count(); ++i) { - sum += top_data[i]; + sum += top_data[i]; } EXPECT_GE(sum, 0); EXPECT_LE(sum, this->blob_top_->count()*scale); @@ -127,12 +129,12 @@ TYPED_TEST(MaxPoolingDropoutTest, CPUBackward) { layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); for (int i = 0; i < this->blob_top_->count(); ++i) { this->blob_top_->mutable_cpu_diff()[i] = 1.; - } + } layer.Backward(this->blob_top_vec_, true, &(this->blob_bottom_vec_)); const TypeParam* bottom_diff = this->blob_bottom_->cpu_diff(); TypeParam sum = 0.; for (int i = 0; i < this->blob_bottom_->count(); ++i) { - sum += bottom_diff[i]; + sum += bottom_diff[i]; } EXPECT_EQ(sum, this->blob_top_->count()); // Dropout in-place @@ -144,7 +146,7 @@ TYPED_TEST(MaxPoolingDropoutTest, CPUBackward) { TypeParam sum_with_dropout = 0.; bottom_diff = this->blob_bottom_->cpu_diff(); for (int i = 0; i < this->blob_bottom_->count(); ++i) { - sum_with_dropout += bottom_diff[i]; + sum_with_dropout += bottom_diff[i]; } EXPECT_GE(sum_with_dropout, sum); } @@ -161,12 +163,12 @@ TYPED_TEST(MaxPoolingDropoutTest, GPUBackward) { layer.Forward(this->blob_bottom_vec_, &(this->blob_top_vec_)); for (int i = 0; i < this->blob_top_->count(); ++i) { this->blob_top_->mutable_cpu_diff()[i] = 1.; - } + } layer.Backward(this->blob_top_vec_, true, &(this->blob_bottom_vec_)); const TypeParam* bottom_diff = this->blob_bottom_->cpu_diff(); TypeParam sum = 0.; for (int i = 0; i < this->blob_bottom_->count(); ++i) { - sum += bottom_diff[i]; + sum += bottom_diff[i]; } EXPECT_EQ(sum, this->blob_top_->count()); // Dropout in-place @@ -178,9 +180,9 @@ TYPED_TEST(MaxPoolingDropoutTest, GPUBackward) { TypeParam sum_with_dropout = 0.; bottom_diff = this->blob_bottom_->cpu_diff(); for (int i = 0; i < this->blob_bottom_->count(); ++i) { - sum_with_dropout += bottom_diff[i]; + sum_with_dropout += bottom_diff[i]; } EXPECT_GE(sum_with_dropout, sum); } -} \ No newline at end of file +} // namespace caffe -- 2.7.4