yolo_box_op.cu 4.5 KB
Newer Older
D
dengkaipeng 已提交
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License. */

#include "paddle/fluid/operators/detection/yolo_box_op.h"
D
dengkaipeng 已提交
16
#include "paddle/fluid/operators/math/math_function.h"
D
dengkaipeng 已提交
17 18 19 20 21 22 23

namespace paddle {
namespace operators {

using Tensor = framework::Tensor;

template <typename T>
D
dengkaipeng 已提交
24
__global__ void KeYoloBoxFw(const T* input, const int* imgsize, T* boxes,
25 26
                            T* scores, const float conf_thresh, const int* anchors, 
			    const int n, const int h, const int w,
D
dengkaipeng 已提交
27
                            const int an_num, const int class_num,
D
dengkaipeng 已提交
28
                            const int box_num, int input_size) {
D
dengkaipeng 已提交
29 30
  int tid = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
D
dengkaipeng 已提交
31
  T box[4];
32
  for (; tid < n * box_num; tid += stride) {
D
dengkaipeng 已提交
33 34 35 36 37 38
    int grid_num = h * w;
    int i = tid / box_num;
    int j = (tid % box_num) / grid_num;
    int k = (tid % grid_num) / w;
    int l = tid % w;

39
    int an_stride = (5 + class_num) * grid_num;
D
dengkaipeng 已提交
40 41 42 43 44 45 46 47 48 49 50 51
    int img_height = imgsize[2 * i];
    int img_width = imgsize[2 * i + 1];

    int obj_idx =
        GetEntryIndex(i, j, k * w + l, an_num, an_stride, grid_num, 4);
    T conf = sigmoid<T>(input[obj_idx]);
    if (conf < conf_thresh) {
      continue;
    }

    int box_idx =
        GetEntryIndex(i, j, k * w + l, an_num, an_stride, grid_num, 0);
D
dengkaipeng 已提交
52
    GetYoloBox<T>(box, input, anchors, l, k, j, h, input_size, box_idx,
D
dengkaipeng 已提交
53 54
                                grid_num, img_height, img_width);
    box_idx = (i * box_num + j * grid_num + k * w + l) * 4;
D
dengkaipeng 已提交
55
    CalcDetectionBox<T>(boxes, box, box_idx, img_height, img_width);
D
dengkaipeng 已提交
56 57 58

    int label_idx =
        GetEntryIndex(i, j, k * w + l, an_num, an_stride, grid_num, 5);
59
    int score_idx = (i * box_num + j * grid_num + k * w + l) * class_num;
D
dengkaipeng 已提交
60 61 62
    CalcLabelScore<T>(scores, input, label_idx, score_idx, class_num, conf,
                      grid_num);
  }
D
dengkaipeng 已提交
63 64 65 66 67 68
}

template <typename T>
class YoloBoxOpCUDAKernel : public framework::OpKernel<T> {
 public:
  void Compute(const framework::ExecutionContext& ctx) const override {
D
dengkaipeng 已提交
69
    auto* input = ctx.Input<Tensor>("X");
D
dengkaipeng 已提交
70
    auto* img_size = ctx.Input<Tensor>("ImgSize");
D
dengkaipeng 已提交
71 72 73 74 75 76 77 78 79 80 81 82 83 84 85
    auto* boxes = ctx.Output<Tensor>("Boxes");
    auto* scores = ctx.Output<Tensor>("Scores");

    auto anchors = ctx.Attr<std::vector<int>>("anchors");
    int class_num = ctx.Attr<int>("class_num");
    float conf_thresh = ctx.Attr<float>("conf_thresh");
    int downsample_ratio = ctx.Attr<int>("downsample_ratio");

    const int n = input->dims()[0];
    const int h = input->dims()[2];
    const int w = input->dims()[3];
    const int box_num = boxes->dims()[1];
    const int an_num = anchors.size() / 2;
    int input_size = downsample_ratio * h;

D
dengkaipeng 已提交
86 87 88 89 90 91
    Tensor anchors_t, cpu_anchors_t;
    auto cpu_anchors_data = cpu_anchors_t.mutable_data<int>({an_num*2}, platform::CPUPlace());
    std::copy(anchors.begin(), anchors.end(), cpu_anchors_data);
    TensorCopySync(cpu_anchors_t, ctx.GetPlace(), &anchors_t);
    auto anchors_data = anchors_t.data<int>();

D
dengkaipeng 已提交
92
    const T* input_data = input->data<T>();
D
dengkaipeng 已提交
93
    const int* imgsize_data = img_size->data<int>();
D
dengkaipeng 已提交
94 95 96
    T* boxes_data = boxes->mutable_data<T>({n, box_num, 4}, ctx.GetPlace());
    T* scores_data =
        scores->mutable_data<T>({n, box_num, class_num}, ctx.GetPlace());
D
dengkaipeng 已提交
97 98 99 100
    math::SetConstant<platform::CUDADeviceContext, T> set_zero;
    auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
    set_zero(dev_ctx, boxes, static_cast<T>(0));
    set_zero(dev_ctx, scores, static_cast<T>(0));
D
dengkaipeng 已提交
101

102 103
    int grid_dim = (n * box_num + 512 - 1) / 512;
    grid_dim = grid_dim > 8 ? 8 : grid_dim;
D
dengkaipeng 已提交
104

105
    KeYoloBoxFw<T><<<grid_dim, 512, 0, ctx.cuda_device_context().stream()>>>(
D
dengkaipeng 已提交
106
	input_data, imgsize_data, boxes_data, scores_data, conf_thresh, 
107
	anchors_data, n, h, w, an_num, class_num, box_num, input_size);	
D
dengkaipeng 已提交
108
  }
D
dengkaipeng 已提交
109
};
D
dengkaipeng 已提交
110 111 112 113 114

}  // namespace operators
}  // namespace paddle

namespace ops = paddle::operators;
D
dengkaipeng 已提交
115 116 117
REGISTER_OP_CUDA_KERNEL(yolo_box,
                        ops::YoloBoxOpCUDAKernel<float>,
                        ops::YoloBoxOpCUDAKernel<double>);