diff options
| author | Gian Marco Iodice <gianmarco.iodice@arm.com> | 2017-07-03 12:33:49 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-09-17 14:15:39 +0100 |
| commit | 368da83fdd7406d629e8cca64f3eb0af05437419 (patch) | |
| tree | fadac4142651cb0f86b997c06cbabb1bec622aae /src/core/CL/cl_kernels/gemm.cl | |
| parent | adffa30de9292c96bf29ff0697ac573270046612 (diff) | |
| download | ComputeLibrary-368da83fdd7406d629e8cca64f3eb0af05437419.tar.gz | |
COMPMID-420, COMPMID-414 - Port CLConvolutionLayer and CLFullyConnectedLayer to use 8 bit fixed point
Change-Id: I1cb1b4d7711ad7b569ee691e13a5df1b3430292b
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/79565
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/gemm.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/gemm.cl | 11 |
1 files changed, 9 insertions, 2 deletions
diff --git a/src/core/CL/cl_kernels/gemm.cl b/src/core/CL/cl_kernels/gemm.cl index 46f1645aa7..db15720ad0 100644 --- a/src/core/CL/cl_kernels/gemm.cl +++ b/src/core/CL/cl_kernels/gemm.cl @@ -21,9 +21,12 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ -#include "fixed_point.h" #include "helpers.h" +#ifdef FIXED_POINT_POSITION +#include "fixed_point.h" +#endif // FIXED_POINT_POSITION + /** This OpenCL kernel computes the "vector" 1x4 transposition of input matrix * * @param[in] src_ptr Pointer to the source matrix. Supported data types: U32/S32/F32 @@ -274,7 +277,11 @@ __kernel void gemm_accumulate_biases( accum_value = vload16(0, (__global DATA_TYPE *)accum.ptr); VEC_DATA_TYPE(DATA_TYPE, 16) biases_value = vload16(0, (__global DATA_TYPE *)biases.ptr); - accum_value = biases_value + accum_value; +#ifdef FIXED_POINT_POSITION + accum_value = ADD_SAT_OP_EXPAND(biases_value, accum_value, DATA_TYPE, 16); +#else // FIXED_POINT_POSITION + accum_value = biases_value + accum_value; +#endif // FIXED_POINT_POSITION // Store result in the accummulate buffer vstore16(accum_value, 0, (__global DATA_TYPE *)accum.ptr); |