aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/gemm.cl
diff options
context:
space:
mode:
authorGeorgios Pinitas <georgios.pinitas@arm.com>2020-12-01 17:41:34 +0000
committerGeorgios Pinitas <georgios.pinitas@arm.com>2020-12-02 15:21:11 +0000
commit96b16b65dd96351b8af1b2a785856ce13cc8ba84 (patch)
treed05ae4c07753e9e15d9861eb36df783dbe7a00a0 /src/core/CL/cl_kernels/gemm.cl
parentd308df3186b4f6057f94b45b7bed7935c618ea80 (diff)
downloadComputeLibrary-96b16b65dd96351b8af1b2a785856ce13cc8ba84.tar.gz
Remove support for (NE/CL)LocallyConnectedLayer
Remove out-of-date and unmaintained LocallyConnectedLayer for both NEON and OpenCL. Resolves: COMPMID-3924 Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com> Change-Id: Ia61398ed8cfa3876f41c1b342c4a80d1cca0ca83 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4634 Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/gemm.cl')
-rw-r--r--src/core/CL/cl_kernels/gemm.cl70
1 files changed, 1 insertions, 69 deletions
diff --git a/src/core/CL/cl_kernels/gemm.cl b/src/core/CL/cl_kernels/gemm.cl
index b6afb85aa4..6883aafee5 100644
--- a/src/core/CL/cl_kernels/gemm.cl
+++ b/src/core/CL/cl_kernels/gemm.cl
@@ -4379,72 +4379,4 @@ __kernel void gemm_ma_f16(TENSOR3D_DECLARATION(src),
vstore8(out, 0, (__global half *)dst.ptr);
}
#endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED)
-#endif // defined(BETA)
-
-#if defined(WIDTH_VECTOR_A)
-/** This OpenCL kernel computes the vector by matrix multiplication between each row of A (src0) and matrix B (src1) used for locally connected layer
- *
- * @note The width of A need to be passed at compile time using -DWIDTH_VECTOR_A
- *
- * @note The input A and matrix B must not be reshaped
- *
- * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F32
- * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
- * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
- * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
- * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
- * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
- * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
- * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
- * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
- * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
- * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
- * @param[in] src1_stride_z Stride of the source matrix in Z dimension (in bytes)
- * @param[in] src1_step_z src_stride_z * number of elements along Z processed per workitem(in bytes)
- * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
- * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
- * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
- * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
- * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
- * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
- * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
- */
-__kernel void gemm_lc_vm_f32(IMAGE_DECLARATION(src0),
- TENSOR3D_DECLARATION(src1),
- IMAGE_DECLARATION(dst))
-{
- int idx = get_global_id(0) * 4;
- int idy = get_global_id(1);
-
- // Compute the address for the vector A and matrix B
- int2 src_addr = ((int2)(src0_offset_first_element_in_bytes + src0_stride_y * idy, src1_offset_first_element_in_bytes + src1_stride_z * idy));
- src_addr.s1 += idx * sizeof(float);
-
- int end_row_vec_a = src_addr.s0 + (WIDTH_VECTOR_A * sizeof(float));
-
- float4 acc = 0.0f;
-
- for(; src_addr.s0 <= (end_row_vec_a - 2 * (int)sizeof(float)); src_addr += (int2)(2 * sizeof(float), 2 * src1_stride_y))
- {
- float2 a0 = vload2(0, (__global float *)(src0_ptr + src_addr.s0));
- float4 b0 = vload4(0, (__global float *)(src1_ptr + src_addr.s1));
- float4 b1 = vload4(0, (__global float *)(src1_ptr + src_addr.s1 + src1_stride_y));
-
- acc += b0 * (float4)a0.s0;
- acc += b1 * (float4)a0.s1;
- }
-
- for(; src_addr.s0 < end_row_vec_a; src_addr += (int2)(sizeof(float), src1_stride_y))
- {
- float a0 = *((__global float *)(src0_ptr + src_addr.s0));
- float4 b0 = vload4(0, (__global float *)(src1_ptr + src_addr.s1));
-
- acc += b0 * (float4)a0;
- }
-
- // Compute destination address
- Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
-
- vstore4(acc, 0, (__global float *)(offset(&dst, 0, 0)));
-}
-#endif // defined(WIDTH_VECTOR_A)
+#endif // defined(BETA) \ No newline at end of file