/* * 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)