diff options
| author | Georgios Pinitas <georgios.pinitas@arm.com> | 2019-10-11 17:38:50 +0100 |
|---|---|---|
| committer | Michalis Spyrou <michalis.spyrou@arm.com> | 2019-10-16 10:49:36 +0000 |
| commit | a07ce151674e28a3e755f1c48785b599f1d34827 (patch) | |
| tree | 7c4f699f990f1b58975867bf5f693b681e3efd76 | |
| parent | db8485ac24135f17e9882c76196924435abc064f (diff) | |
| download | ComputeLibrary-a07ce151674e28a3e755f1c48785b599f1d34827.tar.gz | |
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 <georgios.pinitas@arm.com>
Reviewed-on: https://review.mlplatform.org/c/2080
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/gemm.cl | 8 |
1 files changed, 8 insertions, 0 deletions
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 |