modules/dnn/src/cuda/region.cu

   1 // This file is part of OpenCV project.
   2 // It is subject to the license terms in the LICENSE file found in the top-level directory
   3 // of this distribution and at http://opencv.org/license.html.
   4
   5 #include <cuda_runtime.h>
   6 #include <cuda_fp16.h>
   7
   8 #include "math.hpp"
   9 #include "grid_stride_range.hpp"
  10 #include "execution.hpp"
  11 #include "limits.hpp"
  12 #include "vector_traits.hpp"
  13
  14 #include "../cuda4dnn/csl/stream.hpp"
  15 #include "../cuda4dnn/csl/span.hpp"
  16
  17 #include <opencv2/core.hpp>
  18
  19 #include <cstddef>
  20
  21 using namespace cv::dnn::cuda4dnn::csl;
  22 using namespace cv::dnn::cuda4dnn::csl::device;
  23
  24 namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
  25
  26     namespace raw {
  27         template <class T>
  28         __global__ void sigmoid_strided(Span<T> output, View<T> input, size_type n, size_type stride, size_type offset) {
  29             /* - the input is divided into equal blocks strided by `stride`
  30              * - we must apply sigmoid to a continuous range of `n` values starting from `offset` in every block
  31              */
  32             for (auto i : grid_stride_range(n * output.size() / stride)) {
  33                 auto block_idx = i / n;
  34                 auto index = block_idx * stride + offset + (i % n);
  35
  36                 using device::sigmoid;
  37                 output[index] = sigmoid(input[index]);
  38             }
  39         }
  40
  41         template <class T>
  42         __global__ void softmax_strided(Span<T> output, View<T> input, size_type n, size_type stride, size_type offset_) {
  43             for (auto idx : grid_stride_range(output.size() / stride)) {
  44                 index_type offset = idx * stride + offset_;
  45
  46                 auto largest = numeric_limits<T>::lowest();
  47                 for (int i = 0; i < n; i++) {
  48                     using device::max;
  49                     largest = max(largest, output[offset + i]);
  50                 }
  51
  52                 auto sum = T(0);
  53                 for (int i = 0; i < n; i++) {
  54                     using device::exp;
  55                     auto temp = exp(output[offset + i] - largest);
  56                     sum += temp;
  57                     output[offset + i] = temp;
  58                 }
  59
  60                 for (int i = 0; i < n; i++) {
  61                     output[offset + i] /= sum;
  62                 }
  63             }
  64         }
  65
  66         template <class T>
  67         __global__ void region_finalize(Span<T> output, View<T> input, View<T> bias,
  68             T object_prob_cutoff, T class_prob_cutoff,
  69             size_type height_norm, size_type width_norm,
  70             size_type rows, size_type cols,
  71             size_type boxes_per_cell,
  72             size_type box_size,
  73             size_type classes)
  74         {
  75             for (auto box_index : grid_stride_range(output.size() / box_size)) {
  76                 auto box_of_the_cell = box_index % boxes_per_cell; /* box number within a cell */
  77                 auto box_offset = box_index * box_size;
  78
  79                 auto batch_inner_size = rows * cols * boxes_per_cell;
  80                 auto row_inner_size = cols * boxes_per_cell;
  81                 auto col_inner_size = boxes_per_cell;
  82
  83                 auto y = (box_index % batch_inner_size) / row_inner_size;
  84                 auto x = (box_index % row_inner_size) / col_inner_size;
  85
  86                 using device::sigmoid;
  87                 using device::exp;
  88                 output[box_offset + 0] = (T(x) + sigmoid(input[box_offset + 0])) / T(cols);
  89                 output[box_offset + 1] = (T(y) + sigmoid(input[box_offset + 1])) / T(rows);
  90                 output[box_offset + 2] = exp(input[box_offset + 2]) * bias[2 * box_of_the_cell + 0] / T(width_norm);
  91                 output[box_offset + 3] = exp(input[box_offset + 3]) * bias[2 * box_of_the_cell + 1] / T(height_norm);
  92
  93                 /* squash objectness score into a probability */
  94                 using device::sigmoid;
  95                 T objectness_prob = sigmoid(output[box_offset + 4]);
  96                 output[box_offset + 4] = objectness_prob;
  97
  98                 /* ignore prediction if the objectness probability is less than the cutoff */
  99                 if (objectness_prob < object_prob_cutoff)
 100                     objectness_prob = 0;
 101
 102                 /* the class probabilities we have currently are conditional class probabilities
 103                  * given the object
 104                  *
 105                  * to obtain the actual class probability, we multiply the conditional probability
 106                  * with the object probability
 107                  */
 108                 const index_type class_begin = box_offset + 5; /* 4 box coordinates, 1 obj prob, class probs... */
 109                 const index_type class_end = class_begin + classes;
 110                 index_type offset = class_begin;
 111
 112                 using vector_type = get_vector_type_t<T, 4>;
 113
 114                 /* process each class independently until the offset is aligned to an n-element boundary */
 115                 while (offset % vector_type::size() != 0 && offset < class_end) {
 116                     T actual_class_prob = objectness_prob * output[offset];
 117                     if (actual_class_prob <= class_prob_cutoff)
 118                         actual_class_prob = T(0);
 119                     output[offset] = actual_class_prob;
 120                     offset++;
 121                 }
 122
 123                 auto output_vPtr = vector_type::get_pointer(output.data() + offset);
 124                 auto input_vPtr = vector_type::get_pointer(input.data() + offset);
 125                 for (int i = 0; (offset + vector_type::size()) < class_end; i++) {
 126                     vector_type vec;
 127                     v_load(vec, output_vPtr[i]);
 128                     for (int j = 0; j < vector_type::size(); j++) {
 129                         T actual_class_prob = objectness_prob * vec.data[j];
 130                         if (actual_class_prob <= class_prob_cutoff)
 131                             actual_class_prob = T(0);
 132                         vec.data[j] = actual_class_prob;
 133                     }
 134                     v_store(output_vPtr[i], vec);
 135                     offset += vector_type::size();
 136                 }
 137
 138                 /* process the remaining classes */
 139                 while (offset < class_end) {
 140                     T actual_class_prob = objectness_prob * output[offset];
 141                     if (actual_class_prob <= class_prob_cutoff)
 142                         actual_class_prob = T(0);
 143                     output[offset] = actual_class_prob;
 144                     offset++;
 145                 }
 146             }
 147         }
 148     }
 149
 150     template <class T>
 151     void sigmoid_strided(const Stream& stream, Span<T> output, View<T> input, std::size_t n, std::size_t stride, std::size_t offset) {
 152         CV_Assert(output.size() % stride == 0);
 153
 154         auto kernel = raw::sigmoid_strided<T>;
 155         auto policy = make_policy(kernel, n * output.size() / stride, 0, stream);
 156         launch_kernel(kernel, policy, output, input, n, stride, offset);
 157     }
 158
 159     template void sigmoid_strided(const Stream&, Span<__half>, View<__half>, std::size_t, std::size_t, std::size_t);
 160     template void sigmoid_strided(const Stream&, Span<float>, View<float>, std::size_t, std::size_t, std::size_t);
 161
 162     template <class T>
 163     void softmax_strided(const Stream& stream, Span<T> output, View<T> input, std::size_t n, std::size_t stride, std::size_t offset) {
 164         CV_Assert(output.size() % stride == 0);
 165
 166         auto kernel = raw::softmax_strided<T>;
 167         auto policy = make_policy(kernel, output.size() / stride, 0, stream);
 168         launch_kernel(kernel, policy, output, input, n, stride, offset);
 169     }
 170
 171     template void softmax_strided(const Stream&, Span<__half>, View<__half>, std::size_t, std::size_t, std::size_t);
 172     template void softmax_strided(const Stream&, Span<float>, View<float>, std::size_t, std::size_t, std::size_t);
 173
 174     template <class T>
 175     void region_finalize(const Stream& stream, Span<T> output, View<T> input, View<T> bias,
 176         T object_prob_cutoff, T class_prob_cutoff,
 177         std::size_t height_norm, std::size_t width_norm,
 178         std::size_t rows, std::size_t cols,
 179         std::size_t boxes_per_cell,
 180         std::size_t box_size,
 181         std::size_t classes)
 182     {
 183         CV_Assert(output.size() % box_size == 0);
 184
 185         auto kernel = raw::region_finalize<T>;
 186         auto policy = make_policy(kernel, output.size() / box_size, 0, stream);
 187         launch_kernel(kernel, policy, output, input, bias,
 188             object_prob_cutoff, class_prob_cutoff,
 189             height_norm, width_norm,
 190             rows, cols, boxes_per_cell, box_size, classes);
 191     }
 192
 193     template void region_finalize(const Stream&, Span<__half>, View<__half>, View<__half>,
 194         __half, __half, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t);
 195
 196     template void region_finalize(const Stream&, Span<float>, View<float>, View<float>,
 197         float, float, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t);
 198
 199 }}}} /* namespace cv::dnn::cuda4dnn::kernels */