diff options
author | Giorgio Arena <giorgio.arena@arm.com> | 2020-10-12 11:53:51 +0100 |
---|---|---|
committer | Giorgio Arena <giorgio.arena@arm.com> | 2020-10-12 14:06:24 +0000 |
commit | d056e574f60ca731b2d078e56c6baca5a6c642ac (patch) | |
tree | 30a47f1df1341b0462f344f4234bcf5dbf5d4360 /src/core/CL/cl_kernels/yolo_layer.cl | |
parent | ff9612cad3288ab96f94e86485e591401753f56b (diff) | |
download | ComputeLibrary-d056e574f60ca731b2d078e56c6baca5a6c642ac.tar.gz |
COMPMID-3826 ArmNN Nightly failing for CL
Change-Id: I09f557b5cecafc669e12764e8592457212168d62
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4131
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/yolo_layer.cl')
-rw-r--r-- | src/core/CL/cl_kernels/yolo_layer.cl | 14 |
1 files changed, 7 insertions, 7 deletions
diff --git a/src/core/CL/cl_kernels/yolo_layer.cl b/src/core/CL/cl_kernels/yolo_layer.cl index 2a15a32e2a..fe7b5cbb55 100644 --- a/src/core/CL/cl_kernels/yolo_layer.cl +++ b/src/core/CL/cl_kernels/yolo_layer.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 Arm Limited. + * Copyright (c) 2018-2020 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -21,13 +21,14 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ -#if defined(DATA_TYPE) && defined(SELECT_DATA_TYPE) && defined(ACTIVATION_TYPE) && defined(NUM_CLASSES) && defined(VEC_SIZE) +#if defined(DATA_TYPE) && defined(ACTIVATION_TYPE) && defined(NUM_CLASSES) && defined(VEC_SIZE) #include "activation_float_helpers.h" +#define SELECT_TYPE SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE) + #if VEC_SIZE != 1 #define TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) -#define SELECT_TYPE VEC_DATA_TYPE(SELECT_DATA_TYPE, VEC_SIZE) /** This performs a YOLO partial activation function for NCHW data layout * @@ -79,7 +80,7 @@ __kernel void yolo_layer_nchw( { // Load data TYPE data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input.ptr); - data = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, data, A_VAL, B_VAL); // select(1.0f, ACTIVATION_OP(ACTIVATION_TYPE, data), (SELECT_TYPE)activate); + data = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, data, A_VAL, B_VAL); // select(1.0f, ACTIVATION_OP(ACTIVATION_TYPE, data), (SELECT_TYPE)activate); // Store result VSTORE(VEC_SIZE) @@ -100,7 +101,6 @@ __kernel void yolo_layer_nchw( #else // VEC_SIZE != 1 -#define SELECT_TYPE SELECT_DATA_TYPE /** This performs a YOLO partial activation function for NCHW data layout * * @note In order to perform the activation function "in-place", the pre-processor -DIN_PLACE must be passed at compile time @@ -151,7 +151,7 @@ __kernel void yolo_layer_nhwc( { // Load data DATA_TYPE data = *((__global DATA_TYPE *)input.ptr); - data = select(data, ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, data, A_VAL, B_VAL), (SELECT_TYPE)activate); + data = select(data, ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, data, A_VAL, B_VAL), (SELECT_TYPE)activate); // Store result *((__global DATA_TYPE *)output.ptr) = data; @@ -169,4 +169,4 @@ __kernel void yolo_layer_nhwc( } #endif // VEC_SIZE != 1 -#endif // defined(DATA_TYPE) && defined(SELECT_DATA_TYPE) && defined(ACTIVATION_TYPE) && defined(NUM_CLASSES) && defined(VEC_SIZE) +#endif // defined(DATA_TYPE) && defined(ACTIVATION_TYPE) && defined(NUM_CLASSES) && defined(VEC_SIZE) |