From a07ce151674e28a3e755f1c48785b599f1d34827 Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Fri, 11 Oct 2019 17:38:50 +0100 Subject: COMPMID-2741: [CL] GEMMMatrixMultiplyReshaped clCreateKernel Error Using the accumulator type during the activation in gemm for mixed precision operations as some OpenCL compiler versions complain for mixing float data types for primitive built-ins. Change-Id: I1c4b394c57962aeebb541de572614a7e3ee3f223 Signed-off-by: Georgios Pinitas Reviewed-on: https://review.mlplatform.org/c/2080 Comments-Addressed: Arm Jenkins Reviewed-by: Michele Di Giorgio Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/gemm.cl | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/src/core/CL/cl_kernels/gemm.cl b/src/core/CL/cl_kernels/gemm.cl index 57a5af8ec2..66d0e10b71 100644 --- a/src/core/CL/cl_kernels/gemm.cl +++ b/src/core/CL/cl_kernels/gemm.cl @@ -2094,7 +2094,11 @@ __kernel void gemm_mm_reshaped_lhs_nt_rhs_t(IMAGE_DECLARATION(lhs), #endif // defined(BETA) #if defined(ACTIVATION_TYPE) +#if defined(MIXED_PRECISION) ACTIVATION_BLOCK(M0, ACTIVATION_TYPE, DATA_TYPE_ACCUMULATOR, c, A_VAL, B_VAL); +#else // defined(MIXED_PRECISION) + ACTIVATION_BLOCK(M0, ACTIVATION_TYPE, DATA_TYPE, c, A_VAL, B_VAL); +#endif // defined(MIXED_PRECISION) #endif // defined(ACTIVATION_TYPE) // Store output block @@ -2570,7 +2574,11 @@ __kernel void gemm_mm_reshaped_lhs_t_rhs_nt(IMAGE_DECLARATION(lhs), #endif // defined(BETA) #if defined(ACTIVATION_TYPE) +#if defined(MIXED_PRECISION) + ACTIVATION_BLOCK(M0, ACTIVATION_TYPE, DATA_TYPE_ACCUMULATOR, c, A_VAL, B_VAL); +#else // defined(MIXED_PRECISION) ACTIVATION_BLOCK(M0, ACTIVATION_TYPE, DATA_TYPE, c, A_VAL, B_VAL); +#endif // defined(MIXED_PRECISION) #endif // defined(ACTIVATION_TYPE) // Store output block -- cgit v1.2.1