diff options
Diffstat (limited to 'src/core/CL/cl_kernels/common/arg_min_max.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/common/arg_min_max.cl | 388 |
1 files changed, 388 insertions, 0 deletions
diff --git a/src/core/CL/cl_kernels/common/arg_min_max.cl b/src/core/CL/cl_kernels/common/arg_min_max.cl new file mode 100644 index 0000000000..413fcf5333 --- /dev/null +++ b/src/core/CL/cl_kernels/common/arg_min_max.cl @@ -0,0 +1,388 @@ +/* + * Copyright (c) 2019-2021, 2023 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" +#include "tile_helpers.h" + +#if defined(VEC_SIZE) && defined(DATA_TYPE) && defined(DATA_TYPE_OUTPUT) + +#define VEC_TYPE_IN VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) +#define VEC_TYPE_OUT VEC_DATA_TYPE(DATA_TYPE_OUTPUT, VEC_SIZE) +#define VEC_SELECT_IN SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) +#define VEC_SIGNED_INT_IN SIGNED_INT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + +#if defined(FLOAT_DATA_TYPE) +#define ISGREATER(x, y) (VEC_SELECT_IN) isgreater(x, y) +#define ISLESS(x, y) (VEC_SELECT_IN) isless(x, y) +#else // !FLOAT_DATA_TYPE +#if defined(WIDTH) +#define ISGREATER(x, y) (x > y) ? 1 : 0 +#define ISLESS(x, y) (x < y) ? 1 : 0 +#else // !defined(WIDTH) +#define ISGREATER(x, y) select((VEC_SIGNED_INT_IN)0, (VEC_SIGNED_INT_IN)-1, (VEC_SIGNED_INT_IN)(x > y)) +#define ISLESS(x, y) select((VEC_SIGNED_INT_IN)0, (VEC_SIGNED_INT_IN)-1, (VEC_SIGNED_INT_IN)(x < y)) +#endif // defined(WIDTH) +#endif // defined(FLOAT_DATA_TYPE) + +#if defined(ARG_MAX) +#define CONDITION_TO_USE(x, y) ISGREATER(x, y) +#elif defined(ARG_MIN) +#define CONDITION_TO_USE(x, y) ISLESS(x, y) +#else // !(defined(ARG_MAX) || defined(ARG_MIN)) +#error "Unsupported reduction operation!" +#endif // defined(ARG_MAX) + +#if defined(WIDTH) + +#if defined(ARG_MAX) +#define VECTOR_PREDICATE_EQ(x, y) ((x) >= (y)) +#define VECTOR_PREDICATE(x, y) ((x) > (y)) +#define SCALAR_SELECT_OP(x, y) ((x) > (y)) ? (x) : (y); +#elif defined(ARG_MIN) +#define VECTOR_PREDICATE_EQ(x, y) ((x) <= (y)) +#define VECTOR_PREDICATE(x, y) ((x) < (y)) +#define SCALAR_SELECT_OP(x, y) ((x) < (y)) ? (x) : (y); +#else // !(defined(ARG_MAX) || defined(ARG_MIN)) +#error "Unsupported reduction operation!" +#endif // defined(ARG_MAX) + +inline DATA_TYPE_OUTPUT vectorized_compute_arg_min_max_2(DATA_TYPE *min_max_val, DATA_TYPE_OUTPUT *min_max_idx, VEC_DATA_TYPE(DATA_TYPE, 2) in, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 2) res) +{ + if( VECTOR_PREDICATE_EQ(in.s0,in.s1) ) + { + *min_max_val = in.s0; + *min_max_idx = res.s0; + } + else + { + *min_max_val = in.s1; + *min_max_idx = res.s1; + } +} + +inline DATA_TYPE_OUTPUT vectorized_compute_arg_min_max_4(DATA_TYPE *min_max_val, DATA_TYPE_OUTPUT *min_max_idx, VEC_DATA_TYPE(DATA_TYPE, 4) in, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 4) res) +{ + VEC_DATA_TYPE(COND_DATA_TYPE, 2) + idx_sel = VECTOR_PREDICATE_EQ(in.s01, in.s23); + in.s01 = select(in.s23, in.s01, idx_sel); + res.s01 = select(res.s23, res.s01, CONVERT(idx_sel, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 2) )); + idx_sel.s0 = VECTOR_PREDICATE(in.s0, in.s1) || (in.s0 == in.s1 && CONVERT((res.s0 < res.s1), COND_DATA_TYPE)); + res.s0 = select(res.s1, res.s0, CONVERT(idx_sel.s0, DATA_TYPE_OUTPUT)); + *min_max_val = SCALAR_SELECT_OP(in.s0, in.s1); + *min_max_idx = res.s0; +} + +inline DATA_TYPE_OUTPUT vectorized_compute_arg_min_max_8(DATA_TYPE *min_max_val, DATA_TYPE_OUTPUT *min_max_idx, VEC_DATA_TYPE(DATA_TYPE, 8) in, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 8) res) +{ + VEC_DATA_TYPE(COND_DATA_TYPE, 4) + idx_sel = VECTOR_PREDICATE_EQ(in.s0123, in.s4567); + in.s0123 = select(in.s4567, in.s0123, idx_sel); + res.s0123 = select(res.s4567, res.s0123, CONVERT(idx_sel, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 4) )); + idx_sel.s01 = (VECTOR_PREDICATE(in.s01, in.s23)) || (in.s01 == in.s23 && CONVERT(((res.s01 < res.s23)), VEC_DATA_TYPE(COND_DATA_TYPE, 2))); + in.s01 = select(in.s23, in.s01, idx_sel.s01); + res.s01 = select(res.s23, res.s01, CONVERT(idx_sel.s01, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 2) )); + idx_sel.s0 = VECTOR_PREDICATE(in.s0, in.s1) || (in.s0 == in.s1 && CONVERT((res.s0 < res.s1), COND_DATA_TYPE)); + res.s0 = select(res.s1, res.s0, CONVERT(idx_sel.s0, DATA_TYPE_OUTPUT)); + *min_max_val = SCALAR_SELECT_OP(in.s0, in.s1); + *min_max_idx = res.s0; +} + +inline DATA_TYPE_OUTPUT vectorized_compute_arg_min_max_16(DATA_TYPE *min_max_val, DATA_TYPE_OUTPUT *min_max_idx, VEC_DATA_TYPE(DATA_TYPE, 16) in, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16) res) +{ + VEC_DATA_TYPE(COND_DATA_TYPE, 8) + idx_sel = VECTOR_PREDICATE_EQ(in.s01234567, in.s89abcdef); + in.s01234567 = select(in.s89abcdef, in.s01234567, idx_sel); + res.s01234567 = select(res.s89abcdef, res.s01234567, CONVERT(idx_sel, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 8) )); + idx_sel.s0123 = VECTOR_PREDICATE(in.s0123, in.s4567) || (in.s0123 == in.s4567 && CONVERT(((res.s0123 < res.s4567)), VEC_DATA_TYPE(COND_DATA_TYPE, 4))); + in.s0123 = select(in.s4567, in.s0123, idx_sel.s0123); + res.s0123 = select(res.s4567, res.s0123, CONVERT(idx_sel.s0123, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 4) )); + idx_sel.s01 = (VECTOR_PREDICATE(in.s01, in.s23)) || (in.s01 == in.s23 && CONVERT(((res.s01 < res.s23)), VEC_DATA_TYPE(COND_DATA_TYPE, 2))); + in.s01 = select(in.s23, in.s01, idx_sel.s01); + res.s01 = select(res.s23, res.s01, CONVERT(idx_sel.s01, VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 2) )); + idx_sel.s0 = VECTOR_PREDICATE(in.s0, in.s1) || (in.s0 == in.s1 && CONVERT((res.s0 < res.s1), COND_DATA_TYPE)); + res.s0 = select(res.s1, res.s0, CONVERT(idx_sel.s0, DATA_TYPE_OUTPUT)); + *min_max_val = SCALAR_SELECT_OP(in.s0, in.s1); + *min_max_idx = res.s0; +} + + + +inline void scalar_compute_global_min_max(DATA_TYPE in_val, int idx, DATA_TYPE *out_min_max_val, DATA_TYPE_OUTPUT *out_idx) +{ +#if defined(ARG_MAX) + if(in_val > *out_min_max_val) +#else // defined(ARG_MAX) + if(in_val < *out_min_max_val) +#endif // defined(ARG_MAX) + { + *out_min_max_val = in_val; + *out_idx = idx; + } +} + +#if VEC_SIZE > 1 +#if VEC_SIZE == 16 + #define VECTORIZED_OP(min_max_val,min_max_idx,in,res) vectorized_compute_arg_min_max_16(min_max_val,min_max_idx,in,res) +#elif VEC_SIZE == 8 // #if VEC_SIZE == 16 + #define VECTORIZED_OP(min_max_val,min_max_idx,in,res) vectorized_compute_arg_min_max_8(min_max_val,min_max_idx,in,res) +#elif VEC_SIZE == 4 // # elif VEC_SIZE == 8 + #define VECTORIZED_OP(min_max_val,min_max_idx,in,res) vectorized_compute_arg_min_max_4(min_max_val,min_max_idx,in,res) +#elif VEC_SIZE == 2 // elif VEC_SIZE == 4 + #define VECTORIZED_OP(min_max_val,min_max_idx,in,res) vectorized_compute_arg_min_max_2(min_max_val,min_max_idx,in,res) +#else // elif VEC_SIZE == 2 + #error "Not supported" +#endif // #if VEC_SIZE == 16 + +inline VEC_DATA_TYPE(DATA_TYPE_OUTPUT, VEC_SIZE) init_idx_vector() +{ +#if VEC_SIZE == 16 + VEC_DATA_TYPE(DATA_TYPE_OUTPUT, VEC_SIZE) + vidx = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }; +#elif VEC_SIZE == 8 // #if VEC_SIZE == 16 + VEC_DATA_TYPE(DATA_TYPE_OUTPUT, VEC_SIZE) + vidx = { 0, 1, 2, 3, 4, 5, 6, 7 }; +#elif VEC_SIZE == 4 // elif VEC_SIZE == 8 + VEC_DATA_TYPE(DATA_TYPE_OUTPUT, VEC_SIZE) + vidx = { 0, 1, 2, 3 }; +#elif VEC_SIZE == 2 // elif VEC_SIZE == 4 + VEC_DATA_TYPE(DATA_TYPE_OUTPUT, VEC_SIZE) + vidx = { 0, 1 }; +#else // elif VEC_SIZE == 2 +#error "Not supported" +#endif // #if VEC_SIZE == 16 + return vidx; +} +#endif // VEC_SIZE > 1 + +/** This kernel performs reduction on x-axis. + * + * @note The input data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float + * @note The data type of the output must be passed at compile time using -DDATA_TYPE_OUTPUT: e.g. -DDATA_TYPE_OUTPUT=uint + * @note The data type used for the comparing indexe must be passed at compile type using -DCOND_DATA_TYPE: e.g -DCOND_DATA_TYPE=uint + * @note The height size must be passed at compile time using -DHEIGHT e.g. -DHEIGHT=128 + * + * @param[in] input_ptr Pointer to the source tensor. Supported data types: QASYMM8/QASYMM8_SIGNED/S32/F16/F32 + * @param[in] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] output_ptr The local buffer to hold sumed values. Supported data types: U32/S32 + * @param[in] output_stride_x Stride of the output 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 output 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 source tensor + */ +__kernel void arg_min_max_x( + IMAGE_DECLARATION(input), + IMAGE_DECLARATION(output)) +{ + __global DATA_TYPE *input_addr = (__global DATA_TYPE *)(input_ptr + input_offset_first_element_in_bytes + get_global_id(1) * input_stride_y); + __global DATA_TYPE_OUTPUT *output_addr = (__global DATA_TYPE_OUTPUT *)(output_ptr + output_offset_first_element_in_bytes + get_global_id(1) * output_stride_y); + + DATA_TYPE final_value = input_addr[0]; + DATA_TYPE_OUTPUT final_idx = 0; + +#if VEC_SIZE > 1 + VEC_DATA_TYPE(DATA_TYPE_OUTPUT, VEC_SIZE) + vidx = init_idx_vector(); + + int x = 0; + for(; x <= (WIDTH - VEC_SIZE); x += VEC_SIZE) + { + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + vals = VLOAD(VEC_SIZE)(0, (input_addr + x)); + DATA_TYPE local_min_max_value; + DATA_TYPE_OUTPUT local_min_max_idx; + + VECTORIZED_OP(&local_min_max_value, &local_min_max_idx, vals, vidx); + local_min_max_idx += x; + scalar_compute_global_min_max(local_min_max_value, local_min_max_idx, &final_value, &final_idx); + } +#endif // VEC_SIZE > 1 + +#if(WIDTH % VEC_SIZE) + LOOP_UNROLLING(int, j, 0, 1, WIDTH % VEC_SIZE, + { + scalar_compute_global_min_max(*(input_addr + j + x), j + x, &final_value, &final_idx); + }) +#endif // (WIDTH % VEC_SIZE) + + output_addr[0] = final_idx; +} +#endif // defined(WIDTH) + +#if defined(HEIGHT) +/** This kernel performs reduction on y-axis. + * + * @note The input data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float + * @note Leftover vector size has to be passed at compile time using -DVEC_SIZE_LEFTOVER. e.g. -DVEC_SIZE_LEFTOVER=3. It is defined as the remainder between the input's first dimension and VEC_SIZE + * @note The data type of the output must be passed at compile time using -DDATA_TYPE_OUTPUT: e.g. -DDATA_TYPE_OUTPUT=uint + * @note The height size must be passed at compile time using -DHEIGHT e.g. -DHEIGHT=128 + * + * @param[in] input_ptr Pointer to the source tensor. Supported data types: QASYMM8/QASYMM8_SIGNED/S32/F16/F32 + * @param[in] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] output_ptr The local buffer to hold sumed values. Supported data types: U32/S32 + * @param[in] output_stride_x Stride of the output 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 output 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 source tensor + */ +__kernel void arg_min_max_y( + IMAGE_DECLARATION(input), + IMAGE_DECLARATION(output)) +{ + const int x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); + __global uchar *input_addr = input_ptr + input_offset_first_element_in_bytes + x_offs * sizeof(DATA_TYPE) + get_global_id(1) * input_stride_y; + __global uchar *output_addr = output_ptr + output_offset_first_element_in_bytes + x_offs * sizeof(DATA_TYPE_OUTPUT) + get_global_id(1) * output_stride_y; + + VEC_TYPE_IN res = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input_addr), VEC_TYPE_IN); + + VEC_TYPE_OUT indx0 = 0; + for(DATA_TYPE_OUTPUT y = 1; y < HEIGHT; ++y) + { + VEC_TYPE_IN in = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(input_addr + y * input_stride_y)), VEC_TYPE_IN); + + VEC_TYPE_OUT cond_conv = CONVERT(CONDITION_TO_USE(in, res), VEC_TYPE_OUT); + indx0 = select(indx0, (VEC_TYPE_OUT)y, cond_conv); + res = select(res, in, CONDITION_TO_USE(in, res)); + } + + // Store result + STORE_VECTOR_SELECT(indx, DATA_TYPE_OUTPUT, output_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); +} +#endif // defined(HEIGHT) + +#if defined(DEPTH) && !defined(BATCH) +/** This kernel performs reduction on z-axis. + * + * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float + * @note Leftover vector size has to be passed at compile time using -DVEC_SIZE_LEFTOVER. e.g. -DVEC_SIZE_LEFTOVER=3. It is defined as the remainder between the input's first dimension and VEC_SIZE + * @note The depth size must be passed at compile time using -DDEPTH e.g. -DDEPTH=128 + * + * @param[in] input_ptr Pointer to the source tensor. Supported data types: QASYMM8/QASYMM8_SIGNED/S32/F16/F32 + * @param[in] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] output_ptr The local buffer to hold sumed values. Supported data types: U32/S32 + * @param[in] output_stride_x Stride of the output 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 output 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_stride_z Stride of the output tensor in Z dimension (in bytes) + * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes The offset of the first element in the source tensor + */ +__kernel void arg_min_max_z( + TENSOR3D_DECLARATION(input), + TENSOR3D_DECLARATION(output)) +{ + const int x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); + + __global uchar *input_addr = input_ptr + input_offset_first_element_in_bytes + x_offs * sizeof(DATA_TYPE) + get_global_id(1) * input_stride_y + get_global_id(2) * input_stride_z; + __global uchar *output_addr = output_ptr + output_offset_first_element_in_bytes + x_offs * sizeof(DATA_TYPE_OUTPUT) + get_global_id(1) * output_stride_y + get_global_id(2) * output_stride_z; + + VEC_TYPE_IN res = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input_addr), VEC_TYPE_IN); + + VEC_TYPE_OUT indx0 = 0; + for(DATA_TYPE_OUTPUT z = 1; z < DEPTH; ++z) + { + VEC_TYPE_IN in = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(input_addr + z * input_stride_z)), VEC_TYPE_IN); + + VEC_TYPE_OUT cond_conv = CONVERT(CONDITION_TO_USE(in, res), VEC_TYPE_OUT); + indx0 = select(indx0, (VEC_TYPE_OUT)z, cond_conv); + res = select(res, in, CONDITION_TO_USE(in, res)); + } + + // Store result + STORE_VECTOR_SELECT(indx, DATA_TYPE_OUTPUT, output_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); +} +#endif /* defined(DEPTH) && !defined(BATCH) */ + +#if defined(BATCH) && defined(DEPTH) +/** This kernel performs reduction on w-axis. + * + * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float + * @note Leftover vector size has to be passed at compile time using -DVEC_SIZE_LEFTOVER. e.g. -DVEC_SIZE_LEFTOVER=3. It is defined as the remainder between the input's first dimension and VEC_SIZE + * @note The batch size must be passed at compile time using -DBATCH e.g. -DBATCH=128 + * @note The depth size must be passed at compile time using -DBATCH e.g. -DDEPTH=128 + * + * @param[in] input_ptr Pointer to the source tensor. Supported data types: QASYMM8/QASYMM8_SIGNED/S32/F16/F32 + * @param[in] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] input_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] input_step_w input_stride_w * number of elements along W processed per workitem(in bytes) + * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] output_ptr The local buffer to hold sumed values. Supported data types: U32/S32 + * @param[in] output_stride_x Stride of the output 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 output 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_stride_z Stride of the output tensor in Z dimension (in bytes) + * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_stride_w Stride of the output tensor in W dimension (in bytes) + * @param[in] output_step_w output_stride_w * number of elements along W processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes The offset of the first element in the source tensor + */ +__kernel void arg_min_max_w( + TENSOR4D_DECLARATION(input), + TENSOR4D_DECLARATION(output)) +{ + const int x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); + + __global uchar *input_addr = input_ptr + input_offset_first_element_in_bytes + x_offs * sizeof(DATA_TYPE) + get_global_id(1) * input_stride_y + (get_global_id(2) % DEPTH) * input_stride_z + + (get_global_id(2) / DEPTH) * input_stride_w; + __global uchar *output_addr = output_ptr + output_offset_first_element_in_bytes + x_offs * sizeof(DATA_TYPE_OUTPUT) + get_global_id(1) * output_stride_y + (get_global_id( + 2) % DEPTH) * output_stride_z + (get_global_id(2) / DEPTH) * output_stride_w; + + VEC_TYPE_IN res = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input_addr), VEC_TYPE_IN); + + VEC_TYPE_OUT indx0 = 0; + for(DATA_TYPE_OUTPUT w = 1; w < BATCH; ++w) + { + VEC_TYPE_IN in = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(input_addr + w * input_stride_w)), VEC_TYPE_IN); + + VEC_TYPE_OUT cond_conv = CONVERT(CONDITION_TO_USE(in, res), VEC_TYPE_OUT); + indx0 = select(indx0, (VEC_TYPE_OUT)w, cond_conv); + res = select(res, in, CONDITION_TO_USE(in, res)); + } + + // Store result + STORE_VECTOR_SELECT(indx, DATA_TYPE_OUTPUT, output_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); +} +#endif /* defined(BATCH) && defined(DEPTH) */ +#endif // defined(VEC_SIZE) && defined(DATA_TYPE) && defined(DATA_TYPE_OUTPUT) |