compute/ARMComputeEx/src/core/CL/cl_kernels/hashtable_lookup.cl

   1 /*
   2  * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
   3  *
   4  * Licensed under the Apache License, Version 2.0 (the "License");
   5  * you may not use this file except in compliance with the License.
   6  * You may obtain a copy of the License at
   7  *
   8  *      http://www.apache.org/licenses/LICENSE-2.0
   9  *
  10  * Unless required by applicable law or agreed to in writing, software
  11  * distributed under the License is distributed on an "AS IS" BASIS,
  12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  13  * See the License for the specific language governing permissions and
  14  * limitations under the License.
  15  */
  16
  17 /*
  18  * Copyright (c) 2017 ARM Limited.
  19  *
  20  * SPDX-License-Identifier: MIT
  21  *
  22  * Permission is hereby granted, free of charge, to any person obtaining a copy
  23  * of this software and associated documentation files (the "Software"), to
  24  * deal in the Software without restriction, including without limitation the
  25  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
  26  * sell copies of the Software, and to permit persons to whom the Software is
  27  * furnished to do so, subject to the following conditions:
  28  *
  29  * The above copyright notice and this permission notice shall be included in all
  30  * copies or substantial portions of the Software.
  31  *
  32  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  33  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  34  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
  35  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
  36  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
  37  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  38  * SOFTWARE.
  39  */
  40
  41 #include "helpers.h"
  42
  43 #ifndef VEC_SIZE
  44 #define VEC_SIZE 1
  45 #endif
  46
  47 #if defined(DATA_TYPE) && defined(DEPTH_OUT) && defined(NUM_DIMS)
  48 /** Perform hashtable_lookup of input tensor
  49  *
  50  * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g.
  51  *       -DDATA_TYPE=short
  52  * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g.
  53  *            -DVEC_SIZE=16
  54  * @attention Output tensor depth should be given as a preprocessor argument using
  55  *            -DDEPTH_OUT=depth. e.g. -DDEPTH_OUT=16
  56  * @attention Number of input dimensions are passed as a preprocessor argument using
  57  *            -DNUM_DIMS=size, e.g. -DNUM_DIMS=4
  58  *
  59  * @param[in]  input_ptr                             Pointer to the source tensor. Supported data
  60  *                                                   types: U8/S8/QASYMM8/U16/S16/F16/U32/S32/F32
  61  * @param[in]  input_stride_x                        Stride of the source tensor in X dimension (in
  62  *                                                   bytes)
  63  * @param[in]  input_step_x                          input_stride_x * number of elements along X
  64  *                                                   processed per workitem(in bytes)
  65  * @param[in]  input_stride_y                        Stride of the source tensor in Y dimension (in
  66  *                                                   bytes)
  67  * @param[in]  input_step_y                          input_stride_y * number of elements along Y
  68  *                                                   processed per workitem(in bytes)
  69  * @param[in]  input_stride_z                        Stride of the source tensor in Z dimension (in
  70  *                                                   bytes)
  71  * @param[in]  input_step_z                          input_stride_z * number of elements along Z
  72  *                                                   processed per workitem(in bytes)
  73  * @param[in]  input_offset_first_element_in_bytes   The offset of the first element in the source
  74  *                                                   tensor
  75  * @param[in]  input_stride_w                        Stride of the source tensor in W dimension (in
  76  *                                                   bytes)
  77  * @param[in]  input_step_w                          output_stride_w * number of elements along W
  78  *                                                   processed per workitem(in bytes)
  79  * @param[out] output_ptr                            Pointer to the destination tensor. Supported
  80  *                                                   data types: same as @p input_ptr
  81  * @param[in]  output_stride_x                       Stride of the destination tensor in X dimension
  82  *                                                   (in bytes)
  83  * @param[in]  output_step_x                         output_stride_x * number of elements along X
  84  *                                                   processed per workitem(in bytes)
  85  * @param[in]  output_stride_y                       Stride of the destination tensor in Y dimension
  86  *                                                   (in bytes)
  87  * @param[in]  output_step_y                         output_stride_y * number of elements along Y
  88  *                                                   processed per workitem(in bytes)
  89  * @param[in]  output_stride_z                       Stride of the source tensor in Z dimension (in
  90  *                                                   bytes)
  91  * @param[in]  output_step_z                         output_stride_z * number of elements along Z
  92  *                                                   processed per workitem(in bytes)
  93  * @param[in]  output_stride_w                       Stride of the source tensor in W dimension (in
  94  *                                                   bytes)
  95  * @param[in]  output_step_w                         output_stride_w * number of elements along W
  96  *                                                   processed per workitem(in bytes)
  97  * @param[in]  output_offset_first_element_in_bytes  The offset of the first element in the
  98  *                                                   destination tensor
  99  * @param[in]  lookups_ptr                           Pointer to the lookups vector. Supported data
 100  *                                                   types: S32
 101  * @param[in]  lookups_stride_x                      Stride of the lookups vector in X dimension (in
 102  *                                                   bytes)
 103  * @param[in]  lookups_step_x                        lookups_stride_x * number of elements along X
 104  *                                                   processed per workitem(in bytes)
 105  * @param[in]  lookups_offset_first_element_in_bytes The offset of the first element in the lookups
 106  *                                                   vector
 107  */
 108 __kernel void hashtable_lookup(TENSOR4D_DECLARATION(input), TENSOR4D_DECLARATION(output),
 109                                VECTOR_DECLARATION(lookups))
 110 {
 111   Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT(output, DEPTH_OUT);
 112   Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input, DEPTH_OUT);
 113
 114   Vector lups = CONVERT_TO_VECTOR_STRUCT_NO_STEP(lookups);
 115
 116   int lup_id[4] = {0};
 117
 118   lup_id[0] = (NUM_DIMS == 1) ? *((__global int *)vector_offset(&lups, get_global_id(0)))
 119                               : get_global_id(0);
 120   lup_id[1] = (NUM_DIMS == 2) ? *((__global int *)vector_offset(&lups, get_global_id(1)))
 121                               : get_global_id(1);
 122   lup_id[2] = (NUM_DIMS == 3) ? *((__global int *)vector_offset(&lups, get_global_id(2)))
 123                               : get_global_id(2) % DEPTH_OUT;
 124   lup_id[3] = (NUM_DIMS == 4)
 125                   ? *((__global int *)vector_offset(&lups, get_global_id(2) / DEPTH_OUT))
 126                   : get_global_id(2) / DEPTH_OUT;
 127
 128   if (lup_id[NUM_DIMS - 1] < 0)
 129   {
 130     VSTORE(VEC_SIZE)((VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0, 0, (__global DATA_TYPE *)out.ptr);
 131     return;
 132   }
 133
 134   in.ptr += input_offset_first_element_in_bytes + lup_id[0] * input_step_x +
 135             lup_id[1] * input_step_y + lup_id[2] * input_step_z + lup_id[3] * input_step_w;
 136
 137   VSTORE(VEC_SIZE)
 138   (CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in.ptr), VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)), 0,
 139    (__global DATA_TYPE *)out.ptr);
 140 }
 141 #endif // defined(DATA_TYPE) && defined(DEPTH_OUT) && defined(NUM_DIMS)