diff options
author | Georgios Pinitas <georgios.pinitas@arm.com> | 2020-12-01 17:41:34 +0000 |
---|---|---|
committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2020-12-02 15:21:11 +0000 |
commit | 96b16b65dd96351b8af1b2a785856ce13cc8ba84 (patch) | |
tree | d05ae4c07753e9e15d9861eb36df783dbe7a00a0 /src/core/CL/cl_kernels/gemm.cl | |
parent | d308df3186b4f6057f94b45b7bed7935c618ea80 (diff) | |
download | ComputeLibrary-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.cl | 70 |
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 |