summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--src/caffe/test/test_convolution_layer.cpp82
-rw-r--r--src/caffe/util/im2col.cu2
2 files changed, 66 insertions, 18 deletions
diff --git a/src/caffe/test/test_convolution_layer.cpp b/src/caffe/test/test_convolution_layer.cpp
index c8d79083..b08486e1 100644
--- a/src/caffe/test/test_convolution_layer.cpp
+++ b/src/caffe/test/test_convolution_layer.cpp
@@ -24,7 +24,7 @@ class ConvolutionLayerTest : public ::testing::Test {
: blob_bottom_(new Blob<Dtype>()),
blob_top_(new Blob<Dtype>()) {}
virtual void SetUp() {
- blob_bottom_->Reshape(2, 3, 6, 5);
+ blob_bottom_->Reshape(2, 3, 6, 4);
// fill the values
FillerParameter filler_param;
filler_param.set_value(1.);
@@ -57,7 +57,7 @@ TYPED_TEST(ConvolutionLayerTest, TestSetup) {
EXPECT_EQ(this->blob_top_->num(), 2);
EXPECT_EQ(this->blob_top_->channels(), 4);
EXPECT_EQ(this->blob_top_->height(), 2);
- EXPECT_EQ(this->blob_top_->width(), 2);
+ EXPECT_EQ(this->blob_top_->width(), 1);
// setting group should not change the shape
convolution_param->set_num_output(3);
convolution_param->set_group(3);
@@ -66,10 +66,10 @@ TYPED_TEST(ConvolutionLayerTest, TestSetup) {
EXPECT_EQ(this->blob_top_->num(), 2);
EXPECT_EQ(this->blob_top_->channels(), 3);
EXPECT_EQ(this->blob_top_->height(), 2);
- EXPECT_EQ(this->blob_top_->width(), 2);
+ EXPECT_EQ(this->blob_top_->width(), 1);
}
-TYPED_TEST(ConvolutionLayerTest, TestSimpleConvolution) {
+TYPED_TEST(ConvolutionLayerTest, TestCPUSimpleConvolution) {
// We will simply see if the convolution layer carries out averaging well.
FillerParameter filler_param;
filler_param.set_value(1.);
@@ -93,21 +93,39 @@ TYPED_TEST(ConvolutionLayerTest, TestSimpleConvolution) {
// After the convolution, the output should all have output values 27.1
const TypeParam* top_data = this->blob_top_->cpu_data();
for (int i = 0; i < this->blob_top_->count(); ++i) {
- EXPECT_GE(top_data[i], 27.1 - 1e-4);
- EXPECT_LE(top_data[i], 27.1 + 1e-4);
+ EXPECT_NEAR(top_data[i], 27.1, 1e-4);
}
- // Test GPU
+}
+
+TYPED_TEST(ConvolutionLayerTest, TestGPUSimpleConvolution) {
+ // We will simply see if the convolution layer carries out averaging well.
+ FillerParameter filler_param;
+ filler_param.set_value(1.);
+ ConstantFiller<TypeParam> filler(filler_param);
+ filler.Fill(this->blob_bottom_);
+ LayerParameter layer_param;
+ ConvolutionParameter* convolution_param =
+ layer_param.mutable_convolution_param();
+ convolution_param->set_kernel_size(3);
+ convolution_param->set_stride(2);
+ convolution_param->set_num_output(4);
+ convolution_param->mutable_weight_filler()->set_type("constant");
+ convolution_param->mutable_weight_filler()->set_value(1);
+ convolution_param->mutable_bias_filler()->set_type("constant");
+ convolution_param->mutable_bias_filler()->set_value(0.1);
+ shared_ptr<Layer<TypeParam> > layer(
+ new ConvolutionLayer<TypeParam>(layer_param));
+ layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
Caffe::set_mode(Caffe::GPU);
layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_));
// After the convolution, the output should all have output values 27.1
- top_data = this->blob_top_->cpu_data();
+ const TypeParam* top_data = this->blob_top_->cpu_data();
for (int i = 0; i < this->blob_top_->count(); ++i) {
- EXPECT_GE(top_data[i], 27.1 - 1e-4);
- EXPECT_LE(top_data[i], 27.1 + 1e-4);
+ EXPECT_NEAR(top_data[i], 27.1, 1e-4);
}
}
-TYPED_TEST(ConvolutionLayerTest, TestSimpleConvolutionGroup) {
+TYPED_TEST(ConvolutionLayerTest, TestCPUSimpleConvolutionGroup) {
// We will simply see if the convolution layer carries out averaging well.
FillerParameter filler_param;
filler_param.set_value(1.);
@@ -146,24 +164,54 @@ TYPED_TEST(ConvolutionLayerTest, TestSimpleConvolutionGroup) {
for (int h = 0; h < this->blob_top_->height(); ++h) {
for (int w = 0; w < this->blob_top_->width(); ++w) {
TypeParam data = top_data[this->blob_top_->offset(n, c, h, w)];
- EXPECT_GE(data, c * 9 + 0.1 - 1e-4);
- EXPECT_LE(data, c * 9 + 0.1 + 1e-4);
+ EXPECT_NEAR(data, c * 9 + 0.1, 1e-4);
+ }
+ }
+ }
+ }
+}
+
+
+TYPED_TEST(ConvolutionLayerTest, TestGPUSimpleConvolutionGroup) {
+ // We will simply see if the convolution layer carries out averaging well.
+ FillerParameter filler_param;
+ filler_param.set_value(1.);
+ ConstantFiller<TypeParam> filler(filler_param);
+ filler.Fill(this->blob_bottom_);
+ TypeParam* bottom_data = this->blob_bottom_->mutable_cpu_data();
+ for (int n = 0; n < this->blob_bottom_->num(); ++n) {
+ for (int c = 0; c < this->blob_bottom_->channels(); ++c) {
+ for (int h = 0; h < this->blob_bottom_->height(); ++h) {
+ for (int w = 0; w < this->blob_bottom_->width(); ++w) {
+ bottom_data[this->blob_bottom_->offset(n, c, h, w)] = c;
}
}
}
}
- // Test GPU
+ LayerParameter layer_param;
+ ConvolutionParameter* convolution_param =
+ layer_param.mutable_convolution_param();
+ convolution_param->set_kernel_size(3);
+ convolution_param->set_stride(2);
+ convolution_param->set_num_output(3);
+ convolution_param->set_group(3);
+ convolution_param->mutable_weight_filler()->set_type("constant");
+ convolution_param->mutable_weight_filler()->set_value(1);
+ convolution_param->mutable_bias_filler()->set_type("constant");
+ convolution_param->mutable_bias_filler()->set_value(0.1);
+ shared_ptr<Layer<TypeParam> > layer(
+ new ConvolutionLayer<TypeParam>(layer_param));
+ layer->SetUp(this->blob_bottom_vec_, &(this->blob_top_vec_));
Caffe::set_mode(Caffe::GPU);
layer->Forward(this->blob_bottom_vec_, &(this->blob_top_vec_));
// After the convolution, the output should all have output values 9.1
- top_data = this->blob_top_->cpu_data();
+ const TypeParam* top_data = this->blob_top_->cpu_data();
for (int n = 0; n < this->blob_top_->num(); ++n) {
for (int c = 0; c < this->blob_top_->channels(); ++c) {
for (int h = 0; h < this->blob_top_->height(); ++h) {
for (int w = 0; w < this->blob_top_->width(); ++w) {
TypeParam data = top_data[this->blob_top_->offset(n, c, h, w)];
- EXPECT_GE(data, c * 9 + 0.1 - 1e-4);
- EXPECT_LE(data, c * 9 + 0.1 + 1e-4);
+ EXPECT_NEAR(data, c * 9 + 0.1, 1e-4);
}
}
}
diff --git a/src/caffe/util/im2col.cu b/src/caffe/util/im2col.cu
index 28325fdb..6aecb0e5 100644
--- a/src/caffe/util/im2col.cu
+++ b/src/caffe/util/im2col.cu
@@ -29,7 +29,7 @@ __global__ void im2col_gpu_kernel(const int n, const Dtype* data_im,
for (int j = 0; j < ksize; ++j) {
int h = h_in + i;
int w = w_in + j;
- *data_col = (h >= 0 && w >= 0 && h < width && w < height) ?
+ *data_col = (h >= 0 && w >= 0 && h < height && w < width) ?
data_im[i * width + j] : 0;
data_col += height_col * width_col;
}