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

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,
D
dengkaipeng 已提交
25 26 27
                            T* scores, const float conf_thresh,
                            const int* anchors, const int n, const int h,
                            const int w, 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
                  grid_num, img_height, img_width);
D
dengkaipeng 已提交
54
    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
    auto& dev_ctx = ctx.cuda_device_context();
D
dengkaipeng 已提交
87 88
    auto& allocator =
        platform::DeviceTemporaryAllocator::Instance().Get(dev_ctx);
D
dengkaipeng 已提交
89 90 91 92 93 94
    int bytes = sizeof(int) * anchors.size();
    auto anchors_ptr = allocator.Allocate(sizeof(int) * anchors.size());
    int* anchors_data = reinterpret_cast<int*>(anchors_ptr->ptr());
    const auto gplace = boost::get<platform::CUDAPlace>(ctx.GetPlace());
    const auto cplace = platform::CPUPlace();
    memory::Copy(gplace, anchors_data, cplace, anchors.data(), bytes,
D
dengkaipeng 已提交
95
                 dev_ctx.stream());
D
dengkaipeng 已提交
96

D
dengkaipeng 已提交
97
    const T* input_data = input->data<T>();
D
dengkaipeng 已提交
98
    const int* imgsize_data = img_size->data<int>();
D
dengkaipeng 已提交
99 100 101
    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 已提交
102 103 104
    math::SetConstant<platform::CUDADeviceContext, T> set_zero;
    set_zero(dev_ctx, boxes, static_cast<T>(0));
    set_zero(dev_ctx, scores, static_cast<T>(0));
D
dengkaipeng 已提交
105

106 107
    int grid_dim = (n * box_num + 512 - 1) / 512;
    grid_dim = grid_dim > 8 ? 8 : grid_dim;
D
dengkaipeng 已提交
108

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

}  // namespace operators
}  // namespace paddle

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