mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-08-06 06:26:29 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

optimize region kernels

Browse Source
This commit is contained in:
Yashas 2019-12-08 21:03:30 +05:30
parent 65d606630d
commit dd3f517fe9
3 changed files with 115 additions and 144 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

228
modules/dnn/src/cuda/region.cu
View File

@ -24,176 +24,154 @@ using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw { namespace raw {
template <class T>
__global__ void sigmoid_strided(Span<T> output, View<T> input, size_type n, size_type stride, size_type offset) {
/* - the input is divided into equal blocks strided by `stride`
* - we must apply sigmoid to a continuous range of `n` values starting from `offset` in every block
*/
for (auto i : grid_stride_range(n * output.size() / stride)) {
auto block_idx = i / n;
auto index = block_idx * stride + offset + (i % n);
using device::sigmoid;
output[index] = sigmoid(input[index]);
}
}
template <class T> template <class T>
__global__ void softmax_strided(Span<T> output, View<T> input, size_type n, size_type stride, size_type offset_) { __global__ void region_box(
for (auto idx : grid_stride_range(output.size() / stride)) { Span<T> output, View<T> input, View<T> bias,
index_type offset = idx * stride + offset_; size_type boxes_per_cell, size_type box_size,
auto largest = numeric_limits<T>::lowest();
for (int i = 0; i < n; i++) {
using device::max;
largest = max(largest, output[offset + i]);
}
auto sum = T(0);
for (int i = 0; i < n; i++) {
using device::exp;
auto temp = exp(output[offset + i] - largest);
sum += temp;
output[offset + i] = temp;
}
for (int i = 0; i < n; i++) {
output[offset + i] /= sum;
}
}
}
template <class T>
__global__ void region_finalize(Span<T> output, View<T> input, View<T> bias,
T object_prob_cutoff, T class_prob_cutoff,
size_type height_norm, size_type width_norm,
size_type rows, size_type cols, size_type rows, size_type cols,
size_type boxes_per_cell, size_type height_norm, size_type width_norm,
size_type box_size, T object_prob_cutoff)
size_type classes)
{ {
using vector2_type = get_vector_type_t<T, 2>;
auto bias_vPtr = vector2_type::get_pointer(bias.data());
for (auto box_index : grid_stride_range(output.size() / box_size)) { for (auto box_index : grid_stride_range(output.size() / box_size)) {
auto box_of_the_cell = box_index % boxes_per_cell; /* box number within a cell */ const auto box_of_the_cell = box_index % boxes_per_cell; /* box number within a cell */
auto box_offset = box_index * box_size; const auto box_offset = box_index * box_size;
auto batch_inner_size = rows * cols * boxes_per_cell; const auto batch_inner_size = rows * cols * boxes_per_cell;
auto row_inner_size = cols * boxes_per_cell; const auto row_inner_size = cols * boxes_per_cell;
auto col_inner_size = boxes_per_cell; const auto col_inner_size = boxes_per_cell;
auto y = (box_index % batch_inner_size) / row_inner_size; const auto y = (box_index % batch_inner_size) / row_inner_size;
auto x = (box_index % row_inner_size) / col_inner_size; const auto x = (box_index % row_inner_size) / col_inner_size;
using device::sigmoid; using device::sigmoid;
using device::exp;
output[box_offset + 0] = (T(x) + sigmoid(input[box_offset + 0])) / T(cols); output[box_offset + 0] = (T(x) + sigmoid(input[box_offset + 0])) / T(cols);
output[box_offset + 1] = (T(y) + sigmoid(input[box_offset + 1])) / T(rows); output[box_offset + 1] = (T(y) + sigmoid(input[box_offset + 1])) / T(rows);
output[box_offset + 2] = exp(input[box_offset + 2]) * bias[2 * box_of_the_cell + 0] / T(width_norm);
output[box_offset + 3] = exp(input[box_offset + 3]) * bias[2 * box_of_the_cell + 1] / T(height_norm); vector2_type bias_xy;
v_load(bias_xy, bias_vPtr[box_of_the_cell]);
using device::exp;
output[box_offset + 2] = exp(input[box_offset + 2]) * bias_xy.data[0] / T(width_norm);
output[box_offset + 3] = exp(input[box_offset + 3]) * bias_xy.data[1] / T(height_norm);
/* squash objectness score into a probability */ /* squash objectness score into a probability */
using device::sigmoid; using device::sigmoid;
T objectness_prob = sigmoid(output[box_offset + 4]); T objectness_prob = sigmoid(input[box_offset + 4]);
output[box_offset + 4] = objectness_prob;
/* ignore prediction if the objectness probability is less than the cutoff */ /* ignore prediction if the objectness probability is less than the cutoff */
if (objectness_prob < object_prob_cutoff) if (objectness_prob < object_prob_cutoff)
objectness_prob = 0; objectness_prob = 0;
/* the class probabilities we have currently are conditional class probabilities output[box_offset + 4] = objectness_prob;
}
}
template <class T>
__global__ void region_sigmoid_class_score(Span<T> output, View<T> input, T class_prob_cutoff, size_type box_size)
{
for (auto idx : grid_stride_range(output.size())) {
const index_type box_no = idx / box_size;
const index_type start_of_box = box_no * box_size;
const index_type box_offset = idx % box_size;
if (box_offset < 5) {
/* continue as we have already processed these in region_box */
continue;
}
auto objectness_prob = output[start_of_box + 4];
/* the class probabilities we currently have are conditional class probabilities
* given the object * given the object
* *
* to obtain the actual class probability, we multiply the conditional probability * to obtain the actual class probability, we multiply the conditional probability
* with the object probability * with the object probability
*/ */
const index_type class_begin = box_offset + 5; /* 4 box coordinates, 1 obj prob, class probs... */ auto actual_class_prob = objectness_prob * sigmoid(input[idx]);
const index_type class_end = class_begin + classes; if (actual_class_prob <= class_prob_cutoff)
index_type offset = class_begin; actual_class_prob = T(0);
output[idx] = actual_class_prob;
}
}
using vector_type = get_vector_type_t<T, 4>; template <class T>
__global__ void region_softmax_class_score(Span<T> output, View<T> input, T class_prob_cutoff, size_type box_size) {
for (auto box_no : grid_stride_range(output.size() / box_size)) {
const index_type start_of_box = box_no * box_size;
const index_type start_idx = start_of_box + 5;
const index_type end_idx = start_of_box + box_size;
/* process each class independently until the offset is aligned to an n-element boundary */ auto largest = numeric_limits<T>::lowest();
while (offset % vector_type::size() != 0 && offset < class_end) { for (int idx = start_idx; idx < end_idx; idx++) {
T actual_class_prob = objectness_prob * output[offset]; using device::max;
if (actual_class_prob <= class_prob_cutoff) largest = max(largest, input[idx]);
actual_class_prob = T(0);
output[offset] = actual_class_prob;
offset++;
} }
auto output_vPtr = vector_type::get_pointer(output.data() + offset); auto sum = T(0);
auto input_vPtr = vector_type::get_pointer(input.data() + offset); for (int idx = start_idx; idx < end_idx; idx++) {
for (int i = 0; (offset + vector_type::size()) < class_end; i++) { using device::exp;
vector_type vec; auto temp = exp(input[idx] - largest);
v_load(vec, output_vPtr[i]); sum += temp;
for (int j = 0; j < vector_type::size(); j++) { output[idx] = temp;
T actual_class_prob = objectness_prob * vec.data[j];
if (actual_class_prob <= class_prob_cutoff)
actual_class_prob = T(0);
vec.data[j] = actual_class_prob;
}
v_store(output_vPtr[i], vec);
offset += vector_type::size();
} }
/* process the remaining classes */ for (int idx = start_idx; idx < end_idx; idx++) {
while (offset < class_end) { auto softmax_score = output[idx] / sum;
T actual_class_prob = objectness_prob * output[offset];
/* the class probabilities we currently have are conditional class probabilities
* given the object
*
* to obtain the actual class probability, we multiply the conditional probability
* with the object probability
*/
auto objectness_prob = output[start_of_box + 4];
auto actual_class_prob = objectness_prob * softmax_score;
if (actual_class_prob <= class_prob_cutoff) if (actual_class_prob <= class_prob_cutoff)
actual_class_prob = T(0); actual_class_prob = T(0);
output[offset] = actual_class_prob; output[idx] = actual_class_prob;
offset++;
} }
} }
} }
} }
template <class T> template <class T>
void sigmoid_strided(const Stream& stream, Span<T> output, View<T> input, std::size_t n, std::size_t stride, std::size_t offset) { void region(const Stream& stream, Span<T> output, View<T> input, View<T> bias,
CV_Assert(output.size() % stride == 0);
auto kernel = raw::sigmoid_strided<T>;
auto policy = make_policy(kernel, n * output.size() / stride, 0, stream);
launch_kernel(kernel, policy, output, input, n, stride, offset);
}
template void sigmoid_strided(const Stream&, Span<__half>, View<__half>, std::size_t, std::size_t, std::size_t);
template void sigmoid_strided(const Stream&, Span<float>, View<float>, std::size_t, std::size_t, std::size_t);
template <class T>
void softmax_strided(const Stream& stream, Span<T> output, View<T> input, std::size_t n, std::size_t stride, std::size_t offset) {
CV_Assert(output.size() % stride == 0);
auto kernel = raw::softmax_strided<T>;
auto policy = make_policy(kernel, output.size() / stride, 0, stream);
launch_kernel(kernel, policy, output, input, n, stride, offset);
}
template void softmax_strided(const Stream&, Span<__half>, View<__half>, std::size_t, std::size_t, std::size_t);
template void softmax_strided(const Stream&, Span<float>, View<float>, std::size_t, std::size_t, std::size_t);
template <class T>
void region_finalize(const Stream& stream, Span<T> output, View<T> input, View<T> bias,
T object_prob_cutoff, T class_prob_cutoff, T object_prob_cutoff, T class_prob_cutoff,
std::size_t height_norm, std::size_t width_norm, std::size_t boxes_per_cell, std::size_t box_size,
std::size_t rows, std::size_t cols, std::size_t rows, std::size_t cols,
std::size_t boxes_per_cell, std::size_t height_norm, std::size_t width_norm,
std::size_t box_size, bool if_true_sigmoid_else_softmax /* true = sigmoid, false = softmax */)
std::size_t classes)
{ {
CV_Assert(output.size() == input.size());
CV_Assert(output.size() % box_size == 0); CV_Assert(output.size() % box_size == 0);
CV_Assert(is_fully_aligned(bias, 2));
auto kernel = raw::region_finalize<T>; auto box_kernel = raw::region_box<T>;
auto policy = make_policy(kernel, output.size() / box_size, 0, stream); auto box_policy = make_policy(box_kernel, output.size() / box_size, 0, stream);
launch_kernel(kernel, policy, output, input, bias, launch_kernel(box_kernel, box_policy,
object_prob_cutoff, class_prob_cutoff, output, input, bias, boxes_per_cell, box_size,
height_norm, width_norm, rows, cols, height_norm, width_norm,
rows, cols, boxes_per_cell, box_size, classes); object_prob_cutoff);
if (if_true_sigmoid_else_softmax) {
auto kernel_score = raw::region_sigmoid_class_score<T>;
auto policy_score = make_policy(kernel_score, output.size(), 0, stream);
launch_kernel(kernel_score, policy_score, output, input, class_prob_cutoff, box_size);
} else {
auto kernel_score = raw::region_softmax_class_score<T>;
auto policy_score = make_policy(kernel_score, output.size(), 0, stream);
launch_kernel(kernel_score, policy_score, output, input, class_prob_cutoff, box_size);
}
} }
template void region_finalize(const Stream&, Span<__half>, View<__half>, View<__half>, template void region(const Stream&, Span<__half>, View<__half>, View<__half>,
__half, __half, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t); __half, __half, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, bool);
template void region_finalize(const Stream&, Span<float>, View<float>, View<float>, template void region(const Stream&, Span<float>, View<float>, View<float>,
float, float, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t); float, float, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, std::size_t, bool);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */ }}}} /* namespace cv::dnn::cuda4dnn::kernels */

15
modules/dnn/src/cuda4dnn/kernels/region.hpp
View File

@ -13,19 +13,12 @@
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T> template <class T>
void sigmoid_strided(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input, std::size_t n, std::size_t stride, std::size_t offset); void region(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input, csl::View<T> bias,
template <class T>
void softmax_strided(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input, std::size_t n, std::size_t stride, std::size_t offset);
template <class T>
void region_finalize(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input, csl::View<T> bias,
T object_prob_cutoff, T class_prob_cutoff, T object_prob_cutoff, T class_prob_cutoff,
std::size_t height_norm, std::size_t width_norm, std::size_t boxes_per_cell, std::size_t box_size,
std::size_t rows, std::size_t cols, std::size_t rows, std::size_t cols,
std::size_t boxes_per_cell, std::size_t height_norm, std::size_t width_norm,
std::size_t box_size, bool if_true_sigmoid_else_softmax);
std::size_t classes);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */ }}}} /* namespace cv::dnn::cuda4dnn::kernels */

16
modules/dnn/src/cuda4dnn/primitives/region.hpp
View File

@ -102,21 +102,21 @@ namespace cv { namespace dnn { namespace cuda4dnn {
auto output_wrapper = outputs[0].dynamicCast<wrapper_type>(); auto output_wrapper = outputs[0].dynamicCast<wrapper_type>();
auto output = output_wrapper->getSpan(); auto output = output_wrapper->getSpan();
csl::memcpy<T>(output.get(), input.get(), output.size(), stream);
auto rows = input.get_axis_size(1); auto rows = input.get_axis_size(1);
auto cols = input.get_axis_size(2); auto cols = input.get_axis_size(2);
auto cell_box_size = classes + 4 + 1; auto cell_box_size = classes + 4 + 1;
/* we squash class scores into probabilities using softmax or sigmoid */ /* we squash class scores into probabilities using softmax or sigmoid */
if (squash_type == SquashMethod::SOFTMAX) bool if_true_sigmoid_else_softmax = (squash_type == SquashMethod::SIGMOID);
kernels::softmax_strided<T>(stream, output, input, classes, cell_box_size, 5);
else if (squash_type == SquashMethod::SIGMOID)
kernels::sigmoid_strided<T>(stream, output, input, classes, cell_box_size, 5);
kernels::region_finalize<T>(stream, output, input, biasTensor, object_prob_cutoff, class_prob_cutoff, kernels::region<T>(stream, output, input, biasTensor,
height_norm, width_norm, rows, cols, boxes_per_cell, cell_box_size, classes); object_prob_cutoff, class_prob_cutoff,
boxes_per_cell, cell_box_size,
rows, cols,
height_norm, width_norm,
if_true_sigmoid_else_softmax
);
if (nms_iou_threshold > 0) { if (nms_iou_threshold > 0) {
auto output_mat = output_wrapper->getMutableHostMat(); auto output_mat = output_wrapper->getMutableHostMat();