From b93f5de2676ce0608b283a8f10213ff1c0d2489e Mon Sep 17 00:00:00 2001 From: Gian Marco Iodice Date: Wed, 5 Jul 2017 15:48:39 +0100 Subject: COMPMID-417 - Fixed bug in gemm_interleave_16bit and gemm_interleave_32_bit due to the non non representable numbers in half and float Change-Id: I1590699e788a2cd3ed6e72de4afff75e6346f6d8 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/79800 Reviewed-by: Moritz Pflanzer Reviewed-by: Georgios Pinitas Tested-by: Kaizen --- src/core/CL/cl_kernels/gemm.cl | 48 +++++++++++++++++++++--------------------- 1 file changed, 24 insertions(+), 24 deletions(-) diff --git a/src/core/CL/cl_kernels/gemm.cl b/src/core/CL/cl_kernels/gemm.cl index 7ac421b7b6..46f1645aa7 100644 --- a/src/core/CL/cl_kernels/gemm.cl +++ b/src/core/CL/cl_kernels/gemm.cl @@ -143,22 +143,22 @@ __kernel void gemm_interleave4x4_32bit(IMAGE_DECLARATION(src), Image dst = CONVERT_TO_IMAGE_STRUCT(dst); /* Load values from Matrix A */ - float4 a0 = vload4(0, (__global float *)(offset(&src, 0, 0))); - float4 a1 = vload4(0, (__global float *)(offset(&src, 0, 1))); - float4 a2 = vload4(0, (__global float *)(offset(&src, 0, 2))); - float4 a3 = vload4(0, (__global float *)(offset(&src, 0, 3))); + uint4 a0 = vload4(0, (__global uint *)(offset(&src, 0, 0))); + uint4 a1 = vload4(0, (__global uint *)(offset(&src, 0, 1))); + uint4 a2 = vload4(0, (__global uint *)(offset(&src, 0, 2))); + uint4 a3 = vload4(0, (__global uint *)(offset(&src, 0, 3))); - float4 val0 = (float4)(a0.s0, a1.s0, a2.s0, a3.s0); - vstore4(val0, 0, ((__global float *)dst.ptr) + 0); + uint4 val0 = (uint4)(a0.s0, a1.s0, a2.s0, a3.s0); + vstore4(val0, 0, ((__global uint *)dst.ptr) + 0); - val0 = (float4)(a0.s1, a1.s1, a2.s1, a3.s1); - vstore4(val0, 0, ((__global float *)dst.ptr) + 4); + val0 = (uint4)(a0.s1, a1.s1, a2.s1, a3.s1); + vstore4(val0, 0, ((__global uint *)dst.ptr) + 4); - val0 = (float4)(a0.s2, a1.s2, a2.s2, a3.s2); - vstore4(val0, 0, ((__global float *)dst.ptr) + 8); + val0 = (uint4)(a0.s2, a1.s2, a2.s2, a3.s2); + vstore4(val0, 0, ((__global uint *)dst.ptr) + 8); - val0 = (float4)(a0.s3, a1.s3, a2.s3, a3.s3); - vstore4(val0, 0, ((__global float *)dst.ptr) + 12); + val0 = (uint4)(a0.s3, a1.s3, a2.s3, a3.s3); + vstore4(val0, 0, ((__global uint *)dst.ptr) + 12); } /** This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values @@ -184,22 +184,22 @@ __kernel void gemm_interleave4x4_16bit(IMAGE_DECLARATION(src), Image dst = CONVERT_TO_IMAGE_STRUCT(dst); /* Load values from Matrix A */ - half8 a0 = vload8(0, (__global half *)(offset(&src, 0, 0))); - half8 a1 = vload8(0, (__global half *)(offset(&src, 0, 1))); - half8 a2 = vload8(0, (__global half *)(offset(&src, 0, 2))); - half8 a3 = vload8(0, (__global half *)(offset(&src, 0, 3))); + ushort8 a0 = vload8(0, (__global ushort *)(offset(&src, 0, 0))); + ushort8 a1 = vload8(0, (__global ushort *)(offset(&src, 0, 1))); + ushort8 a2 = vload8(0, (__global ushort *)(offset(&src, 0, 2))); + ushort8 a3 = vload8(0, (__global ushort *)(offset(&src, 0, 3))); - half8 val0 = (half8)((half4)(a0.s0, a1.s0, a2.s0, a3.s0), (half4)(a0.s1, a1.s1, a2.s1, a3.s1)); - vstore8(val0, 0, ((__global half *)dst.ptr) + 0); + ushort8 val0 = (ushort8)((ushort4)(a0.s0, a1.s0, a2.s0, a3.s0), (ushort4)(a0.s1, a1.s1, a2.s1, a3.s1)); + vstore8(val0, 0, ((__global ushort *)dst.ptr) + 0); - val0 = (half8)((half4)(a0.s2, a1.s2, a2.s2, a3.s2), (half4)(a0.s3, a1.s3, a2.s3, a3.s3)); - vstore8(val0, 0, ((__global half *)dst.ptr) + 8); + val0 = (ushort8)((ushort4)(a0.s2, a1.s2, a2.s2, a3.s2), (ushort4)(a0.s3, a1.s3, a2.s3, a3.s3)); + vstore8(val0, 0, ((__global ushort *)dst.ptr) + 8); - val0 = (half8)((half4)(a0.s4, a1.s4, a2.s4, a3.s4), (half4)(a0.s5, a1.s5, a2.s5, a3.s5)); - vstore8(val0, 0, ((__global half *)dst.ptr) + 16); + val0 = (ushort8)((ushort4)(a0.s4, a1.s4, a2.s4, a3.s4), (ushort4)(a0.s5, a1.s5, a2.s5, a3.s5)); + vstore8(val0, 0, ((__global ushort *)dst.ptr) + 16); - val0 = (half8)((half4)(a0.s6, a1.s6, a2.s6, a3.s6), (half4)(a0.s7, a1.s7, a2.s7, a3.s7)); - vstore8(val0, 0, ((__global half *)dst.ptr) + 24); + val0 = (ushort8)((ushort4)(a0.s6, a1.s6, a2.s6, a3.s6), (ushort4)(a0.s7, a1.s7, a2.s7, a3.s7)); + vstore8(val0, 0, ((__global ushort *)dst.ptr) + 24); } /** This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values -- cgit v1.2.1