aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/common/gemm.cl
diff options
context:
space:
mode:
authorGian Marco Iodice <gianmarco.iodice@arm.com>2021-10-15 10:23:24 +0100
committerGian Marco Iodice <gianmarco.iodice@arm.com>2021-10-18 13:26:35 +0000
commitc9cecc0e565e7b4978cecc92e03e6c93bb8d0cb9 (patch)
treee5e26af2cd1d8537c528a401ddbc729ae7beb2bb /src/core/CL/cl_kernels/common/gemm.cl
parent36118524d2f387be53dc95e5eebabfcb3ec21f31 (diff)
downloadComputeLibrary-c9cecc0e565e7b4978cecc92e03e6c93bb8d0cb9.tar.gz
Remove legacy GeMM kernels on OpenCL
Resolves COMPMID-4446 Change-Id: I1d3c2391b67681f4d3af440826aa95b47a1288a6 Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/6444 Reviewed-by: Giorgio Arena <giorgio.arena@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/common/gemm.cl')
-rw-r--r--src/core/CL/cl_kernels/common/gemm.cl7
1 files changed, 2 insertions, 5 deletions
diff --git a/src/core/CL/cl_kernels/common/gemm.cl b/src/core/CL/cl_kernels/common/gemm.cl
index 87921f51fd..431c97becc 100644
--- a/src/core/CL/cl_kernels/common/gemm.cl
+++ b/src/core/CL/cl_kernels/common/gemm.cl
@@ -4141,6 +4141,7 @@ __kernel void gemm_mm_native(IMAGE_DECLARATION(lhs),
REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE, N0), c, 0); //VEC_DATA_TYPE(DATA_TYPE, N0) c0=0,c1=0,c2=0,... c(M0-1)=0;
int i = 0;
+#if K0 > 1
for(; i <= (K - K0); i += K0)
{
// Supported cases (M0, K0):
@@ -4186,7 +4187,7 @@ __kernel void gemm_mm_native(IMAGE_DECLARATION(lhs),
lhs_offset += K0 * sizeof(DATA_TYPE);
rhs_offset += K0 * rhs_stride_y;
}
-
+#endif // K0 > 1
// Left-over accumulations
for(; i < K; ++i)
{
@@ -4292,10 +4293,6 @@ __kernel void gemm_mm_native(IMAGE_DECLARATION(lhs),
// Store output block
STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x);
-
-#undef RHS_BLOCK_SIZE
-#undef RHS_OFFSET_X
-#undef RHS_STEP_X
}
#endif // defined(M0) && defined(N0) && defined(K0) && defined(K) && defined(DATA_TYPE)