src/core/CL/cl_kernels/prior_box_layer.cl

/*
 * Copyright (c) 2018 ARM Limited.
 *
 * SPDX-License-Identifier: MIT
 *
 * Permission is hereby granted, free of charge, to any person obtaining a copy
 * of this software and associated documentation files (the "Software"), to
 * deal in the Software without restriction, including without limitation the
 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
 * sell copies of the Software, and to permit persons to whom the Software is
 * furnished to do so, subject to the following conditions:
 *
 * The above copyright notice and this permission notice shall be included in all
 * copies or substantial portions of the Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 * SOFTWARE.
 */
#include "helpers.h"

#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)

/**  Compute prior boxes and clip (NCHW)
 *
 * @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_nchw(Image *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)

    // Store result
    vstore4((VEC_DATA_TYPE(DATA_TYPE, 4))(xmin, ymin, xmax, ymax), 0, ((__global DATA_TYPE *)offset(out, idx + 0, 0)));
}

/** Compute prior boxes (NCHW)
 *
 * @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_size                             Prior box min 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_nchw(Image *out, float *max, 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(0) % LAYER_WIDTH) + OFFSET) * STEP_X;
    const float center_y = ((float)(get_global_id(0) / LAYER_WIDTH) + OFFSET) * STEP_Y;

    float box_width  = min_size;
    float box_height = min_size;
    calculate_xy_min_max_nchw(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_nchw(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_nchw(out, idx, center_x, center_y, box_width, box_height);
        idx += 4;
    }

    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(Image *out, float *max, 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
 * @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_nchw(IMAGE_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)
{
    Image out = CONVERT_TO_IMAGE_STRUCT(output);

    int idx = 0;
    for(unsigned int i = 0; i < min_size; ++i)
    {
        idx = calculate_min_nchw(&out, max, aspect_ratios, max_size, aspect_ratios_size, min[i], i, idx);
    }

    // Store variances
    for(int i = 0; i < (NUM_PRIORS * 4); i += 4)
    {
        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) */