FazBrowse GitHub Viewer | Trending |
URL:
| Home
Tools: [Download Repo ZIP]   [Original HTTPS Page]

Extend Crop to N-D, changed CropParameter. · OAID/Caffe-HRT@952fd17 · GitHub

  1. OAID
  2. Caffe-HRT

Repository navigation

Commit 952fd17

Browse files
BlGene
committed
Extend Crop to N-D, changed CropParameter.
1 parent 64e78bd commit 952fd17

4 files changed

Lines changed: 210 additions & 52 deletions

File tree

‎include/caffe/layers/crop_layer.hpp‎

Lines changed: 20 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -41,9 +41,27 @@ class CropLayer : public Layer<Dtype> {
4141
virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
4242
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom);
4343

44-
int crop_h_, crop_w_;
45-
};
44+
vector<int> offsets;
45+
46+
private:
47+
void crop_copy(const vector<Blob<Dtype>*>& bottom,
48+
const vector<Blob<Dtype>*>& top,
49+
const vector<int>& offsets,
50+
vector<int> indices,
51+
int cur_dim,
52+
const Dtype* src_data,
53+
Dtype* dest_data,
54+
bool is_forward);
4655

56+
void crop_copy_gpu(const vector<Blob<Dtype>*>& bottom,
57+
const vector<Blob<Dtype>*>& top,
58+
const vector<int>& offsets,
59+
vector<int> indices,
60+
int cur_dim,
61+
const Dtype* src_data,
62+
Dtype* dest_data,
63+
bool is_forward);
64+
};
4765
} // namespace caffe
4866

4967
#endif // CAFFE_CROP_LAYER_HPP_

‎src/caffe/layers/crop_layer.cpp‎

Lines changed: 94 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,10 @@
11
#include <algorithm>
2+
#include <functional>
23
#include <map>
34
#include <set>
45
#include <vector>
56

7+
68
#include "caffe/layer.hpp"
79
#include "caffe/layers/crop_layer.hpp"
810
#include "caffe/net.hpp"
@@ -13,58 +15,120 @@ namespace caffe {
1315
template <typename Dtype>
1416
void CropLayer<Dtype>::LayerSetUp(const vector<Blob<Dtype>*>& bottom,
1517
const vector<Blob<Dtype>*>& top) {
16-
const CropParameter& param = this->layer_param_.crop_param();
1718
CHECK_EQ(bottom.size(), 2) << "Wrong number of bottom blobs.";
18-
CHECK_EQ(bottom[0]->num_axes(), 4) << "Only works with 4D blobs.";
19-
CHECK_EQ(bottom[1]->num_axes(), 4) << "Only works with 4D blobs.";
20-
crop_h_ = param.offset_height();
21-
crop_w_ = param.offset_width();
19+
// parameter setup moved to Reshape because it depends on size.
2220
}
2321

2422
template <typename Dtype>
2523
void CropLayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
2624
const vector<Blob<Dtype>*>& top) {
27-
// Check that the image we are cropping minus the margin is bigger than the
28-
// destination image.
29-
CHECK_GT(bottom[0]->height()-crop_h_, bottom[1]->height())
30-
<< "invalid offset";
31-
CHECK_GT(bottom[0]->width()-crop_w_, bottom[1]->width()) << "invalid offset";
32-
top[0]->Reshape(bottom[0]->num(), bottom[0]->channels(), bottom[1]->height(),
33-
bottom[1]->width());
25+
const CropParameter& param = this->layer_param_.crop_param();
26+
// bottom[0] supplies the data
27+
// bottom[1] supplies the size
28+
int input_dim = bottom[0]->num_axes();
29+
CHECK_LT(param.axis(), input_dim) << "crop axis bigger than input dim";
30+
// initialize all offsets to 0
31+
offsets = vector<int>(input_dim, 0);
32+
// initialize new shape to bottom[0]
33+
vector<int> new_shape(bottom[0]->shape());
34+
35+
if (param.offset_size() > 1) {
36+
// the number of crop values specified must be equal to the number
37+
// of dimensions following axis
38+
CHECK_EQ(param.axis() + param.offset_size(), input_dim)
39+
<< "number of crop values specified must be equal to the number of "
40+
<< "dimensions following axis.";
41+
}
42+
// apply crops
43+
for (int i = 0; i < input_dim; ++i) {
44+
int crop_offset = 0;
45+
int new_size = bottom[0]->shape(i);
46+
if (i >= param.axis() && param.offset_size() == 1) {
47+
// if only one crop value is supplied, crop all dimensions after axis
48+
// by this crop value
49+
crop_offset = param.offset(0);
50+
new_size = bottom[1]->shape(i);
51+
} else if (i >= param.axis() && param.offset_size() > 1) {
52+
// crop values specified must be equal to the number of dimensions
53+
// following axis
54+
crop_offset = param.offset(i - param.axis());
55+
new_size = bottom[1]->shape(i);
56+
}
57+
// Check that the image we are cropping minus the margin is bigger
58+
// than the destination image.
59+
CHECK_GE(bottom[0]->shape(i) - crop_offset,
60+
bottom[1]->shape(i))
61+
<< "invalid crop parameters in dimension: " << i;
62+
// Now set new size and offsets
63+
new_shape[i] = new_size;
64+
offsets[i] = crop_offset;
65+
}
66+
top[0]->Reshape(new_shape);
67+
}
68+
69+
// recursive copy function
70+
template <typename Dtype>
71+
void CropLayer<Dtype>::crop_copy(const vector<Blob<Dtype>*>& bottom,
72+
const vector<Blob<Dtype>*>& top,
73+
const vector<int>& offsets,
74+
vector<int> indices,
75+
int cur_dim,
76+
const Dtype* src_data,
77+
Dtype* dest_data,
78+
bool is_forward) {
79+
if (cur_dim + 1 < top[0]->num_axes()) {
80+
// We are not yet at the final dimension, call copy recursivley
81+
for (int i = 0; i < top[0]->shape(cur_dim); ++i) {
82+
indices[cur_dim] = i;
83+
crop_copy(bottom, top, offsets, indices, cur_dim+1,
84+
src_data, dest_data, is_forward);
85+
}
86+
} else {
87+
// We are at the last dimensions, which is stored continously in memory
88+
for (int i = 0; i < top[0]->shape(cur_dim); ++i) {
89+
// prepare index vector reduced(red) and with offsets(off)
90+
std::vector<int> ind_red(cur_dim, 0);
91+
std::vector<int> ind_off(cur_dim+1, 0);
92+
for (int j = 0; j < cur_dim; ++j) {
93+
ind_red[j] = indices[j];
94+
ind_off[j] = indices[j] + offsets[j];
95+
}
96+
ind_off[cur_dim] = offsets[cur_dim];
97+
// do the copy
98+
if (is_forward) {
99+
caffe_copy(top[0]->shape(cur_dim),
100+
src_data + bottom[0]->offset(ind_off),
101+
dest_data + top[0]->offset(ind_red));
102+
} else {
103+
// in the backwards pass the src_data is top_diff
104+
// and the dest_data is bottom_diff
105+
caffe_copy(top[0]->shape(cur_dim),
106+
src_data + top[0]->offset(ind_red),
107+
dest_data + bottom[0]->offset(ind_off));
108+
}
109+
}
110+
}
34111
}
35112

36113
template <typename Dtype>
37114
void CropLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
38115
const vector<Blob<Dtype>*>& top) {
116+
std::vector<int> indices(top[0]->num_axes(), 0);
39117
const Dtype* bottom_data = bottom[0]->cpu_data();
40118
Dtype* top_data = top[0]->mutable_cpu_data();
41-
for (int n = 0; n < top[0]->num(); ++n) {
42-
for (int c = 0; c < top[0]->channels(); ++c) {
43-
for (int h = 0; h < top[0]->height(); ++h) {
44-
caffe_copy(top[0]->width(),
45-
bottom_data + bottom[0]->offset(n, c, crop_h_ + h, crop_w_),
46-
top_data + top[0]->offset(n, c, h));
47-
}
48-
}
49-
}
119+
crop_copy(bottom, top, offsets, indices, 0, bottom_data, top_data, true);
50120
}
51121

52122
template <typename Dtype>
53123
void CropLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
54124
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom) {
55125
const Dtype* top_diff = top[0]->cpu_diff();
56126
Dtype* bottom_diff = bottom[0]->mutable_cpu_diff();
127+
57128
if (propagate_down[0]) {
58129
caffe_set(bottom[0]->count(), static_cast<Dtype>(0), bottom_diff);
59-
for (int n = 0; n < top[0]->num(); ++n) {
60-
for (int c = 0; c < top[0]->channels(); ++c) {
61-
for (int h = 0; h < top[0]->height(); ++h) {
62-
caffe_copy(top[0]->width(),
63-
top_diff + top[0]->offset(n, c, h),
64-
bottom_diff + bottom[0]->offset(n, c, crop_h_ + h, crop_w_));
65-
}
66-
}
67-
}
130+
std::vector<int> indices(top[0]->num_axes(), 0);
131+
crop_copy(bottom, top, offsets, indices, 0, top_diff, bottom_diff, false);
68132
}
69133
}
70134

‎src/caffe/layers/crop_layer.cu‎

Lines changed: 82 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -22,36 +22,103 @@ __global__ void copy_kernel(const int n, const int height, const int width,
2222
}
2323
}
2424

25+
// recursive copy function, this function is similar to crop_copy but loops
26+
// over all but the last two dimensions. It is implemented this way to allow
27+
// for ND cropping while still relying on a CUDA kernel for the innermost
28+
// two dimensions for performance reasons.
29+
// An alternative way to implement ND cropping relying more on the kernel
30+
// would require passing offsets to the kernel, which is a bit problematic
31+
// because it is of variable length. Since in the standard (N,C,W,H) case
32+
// N,C are usually not cropped a speedup could be achieved by not looping
33+
// the application of the copy_kernel around these dimensions.
34+
template <typename Dtype>
35+
void CropLayer<Dtype>::crop_copy_gpu(const vector<Blob<Dtype>*>& bottom,
36+
const vector<Blob<Dtype>*>& top,
37+
const vector<int>& offsets,
38+
vector<int> indices,
39+
int cur_dim,
40+
const Dtype* src_data,
41+
Dtype* dest_data,
42+
bool is_forward) {
43+
if (cur_dim + 2 < top[0]->num_axes()) {
44+
// We are not yet at the final dimension, call copy recursivley
45+
for (int i = 0; i < top[0]->shape(cur_dim); ++i) {
46+
indices[cur_dim] = i;
47+
crop_copy_gpu(bottom, top, offsets, indices, cur_dim+1,
48+
src_data, dest_data, is_forward);
49+
}
50+
} else {
51+
// We are at the last two dimensions, which are stored continously in memory
52+
// With (N,C,H,W)
53+
// (0,1,2,3) cur_dim -> H
54+
// cur_dim+1 -> W
55+
const int lines = top[0]->shape(cur_dim);
56+
const int height = top[0]->shape(cur_dim);
57+
const int width = top[0]->shape(cur_dim+1);
58+
std::vector<int> ind_off(cur_dim+2, 0);
59+
for (int j = 0; j < cur_dim; ++j) {
60+
ind_off[j] = indices[j] + offsets[j];
61+
}
62+
ind_off[cur_dim] = offsets[cur_dim];
63+
ind_off[cur_dim+1] = offsets[cur_dim+1];
64+
// Compute copy strides
65+
const int src_outer_stride =
66+
bottom[0]->shape(cur_dim)*bottom[0]->shape(cur_dim+1);
67+
const int src_inner_stride = bottom[0]->shape(cur_dim+1);
68+
const int dest_outer_stride =
69+
top[0]->shape(cur_dim)*top[0]->shape(cur_dim+1);
70+
const int dest_inner_stride = top[0]->shape(cur_dim+1);
71+
72+
if (is_forward) {
73+
const Dtype* bottom_data = bottom[0]->gpu_data() +
74+
bottom[0]->offset(ind_off);
75+
Dtype* top_data = top[0]->mutable_gpu_data() +
76+
top[0]->offset(indices);
77+
// NOLINT_NEXT_LINE(whitespace/operators)
78+
copy_kernel<<<CAFFE_GET_BLOCKS(lines), CAFFE_CUDA_NUM_THREADS>>>(
79+
lines, height, width,
80+
src_outer_stride, src_inner_stride,
81+
dest_outer_stride, dest_inner_stride,
82+
bottom_data, top_data);
83+
84+
} else {
85+
const Dtype* top_diff = top[0]->gpu_diff() +
86+
top[0]->offset(indices);
87+
Dtype* bottom_diff = bottom[0]->mutable_gpu_diff() +
88+
bottom[0]->offset(ind_off);
89+
// NOLINT_NEXT_LINE(whitespace/operators)
90+
copy_kernel<<<CAFFE_GET_BLOCKS(lines), CAFFE_CUDA_NUM_THREADS>>>(
91+
lines, height, width,
92+
dest_outer_stride, dest_inner_stride,
93+
src_outer_stride, src_inner_stride,
94+
top_diff, bottom_diff);
95+
}
96+
}
97+
}
98+
2599
template <typename Dtype>
26100
void CropLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
27101
const vector<Blob<Dtype>*>& top) {
102+
std::vector<int> indices(top[0]->num_axes(), 0);
103+
// This works because crop_copy uses caffe_copy which calls cudaMemcpy.
104+
// My intuition is that calling this thousands of times is probably less
105+
// efficient than writing a custom kernel.
28106
const Dtype* bottom_data = bottom[0]->gpu_data();
29107
Dtype* top_data = top[0]->mutable_gpu_data();
30-
const int lines = top[0]->count() / top[0]->width();
31-
32-
// NOLINT_NEXT_LINE(whitespace/operators)
33-
copy_kernel<<<CAFFE_GET_BLOCKS(lines), CAFFE_CUDA_NUM_THREADS>>>(
34-
lines, top[0]->height(), top[0]->width(),
35-
bottom[0]->height() * bottom[0]->width(), bottom[0]->width(),
36-
top[0]->height() * top[0]->width(), top[0]->width(),
37-
bottom_data + bottom[0]->offset(0, 0, crop_h_, crop_w_), top_data);
108+
crop_copy_gpu(bottom, top, offsets, indices, 0, bottom_data, top_data, true);
38109
}
39110

40111
template <typename Dtype>
41112
void CropLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
42113
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom) {
43114
const Dtype* top_diff = top[0]->gpu_diff();
44115
Dtype* bottom_diff = bottom[0]->mutable_gpu_diff();
45-
const int lines = top[0]->count() / top[0]->width();
46116

47117
if (propagate_down[0]) {
48118
caffe_gpu_set(bottom[0]->count(), static_cast<Dtype>(0), bottom_diff);
49-
// NOLINT_NEXT_LINE(whitespace/operators)
50-
copy_kernel<<<CAFFE_GET_BLOCKS(lines), CAFFE_CUDA_NUM_THREADS>>>(
51-
lines, top[0]->height(), top[0]->width(),
52-
top[0]->height() * top[0]->width(), top[0]->width(),
53-
bottom[0]->height() * bottom[0]->width(), bottom[0]->width(),
54-
top_diff, bottom_diff + bottom[0]->offset(0, 0, crop_h_, crop_w_));
119+
std::vector<int> indices(top[0]->num_axes(), 0);
120+
crop_copy_gpu(bottom, top, offsets, indices, 0, top_diff, bottom_diff,
121+
false);
55122
}
56123
}
57124

‎src/caffe/proto/caffe.proto‎

Lines changed: 14 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -600,10 +600,19 @@ message ConvolutionParameter {
600600
}
601601

602602
message CropParameter {
603-
// Assumes standard dimensions: ( N,C,H,W )
604-
// This could possibly be extended to use "optional BlobShape offsets"
605-
optional uint32 offset_height = 1[default = 0];
606-
optional uint32 offset_width = 2[default = 0];
603+
// To crop, elements of the first bottom are selected to fit the dimensions
604+
// of the second, reference bottom. The crop is configured by
605+
// - the crop `axis` to pick the dimensions for cropping
606+
// - the crop `offset` to set the shift for all/each dimension
607+
// to align the cropped bottom with the reference bottom.
608+
// All dimensions up to but excluding `axis` are preserved, while
609+
// the dimensions including and trailing `axis` are cropped.
610+
// If only one `offset` is set, then all dimensions are offset by this amount.
611+
// Otherwise, the number of offsets must equal the number of cropped axes to
612+
// shift the crop in each dimension accordingly.
613+
// Note: standard dimensions are N,C,H,W so the default is a spatial crop.
614+
optional uint32 axis = 1 [default = 2];
615+
repeated uint32 offset = 2;
607616
}
608617

609618
message DataParameter {
@@ -680,7 +689,7 @@ message EltwiseParameter {
680689
// Message that stores parameters used by ELULayer
681690
message ELUParameter {
682691
// Described in:
683-
// Clevert, D.-A., Unterthiner, T., & Hochreiter, S. (2015). Fast and Accurate
692+
// Clevert, D.-A., Unterthiner, T., & Hochreiter, S. (2015). Fast and Accurate
684693
// Deep Network Learning by Exponential Linear Units (ELUs). arXiv
685694
optional float alpha = 1 [default = 1];
686695
}

0 commit comments

Comments
 (0)
Back | FazBrowse Home | New Git URL