Skip to content

Commit 9a14ca9

Browse files
jerrywgzjerrywgz
committed
test=develop
1 parent 4c9884e commit 9a14ca9

File tree

5 files changed

+118
-136
lines changed

5 files changed

+118
-136
lines changed

‎paddle/fluid/API.spec

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -114,7 +114,7 @@ paddle.fluid.layers.pad ArgSpec(args=['x', 'paddings', 'pad_value', 'name'], var
114114
paddle.fluid.layers.pad_constant_like ArgSpec(args=['x', 'y', 'pad_value', 'name'], varargs=None, keywords=None, defaults=(0.0, None))
115115
paddle.fluid.layers.label_smooth ArgSpec(args=['label', 'prior_dist', 'epsilon', 'dtype', 'name'], varargs=None, keywords=None, defaults=(None, 0.1, 'float32', None))
116116
paddle.fluid.layers.roi_pool ArgSpec(args=['input', 'rois', 'pooled_height', 'pooled_width', 'spatial_scale'], varargs=None, keywords=None, defaults=(1, 1, 1.0))
117-
paddle.fluid.layers.roi_align ArgSpec(args=['input', 'rois', 'pooled_height', 'pooled_width', 'spatial_scale', 'sampling_ratio'], varargs=None, keywords=None, defaults=(1, 1, 1.0, -1))
117+
paddle.fluid.layers.roi_align ArgSpec(args=['input', 'rois', 'pooled_height', 'pooled_width', 'spatial_scale', 'sampling_ratio', 'name'], varargs=None, keywords=None, defaults=(1, 1, 1.0, -1, None))
118118
paddle.fluid.layers.dice_loss ArgSpec(args=['input', 'label', 'epsilon'], varargs=None, keywords=None, defaults=(1e-05,))
119119
paddle.fluid.layers.image_resize ArgSpec(args=['input', 'out_shape', 'scale', 'name', 'resample'], varargs=None, keywords=None, defaults=(None, None, None, 'BILINEAR'))
120120
paddle.fluid.layers.image_resize_short ArgSpec(args=['input', 'out_short_len', 'resample'], varargs=None, keywords=None, defaults=('BILINEAR',))

‎paddle/fluid/operators/roi_align_op.cc

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -94,7 +94,7 @@ class ROIAlignOpMaker : public framework::OpProtoAndCheckerMaker {
9494
void Make() override {
9595
AddInput("X",
9696
"(Tensor), "
97-
"the input of ROIAlignOp. "
97+
"The input of ROIAlignOp. "
9898
"The format of input tensor is NCHW. Where N is batch size, "
9999
"C is the number of input channels, "
100100
"H is the height of the feature, and "
@@ -104,7 +104,6 @@ class ROIAlignOpMaker : public framework::OpProtoAndCheckerMaker {
104104
"ROIs (Regions of Interest) to pool over. "
105105
"should be a 2-D LoDTensor of shape (num_rois, 4)"
106106
"given as [[x1, y1, x2, y2], …]. "
107-
"Where batch_id is the id of the data, "
108107
"(x1, y1) is the top left coordinates, and "
109108
"(x2, y2) is the bottom right coordinates.");
110109
AddOutput("Out",

‎paddle/fluid/operators/roi_align_op.cu

Lines changed: 47 additions & 54 deletions
Original file line numberDiff line numberDiff line change
@@ -34,17 +34,13 @@ static inline int NumBlocks(const int N) {
3434
i += blockDim.x * gridDim.x)
3535

3636
template <class T>
37-
__device__ T bilinear_interpolate(const T* input_data, const int height,
38-
const int width, T y, T x) {
37+
__device__ T BilinearInterpolate(const T* input_data, const int height,
38+
const int width, T y, T x) {
3939
if (y < -1.0 || y > height || x < -1.0 || x > width) {
4040
return 0;
4141
}
42-
if (y <= 0) {
43-
y = 0;
44-
}
45-
if (x <= 0) {
46-
x = 0;
47-
}
42+
y = y <= 0 ? 0 : y;
43+
x = x <= 0 ? 0 : x;
4844
int y_low = static_cast<int>(y);
4945
int x_low = static_cast<int>(x);
5046
int y_high;
@@ -75,20 +71,16 @@ __device__ T bilinear_interpolate(const T* input_data, const int height,
7571
}
7672

7773
template <class T>
78-
__device__ void bilinear_interpolate_gradient(const int height, const int width,
79-
T y, T x, T* w1, T* w2, T* w3,
80-
T* w4, int* x_low, int* x_high,
81-
int* y_low, int* y_high) {
74+
__device__ void BilinearInterpolateGradient(const int height, const int width,
75+
T y, T x, T* w1, T* w2, T* w3,
76+
T* w4, int* x_low, int* x_high,
77+
int* y_low, int* y_high) {
8278
if (y < -1.0 || y > height || x < -1.0 || x > width) {
8379
return;
8480
}
8581

86-
if (y <= 0) {
87-
y = 0;
88-
}
89-
if (x <= 0) {
90-
x = 0;
91-
}
82+
y = y <= 0 ? 0 : y;
83+
x = x <= 0 ? 0 : x;
9284
*y_low = static_cast<int>(y);
9385
*x_low = static_cast<int>(x);
9486
if (*y_low >= height - 1) {
@@ -153,7 +145,7 @@ __global__ void GPUROIAlignForward(
153145
const T x = roi_xmin + pw * bin_size_w +
154146
static_cast<T>(ix + .5f) * bin_size_w /
155147
static_cast<T>(roi_bin_grid_w);
156-
T val = bilinear_interpolate(offset_input_data, height, width, y, x);
148+
T val = BilinearInterpolate(offset_input_data, height, width, y, x);
157149
output_val += val;
158150
}
159151
}
@@ -213,8 +205,8 @@ __global__ void GPUROIAlignBackward(const int nthreads, const T* input_rois,
213205
static_cast<T>(roi_bin_grid_w);
214206
T w1 = 0, w2 = 0, w3 = 0, w4 = 0;
215207
int x_low = -1, x_high = -1, y_low = -1, y_high = -1;
216-
bilinear_interpolate_gradient(height, width, y, x, &w1, &w2, &w3, &w4,
217-
&x_low, &x_high, &y_low, &y_high);
208+
BilinearInterpolateGradient(height, width, y, x, &w1, &w2, &w3, &w4,
209+
&x_low, &x_high, &y_low, &y_high);
218210
T diff1 = out_grad_this_bin * w1 / count;
219211
T diff2 = out_grad_this_bin * w2 / count;
220212
T diff3 = out_grad_this_bin * w3 / count;
@@ -279,8 +271,8 @@ class GPUROIAlignOpKernel : public framework::OpKernel<T> {
279271
}
280272
}
281273
Tensor roi_batch_id_list_gpu;
282-
framework::TensorCopy(roi_batch_id_list, ctx.GetPlace(),
283-
ctx.device_context(), &roi_batch_id_list_gpu);
274+
framework::TensorCopySync(roi_batch_id_list, ctx.GetPlace(),
275+
&roi_batch_id_list_gpu);
284276
GPUROIAlignForward<
285277
T><<<blocks, threads, 0, ctx.cuda_device_context().stream()>>>(
286278
output_size, in->data<T>(), rois->data<T>(), spatial_scale, channels,
@@ -310,39 +302,40 @@ class GPUROIAlignGradOpKernel : public framework::OpKernel<T> {
310302
int height = in->dims()[2];
311303
int width = in->dims()[3];
312304

313-
if (in_grad) {
314-
Tensor roi_batch_id_list;
315-
roi_batch_id_list.Resize({rois_num});
316-
int* roi_batch_id_data =
317-
roi_batch_id_list.mutable_data<int>(platform::CPUPlace());
318-
auto rois_lod = rois->lod().back();
319-
int rois_batch_size = rois_lod.size() - 1;
320-
for (int n = 0; n < rois_batch_size; ++n) {
321-
for (size_t i = rois_lod[n]; i < rois_lod[n + 1]; ++i) {
322-
roi_batch_id_data[i] = n;
323-
}
324-
}
325-
Tensor roi_batch_id_list_gpu;
326-
framework::TensorCopy(roi_batch_id_list, ctx.GetPlace(),
327-
ctx.device_context(), &roi_batch_id_list_gpu);
328-
329-
in_grad->mutable_data<T>(ctx.GetPlace());
330-
math::SetConstant<Place, T> set_zero;
331-
set_zero(ctx.cuda_device_context(), in_grad, static_cast<T>(0));
332-
333-
int output_grad_size = out_grad->numel();
334-
int blocks = NumBlocks(output_grad_size);
335-
int threads = kNumCUDAThreads;
336-
337-
if (output_grad_size > 0) {
338-
GPUROIAlignBackward<
339-
T><<<blocks, threads, 0, ctx.cuda_device_context().stream()>>>(
340-
output_grad_size, rois->data<T>(), out_grad->data<T>(), rois_num,
341-
spatial_scale, channels, height, width, pooled_height, pooled_width,
342-
sampling_ratio, roi_batch_id_list_gpu.data<int>(),
343-
in_grad->mutable_data<T>(ctx.GetPlace()));
305+
if (!in_grad) {
306+
return;
307+
}
308+
Tensor roi_batch_id_list;
309+
roi_batch_id_list.Resize({rois_num});
310+
int* roi_batch_id_data =
311+
roi_batch_id_list.mutable_data<int>(platform::CPUPlace());
312+
auto rois_lod = rois->lod().back();
313+
int rois_batch_size = rois_lod.size() - 1;
314+
for (int n = 0; n < rois_batch_size; ++n) {
315+
for (size_t i = rois_lod[n]; i < rois_lod[n + 1]; ++i) {
316+
roi_batch_id_data[i] = n;
344317
}
345318
}
319+
Tensor roi_batch_id_list_gpu;
320+
framework::TensorCopySync(roi_batch_id_list, ctx.GetPlace(),
321+
&roi_batch_id_list_gpu);
322+
323+
in_grad->mutable_data<T>(ctx.GetPlace());
324+
math::SetConstant<Place, T> set_zero;
325+
set_zero(ctx.cuda_device_context(), in_grad, static_cast<T>(0));
326+
327+
int output_grad_size = out_grad->numel();
328+
int blocks = NumBlocks(output_grad_size);
329+
int threads = kNumCUDAThreads;
330+
331+
if (output_grad_size > 0) {
332+
GPUROIAlignBackward<
333+
T><<<blocks, threads, 0, ctx.cuda_device_context().stream()>>>(
334+
output_grad_size, rois->data<T>(), out_grad->data<T>(), rois_num,
335+
spatial_scale, channels, height, width, pooled_height, pooled_width,
336+
sampling_ratio, roi_batch_id_list_gpu.data<int>(),
337+
in_grad->mutable_data<T>(ctx.GetPlace()));
338+
}
346339
}
347340
};
348341

‎paddle/fluid/operators/roi_align_op.h

Lines changed: 67 additions & 78 deletions
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@ using LoDTensor = framework::LoDTensor;
2424
static constexpr int kROISize = 4;
2525

2626
template <class T>
27-
void pre_calc_for_bilinear_interpolate(
27+
void PreCalcForBilinearInterpolate(
2828
const platform::DeviceContext& ctx, const int height, const int width,
2929
const int pooled_height, const int pooled_width, const int iy_upper,
3030
const int ix_upper, T roi_ymin, T roi_xmin, T bin_size_h, T bin_size_w,
@@ -53,12 +53,8 @@ void pre_calc_for_bilinear_interpolate(
5353
pre_calc_index += 1;
5454
continue;
5555
}
56-
if (y <= 0) {
57-
y = 0;
58-
}
59-
if (x <= 0) {
60-
x = 0;
61-
}
56+
y = y <= 0 ? 0 : y;
57+
x = x <= 0 ? 0 : x;
6258

6359
int y_low = static_cast<int>(y);
6460
int x_low = static_cast<int>(x);
@@ -104,12 +100,8 @@ void bilinear_interpolate_gradient(const int height, const int width, T y, T x,
104100
x_low = x_high = y_low = y_high = -1;
105101
return;
106102
}
107-
if (y <= 0) {
108-
y = 0;
109-
}
110-
if (x <= 0) {
111-
x = 0;
112-
}
103+
y = y <= 0 ? 0 : y;
104+
x = x <= 0 ? 0 : x;
113105
y_low = static_cast<int>(y);
114106
x_low = static_cast<int>(x);
115107
if (y_low >= height - 1) {
@@ -139,7 +131,6 @@ void bilinear_interpolate_gradient(const int height, const int width, T y, T x,
139131
*(batch_grad_data + y_high * width + x_low) += diff3;
140132
*(batch_grad_data + y_high * width + x_high) += diff4;
141133
}
142-
return;
143134
}
144135

145136
template <typename DeviceContext, typename T>
@@ -214,7 +205,7 @@ class CPUROIAlignOpKernel : public framework::OpKernel<T> {
214205
pre_pos.Resize({pre_size, kROISize});
215206
pre_w.Resize({pre_size, kROISize});
216207

217-
pre_calc_for_bilinear_interpolate(
208+
PreCalcForBilinearInterpolate(
218209
dev_ctx, height, width, pooled_height, pooled_width, roi_bin_grid_h,
219210
roi_bin_grid_w, roi_ymin, roi_xmin, bin_size_h, bin_size_w,
220211
roi_bin_grid_h, roi_bin_grid_w, &pre_pos, &pre_w);
@@ -245,7 +236,6 @@ class CPUROIAlignOpKernel : public framework::OpKernel<T> {
245236
}
246237
rois_data += roi_stride[0];
247238
}
248-
return;
249239
}
250240
};
251241

@@ -264,79 +254,78 @@ class CPUROIAlignGradOpKernel : public framework::OpKernel<T> {
264254
auto spatial_scale = ctx.Attr<float>("spatial_scale");
265255
auto sampling_ratio = ctx.Attr<int>("sampling_ratio");
266256
auto in_dims = in->dims();
267-
if (in_grad) {
268-
int channels = in_dims[1];
269-
int height = in_dims[2];
270-
int width = in_dims[3];
271-
int rois_num = rois->dims()[0];
272-
Tensor roi_batch_id_list;
273-
roi_batch_id_list.Resize({rois_num});
274-
int* roi_batch_id_data =
275-
roi_batch_id_list.mutable_data<int>(ctx.GetPlace());
257+
if (!in_grad) {
258+
return;
259+
}
260+
int channels = in_dims[1];
261+
int height = in_dims[2];
262+
int width = in_dims[3];
263+
int rois_num = rois->dims()[0];
264+
Tensor roi_batch_id_list;
265+
roi_batch_id_list.Resize({rois_num});
266+
int* roi_batch_id_data =
267+
roi_batch_id_list.mutable_data<int>(ctx.GetPlace());
276268

277-
auto rois_lod = rois->lod().back();
278-
int rois_batch_size = rois_lod.size() - 1;
279-
for (int n = 0; n < rois_batch_size; ++n) {
280-
for (size_t i = rois_lod[n]; i < rois_lod[n + 1]; ++i) {
281-
roi_batch_id_data[i] = n;
282-
}
269+
auto rois_lod = rois->lod().back();
270+
int rois_batch_size = rois_lod.size() - 1;
271+
for (int n = 0; n < rois_batch_size; ++n) {
272+
for (size_t i = rois_lod[n]; i < rois_lod[n + 1]; ++i) {
273+
roi_batch_id_data[i] = n;
283274
}
275+
}
284276

285-
const T* rois_data = rois->data<T>();
286-
const T* out_grad_data = out_grad->data<T>();
287-
T* in_grad_data = in_grad->mutable_data<T>(ctx.GetPlace());
277+
const T* rois_data = rois->data<T>();
278+
const T* out_grad_data = out_grad->data<T>();
279+
T* in_grad_data = in_grad->mutable_data<T>(ctx.GetPlace());
288280

289-
auto in_stride = framework::stride(in->dims());
290-
auto roi_stride = framework::stride(rois->dims());
291-
auto out_stride = framework::stride(out_grad->dims());
281+
auto in_stride = framework::stride(in->dims());
282+
auto roi_stride = framework::stride(rois->dims());
283+
auto out_stride = framework::stride(out_grad->dims());
292284

293-
for (int n = 0; n < rois_num; ++n) {
294-
int roi_batch_idx = roi_batch_id_data[n];
295-
T roi_xmin = rois_data[0] * spatial_scale;
296-
T roi_ymin = rois_data[1] * spatial_scale;
297-
T roi_xmax = rois_data[2] * spatial_scale;
298-
T roi_ymax = rois_data[3] * spatial_scale;
299-
T roi_width = std::max(roi_xmax - roi_xmin, static_cast<T>(1.));
300-
T roi_height = std::max(roi_ymax - roi_ymin, static_cast<T>(1.));
301-
T bin_size_h =
302-
static_cast<T>(roi_height) / static_cast<T>(pooled_height);
303-
T bin_size_w = static_cast<T>(roi_width) / static_cast<T>(pooled_width);
304-
for (int c = 0; c < channels; ++c) {
305-
T* batch_grad_data =
306-
in_grad_data + roi_batch_idx * in_stride[0] + c * in_stride[1];
307-
const T* batch_out_grad_data =
308-
out_grad_data + n * out_stride[0] + c * out_stride[1];
309-
for (int ph = 0; ph < pooled_height; ++ph) {
310-
for (int pw = 0; pw < pooled_width; ++pw) {
311-
int pool_index = ph * pooled_width + pw;
312-
T out_grad_this_bin = batch_out_grad_data[pool_index];
313-
int roi_bin_grid_h = (sampling_ratio > 0)
314-
? sampling_ratio
315-
: ceil(roi_height / pooled_height);
316-
int roi_bin_grid_w = (sampling_ratio > 0)
317-
? sampling_ratio
318-
: ceil(roi_width / pooled_width);
319-
T count = roi_bin_grid_h * roi_bin_grid_w;
320-
for (int iy = 0; iy < roi_bin_grid_h; iy++) {
321-
const T y = roi_ymin + ph * bin_size_h +
322-
static_cast<T>(iy + .5f) * bin_size_h /
323-
static_cast<T>(roi_bin_grid_h);
324-
for (int ix = 0; ix < roi_bin_grid_w; ix++) {
325-
const T x = roi_xmin + pw * bin_size_w +
326-
static_cast<T>(ix + .5f) * bin_size_w /
327-
static_cast<T>(roi_bin_grid_w);
328-
bilinear_interpolate_gradient(height, width, y, x,
329-
out_grad_this_bin, count,
330-
batch_grad_data);
331-
}
285+
for (int n = 0; n < rois_num; ++n) {
286+
int roi_batch_idx = roi_batch_id_data[n];
287+
T roi_xmin = rois_data[0] * spatial_scale;
288+
T roi_ymin = rois_data[1] * spatial_scale;
289+
T roi_xmax = rois_data[2] * spatial_scale;
290+
T roi_ymax = rois_data[3] * spatial_scale;
291+
T roi_width = std::max(roi_xmax - roi_xmin, static_cast<T>(1.));
292+
T roi_height = std::max(roi_ymax - roi_ymin, static_cast<T>(1.));
293+
T bin_size_h = static_cast<T>(roi_height) / static_cast<T>(pooled_height);
294+
T bin_size_w = static_cast<T>(roi_width) / static_cast<T>(pooled_width);
295+
for (int c = 0; c < channels; ++c) {
296+
T* batch_grad_data =
297+
in_grad_data + roi_batch_idx * in_stride[0] + c * in_stride[1];
298+
const T* batch_out_grad_data =
299+
out_grad_data + n * out_stride[0] + c * out_stride[1];
300+
for (int ph = 0; ph < pooled_height; ++ph) {
301+
for (int pw = 0; pw < pooled_width; ++pw) {
302+
int pool_index = ph * pooled_width + pw;
303+
T out_grad_this_bin = batch_out_grad_data[pool_index];
304+
int roi_bin_grid_h = (sampling_ratio > 0)
305+
? sampling_ratio
306+
: ceil(roi_height / pooled_height);
307+
int roi_bin_grid_w = (sampling_ratio > 0)
308+
? sampling_ratio
309+
: ceil(roi_width / pooled_width);
310+
T count = roi_bin_grid_h * roi_bin_grid_w;
311+
for (int iy = 0; iy < roi_bin_grid_h; iy++) {
312+
const T y = roi_ymin + ph * bin_size_h +
313+
static_cast<T>(iy + .5f) * bin_size_h /
314+
static_cast<T>(roi_bin_grid_h);
315+
for (int ix = 0; ix < roi_bin_grid_w; ix++) {
316+
const T x = roi_xmin + pw * bin_size_w +
317+
static_cast<T>(ix + .5f) * bin_size_w /
318+
static_cast<T>(roi_bin_grid_w);
319+
bilinear_interpolate_gradient(height, width, y, x,
320+
out_grad_this_bin, count,
321+
batch_grad_data);
332322
}
333323
}
334324
}
335325
}
336-
rois_data += roi_stride[0];
337326
}
327+
rois_data += roi_stride[0];
338328
}
339-
return;
340329
}
341330
};
342331
} // namespace operators

‎python/paddle/fluid/layers/nn.py

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5184,7 +5184,8 @@ def roi_align(input,
51845184
pooled_height=1,
51855185
pooled_width=1,
51865186
spatial_scale=1.0,
5187-
sampling_ratio=-1):
5187+
sampling_ratio=-1,
5188+
name=None):
51885189
"""
51895190
${comment}
51905191

0 commit comments

Comments
 (0)