diff options
| author | SiCongLi <sicong.li@arm.com> | 2021-11-03 19:01:22 +0000 |
|---|---|---|
| committer | SiCong Li <sicong.li@arm.com> | 2021-11-04 14:03:19 +0000 |
| commit | d928735fee6baefdb74325c05d8152dd13044f32 (patch) | |
| tree | 6fb702e36da2863639149995e2df1cfe70905fc7 /src/core | |
| parent | 6049edadf0c89a026b3fcd1927ee7531d3c40278 (diff) | |
| download | ComputeLibrary-d928735fee6baefdb74325c05d8152dd13044f32.tar.gz | |
Add validate tests for CLConvolutionLayer and CLGEMMConvolutionLayer with post ops
* Add validate tests
* Restrict post ops support in ClGemmConv2d to only those that do not
need im2col or col2im. In practice this means we only support post ops
in conv1x1 with stride = 1, dilation = 1 and data layout = NHWC
Resolves COMPMID-4435
Change-Id: I1fdf0c5d565a4624857250075ac76db35c2f383b
Signed-off-by: SiCongLi <sicong.li@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/6573
Reviewed-by: Gunes Bayir <gunes.bayir@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core')
4 files changed, 12 insertions, 12 deletions
diff --git a/src/core/CL/CLUtils.cpp b/src/core/CL/CLUtils.cpp index 88b31c8349..8dab8aa876 100644 --- a/src/core/CL/CLUtils.cpp +++ b/src/core/CL/CLUtils.cpp @@ -96,9 +96,9 @@ bool PostOpCLKernelUtils::are_post_op_shapes_compliant(const ITensorInfo *dst, c return false; } // NOTE: Kernel limitation: currently only the following broadcasting types are supported: - // 1. Post op arg is scalar, broadcast in both X and Y - // 2. Post op arg is of shape: Y=1, X=N, broadcast only in Y - // This means this case: Post op arg is of shape: Y=M, X=1, broadcast only in X, is NOT supported + // 1. Post op arg is scalar, broadcast in both first and second dims + // 2. Post op arg is of shape: second dim=1, first dim=N, broadcast only in second dim + // This means this case: Post op arg is of shape: second dim=M, first dim=1, broadcast only in first dim, is NOT supported if(dst->dimension(0) > 1 && dst->dimension(1) > 1 && (*tensor)->dimension(0) == 1 && (*tensor)->dimension(1) > 1) { return false; diff --git a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_native.cl b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_native.cl index bbe97b2781..4665d612f5 100644 --- a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_native.cl +++ b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_native.cl @@ -133,7 +133,7 @@ __kernel void gemm_mm_native_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), IMAGE_DECLARATION(bias), #endif // defined(BETA) IMAGE_DECLARATION(dst), - // Post-Op arguments + // Post Op arguments IMAGE_DECLARATION(eltwise_operand), uint lhs_stride_z, uint rhs_stride_z, diff --git a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl index 9e9a73ccf6..32186c359b 100644 --- a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl +++ b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl @@ -233,7 +233,7 @@ __kernel void gemm_mm_reshaped_lhs_nt_rhs_t_post_act_eltwise_op_act(IMAGE_DECLAR IMAGE_DECLARATION(bias), #endif // defined(BETA) IMAGE_DECLARATION(dst), - // Post-Op arguments + // Post Op arguments IMAGE_DECLARATION(eltwise_operand), uint k, uint lhs_stride_z, @@ -453,7 +453,7 @@ __kernel void gemm_mm_reshaped_lhs_nt_rhs_t_texture_post_act_eltwise_op_act(IMAG IMAGE_DECLARATION(bias), #endif // defined(BETA) IMAGE_DECLARATION(dst), - // Post-Op arguments + // Post Op arguments IMAGE_DECLARATION(eltwise_operand), uint k, uint lhs_stride_z, @@ -781,7 +781,7 @@ __kernel void gemm_mm_reshaped_lhs_t_rhs_nt_post_act_eltwise_op_act(IMAGE_DECLAR IMAGE_DECLARATION(bias), #endif // defined(BETA) IMAGE_DECLARATION(dst), - // Post-Op arguments + // Post Op arguments IMAGE_DECLARATION(eltwise_operand), uint k, uint lhs_stride_z, @@ -1110,7 +1110,7 @@ __kernel void gemm_mm_reshaped_lhs_t_rhs_nt_texture_post_act_eltwise_op_act(IMAG IMAGE_DECLARATION(bias), #endif // defined(BETA) IMAGE_DECLARATION(dst), - // Post-Op arguments + // Post Op arguments IMAGE_DECLARATION(eltwise_operand), uint k, uint lhs_stride_z, diff --git a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl index fe2d103de5..e96aba613b 100644 --- a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl +++ b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl @@ -177,7 +177,7 @@ __kernel void gemm_mm_reshaped_only_rhs_t_post_act_eltwise_op_act(IMAGE_DECLARAT IMAGE_DECLARATION(bias), #endif // defined(BETA) IMAGE_DECLARATION(dst), - // Post-Op arguments + // Post Op arguments IMAGE_DECLARATION(eltwise_operand), uint lhs_stride_z, uint rhs_stride_z, @@ -437,7 +437,7 @@ __kernel void gemm_mm_reshaped_only_rhs_t_texture_post_act_eltwise_op_act(IMAGE_ IMAGE_DECLARATION(bias), #endif // defined(BETA) IMAGE_DECLARATION(dst), - // Post-Op arguments + // Post Op arguments IMAGE_DECLARATION(eltwise_operand), uint lhs_stride_z, uint rhs_stride_z, @@ -831,7 +831,7 @@ __kernel void gemm_mm_reshaped_only_rhs_nt_post_act_eltwise_op_act(IMAGE_DECLARA IMAGE_DECLARATION(bias), #endif // defined(BETA) IMAGE_DECLARATION(dst), - // Post-Op arguments + // Post Op arguments IMAGE_DECLARATION(eltwise_operand), uint lhs_stride_z, uint rhs_stride_z, @@ -1116,7 +1116,7 @@ __kernel void gemm_mm_reshaped_only_rhs_nt_texture_post_act_eltwise_op_act(IMAGE IMAGE_DECLARATION(bias), #endif // defined(BETA) IMAGE_DECLARATION(dst), - // Post-Op arguments + // Post Op arguments IMAGE_DECLARATION(eltwise_operand), uint lhs_stride_z, uint rhs_stride_z, |