diff options
| author | SiCong Li <sicong.li@arm.com> | 2020-10-20 18:07:27 +0100 |
|---|---|---|
| committer | SiCong Li <sicong.li@arm.com> | 2020-10-21 11:54:24 +0000 |
| commit | ed5fb39d1d9d3e56d26b621cd1d56ceb39270701 (patch) | |
| tree | 3a1b417e67a8ebc88ff3fdff5ff7fce78d6808a8 | |
| parent | 9f8937568c65863e2bc080a1f586068153289b83 (diff) | |
| download | ComputeLibrary-ed5fb39d1d9d3e56d26b621cd1d56ceb39270701.tar.gz | |
COMPMID-3599 Fix CLGEMMLowpMatrixMultiplyNativeKernel nightly failures
Guard the kernel with all required compile-time arguments, otherwise
the kernel might be wrongly included when compiling for other kernels
which don't have the required compile-time arguments, resulting in
mysterious kernel build errors.
Change-Id: Ib45b46a5ab14e6dc6a415c0466cf9a5963452364
Signed-off-by: SiCong Li <sicong.li@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4224
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/gemmlowp.cl | 9 |
1 files changed, 3 insertions, 6 deletions
diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl index c962d3c20e..da92548634 100644 --- a/src/core/CL/cl_kernels/gemmlowp.cl +++ b/src/core/CL/cl_kernels/gemmlowp.cl @@ -916,7 +916,7 @@ __kernel void gemmlowp_mm_reshaped_only_rhs_t_fused_output_stage_fixedpoint(IMAG #endif // defined(RESULT_OFFSET) && defined(RESULT_SHIFT) && defined(RESULT_MULTIPLIER) #endif // defined(M0) && defined(N0) && defined(K0) && defined(H0) && defined(DATA_TYPE) && defined(K) -#if defined(M0) && defined(N0) && defined(K0) && defined(K) +#if defined(M0) && defined(N0) && defined(K0) && defined(K) && defined(PARTIAL_STORE_M0) && defined(PARTIAL_STORE_N0) /** This OpenCL kernel computes the matrix multiplication between 2 matrices. * The LHS matrix is NOT reshaped @@ -997,7 +997,6 @@ __kernel void gemmlowp_mm_native(IMAGE_DECLARATION(lhs), // Compute RHS matrix address uint rhs_offset = rhs_offset_first_element_in_bytes + x * N0 * sizeof(DATA_TYPE); - #if defined(MATRIX_B_DEPTH) // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 rhs_offset += (z % MATRIX_B_DEPTH) * rhs_stride_z; @@ -1077,7 +1076,6 @@ __kernel void gemmlowp_mm_native(IMAGE_DECLARATION(lhs), __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(int)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * dst_stride_y); - REPEAT_VAR_INIT_TO_CONST(M0, uint, zout, 0); //uint zout0=0,zout1=0,zout2=0,... zout7=0; #if defined(REINTERPRET_OUTPUT_AS_3D) @@ -1097,11 +1095,10 @@ __kernel void gemmlowp_mm_native(IMAGE_DECLARATION(lhs), const bool cond_y = y == 0; const bool cond_x = ((x + 1) * N0 >= N); - - // Store output block + // Store output block STORE_BLOCK_BOUNDARY_AWARE(M0, N0, int, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, N, cond_y, cond_x); } -#endif // defined(M0) && defined(N0) && defined(K0) && defined(K) +#endif // defined(M0) && defined(N0) && defined(K0) && defined(K) && defined(PARTIAL_STORE_M0) && defined(PARTIAL_STORE_N0) #if defined(COLS_A) /** OpenCL kernel used to compute the row-vectors of sums of all the entries in each row of Matrix A. |