1/* 2 * Copyright (c) 2018-2021 Arm Limited. 3 * 4 * SPDX-License-Identifier: MIT 5 * 6 * Permission is hereby granted, free of charge, to any person obtaining a copy 7 * of this software and associated documentation files (the "Software"), to 8 * deal in the Software without restriction, including without limitation the 9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or 10 * sell copies of the Software, and to permit persons to whom the Software is 11 * furnished to do so, subject to the following conditions: 12 * 13 * The above copyright notice and this permission notice shall be included in all 14 * copies or substantial portions of the Software. 15 * 16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE 22 * SOFTWARE. 23 */ 24#include "helpers.h" 25 26#if defined(DATA_TYPE) && defined(WIDTH) && defined(HEIGHT) && defined(LAYER_WIDTH) && defined(LAYER_HEIGHT) && defined(OFFSET) && defined(STEP_X) && defined(STEP_Y) && defined(NUM_PRIORS) && defined(VARIANCE_0) && defined(VARIANCE_1) && defined(VARIANCE_2) && defined(VARIANCE_3) 27 28/** Compute prior boxes and clip (NCHW) 29 * 30 * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32 31 * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes) 32 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) 33 * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes) 34 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) 35 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor 36 * @param[in] idx Index to write to 37 * @param[in] center_x Center value of the x axis 38 * @param[in] center_y Center value of the y axis 39 * @param[in] box_width Prior box width 40 * @param[in] box_height Prior box height 41 * 42 */ 43inline void calculate_xy_min_max_nchw(Image *out, int idx, float center_x, float center_y, float box_width, float box_height) 44{ 45 float xmin = (center_x - box_width / 2.f) / WIDTH; 46 float ymin = (center_y - box_height / 2.f) / HEIGHT; 47 float xmax = (center_x + box_width / 2.f) / WIDTH; 48 float ymax = (center_y + box_height / 2.f) / HEIGHT; 49 50#if defined(CLIP) 51 xmin = clamp(xmin, 0.f, 1.f); 52 ymin = clamp(ymin, 0.f, 1.f); 53 xmax = clamp(xmax, 0.f, 1.f); 54 ymax = clamp(ymax, 0.f, 1.f); 55#endif // defined(CLIP) 56 57 // Store result 58 vstore4((VEC_DATA_TYPE(DATA_TYPE, 4))(xmin, ymin, xmax, ymax), 0, ((__global DATA_TYPE *)offset(out, idx + 0, 0))); 59} 60 61/** Compute prior boxes (NCHW) 62 * 63 * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32 64 * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes) 65 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) 66 * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes) 67 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) 68 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor 69 * @param[in] min_size Prior box min size 70 * @param[in] min_idx Index of the min vector 71 * @param[in] idx Index to write to 72 * 73 * @return The updated index 74 */ 75inline int calculate_min_nchw(Image *out, __global float *max, __global float *aspect_ratios, int max_size, int aspect_ratios_size, float min_size, int min_idx, int idx) 76{ 77 const float center_x = ((float)(get_global_id(0) % LAYER_WIDTH) + OFFSET) * STEP_X; 78 const float center_y = ((float)(get_global_id(0) / LAYER_WIDTH) + OFFSET) * STEP_Y; 79 80 float box_width = min_size; 81 float box_height = min_size; 82 calculate_xy_min_max_nchw(out, idx, center_x, center_y, box_width, box_height); 83 idx += 4; 84 85 if(max_size > 0) 86 { 87 box_width = sqrt(min_size * max[min_idx]); 88 box_height = box_width; 89 calculate_xy_min_max_nchw(out, idx, center_x, center_y, box_width, box_height); 90 idx += 4; 91 } 92 for(unsigned int i = 0; i < aspect_ratios_size; ++i) 93 { 94 if(fabs(aspect_ratios[i] - 1.f) < 1e-6f) 95 { 96 continue; 97 } 98 box_width = min_size * sqrt(aspect_ratios[i]); 99 box_height = min_size * rsqrt(aspect_ratios[i]); 100 101 calculate_xy_min_max_nchw(out, idx, center_x, center_y, box_width, box_height); 102 idx += 4; 103 } 104 105 return idx; 106} 107/** Calculate prior boxes with NCHW format. 108 * 109 * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32 110 * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes) 111 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) 112 * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes) 113 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) 114 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor 115 * @param[in] min The minimum values 116 * @param[in] max The maximum_values 117 * @param[in] aspect_ratios The aspect ratio values 118 * @param[in] min_size The minimum values size 119 * @param[in] max_size The maximum_values values size 120 * @param[in] aspect_ratios_size The aspect ratio values size 121 */ 122__kernel void prior_box_layer_nchw(IMAGE_DECLARATION(output), __global float *min, __global float *max, __global float *aspect_ratios, unsigned int min_size, unsigned int max_size, 123 unsigned int aspect_ratios_size) 124{ 125 Image out = CONVERT_TO_IMAGE_STRUCT(output); 126 127 int idx = 0; 128 for(unsigned int i = 0; i < min_size; ++i) 129 { 130 idx = calculate_min_nchw(&out, max, aspect_ratios, max_size, aspect_ratios_size, min[i], i, idx); 131 } 132 133 // Store variances 134 for(int i = 0; i < (NUM_PRIORS * 4); i += 4) 135 { 136 vstore4((VEC_DATA_TYPE(DATA_TYPE, 4))(VARIANCE_0, VARIANCE_1, VARIANCE_2, VARIANCE_3), 0, ((__global DATA_TYPE *)offset(&out, i, 1))); 137 } 138} 139#endif /* defined(DATA_TYPE) && defined(WIDTH) && defined(HEIGHT) && defined(LAYER_WIDTH) && defined(LAYER_HEIGHT) && defined(OFFSET) && defined(STEP_X) && defined(STEP_Y) && defined(NUM_PRIORS) && defined(VARIANCE_0) && defined(VARIANCE_1) && defined(VARIANCE_2) && defined(VARIANCE_3) */ 140