aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/arg_min_max.cl
diff options
context:
space:
mode:
Diffstat (limited to 'src/core/CL/cl_kernels/arg_min_max.cl')
-rw-r--r--src/core/CL/cl_kernels/arg_min_max.cl65
1 files changed, 39 insertions, 26 deletions
diff --git a/src/core/CL/cl_kernels/arg_min_max.cl b/src/core/CL/cl_kernels/arg_min_max.cl
index 3f75377636..06dcc8ddde 100644
--- a/src/core/CL/cl_kernels/arg_min_max.cl
+++ b/src/core/CL/cl_kernels/arg_min_max.cl
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2019 ARM Limited.
+ * Copyright (c) 2019-2020 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -23,6 +23,19 @@
*/
#include "helpers.h"
+#if defined(FLOAT_DATA_TYPE)
+#define ISGREATER(x, y) isgreater(x, y)
+#define ISLESS(x, y) 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_DATA_TYPE(DATA_TYPE_SELECT, 16))0, (VEC_DATA_TYPE(DATA_TYPE_SELECT, 16)) - 1, x > y)
+#define ISLESS(x, y) select((VEC_DATA_TYPE(DATA_TYPE_SELECT, 16))0, (VEC_DATA_TYPE(DATA_TYPE_SELECT, 16)) - 1, 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)
@@ -31,7 +44,7 @@
#error "Unsupported reduction operation!"
#endif // defined(ARG_MAX)
-#if defined(DATA_TYPE_OUTPUT)
+#if defined(DATA_TYPE_OUTPUT) && defined(DATA_TYPE_SELECT)
#if defined(WIDTH)
#if defined(ARG_MIN)
#if defined(PREV_OUTPUT)
@@ -81,20 +94,20 @@ inline DATA_TYPE_OUTPUT arg_idx_min(__global const DATA_TYPE *input, const int x
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)
+ VEC_DATA_TYPE(DATA_TYPE_SELECT, 8)
idx_sel = (in.s01234567 <= in.s89abcdef);
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)));
+ idx_sel.s0123 = (in.s0123 < in.s4567) || (in.s0123 == in.s4567 && CONVERT((res.s0123 < res.s4567), VEC_DATA_TYPE(DATA_TYPE_SELECT, 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)));
+ idx_sel.s01 = (in.s01 < in.s23) || (in.s01 == in.s23 && CONVERT((res.s01 < res.s23), VEC_DATA_TYPE(DATA_TYPE_SELECT, 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));
+ idx_sel.s0 = (in.s0 < in.s1) || (in.s0 == in.s1 && CONVERT((res.s0 < res.s1), DATA_TYPE_SELECT));
res.s0 = select(res.s1, res.s0, CONVERT(idx_sel.s0, int));
return res.s0 + x_elem;
@@ -150,20 +163,20 @@ inline DATA_TYPE_OUTPUT arg_idx_max(__global const DATA_TYPE *input, const int x
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)
+ VEC_DATA_TYPE(DATA_TYPE_SELECT, 8)
idx_sel = (in.s01234567 >= in.s89abcdef);
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)));
+ idx_sel.s0123 = (in.s0123 > in.s4567) || (in.s0123 == in.s4567 && CONVERT((res.s0123 < res.s4567), VEC_DATA_TYPE(DATA_TYPE_SELECT, 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)));
+ idx_sel.s01 = (in.s01 > in.s23) || (in.s01 == in.s23 && CONVERT((res.s01 < res.s23), VEC_DATA_TYPE(DATA_TYPE_SELECT, 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));
+ idx_sel.s0 = (in.s0 > in.s1) || (in.s0 == in.s1 && CONVERT((res.s0 < res.s1), DATA_TYPE_SELECT));
res.s0 = select(res.s1, res.s0, CONVERT(idx_sel.s0, int));
return res.s0 + x_elem;
@@ -276,7 +289,7 @@ __kernel void arg_min_max_x(
*
* @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 of the intermediate results must be passed at compile time using -DDATA_TYPE_PROMOTED: e.g. -DDATA_TYPE_PROMOTED=uint
+ * @note The data type of the select results must be passed at compile time using -DDATA_TYPE_SELECT: e.g. -DDATA_TYPE_SELECT=int
* @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: S32/F16/F32
@@ -299,15 +312,15 @@ __kernel void arg_min_max_y(
Image src = CONVERT_TO_IMAGE_STRUCT(src);
Image output = CONVERT_TO_IMAGE_STRUCT(output);
- VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)
- res = CONVERT(vload16(0, (__global DATA_TYPE *)offset(&src, 0, 0)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16));
+ VEC_DATA_TYPE(DATA_TYPE, 16)
+ res = CONVERT(vload16(0, (__global DATA_TYPE *)offset(&src, 0, 0)), VEC_DATA_TYPE(DATA_TYPE, 16));
VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)
indx = 0;
for(unsigned int y = 1; y < HEIGHT; ++y)
{
- VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)
- in = CONVERT(vload16(0, (__global DATA_TYPE *)offset(&src, 0, y)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16));
+ VEC_DATA_TYPE(DATA_TYPE, 16)
+ in = CONVERT(vload16(0, (__global DATA_TYPE *)offset(&src, 0, y)), VEC_DATA_TYPE(DATA_TYPE, 16));
VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)
cond_conv = CONVERT(CONDITION_TO_USE(in, res), VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16));
@@ -324,7 +337,7 @@ __kernel void arg_min_max_y(
/** 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 data type of the select results must be passed at compile time using -DDATA_TYPE_SELECT: e.g. -DDATA_TYPE_SELECT=int
* @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: S32/F16/F32
@@ -351,15 +364,15 @@ __kernel void arg_min_max_z(
Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
- VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)
- res = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16));
+ VEC_DATA_TYPE(DATA_TYPE, 16)
+ res = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)), VEC_DATA_TYPE(DATA_TYPE, 16));
VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)
indx = 0;
for(DATA_TYPE_OUTPUT z = 1; z < DEPTH; ++z)
{
- VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)
- in = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, z)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16));
+ VEC_DATA_TYPE(DATA_TYPE, 16)
+ in = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, z)), VEC_DATA_TYPE(DATA_TYPE, 16));
VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)
cond_conv = CONVERT(CONDITION_TO_USE(in, res), VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16));
@@ -376,7 +389,7 @@ __kernel void arg_min_max_z(
/** 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 data type of the select results must be passed at compile time using -DDATA_TYPE_SELECT: e.g. -DDATA_TYPE_SELECT=int
* @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
*
@@ -408,15 +421,15 @@ __kernel void arg_min_max_w(
Tensor4D input = CONVERT_TO_TENSOR4D_STRUCT(input, DEPTH);
Tensor4D output = CONVERT_TO_TENSOR4D_STRUCT(output, DEPTH);
- VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)
- res = CONVERT(vload16(0, (__global DATA_TYPE *)tensor4D_offset(&input, 0, 0, 0, 0)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16));
+ VEC_DATA_TYPE(DATA_TYPE, 16)
+ res = CONVERT(vload16(0, (__global DATA_TYPE *)tensor4D_offset(&input, 0, 0, 0, 0)), VEC_DATA_TYPE(DATA_TYPE, 16));
VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)
indx = 0;
for(DATA_TYPE_OUTPUT w = 1; w < BATCH; ++w)
{
- VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)
- in = CONVERT(vload16(0, (__global DATA_TYPE *)tensor4D_offset(&input, 0, 0, 0, w)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16));
+ VEC_DATA_TYPE(DATA_TYPE, 16)
+ in = CONVERT(vload16(0, (__global DATA_TYPE *)tensor4D_offset(&input, 0, 0, 0, w)), VEC_DATA_TYPE(DATA_TYPE, 16));
VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)
cond_conv = CONVERT(CONDITION_TO_USE(in, res), VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16));
@@ -428,4 +441,4 @@ __kernel void arg_min_max_w(
vstore16(indx, 0, (__global DATA_TYPE_OUTPUT *)output.ptr);
}
#endif /* defined(BATCH) && defined(DEPTH) */
-#endif // defined(DATA_TYPE_OUTPUT) \ No newline at end of file
+#endif /* defined(DATA_TYPE_OUTPUT) && defined(DATA_TYPE_SELECT) */ \ No newline at end of file
generated by cgit v1.2.1 (git 2.23.0) at 2024-05-04 17:07:11 +0000