Compute Library
 22.08
prior_box_layer.cl
Go to the documentation of this file.
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  */
43 inline 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  */
75 inline 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) */
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
Definition: helpers.h:1084
#define CONVERT_TO_IMAGE_STRUCT(name)
Definition: helpers.h:855
#define IMAGE_DECLARATION(name)
Definition: helpers.h:805
DataType clamp(const DataType &n, const DataType &lower=std::numeric_limits< RangeType >::lowest(), const DataType &upper=std::numeric_limits< RangeType >::max())
Performs clamping among a lower and upper value.
Definition: Utility.h:101
Structure to hold Image information.
Definition: helpers.h:897
#define VEC_DATA_TYPE(type, size)
Definition: helpers.h:728