Merge pull request BVLC#2083 from jeffdonahue/tile-layer

TileLayer

include/caffe/common_layers.hpp

@@ -644,6 +644,35 @@ class SliceLayer : public Layer<Dtype> {
   vector<int> slice_point_;
 };
 
+/**
+ * @brief Copy a Blob along specified dimensions.
+ */
+template <typename Dtype>
+class TileLayer : public Layer<Dtype> {
+ public:
+  explicit TileLayer(const LayerParameter& param)
+      : Layer<Dtype>(param) {}
+  virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
+      const vector<Blob<Dtype>*>& top);
+
+  virtual inline const char* type() const { return "Tile"; }
+  virtual inline int ExactNumBottomBlobs() const { return 1; }
+  virtual inline int ExactNumTopBlobs() const { return 1; }
+
+ protected:
+  virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
+      const vector<Blob<Dtype>*>& top);
+  virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
+      const vector<Blob<Dtype>*>& top);
+
+  virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
+      const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom);
+  virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
+      const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom);
+
+  unsigned int axis_, tiles_, outer_dim_, inner_dim_;
+};
+
 }  // namespace caffe
 
 #endif  // CAFFE_COMMON_LAYERS_HPP_

src/caffe/layers/tile_layer.cpp

@@ -0,0 +1,62 @@
+#include <vector>
+
+#include "caffe/common_layers.hpp"
+#include "caffe/layer.hpp"
+#include "caffe/util/math_functions.hpp"
+
+namespace caffe {
+
+template <typename Dtype>
+void TileLayer<Dtype>::Reshape(
+    const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
+  const TileParameter& tile_param = this->layer_param_.tile_param();
+  axis_ = bottom[0]->CanonicalAxisIndex(tile_param.axis());
+  CHECK(tile_param.has_tiles()) << "Number of tiles must be specified";
+  tiles_ = tile_param.tiles();
+  CHECK_GT(tiles_, 0) << "Number of tiles must be positive.";
+  vector<int> top_shape = bottom[0]->shape();
+  top_shape[axis_] = bottom[0]->shape(axis_) * tiles_;
+  top[0]->Reshape(top_shape);
+  outer_dim_ = bottom[0]->count(0, axis_);
+  inner_dim_ = bottom[0]->count(axis_);
+}
+
+template <typename Dtype>
+void TileLayer<Dtype>::Forward_cpu(
+    const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
+  const Dtype* bottom_data = bottom[0]->cpu_data();
+  Dtype* top_data = top[0]->mutable_cpu_data();
+  for (int i = 0; i < outer_dim_; ++i) {
+    for (int t = 0; t < tiles_; ++t) {
+      caffe_copy(inner_dim_, bottom_data, top_data);
+      top_data += inner_dim_;
+    }
+    bottom_data += inner_dim_;
+  }
+}
+
+template <typename Dtype>
+void TileLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
+    const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom) {
+  if (!propagate_down[0]) { return; }
+  const Dtype* top_diff = top[0]->cpu_diff();
+  Dtype* bottom_diff = bottom[0]->mutable_cpu_diff();
+  for (int i = 0; i < outer_dim_; ++i) {
+    caffe_copy(inner_dim_, top_diff, bottom_diff);
+    top_diff += inner_dim_;
+    for (int t = 1; t < tiles_; ++t) {
+      caffe_axpy(inner_dim_, Dtype(1), top_diff, bottom_diff);
+      top_diff += inner_dim_;
+    }
+    bottom_diff += inner_dim_;
+  }
+}
+
+#ifdef CPU_ONLY
+STUB_GPU(TileLayer);
+#endif
+
+INSTANTIATE_CLASS(TileLayer);
+REGISTER_LAYER_CLASS(Tile);
+
+}  // namespace caffe

src/caffe/layers/tile_layer.cu

@@ -0,0 +1,67 @@
+#include <vector>
+
+#include "caffe/common_layers.hpp"
+#include "caffe/layer.hpp"
+#include "caffe/util/math_functions.hpp"
+
+namespace caffe {
+
+template <typename Dtype>
+__global__ void Tile(const int nthreads, const Dtype* bottom_data,
+    const int tile_size, const int num_tiles, const int bottom_tile_axis,
+    Dtype* top_data) {
+  CUDA_KERNEL_LOOP(index, nthreads) {
+    const int d = index % tile_size;
+    const int b = (index / tile_size / num_tiles) % bottom_tile_axis;
+    const int n = index / tile_size / num_tiles / bottom_tile_axis;
+    const int bottom_index = (n * bottom_tile_axis + b) * tile_size + d;
+    top_data[index] = bottom_data[bottom_index];
+  }
+}
+
+template <typename Dtype>
+void TileLayer<Dtype>::Forward_gpu(
+    const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
+  const Dtype* bottom_data = bottom[0]->gpu_data();
+  Dtype* top_data = top[0]->mutable_gpu_data();
+  const int bottom_tile_axis = bottom[0]->shape(axis_);
+  const int nthreads = top[0]->count();
+  Tile<Dtype>  // NOLINT_NEXT_LINE(whitespace/operators)
+      <<<CAFFE_GET_BLOCKS(nthreads), CAFFE_CUDA_NUM_THREADS>>>(
+      nthreads, bottom_data, inner_dim_, tiles_, bottom_tile_axis, top_data);
+}
+
+template <typename Dtype>
+__global__ void TileBackward(const int nthreads, const Dtype* top_diff,
+    const int tile_size, const int num_tiles, const int bottom_tile_axis,
+    Dtype* bottom_diff) {
+  CUDA_KERNEL_LOOP(index, nthreads) {
+    const int d = index % tile_size;
+    const int b = (index / tile_size) % bottom_tile_axis;
+    const int n = index / tile_size / bottom_tile_axis;
+    bottom_diff[index] = 0;
+    int top_index = (n * num_tiles * bottom_tile_axis + b) * tile_size + d;
+    for (int t = 0; t < num_tiles; ++t) {
+      bottom_diff[index] += top_diff[top_index];
+      top_index += bottom_tile_axis * tile_size;
+    }
+  }
+}
+
+template <typename Dtype>
+void TileLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
+    const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom) {
+  if (!propagate_down[0]) { return; }
+  const Dtype* top_diff = top[0]->gpu_diff();
+  Dtype* bottom_diff = bottom[0]->mutable_gpu_diff();
+  const int bottom_tile_axis = bottom[0]->shape(axis_);
+  const int tile_size = inner_dim_ / bottom_tile_axis;
+  const int nthreads = bottom[0]->count();
+  TileBackward<Dtype>  // NOLINT_NEXT_LINE(whitespace/operators)
+      <<<CAFFE_GET_BLOCKS(nthreads), CAFFE_CUDA_NUM_THREADS>>>(
+      nthreads, top_diff, tile_size, tiles_, bottom_tile_axis, bottom_diff);
+}
+
+INSTANTIATE_LAYER_GPU_FUNCS(TileLayer);
+
+}  // namespace caffe

src/caffe/proto/caffe.proto

@@ -301,7 +301,7 @@ message ParamSpec {
 // NOTE
 // Update the next available ID when you add a new LayerParameter field.
 //
-// LayerParameter next available layer-specific ID: 138 (last added: embed_param)
+// LayerParameter next available layer-specific ID: 139 (last added: tile_param)
 message LayerParameter {
   optional string name = 1; // the layer name
   optional string type = 2; // the layer type
@@ -383,6 +383,7 @@ message LayerParameter {
   optional SliceParameter slice_param = 126;
   optional TanHParameter tanh_param = 127;
   optional ThresholdParameter threshold_param = 128;
+  optional TileParameter tile_param = 138;
   optional WindowDataParameter window_data_param = 129;
 }
 
@@ -919,6 +920,16 @@ message TanHParameter {
   optional Engine engine = 1 [default = DEFAULT];
 }
 
+// Message that stores parameters used by TileLayer
+message TileParameter {
+  // The index of the axis to tile.
+  optional int32 axis = 1 [default = 1];
+
+  // The number of copies (tiles) of the blob to output.
+  optional int32 tiles = 2;
+}
+
+// Message that stores parameters used by ThresholdLayer
 message ThresholdParameter {
   optional float threshold = 1 [default = 0]; // Strictly positive values
 }

src/caffe/test/test_tile_layer.cpp

@@ -0,0 +1,162 @@
+#include <cstring>
+#include <vector>
+
+#include "gtest/gtest.h"
+
+#include "caffe/blob.hpp"
+#include "caffe/common.hpp"
+#include "caffe/filler.hpp"
+#include "caffe/vision_layers.hpp"
+
+#include "caffe/test/test_caffe_main.hpp"
+#include "caffe/test/test_gradient_check_util.hpp"
+
+namespace caffe {
+
+template <typename TypeParam>
+class TileLayerTest : public MultiDeviceTest<TypeParam> {
+  typedef typename TypeParam::Dtype Dtype;
+
+ protected:
+  TileLayerTest()
+      : blob_bottom_(new Blob<Dtype>(2, 3, 4, 5)),
+        blob_top_(new Blob<Dtype>()) {}
+  virtual void SetUp() {
+    blob_bottom_vec_.push_back(blob_bottom_);
+    blob_top_vec_.push_back(blob_top_);
+    FillerParameter filler_param;
+    filler_param.set_mean(0.0);
+    filler_param.set_std(1.0);
+    GaussianFiller<Dtype> filler(filler_param);
+    filler.Fill(blob_bottom_);
+  }
+
+  virtual ~TileLayerTest() {
+    delete blob_bottom_;
+    delete blob_top_;
+  }
+
+  Blob<Dtype>* const blob_bottom_;
+  Blob<Dtype>* const blob_top_;
+  vector<Blob<Dtype>*> blob_bottom_vec_;
+  vector<Blob<Dtype>*> blob_top_vec_;
+};
+
+TYPED_TEST_CASE(TileLayerTest, TestDtypesAndDevices);
+
+TYPED_TEST(TileLayerTest, TestTrivialSetup) {
+  typedef typename TypeParam::Dtype Dtype;
+  LayerParameter layer_param;
+  const int kNumTiles = 1;
+  layer_param.mutable_tile_param()->set_tiles(kNumTiles);
+  for (int i = 0; i < this->blob_bottom_->num_axes(); ++i) {
+    layer_param.mutable_tile_param()->set_axis(i);
+    TileLayer<Dtype> layer(layer_param);
+    layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_);
+    ASSERT_EQ(this->blob_top_->num_axes(), this->blob_bottom_->num_axes());
+    for (int j = 0; j < this->blob_bottom_->num_axes(); ++j) {
+      EXPECT_EQ(this->blob_top_->shape(j), this->blob_bottom_->shape(j));
+    }
+  }
+}
+
+TYPED_TEST(TileLayerTest, TestSetup) {
+  typedef typename TypeParam::Dtype Dtype;
+  LayerParameter layer_param;
+  const int kNumTiles = 3;
+  layer_param.mutable_tile_param()->set_tiles(kNumTiles);
+  for (int i = 0; i < this->blob_bottom_->num_axes(); ++i) {
+    layer_param.mutable_tile_param()->set_axis(i);
+    TileLayer<Dtype> layer(layer_param);
+    layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_);
+    ASSERT_EQ(this->blob_top_->num_axes(), this->blob_bottom_->num_axes());
+    for (int j = 0; j < this->blob_bottom_->num_axes(); ++j) {
+      const int top_dim =
+          ((i == j) ? kNumTiles : 1) * this->blob_bottom_->shape(j);
+      EXPECT_EQ(top_dim, this->blob_top_->shape(j));
+    }
+  }
+}
+
+TYPED_TEST(TileLayerTest, TestForwardNum) {
+  typedef typename TypeParam::Dtype Dtype;
+  LayerParameter layer_param;
+  const int kTileAxis = 0;
+  const int kNumTiles = 3;
+  layer_param.mutable_tile_param()->set_axis(kTileAxis);
+  layer_param.mutable_tile_param()->set_tiles(kNumTiles);
+  TileLayer<Dtype> layer(layer_param);
+  layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_);
+  layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_);
+  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) {
+           const int bottom_n = n % this->blob_bottom_->num();
+           EXPECT_EQ(this->blob_bottom_->data_at(bottom_n, c, h, w),
+                     this->blob_top_->data_at(n, c, h, w));
+         }
+       }
+    }
+  }
+}
+
+TYPED_TEST(TileLayerTest, TestForwardChannels) {
+  typedef typename TypeParam::Dtype Dtype;
+  LayerParameter layer_param;
+  const int kNumTiles = 3;
+  layer_param.mutable_tile_param()->set_tiles(kNumTiles);
+  TileLayer<Dtype> layer(layer_param);
+  layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_);
+  layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_);
+  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) {
+           const int bottom_c = c % this->blob_bottom_->channels();
+           EXPECT_EQ(this->blob_bottom_->data_at(n, bottom_c, h, w),
+                     this->blob_top_->data_at(n, c, h, w));
+         }
+       }
+    }
+  }
+}
+
+TYPED_TEST(TileLayerTest, TestTrivialGradient) {
+  typedef typename TypeParam::Dtype Dtype;
+  LayerParameter layer_param;
+  const int kNumTiles = 1;
+  layer_param.mutable_tile_param()->set_tiles(kNumTiles);
+  TileLayer<Dtype> layer(layer_param);
+  GradientChecker<Dtype> checker(1e-2, 1e-2);
+  checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_,
+      this->blob_top_vec_);
+}
+
+TYPED_TEST(TileLayerTest, TestGradientNum) {
+  typedef typename TypeParam::Dtype Dtype;
+  LayerParameter layer_param;
+  const int kTileAxis = 0;
+  const int kNumTiles = 3;
+  layer_param.mutable_tile_param()->set_axis(kTileAxis);
+  layer_param.mutable_tile_param()->set_tiles(kNumTiles);
+  TileLayer<Dtype> layer(layer_param);
+  GradientChecker<Dtype> checker(1e-2, 1e-2);
+  checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_,
+      this->blob_top_vec_);
+}
+
+TYPED_TEST(TileLayerTest, TestGradientChannels) {
+  typedef typename TypeParam::Dtype Dtype;
+  LayerParameter layer_param;
+  const int kTileAxis = 1;
+  const int kNumTiles = 3;
+  layer_param.mutable_tile_param()->set_axis(kTileAxis);
+  layer_param.mutable_tile_param()->set_tiles(kNumTiles);
+  TileLayer<Dtype> layer(layer_param);
+  GradientChecker<Dtype> checker(1e-2, 1e-2);
+  checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_,
+      this->blob_top_vec_);
+}
+
+}  // namespace caffe