From f680f235838a1881f730b8ea2ea9594b77769601 Mon Sep 17 00:00:00 2001 From: Jonathan L Long Date: Sat, 27 Dec 2014 01:44:36 -0800 Subject: [PATCH] add CropLayer for cropping one blob to another using coordinates from paramters --- include/caffe/layers/crop_layer.hpp | 49 ++++++++++++++++++ src/caffe/layers/crop_layer.cpp | 78 +++++++++++++++++++++++++++++ src/caffe/layers/crop_layer.cu | 60 ++++++++++++++++++++++ src/caffe/proto/caffe.proto | 10 +++- 4 files changed, 196 insertions(+), 1 deletion(-) create mode 100644 include/caffe/layers/crop_layer.hpp create mode 100644 src/caffe/layers/crop_layer.cpp create mode 100644 src/caffe/layers/crop_layer.cu diff --git a/include/caffe/layers/crop_layer.hpp b/include/caffe/layers/crop_layer.hpp new file mode 100644 index 00000000000..bab290718e0 --- /dev/null +++ b/include/caffe/layers/crop_layer.hpp @@ -0,0 +1,49 @@ +#ifndef CAFFE_CROP_LAYER_HPP_ +#define CAFFE_CROP_LAYER_HPP_ + +#include +#include + +#include "caffe/blob.hpp" +#include "caffe/layer.hpp" +#include "caffe/proto/caffe.pb.h" + +namespace caffe { + +/** + * @brief Takes a Blob and crop it along either the width or height dimension, + * outputting a cropped Blob. + * + * TODO(dox): thorough documentation for Forward, Backward, and proto params. + */ + +template +class CropLayer : public Layer { + public: + explicit CropLayer(const LayerParameter& param) + : Layer(param) {} + virtual void LayerSetUp(const vector*>& bottom, + const vector*>& top); + virtual void Reshape(const vector*>& bottom, + const vector*>& top); + + virtual inline const char* type() const { return "Crop"; } + virtual inline int ExactNumBottomBlobs() const { return 2; } + virtual inline int ExactNumTopBlobs() const { return 1; } + + protected: + virtual void Forward_cpu(const vector*>& bottom, + const vector*>& top); + virtual void Backward_cpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom); + virtual void Forward_gpu(const vector*>& bottom, + const vector*>& top); + virtual void Backward_gpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom); + + int crop_h_, crop_w_; +}; + +} // namespace caffe + +#endif // CAFFE_CROP_LAYER_HPP_ diff --git a/src/caffe/layers/crop_layer.cpp b/src/caffe/layers/crop_layer.cpp new file mode 100644 index 00000000000..76409bd7408 --- /dev/null +++ b/src/caffe/layers/crop_layer.cpp @@ -0,0 +1,78 @@ +#include +#include +#include +#include + +#include "caffe/layer.hpp" +#include "caffe/layers/crop_layer.hpp" +#include "caffe/net.hpp" + + +namespace caffe { + +template +void CropLayer::LayerSetUp(const vector*>& bottom, + const vector*>& top) { + const CropParameter& param = this->layer_param_.crop_param(); + CHECK_EQ(bottom.size(), 2) << "Wrong number of bottom blobs."; + CHECK_EQ(bottom[0]->num_axes(), 4) << "Only works with 4D blobs."; + CHECK_EQ(bottom[1]->num_axes(), 4) << "Only works with 4D blobs."; + crop_h_ = param.offset_height(); + crop_w_ = param.offset_width(); +} + +template +void CropLayer::Reshape(const vector*>& bottom, + const vector*>& top) { + // Check that the image we are cropping minus the margin is bigger than the + // destination image. + CHECK_GT(bottom[0]->height()-crop_h_, bottom[1]->height()) + << "invalid offset"; + CHECK_GT(bottom[0]->width()-crop_w_, bottom[1]->width()) << "invalid offset"; + top[0]->Reshape(bottom[0]->num(), bottom[0]->channels(), bottom[1]->height(), + bottom[1]->width()); +} + +template +void CropLayer::Forward_cpu(const vector*>& bottom, + const vector*>& top) { + const Dtype* bottom_data = bottom[0]->cpu_data(); + Dtype* top_data = top[0]->mutable_cpu_data(); + for (int n = 0; n < top[0]->num(); ++n) { + for (int c = 0; c < top[0]->channels(); ++c) { + for (int h = 0; h < top[0]->height(); ++h) { + caffe_copy(top[0]->width(), + bottom_data + bottom[0]->offset(n, c, crop_h_ + h, crop_w_), + top_data + top[0]->offset(n, c, h)); + } + } + } +} + +template +void CropLayer::Backward_cpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom) { + const Dtype* top_diff = top[0]->cpu_diff(); + Dtype* bottom_diff = bottom[0]->mutable_cpu_diff(); + if (propagate_down[0]) { + caffe_set(bottom[0]->count(), static_cast(0), bottom_diff); + for (int n = 0; n < top[0]->num(); ++n) { + for (int c = 0; c < top[0]->channels(); ++c) { + for (int h = 0; h < top[0]->height(); ++h) { + caffe_copy(top[0]->width(), + top_diff + top[0]->offset(n, c, h), + bottom_diff + bottom[0]->offset(n, c, crop_h_ + h, crop_w_)); + } + } + } + } +} + +#ifdef CPU_ONLY +STUB_GPU(CropLayer); +#endif + +INSTANTIATE_CLASS(CropLayer); +REGISTER_LAYER_CLASS(Crop); + +} // namespace caffe diff --git a/src/caffe/layers/crop_layer.cu b/src/caffe/layers/crop_layer.cu new file mode 100644 index 00000000000..262f5fa8483 --- /dev/null +++ b/src/caffe/layers/crop_layer.cu @@ -0,0 +1,60 @@ +#include + +#include "caffe/layers/crop_layer.hpp" + +namespace caffe { + +// Copy (one line per thread) from one array to another, with arbitrary +// strides in the last two dimensions. +template +__global__ void copy_kernel(const int n, const int height, const int width, + const int src_outer_stride, const int src_inner_stride, + const int dest_outer_stride, const int dest_inner_stride, + const Dtype* src, Dtype* dest) { + CUDA_KERNEL_LOOP(index, n) { + int src_start = index / height * src_outer_stride + + index % height * src_inner_stride; + int dest_start = index / height * dest_outer_stride + + index % height * dest_inner_stride; + for (int i = 0; i < width; ++i) { + dest[dest_start + i] = src[src_start + i]; + } + } +} + +template +void CropLayer::Forward_gpu(const vector*>& bottom, + const vector*>& top) { + const Dtype* bottom_data = bottom[0]->gpu_data(); + Dtype* top_data = top[0]->mutable_gpu_data(); + const int lines = top[0]->count() / top[0]->width(); + + // NOLINT_NEXT_LINE(whitespace/operators) + copy_kernel<<>>( + lines, top[0]->height(), top[0]->width(), + bottom[0]->height() * bottom[0]->width(), bottom[0]->width(), + top[0]->height() * top[0]->width(), top[0]->width(), + bottom_data + bottom[0]->offset(0, 0, crop_h_, crop_w_), top_data); +} + +template +void CropLayer::Backward_gpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom) { + const Dtype* top_diff = top[0]->gpu_diff(); + Dtype* bottom_diff = bottom[0]->mutable_gpu_diff(); + const int lines = top[0]->count() / top[0]->width(); + + if (propagate_down[0]) { + caffe_gpu_set(bottom[0]->count(), static_cast(0), bottom_diff); + // NOLINT_NEXT_LINE(whitespace/operators) + copy_kernel<<>>( + lines, top[0]->height(), top[0]->width(), + top[0]->height() * top[0]->width(), top[0]->width(), + bottom[0]->height() * bottom[0]->width(), bottom[0]->width(), + top_diff, bottom_diff + bottom[0]->offset(0, 0, crop_h_, crop_w_)); + } +} + +INSTANTIATE_LAYER_GPU_FUNCS(CropLayer); + +} // namespace caffe diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index f873deba10c..fc27b00c827 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -306,7 +306,7 @@ message ParamSpec { // NOTE // Update the next available ID when you add a new LayerParameter field. // -// LayerParameter next available layer-specific ID: 141 (last added: elu_param) +// LayerParameter next available layer-specific ID: 142 (last added: crop_param) message LayerParameter { optional string name = 1; // the layer name optional string type = 2; // the layer type @@ -359,6 +359,7 @@ message LayerParameter { optional ConcatParameter concat_param = 104; optional ContrastiveLossParameter contrastive_loss_param = 105; optional ConvolutionParameter convolution_param = 106; + optional CropParameter crop_param = 141; optional DataParameter data_param = 107; optional DropoutParameter dropout_param = 108; optional DummyDataParameter dummy_data_param = 109; @@ -563,6 +564,13 @@ message ConvolutionParameter { optional bool force_nd_im2col = 17 [default = false]; } +message CropParameter { + // Assumes standard dimensions: ( N,C,H,W ) + // This could possibly be extended to use "optional BlobShape offsets" + optional uint32 offset_height = 1[default = 0]; + optional uint32 offset_width = 2[default = 0]; +} + message DataParameter { enum DB { LEVELDB = 0;