aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/kernels/CLWinogradInputTransformKernel.cpp
AgeCommit message (Collapse)Author
2021-06-15Port CLWinogradConvolutionLayer with ClWinogradConv2dManuel Bottini
Port CLWinogradInputTransformKernel Port CLWinogradFilterTransformKernel Port CLWinogradOutputTransformKernel Resolves: COMPMID-4504 Change-Id: I3177dda0b9c2f56b36cb317027e94abe8d47229e Signed-off-by: Manuel Bottini <manuel.bottini@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5680 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2021-04-08Rework the OpenCL Winograd Input Transformations NHWCGian Marco Iodice
- Rework Winograd Input Transform 3x3 NHWC using the new macros - Rework Winograd Input Transform 5x5 NHWC using the new macros - Rework Winograd Input Transform 7x7 NHWC using the new macros - The new implementation is also faster than before - Winograd Input Transform 5x5/7x7 3x faster Resolves COMPMID-4139 Change-Id: Ia9c8af23a2d47d2db60ec4c44650a63a34ffa0d5 Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5358 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
2020-12-18Add new shapes to WinogradInputTransform dataset and fix border size for ↵Giorgio Arena
NCHW data layout Fix border size for CLWinogradInputTransformKernel with NCHW data layout by setting it to the input's paddings. Add new the new validation shapes to the WinogradInputTransform's dataset Resolves COMPMID-4042 Change-Id: Id93ac86e75c94ea3f2f35edcedebafada928f34a Signed-off-by: Giorgio Arena <giorgio.arena@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4694 Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Manuel Bottini <manuel.bottini@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2020-11-17COMPMID-3979 Sanitise Padding Removal epicSiCong Li
* Add missing padding immutability asserts in all relevant CL kernels * Remove unnecessary zero padding validation tests. Change-Id: If93f9ccbc988e0286f5e7b135f812141476d5da0 Signed-off-by: SiCong Li <sicong.li@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4446 Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2020-11-07COMPMID-3639: (3RDPARTY_UPDATE) Move CL kernels to srcSang-Hoon Park
Change-Id: I10d27db788e5086adae1841e3e2441cd9b76ef84 Signed-off-by: Sang-Hoon Park <sang-hoon.park@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4310 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2020-10-20COMPMID-3637: Move utility headers from arm_compute to srcSang-Hoon Park
Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com> Change-Id: If9d6fa8c900b68c4b6fd373f2fc1f9abb83ea917 Signed-off-by: Michalis Spyrou <michalis.spyrou@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4145 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Sang-Hoon Park <sang-hoon.park@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2020-10-19COMPMID-3740: Remove OpenCL padding: CLWinogradInputTransformKernelGian Marco Iodice
- Remove padding requirement from the OpenCL kernels - Extend test to validate zero padding requirement Change-Id: I1ddf04eba783721858792efb08a2c97f11f7297e Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4206 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2020-07-09COMPMID-3324: Adjusting capitalization of Arm copyright claim to reflect Arm ↵Michele Di Giorgio
preferred presentation Change-Id: Ib7dcfcbb24b408999dfae366b9da396485aacf78 Signed-off-by: Michele Di Giorgio <michele.digiorgio@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3525 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Georgios Pinitas <georgios.pinitas@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2020-04-22COMPMID-3280: Make all ML primitives for CL use the new interface - Part1 - Fix3Manuel Bottini
- const fix in the CLKernels part 3 Change-Id: I9cfb896f334145249a97c9287fa00399b8319a8e Signed-off-by: Manuel Bottini <manuel.bottini@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3075 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2020-04-17COMPMID-3280: Make all ML primitives for CL use the new interface - Part 1Manuel Bottini
- Only CLKernels have been updated Change-Id: Ife55b847c2e39e712a186eb6ca452503d5b66937 Signed-off-by: Manuel Bottini <manuel.bottini@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3001 Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2020-03-10COMPMID-3069: Improve build time by splitting up ToolchainSupport.hMatthew Bentham
Split out the parts of ToolchainSupport coming from <memory> and the parts coming from <string> into their own new header files. This accounts for 99% of uses of ToolchainSupport, which means that expensive header files such as arm_neon.h don't need to be included everywhere. Knocks about 10% of compilation time off kernel files. Signed-off-by: Matthew Bentham <matthew.bentham@arm.com> Change-Id: I2ae718fe766b5ff28608812b0f686f30eeac1b21 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2852 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
2020-01-10COMPMID-2819: Retain layout during configuration for multiple functions.Georgios Pinitas
Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com> Change-Id: Ia528762dc5a93bebfd8fd037bf1f4e75d0b8a6de Reviewed-on: https://review.mlplatform.org/c/2566 Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
2018-11-02COMPMID-1699: Disable arithmetic operations in CLWinogradLayer when no ↵Georgios Pinitas
batches available. Change-Id: Iad83df2a9116a7f350de83ec59b28cd8893c8d3a Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/155716 Tested-by: bsgcomp <bsgcomp@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
2018-11-02COMPMID-1029: Collapse CLWinogradInputTransform/CLWinogradOutputTransformGeorgios Pinitas
Change-Id: I051748502ca24b9952e7313524bbfd708162efb4 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/155166 Tested-by: bsgcomp <bsgcomp@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
2018-11-02COMPMID-1266 : support for FP16 in CLWinogradConvolutionLayerVidhya Sudhan Loganathan
Added support for FP16 in CLWinogradConvolutionLayer: 5x5 kernels and 3x3 kernels(COMPMID-937) Change-Id: I0f394cbdc978dd04176416e9f612aca3986b09e6 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/145537 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Giorgio Arena <giorgio.arena@arm.com>
2018-11-02COMPMID-1188 - Fix CLWinogradConvolutionLayer for NHWCGian Marco Iodice
Change-Id: Ib4abe0388f218276e79f7c4405827e61722f0ef8 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/144240 Reviewed-by: Giorgio Arena <giorgio.arena@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-1478: Stop relying on static default OpenCL objects in cl2.hppAnthony Barbier
This causes problems when ACL is used as a shared library on Android. Fixes some problems related to creation / destruction order between the Graph's CL backend and core / runtime Change-Id: I716d63fd42f4586df1ffbb6fa97e4db06d3a781b Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/143228 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Michele DiGiorgio <michele.digiorgio@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
2018-11-02COMPMID-1337 Implementing Winograd Convolution Layer 1x3 and 3x1 kernels on ↵Giorgio Arena
OpenCL NHWC Change-Id: Ia07e0dfcbcd07366c4bcb956e298369fb12a0369 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/138759 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
2018-11-02COMPMID-1201 - Implementing Winograd Convolution Layer 1x3 and 3x1 kernels ↵Gian Marco Iodice
on OpenCL Change-Id: I39667bab49daa4da009694163274a59fd3574c73 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/137595 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Giorgio Arena <giorgio.arena@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-11-02COMPMID-1204 Add NHWC data format support to Winograd input transform 4x4_5x5Giorgio Arena
Change-Id: I3dffdd1772b78db27a4374f074a24a15a9552189 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/134859 Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-1048 Add NHWC data format support to Winograd input transform 4x4_3x3Giorgio Arena
https://confluence.arm.com/display/MLENG/Winograd+Input+Transform%3A+NCHW+vs+NHWC+on+OpenCL Change-Id: Iac35a54389266701b7d8f5434a7a37df85b7b187 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/133315 Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-1026 - Add support for 4x4 output tile in CLWinogradConvolutionLayerGian Marco Iodice
The performance achieved can be found at the following confluence page: https://confluence.arm.com/display/MLENG/GEMM-based+convolution+vs+Winograd-based+convolution+on+OpenCL Change-Id: I4b690cfdd4eb4ff0cd17b14fdd49ccaa1d1dc85c Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/127729 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-11-02COMPMID-1037 Add support for F(4x4, 5x5) in CLWinogradInputTransformKernelGiorgio Arena
Change-Id: Iac26936f46d0f7cdd9d2f8393b0092cd5a223c45 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/127675 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
2018-11-02COMPMID-1013 - Create WinogradInfo data structureGian Marco Iodice
COMPMID-1014 - Refactoring Winograd's dataset Change-Id: I6abdcbf9a90d663f4db666cd410afece9f1d034d Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/125899 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02COMPMID-935 - Implementing Convolution with Winograd on OpenCL (part 4)Gian Marco Iodice
Implemented Winograd Output Transform (2x2,3x3) on OpenCL Implemented CLWinogradConvolutionLayer on OpenCL Change-Id: I6a113fc5f052ca07f878d2b800d2ab003f84af65 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/125148 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-935 Implementing Convolution with Winograd on OpenCL (part 3)Giorgio Arena
Change-Id: I51f92f30602fb0a02314f344fa67061f448694bf Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/122793 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Giorgio Arena <giorgio.arena@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
generated by cgit v1.2.1 (git 2.23.0) at 2024-05-20 22:22:13 +0000