Revert "COMPMID-1816: Use parallel reduction on 0 axis in CL ARG_MIN/ARG_MAX"
This reverts commit 36debd4472839997fd3b6ec9d58530d95e3c17de.
Change-Id: I45d63a63cebfd4582214b488027c2b14f492fdc1
Reviewed-on: https://review.mlplatform.org/c/2120
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
diff --git a/src/core/CL/cl_kernels/reduction_operation.cl b/src/core/CL/cl_kernels/reduction_operation.cl
index db034f0..5a4bb9f 100644
--- a/src/core/CL/cl_kernels/reduction_operation.cl
+++ b/src/core/CL/cl_kernels/reduction_operation.cl
@@ -36,10 +36,6 @@
#endif // defined(WIDTH)
#endif // FLOAT_DATA_TYPE
-#if defined(DATA_TYPE)
-
-#if defined(OPERATION) && defined(WIDTH)
-
/** Calculate square sum of a vector
*
* @param[in] input Pointer to the first pixel.
@@ -95,112 +91,10 @@
return (in.s0 * in.s1);
}
-
-#if defined(DATA_TYPE_OUTPUT)
-
-#if defined(ARG_MAX)
-/** Find index maximum value of a vector
- *
- * @param[in] input Pointer to the first value.
- *
- * @return index of the vector.
- */
-inline DATA_TYPE_OUTPUT arg_idx_max(__global const DATA_TYPE *input, const int x_idx)
-{
-#if defined(MULTI_ACCESS_X)
-
- int x_elem = x_idx * 16;
- const int x_goback = select(0, 16 - WIDTH % 16, x_elem + 16 > WIDTH);
- x_elem -= x_goback;
-
- VEC_DATA_TYPE(DATA_TYPE, 16)
- in = vload16(0, input - x_goback);
- VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)
- res = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 };
-
- VEC_DATA_TYPE(COND_DATA_TYPE, 8)
- idx_sel = (in.s01234567 > in.s89abcdef) || (in.s01234567 == in.s89abcdef && CONVERT((res.s01234567 < res.s89abcdef), VEC_DATA_TYPE(COND_DATA_TYPE, 8)));
- in.s01234567 = select(in.s89abcdef, in.s01234567, idx_sel);
- res.s01234567 = select(res.s89abcdef, res.s01234567, CONVERT(idx_sel, int8));
-
- idx_sel.s0123 = (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, int4));
-
- idx_sel.s01 = (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, int2));
-
- idx_sel.s0 = (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, int));
-
- return res.s0 + x_elem;
-#else // defined(MULTI_ACCESS_X)
-
- DATA_TYPE_OUTPUT res = 0;
- for(DATA_TYPE_OUTPUT x_v = res + 1; x_v < WIDTH; ++x_v)
- {
- res = select(res, x_v, *(input + x_v) > *(input + res));
- }
-
- return res;
-#endif // defined(MULTI_ACCESS_X)
-}
-#endif // defined(ARG_MAX)
-
-#if defined(ARG_MIN)
-/** Find index minimum value of a vector
- *
- * @param[in] input Pointer to the first value.
- *
- * @return index of the vector.
- */
-inline DATA_TYPE_OUTPUT arg_idx_min(__global const DATA_TYPE *input, const int x_idx)
-{
-#if defined(MULTI_ACCESS_X)
-
- int x_elem = x_idx * 16;
- const int x_goback = select(0, 16 - WIDTH % 16, x_elem + 16 > WIDTH);
- x_elem -= x_goback;
-
- VEC_DATA_TYPE(DATA_TYPE, 16)
- in = vload16(0, input - x_goback);
- VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)
- res = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 };
-
- VEC_DATA_TYPE(COND_DATA_TYPE, 8)
- idx_sel = (in.s01234567 < in.s89abcdef) || (in.s01234567 == in.s89abcdef && CONVERT((res.s01234567 < res.s89abcdef), VEC_DATA_TYPE(COND_DATA_TYPE, 8)));
- in.s01234567 = select(in.s89abcdef, in.s01234567, idx_sel);
- res.s01234567 = select(res.s89abcdef, res.s01234567, CONVERT(idx_sel, int8));
-
- idx_sel.s0123 = (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, int4));
-
- idx_sel.s01 = (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, int2));
-
- idx_sel.s0 = (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, int));
-
- return res.s0 + x_elem;
-#else // defined(MULTI_ACCESS_X)
-
- DATA_TYPE_OUTPUT res = 0;
- for(DATA_TYPE_OUTPUT x_v = res + 1; x_v < WIDTH; ++x_v)
- {
- res = select(res, x_v, *(input + x_v) < * (input + res));
- }
- return res;
-#endif // defined(MULTI_ACCESS_X)
-}
-#endif // defined(ARG_MIN)
-
+#if defined(OPERATION)
/** This kernel performs parallel reduction given an operation on x-axis.
*
* @note The 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 operation we want to perform must be passed at compile time using -DOPERATION e.g. -DOPERATION=square_sum
* @note The mean flag must be passed at compile time using -DMEAN if we want to compute the mean value
* @note The product flag must be passed at compile time using -DPROD if we want to compute the product, otherwise sum will be used
@@ -223,7 +117,7 @@
__kernel void reduction_operation_x(
IMAGE_DECLARATION(src),
IMAGE_DECLARATION(partial_res),
- __local DATA_TYPE_OUTPUT *local_results)
+ __local DATA_TYPE *local_results)
{
Image src = CONVERT_TO_IMAGE_STRUCT(src);
Image partial_res = CONVERT_TO_IMAGE_STRUCT(partial_res);
@@ -231,17 +125,9 @@
unsigned int lsize = get_local_size(0);
unsigned int lid = get_local_id(0);
- const uint x_idx = get_global_id(0);
- const uint y_idx = get_global_id(1);
-
for(unsigned int y = 0; y < get_local_size(1); ++y)
{
-#if defined(ARG_MAX) || defined(ARG_MIN)
- local_results[lid] = OPERATION((__global DATA_TYPE *)offset(&src, 0, y), x_idx);
-#else // defined(ARG_MAX) || defined(ARG_MIN)
- local_results[lid] = OPERATION((__global DATA_TYPE *)offset(&src, 0, y));
-#endif // defined(ARG_MAX) || defined(ARG_MIN)
-
+ local_results[lid] = OPERATION((__global DATA_TYPE *)offset(&src, 0, y));
barrier(CLK_LOCAL_MEM_FENCE);
// Perform parallel reduction
@@ -251,26 +137,9 @@
{
#if defined(PROD)
local_results[lid] *= local_results[lid + i];
-#elif defined(ARG_MAX)
- __global DATA_TYPE *src_in_row = src_ptr + src_offset_first_element_in_bytes + y_idx * src_step_y;
- DATA_TYPE tmp0 = *(src_in_row + local_results[lid]);
- DATA_TYPE tmp1 = *(src_in_row + local_results[lid + i]);
- local_results[lid] = select(
- local_results[lid],
- local_results[lid + i],
- ((tmp0 == tmp1) && (local_results[lid + i] < local_results[lid])) || (tmp0 < tmp1));
-
-#elif defined(ARG_MIN)
- __global DATA_TYPE *src_in_row = src_ptr + src_offset_first_element_in_bytes + y_idx * src_step_y;
- DATA_TYPE tmp0 = *(src_in_row + local_results[lid]);
- DATA_TYPE tmp1 = *(src_in_row + local_results[lid + i]);
- local_results[lid] = select(
- local_results[lid],
- local_results[lid + i],
- ((tmp0 == tmp1) && (local_results[lid + i] < local_results[lid])) || (tmp0 > tmp1));
-#else // !defined(PROD) && !defined(ARG_MAX) && !defined(ARG_MIN)
+#else // !defined(PROD)
local_results[lid] += local_results[lid + i];
-#endif // !defined(PROD) && !defined(ARG_MAX) && !defined(ARG_MIN)
+#endif // defined(PROD)
}
barrier(CLK_LOCAL_MEM_FENCE);
}
@@ -283,22 +152,16 @@
local_results[0] /= WIDTH;
}
#endif // defined(MEAN) && defined(WIDTH)
- ((__global DATA_TYPE_OUTPUT *)offset(&partial_res, get_group_id(0), y))[0] = local_results[0];
+ ((__global DATA_TYPE *)offset(&partial_res, get_group_id(0), y))[0] = local_results[0];
}
}
}
-
-#endif // defined(DATA_TYPE_OUTPUT)
-
-#endif // defined(OPERATION) && defined(WIDTH)
-
-#if defined(DATA_TYPE_PROMOTED)
+#endif // defined(OPERATION)
#if defined(WIDTH)
/** This kernel performs reduction on x-axis. (Non parallel)
*
* @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
- * @note The data type of the intermediate results must be passed at compile time using -DDATA_TYPE_PROMOTED: e.g. -DDATA_TYPE_PROMOTED=uint
* @note The width size must be passed at compile time using -DWIDTH e.g. -DWIDTH=128
* @note The product flag must be passed at compile time using -DPROD if we want to compute the product, otherwise sum will be used
* @note In case of ARG_MIN and ARG_MAX the condition data type must be passed at compile time using -DCOND_DATA_TYPE e.g. -DCOND_DATA_TYPE=short
@@ -328,7 +191,13 @@
for(unsigned int x = 1; x < WIDTH; ++x)
{
DATA_TYPE_PROMOTED in = *((__global DATA_TYPE *)vector_offset(&src, x));
-#if defined(MIN)
+#if defined(ARG_MAX)
+ indx = select(indx, x, ISGREATER(in, res));
+ res = select(res, in, CONVERT(ISGREATER(in, res), COND_DATA_TYPE));
+#elif defined(ARG_MIN)
+ indx = select(indx, x, ISLESS(in, res));
+ res = select(res, in, CONVERT(ISLESS(in, res), COND_DATA_TYPE));
+#elif defined(MIN)
res = select(res, in, CONVERT(ISLESS(in, res), COND_DATA_TYPE));
#elif defined(MAX)
res = select(res, in, CONVERT(ISGREATER(in, res), COND_DATA_TYPE));
@@ -357,7 +226,6 @@
/** 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 The data type of the intermediate results must be passed at compile time using -DDATA_TYPE_PROMOTED: e.g. -DDATA_TYPE_PROMOTED=uint
* @note The height size must be passed at compile time using -DHEIGHT e.g. -DHEIGHT=128
*
* @param[in] src_ptr Pointer to the source tensor. Supported data types: QASYMM8/S32/F16/F32
@@ -435,7 +303,6 @@
/** 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 The data type of the intermediate results must be passed at compile time using -DDATA_TYPE_PROMOTED: e.g. -DDATA_TYPE_PROMOTED=uint
* @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/S32/F16/F32
@@ -533,9 +400,8 @@
/** 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 The data type of the intermediate results must be passed at compile time using -DDATA_TYPE_PROMOTED: e.g. -DDATA_TYPE_PROMOTED=uint
* @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 -DDEPTH e.g. -DDEPTH=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/S32/F16/F32
* @param[in] input_stride_x Stride of the source tensor in X dimension (in bytes)
@@ -616,7 +482,3 @@
#endif // defined(ARG_MAX) || defined(ARG_MIN)
}
#endif /* defined(BATCH) && defined(DEPTH) */
-
-#endif /* defined(DATA_TYPE_PROMOTED) */
-
-#endif /* defined(DATA_TYPE) */
\ No newline at end of file