Skip to content

Bugs fixed in the euclidean loss layer. #137

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
aravindhm wants to merge 4 commits into from
Closed

Bugs fixed in the euclidean loss layer. #137

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
53130b3
Bugs fixed in the euclidean loss layer.
Feb 20, 2014
94dbb0a
Forgot to add the definition for caffe_gpu_sub()
Feb 20, 2014
f4bae0b
Changed loss function to 1/2 ||x-y||^2 and updated the test cases to …
Feb 25, 2014
88f8597
Fixed the bug with EXPECT_NEAR argument.
Feb 26, 2014
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions include/caffe/util/math_functions.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -78,6 +78,9 @@ void caffe_mul(const int N, const Dtype* a, const Dtype* b, Dtype* y);
template <typename Dtype>
void caffe_gpu_mul(const int N, const Dtype* a, const Dtype* b, Dtype* y);

template <typename Dtype>
void caffe_gpu_sub(int N, const Dtype* a, const Dtype* b, Dtype* y);

template <typename Dtype>
void caffe_div(const int N, const Dtype* a, const Dtype* b, Dtype* y);

Expand Down
7 changes: 4 additions & 3 deletions include/caffe/vision_layers.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -467,22 +467,23 @@ template <typename Dtype>
class EuclideanLossLayer : public Layer<Dtype> {
public:
explicit EuclideanLossLayer(const LayerParameter& param)
: Layer<Dtype>(param), difference_() {}
: Layer<Dtype>(param), difference_(), scale_(param.scale()) {}
virtual void SetUp(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top);

protected:
// The loss layer will do nothing during forward - all computation are
// carried out in the backward pass.
virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) { return; }
vector<Blob<Dtype>*>* top);
virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* top) { return; }
vector<Blob<Dtype>*>* top);
virtual Dtype Backward_cpu(const vector<Blob<Dtype>*>& top,
const bool propagate_down, vector<Blob<Dtype>*>* bottom);
// virtual Dtype Backward_gpu(const vector<Blob<Dtype>*>& top,
// const bool propagate_down, vector<Blob<Dtype>*>* bottom);
Blob<Dtype> difference_;
Dtype scale_;
};


Expand Down
43 changes: 38 additions & 5 deletions src/caffe/layers/loss_layer.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -97,14 +97,45 @@ template <typename Dtype>
void EuclideanLossLayer<Dtype>::SetUp(
const vector<Blob<Dtype>*>& bottom, vector<Blob<Dtype>*>* 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 <typename Dtype>
void EuclideanLossLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* 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 / (2*num);
}
}

template <typename Dtype>
void EuclideanLossLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
vector<Blob<Dtype>*>* 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 / (2*num);
}
}

template <typename Dtype>
Expand All @@ -115,11 +146,13 @@ Dtype EuclideanLossLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& 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, 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 <typename Dtype>
Expand Down
64 changes: 62 additions & 2 deletions src/caffe/test/test_euclidean_loss_layer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,29 +22,89 @@ template <typename Dtype>
class EuclideanLossLayerTest : public ::testing::Test {
protected:
EuclideanLossLayerTest()
: blob_bottom_data_(new Blob<Dtype>(10, 5, 1, 1)),
blob_bottom_label_(new Blob<Dtype>(10, 5, 1, 1)) {
: blob_bottom_data_(new Blob<Dtype>(10, 5, 1, 10)),
blob_bottom_label_(new Blob<Dtype>(10, 5, 1, 10)),
blob_top_(new Blob<Dtype>()) {
// fill the values
FillerParameter filler_param;
GaussianFiller<Dtype> filler(filler_param);
filler.Fill(this->blob_bottom_data_);
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<Dtype>* const blob_bottom_data_;
Blob<Dtype>* const blob_bottom_label_;
Blob<Dtype>* blob_top_;
vector<Blob<Dtype>*> blob_bottom_vec_;
vector<Blob<Dtype>*> blob_top_vec_;
};

typedef ::testing::Types<float, double> Dtypes;
TYPED_TEST_CASE(EuclideanLossLayerTest, Dtypes);

TYPED_TEST(EuclideanLossLayerTest, TestSetUp) {
LayerParameter layer_param;
shared_ptr<EuclideanLossLayer<TypeParam> > layer(
new EuclideanLossLayer<TypeParam>(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<EuclideanLossLayer<TypeParam> > layer(
new EuclideanLossLayer<TypeParam>(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*2);
EXPECT_NEAR(this->blob_top_vec_[0]->data_at(0, 0, 0, 0), sum, 1e-4);
}

TYPED_TEST(EuclideanLossLayerTest, TestGPU) {
LayerParameter layer_param;
Caffe::set_mode(Caffe::GPU);
shared_ptr<EuclideanLossLayer<TypeParam> > layer(
new EuclideanLossLayer<TypeParam>(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*2);
EXPECT_NEAR(this->blob_top_vec_[0]->data_at(0, 0, 0, 0), sum, 1e-4);
}

TYPED_TEST(EuclideanLossLayerTest, TestGradientCPU) {
LayerParameter layer_param;
Caffe::set_mode(Caffe::CPU);
Expand Down
33 changes: 33 additions & 0 deletions src/caffe/util/math_functions.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,5 +32,38 @@ void caffe_gpu_mul<double>(const int N, const double* a,
N, a, b, y);
}

/* grid stride kernel */
template <typename Dtype>
__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<float>(const int N, const float* a,
const float* b, float* y) {
int deviceid;
cudaGetDevice(&deviceid);
int numSMs;
cudaDeviceGetAttribute(&numSMs, cudaDevAttrMultiProcessorCount, deviceid);
sub_kernel<float><<<numSMs, CAFFE_CUDA_NUM_THREADS>>>(
N, a, b, y);
}

template <>
void caffe_gpu_sub<double>(const int N, const double* a,
const double* b, double* y) {
int deviceid;
cudaGetDevice(&deviceid);
int numSMs;
cudaDeviceGetAttribute(&numSMs, cudaDevAttrMultiProcessorCount, deviceid);
sub_kernel<double><<<numSMs, CAFFE_CUDA_NUM_THREADS>>>(
N, a, b, y);
}

} // namespace caffe