diff options
| author | Viet-Hoa Do <viet-hoa.do@arm.com> | 2022-10-05 15:04:23 +0100 |
|---|---|---|
| committer | Pablo Marquez Tello <pablo.tello@arm.com> | 2022-10-07 10:08:02 +0000 |
| commit | 2ffab6dc8afb3716c189a30f75c33b2f2c35a6d4 (patch) | |
| tree | 0f951fc4e56b62e244123070e48caa468951f9f9 | |
| parent | 842ad211c11417ba456a2dca7e89988db98eb256 (diff) | |
| download | ComputeLibrary-2ffab6dc8afb3716c189a30f75c33b2f2c35a6d4.tar.gz | |
Workaround CL compiler issue on FP16
Resolves: COMPMID-5600
Signed-off-by: Viet-Hoa Do <viet-hoa.do@arm.com>
Change-Id: I5196d1639c48d0b8a116d47ed1d6c7334dc8f41e
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/8374
Benchmark: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Jakub Sujak <jakub.sujak@arm.com>
Reviewed-by: Pablo Marquez Tello <pablo.tello@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl | 4 |
1 files changed, 4 insertions, 0 deletions
diff --git a/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl b/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl index 8a8458798e..345469063a 100644 --- a/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl +++ b/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl @@ -158,7 +158,11 @@ __kernel void dwc_native_fp_nhwc( { LOOP_UNROLLING(int, xk, 0, 1, _IWEI_WIDTH, { +#if GPU_ARCH == GPU_ARCH_MIDGARD + c[m0].v += a[xk + m0].v * b[xk].v; +#else // GPU_ARCH == GPU_ARCH_MIDGARD c[m0].v = fma(a[xk + m0].v, b[xk].v, c[m0].v); +#endif // GPU_ARCH == GPU_ARCH_MIDGARD }) }) } |