From 3112e33d8d0b987e85107390a0350bd5988f5f01 Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Wed, 21 Nov 2018 15:44:55 +0000 Subject: COMPMID-1451 Change PriorBox output to NCHw Output of Priorbox should be independent of the input data layout and should always be in NCHW format Change-Id: Ie80cd4e51c78945b158c0db1af1923bdf8d7ea7b --- src/core/CL/cl_kernels/prior_box_layer.cl | 117 ------------------------------ 1 file changed, 117 deletions(-) (limited to 'src/core/CL/cl_kernels/prior_box_layer.cl') diff --git a/src/core/CL/cl_kernels/prior_box_layer.cl b/src/core/CL/cl_kernels/prior_box_layer.cl index be072ecb1d..046151bca4 100644 --- a/src/core/CL/cl_kernels/prior_box_layer.cl +++ b/src/core/CL/cl_kernels/prior_box_layer.cl @@ -104,88 +104,6 @@ inline int calculate_min_nchw(Image *out, __global float *max, __global float *a return idx; } - -/** Compute prior boxes and clip (NHWC) - * - * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32 - * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes) - * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes) - * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[in] idx Index to write to - * @param[in] center_x Center value of the x axis - * @param[in] center_y Center value of the y axis - * @param[in] box_width Prior box width - * @param[in] box_height Prior box height - * - */ -inline void calculate_xy_min_max_nhwc(Tensor3D *out, int idx, float center_x, float center_y, float box_width, float box_height) -{ - float xmin = (center_x - box_width / 2.f) / WIDTH; - float ymin = (center_y - box_height / 2.f) / HEIGHT; - float xmax = (center_x + box_width / 2.f) / WIDTH; - float ymax = (center_y + box_height / 2.f) / HEIGHT; - -#if defined(CLIP) - xmin = clamp(xmin, 0.f, 1.f); - ymin = clamp(ymin, 0.f, 1.f); - xmax = clamp(xmax, 0.f, 1.f); - ymax = clamp(ymax, 0.f, 1.f); -#endif // defined(CLIP) - - *((__global DATA_TYPE *)tensor3D_offset(out, 0, idx + 0, 0)) = xmin; - *((__global DATA_TYPE *)tensor3D_offset(out, 0, idx + 1, 0)) = ymin; - *((__global DATA_TYPE *)tensor3D_offset(out, 0, idx + 2, 0)) = xmax; - *((__global DATA_TYPE *)tensor3D_offset(out, 0, idx + 3, 0)) = ymax; -} - -/** Compute prior boxes (NHWC) - * - * @param[in,out] out Tensor output - * @param[in] max The maximum values - * @param[in] aspect_ratios The aspect ratio values - * @param[in] max_size The maximum values values size - * @param[in] aspect_ratios_size The aspect ratio values size - * @param[in] min_size The minimum values size - * @param[in] min_idx Index of the min vector - * @param[in] idx Index to write to - * - * @return The updated index - */ -inline int calculate_min_nhwc(Tensor3D *out, __global float *max, __global float *aspect_ratios, int max_size, int aspect_ratios_size, float min_size, int min_idx, int idx) -{ - const float center_x = ((float)(get_global_id(1) % LAYER_WIDTH) + OFFSET) * STEP_X; - const float center_y = ((float)(get_global_id(1) / LAYER_WIDTH) + OFFSET) * STEP_Y; - - float box_width = min_size; - float box_height = min_size; - - calculate_xy_min_max_nhwc(out, idx, center_x, center_y, box_width, box_height); - idx += 4; - if(max_size > 0) - { - box_width = sqrt(min_size * max[min_idx]); - box_height = box_width; - calculate_xy_min_max_nhwc(out, idx, center_x, center_y, box_width, box_height); - idx += 4; - } - for(unsigned int i = 0; i < aspect_ratios_size; ++i) - { - if(fabs(aspect_ratios[i] - 1.f) < 1e-6f) - { - continue; - } - box_width = min_size * sqrt(aspect_ratios[i]); - box_height = min_size * rsqrt(aspect_ratios[i]); - - calculate_xy_min_max_nhwc(out, idx, center_x, center_y, box_width, box_height); - idx += 4; - } - - return idx; -} - /** Calculate prior boxes with NCHW format. * * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32 @@ -218,39 +136,4 @@ __kernel void prior_box_layer_nchw(IMAGE_DECLARATION(output), __global float *mi vstore4((VEC_DATA_TYPE(DATA_TYPE, 4))(VARIANCE_0, VARIANCE_1, VARIANCE_2, VARIANCE_3), 0, ((__global DATA_TYPE *)offset(&out, i, 1))); } } - -/** Calculate prior boxes with NHWC format. - * - * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F32 - * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes) - * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes) - * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[in] min The minimum values - * @param[in] max The maximum_values - * @param[in] aspect_ratios The aspect ratio values - * @param[in] min_size The minimum values size - * @param[in] max_size The maximum_values values size - * @param[in] aspect_ratios_size The aspect ratio values size - */ -__kernel void prior_box_layer_nhwc(TENSOR3D_DECLARATION(output), __global float *min, __global float *max, __global float *aspect_ratios, unsigned int min_size, unsigned int max_size, - unsigned int aspect_ratios_size) -{ - Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output); - - int idx = 0; - for(unsigned int i = 0; i < min_size; ++i) - { - idx = calculate_min_nhwc(&out, max, aspect_ratios, max_size, aspect_ratios_size, min[i], i, idx); - } - - for(int i = 0; i < (NUM_PRIORS * 4); i += 4) - { - *((__global DATA_TYPE *)tensor3D_offset(&out, 0, i + 0, 1)) = VARIANCE_0; - *((__global DATA_TYPE *)tensor3D_offset(&out, 0, i + 1, 1)) = VARIANCE_1; - *((__global DATA_TYPE *)tensor3D_offset(&out, 0, i + 2, 1)) = VARIANCE_2; - *((__global DATA_TYPE *)tensor3D_offset(&out, 0, i + 3, 1)) = VARIANCE_3; - } -} #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) */ -- cgit v1.2.1