aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/kernels/CLActivationLayerKernel.cpp
AgeCommit message (Collapse)Author
2020-07-20COMPMID-3532: Align data type support between doxygen and implementation - CLMichele Di Giorgio
Also removes some unused code. Change-Id: I85687c40999c3cdf9e6fccfcd020b0901a9515fe Signed-off-by: Michele Di Giorgio <michele.digiorgio@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3581 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: 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-07-07COMPMID-3387: Support memory injection in CLActivationLayerGeorgios Pinitas
Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com> Change-Id: I31f9620607b372fc3340c71e748a5ea177d9da62 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3520 Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2020-04-22COMPMID-3280: Make all ML primitives for CL use the new interface - Part1 - Fix1Manuel Bottini
- const fix in the CLKernels part 1 Change-Id: I17340cb6ff26afd52b14b46645efedbe07aef1b6 Signed-off-by: Manuel Bottini <manuel.bottini@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3067 Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: 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-03-10COMPMID-3080: Implement Hard-Swish activation in CLmorgolock
Change-Id: I5ed35a5e0fba09791e4371b1a74108a1fdfed900 Signed-off-by: morgolock <pablo.tello@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2848 Reviewed-by: Sang-Hoon Park <sang-hoon.park@arm.com> Reviewed-by: Giorgio Arena <giorgio.arena@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2020-01-14COMPMID-2991: Add support for QASYMM8_SIGNED in CL kernels/functions - part 2Michele Di Giorgio
Adding support for QASYMM8_SIGNED to the following CL kernels/functions: - CLActivationLayerKernel/CLActivationLayer - CLComparisonKernel/CLComparison - CLConvertFullyConnectedWeightsKernel/CLConvertFullyConnectedWeights - CLDeconvolutionLayerUpsampleKernel/CLDeconvolutionLayerUpsample - CLDepthToSpaceLayerKernel/CLDepthToSpaceLayer - CLDequantizationLayerKernel/CLDequantizationLayer - CLGEMMMatrixVectorMultiplyKernel - CLNormalizePlanarYUVLayerKernel - CLPReluLayer - CLPixelWiseMultiplicationKernel/CLPixelWiseMultiplication - CLPoolingLayerKernel/CLPoolingLayer Change-Id: I874bbb7c2b08baa9c5ff4c9e6bc8778b42a6bec5 Signed-off-by: Michele Di Giorgio <michele.digiorgio@arm.com> Reviewed-on: https://review.mlplatform.org/c/2539 Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2019-10-15COMPMID-2205: CL runtime context.Pablo Tello
CL Interfaces implemented. Concrete classes implemented. One test (ActivationLayer) ported to the new interface. Change-Id: I283808bec36ccfc2f13fe048c45cbbee698ce525 Signed-off-by: Pablo Tello <pablo.tello@arm.com> Reviewed-on: https://review.mlplatform.org/c/1998 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2019-07-09COMPMID-2411: Add (logistic and tanh) activation support for QSYMM16 for CLManuel Bottini
Change-Id: I8d72490b1cc58563ba7b94664135586bc40e6526 Signed-off-by: Manuel Bottini <manuel.bottini@arm.com> Reviewed-on: https://review.mlplatform.org/c/1466 Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2019-06-05COMPMID-2372: Add support for QASYMM8 for TanhGeorgios Pinitas
-Perform calculations in the floating point domain -Extends checks for Logistic as scale should be 1/256 and offset 0 Change-Id: I90ef4a042f053976936f5d28f8e09b54eec196a2 Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com> Reviewed-on: https://review.mlplatform.org/c/1287 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2019-06-03COMPMID-2225: Add interface support for new quantized data types.Georgios Pinitas
Add support for: -QSYMM8, 8-bit quantized symmetric -QSYMM8_PER_CHANNEL, 8-bit quantized symmetric with per channel quantization Change-Id: I00c4ff98e44af37419470af61419ee95d0de2463 Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com> Reviewed-on: https://review.mlplatform.org/c/1236 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2019-05-13COMPMID-2194: Refactor activation function macro in OpenCL. New Macros for ↵Usama Arif
activation. This commit contains the new macros for activation. Only the activation_layer utilizes the new macros in this commit. Change-Id: I2fa8567cc876e8cb67a1e876652bc348b7ed23ea Signed-off-by: Usama Arif <usama.arif@arm.com> Reviewed-on: https://review.mlplatform.org/c/1104 Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2019-02-15COMPMID-1710: Fix CL logistic activation for QASYMM8Michele Di Giorgio
Logistic activation in QASYMM8 is performed in floating point and then converting back to QASYMM8. However, if input and output have different quantization information, a double conversion is done leading to loss in accuracy. This patch creates a special case kernel for logistic activation. Change-Id: Ia18ee0b2f84701674f88785bcd13753b50d64a08 Signed-off-by: Michele Di Giorgio <michele.digiorgio@arm.com> Reviewed-on: https://review.mlplatform.org/696 Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-1539 Implement YOLOLayer on CLGiorgio Arena
Change-Id: I332c0703e1399fca0c5b724529b54a28f49c88da Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/146842 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@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-1349: Add support for QASYMM8 LOGISTIC activation in CLActivationLayerMichele Di Giorgio
Change-Id: Ibabce61cf5427de80078a6468023bed05f5e7c2c Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/139006 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-11-02COMPMID-970 : Remove QS8 / QS16 supportVidhya Sudhan Loganathan
Removed fixed point related code. Change-Id: I487acf138dace3b0450e0d72ca7071eaec254566 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/137678 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02COMPMID-655 : Check FP16 is supported by the GPUVidhya Sudhan Loganathan
Change-Id: I507b04680a4e88426b682bd0be03bccb560ec78d Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/132589 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02COMPMID-909: Enabling in-place computation for batchnormalization and ↵Michele Di Giorgio
activation at graph level Change-Id: I84d4a212629b21794451ab5fb5c5b187b5e28f98 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/120127 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-841: Add CL QASYMM8 RELU ActivationMichele Di Giorgio
Change-Id: I8e0b7cad2f977942224d0116e8498bf9b2d6014d Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/117229 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-751 QASYMM8 ActivationLayer optimisation: don't requantize if not ↵Giorgio Arena
necessary Change-Id: Iea8a21f7c71025bfde6fdf7c7a7c92ba749b189b Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/116673 Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-617 Add validation methods to KernelsMichalis Spyrou
- NEActivationLayer - NESoftmax - NEDirectConvolutionLayer - NENormalizationLayer - NEPoolingLayer Change-Id: Ib279f1c1b7f9247679b0d6593aed7393da8fe87b Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/111335 Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-11-02COMPMID-556: Rename Error to Status and inverse logicGeorgios Pinitas
Change-Id: Ib57d4f7177cc6179302bda7ad870acb8bd3825f5 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/112115 Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02IVGCVSW-631 Neon support for Softmax beta parameter (F32 only)Pablo Palmier
Change-Id: Ibf6f038b39f1a4e557f5d04feb08e3d5ef54e223 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/112019 Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-11-02COMPMID-556: Added a rounding policy to the quantize functionJaroslaw Rzepecki
Change-Id: I6272a36636c5d9baff6d35dee0a50dc847f65bfa Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/110266 Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02COMPMID-617 CL window validation fixesGiorgio Arena
Change-Id: I479dc1bd3bf9be641cb89cea5085151ba464f798 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/111471 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com>
2018-11-02COMPMID-617 Add validation window to CLActivationLayerGiorgio Arena
Change-Id: I56e0655a2f3448691cfd7bef6911d7959da12fdd Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/111186 Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-11-02COMPMID-556 - Passed _lws_hint in the run() method where was missingGian Marco
_lws_hint was missing in the run() method for some ML functions where _config_id was set Change-Id: I0a092372cccaf21f29e4c78340201de1606caab4 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/110250 Reviewed-by: Anthony Barbier <anthony.barbier@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com>
2018-11-02COMPMID-685 Extend CLTuner support to other DL functionsGiorgio Arena
Change-Id: Ica97857c2145228e4a6088724681ec1c0a138133 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/95918 Reviewed-by: Anthony Barbier <anthony.barbier@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
2018-11-02COMPMID-661: Add avgpool-uint8 support. Optimize avgpool-fp32 for Bifrost. (#13)Anton Lokhmotov
Change-Id: I32ba6afbac6694ffa053dd16f03a1b3d14627a19 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/94857 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02COMPMID-617: Add validation functions.Georgios Pinitas
Added validation routines to the following kernels. -CLActivationLayer -CLBatchNormalizationLayer -CLArithmeticAddition -CLArithmeticSubtraction -CLPixelwiseMultiplication Change-Id: I0f3a03154f9e392279f715af656683cd0ad4cef5 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/94595 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02COMPMID-556: Support beta for all softmax data types.Georgios Pinitas
Change-Id: I4c0ca033dc53829fb7ac3dd7c7469d143be74e73 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/94251 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02COMPMID-556: Rework CLActivationLayerGeorgios Pinitas
Refactoring. Change-Id: I879353299b655ec3026cccdfcfca2ee98abf14ea Reviewed-on: http://mpd-gerrit.cambridge.arm.com/94191 Reviewed-by: Michel Iwaniec <michel.iwaniec@arm.com> Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02IVGCVSW-619: Support for Cl u8 bounded ReluMichel Iwaniec
Change-Id: I3c39ecbd36f06d5376c35ed4eb38dd73533ef97e Reviewed-on: http://mpd-gerrit.cambridge.arm.com/93686 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-09-17COMPMID-406: Port CLActivationLayer to use QS8/QS16.Georgios Pinitas
Change-Id: Ia4114984c38e1d2027ad97335b3c6c11f5754e23 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/78727 Reviewed-by: Anthony Barbier <anthony.barbier@arm.com> Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
2018-09-17COMPMID-443: Collapse higher dimensions for activation layerAnthony Barbier
Change-Id: I5943235aff1bb6440e3ab08e818d53aa5d94143a Reviewed-on: http://mpd-gerrit.cambridge.arm.com/80349 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-09-17COMPMID-417: Remove val_to_stringMoritz Pflanzer
The function was a duplicate of the to_string function we have in support/ToolchainSupport.h Change-Id: If1eff674333248f7d189824ced717310afec19b0 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/79634 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
2018-09-17COMPMID-345 - In-place computation for Activation LayerGian Marco Iodice
Change-Id: I25ebfccc3d3e758cc8164e0b33805c0bb303891a Reviewed-on: http://mpd-gerrit.cambridge.arm.com/78226 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Pablo Tello <pablo.tello@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-09-17COMPMID-344 Updated doxygenAnthony Barbier
Change-Id: I32f7b84daa560e460b77216add529c8fa8b327ae