From d051e97e36b9981f411093904cc019c2c7f9ac75 Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Wed, 20 Jun 2018 11:46:42 +0100 Subject: COMPMID-811 Add NHWC data format support for CL depthwise convolution Change-Id: I574f7945f0be009c638d860028bce8b52b4120fd Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/136484 Tested-by: Jenkins Reviewed-by: Gian Marco Iodice --- src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl | 5 +++++ 1 file changed, 5 insertions(+) (limited to 'src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl') diff --git a/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl index b58dc7af72..ae87420774 100644 --- a/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl +++ b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl @@ -296,7 +296,12 @@ __kernel void output_stage_quantized( #if defined(HAS_BIAS) // Load and add bias +#if defined(NCHW) int bias_value = *((__global int *)(vector_offset(&bias, get_global_id(2)))); +#else // defined(NCHW) + int16 bias_value = vload16(0, ((__global int *)(vector_offset(&bias, get_global_id(0) * 16)))); +#endif // defined(NCHW) + vals += (int16)(bias_value); #endif //defined(HAS_BIAS) -- cgit v1.2.1