diff options
author | Gian Marco Iodice <gianmarco.iodice@arm.com> | 2023-01-10 12:46:29 +0000 |
---|---|---|
committer | Gian Marco Iodice <gianmarco.iodice@arm.com> | 2023-01-12 15:09:56 +0000 |
commit | bc672082ae31778164ed3ec23b7a4a8f1a8dc454 (patch) | |
tree | ed776c6190e57a95d27577ba4da2570a25242a40 /src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl | |
parent | 6bcdc578a388782f5ec80ec348c5dd3f5c1f8363 (diff) | |
download | ComputeLibrary-bc672082ae31778164ed3ec23b7a4a8f1a8dc454.tar.gz |
Update the heuristic for CLDepthwiseConvolutionNative kernel
- Use T_LOAD2D_INDIRECT macro instead of T_LOAD_NHWC_WITH_DILATION in
the depthwise convolution opencl kernels
- Update the heuristic for Arm® Mali™-G77
Resolves COMPMID-5716
Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Change-Id: I32d375b220e04bf05f5d8f0af2231bde600f0665
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/8930
Benchmark: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Jakub Sujak <jakub.sujak@arm.com>
Reviewed-by: Viet-Hoa Do <viet-hoa.do@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl')
-rw-r--r-- | src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl | 18 |
1 files changed, 16 insertions, 2 deletions
diff --git a/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl b/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl index 2d255e5b61..e502d721d5 100644 --- a/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl +++ b/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl @@ -180,8 +180,22 @@ __kernel void dwc_native_quantized_nhwc( a[i].v = ZERO_VALUE; }) - // Load tile from the src tensor (TILE A) - T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, _IM0_A, _IN0_A, SRC_TENSOR_TYPE, src, bout, yi + yk * DILATION_Y, xi, (cout / DEPTH_MULTIPLIER), src_w, src_h, DILATION_X, 1, _IBOUNDARY_CHECK, a); + TILE(int, 1, _IM0_A, my); + + LOOP_UNROLLING(int, xk_i, 0, 1, _IM0_A, + { + int x_s = xi + xk_i * (DILATION_X); + int y_s = yi + yk * (DILATION_Y); + my[0].s[xk_i] = x_s + y_s * SRC_WIDTH; + my[0].s[xk_i] = my[0].s[xk_i] + bout * (int)(SRC_WIDTH * SRC_HEIGHT); + my[0].s[xk_i] = select(-1, my[0].s[xk_i], x_s >= 0); + my[0].s[xk_i] = select(-1, my[0].s[xk_i], x_s < SRC_WIDTH); + my[0].s[xk_i] = select(-1, my[0].s[xk_i], y_s >= 0); + my[0].s[xk_i] = select(-1, my[0].s[xk_i], y_s < SRC_HEIGHT); + }) + + // Load tile from the src tensor + T_LOAD2D_INDIRECT(SRC_DATA_TYPE, _IM0_A, _IN0_A, SRC_TENSOR_TYPE, src, (cout / DEPTH_MULTIPLIER), src_stride_y, my, a); TILE(WEI_DATA_TYPE, _IM0_B, _IN0_B, b); |