aboutsummaryrefslogtreecommitdiff
path: root/src/runtime/CL/CLTuner.cpp
AgeCommit message (Collapse)Author
2023-09-28Apply clang-format on repositoryFelix Thomasmathibalan
Code is formatted as per a revised clang format configuration file(not part of this delivery). Version 14.0.6 is used. Exclusion List: - files with .cl extension - files that are not strictly C/C++ (e.g. Android.bp, Sconscript ...) And the following directories - compute_kernel_writer/validation/ - tests/ - include/ - src/core/NEON/kernels/convolution/ - src/core/NEON/kernels/arm_gemm/ - src/core/NEON/kernels/arm_conv/ - data/ There will be a follow up for formatting of .cl files and the files under tests/ and compute_kernel_writer/validation/. Signed-off-by: Felix Thomasmathibalan <felixjohnny.thomasmathibalan@arm.com> Change-Id: Ib7eb1fcf4e7537b9feaefcfc15098a804a3fde0a Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/10391 Benchmark: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Gunes Bayir <gunes.bayir@arm.com>
2023-03-06Fix LWS search space used by CLTunerSiCong Li
* Ensure CLTuner uses the real GWS used by run(), instead of the static GWS (which is usually changed at run time), by caching GWS in each kernel Note this is a somewhat inelegant workaround. The real issue stems from the fact that execution window and scheduler are very much coupled with our operator run() / run_op() method. (Please see COMPMID-5934) * Restrict LWS values to explore within GWS bound for exhaustive mode * Refactor gws_from_window() to include all the information required to calculate GWS * Log lws search space used for tuning * Fix ClDirectConv2dKernel config id Resolves COMPMID-5892 Signed-off-by: SiCong Li <sicong.li@arm.com> Change-Id: I420490d8b94d13ada2e44eb0a12078f883379334 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/9193 Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Benchmark: Arm Jenkins <bsgcomp@arm.com>
2022-11-22Remove dynamic fusion prototype with tests and examplesSiCong Li
Public headers of the new experimental dynamic fusion can be found in arm_compute/dynamic_fusion/ New examples on how to use the interface can be found in tests/validation/dynamic_fusion/gpu/Integration.cpp Resolves COMPMID-5683 Change-Id: I7ccb902a227fb487562df15fc3c30118d1d95bbd Signed-off-by: SiCong Li <sicong.li@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/8671 Reviewed-by: Jakub Sujak <jakub.sujak@arm.com> Reviewed-by: Gunes Bayir <gunes.bayir@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Benchmark: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2022-05-06Integrate Dynamic Fusion patchesSiCong Li
* Add public interfaces: * OperatorGraph: Describe a workload that could contain fused kernels * IWorkload: Generic interface for workloads built from OperatorGraph * ClWorkload: OpenCL workloads built from OperatorGraph * ClCompositeOperator: Runtime async operator to execute a ClWorkload * DependencyGraph (will likely be deprecated in later iterations) * Add example * cl_fused_conv2d_elementwise_add.cpp to explain how to use the new interfaces * Add internal translation layer * Refactor ClKernelBuildingAPI * Remove non-tile based gemm native kernel component * Minor interface changes * Add integration tests Resolves COMPMID-5161 Signed-off-by: SiCong Li <sicong.li@arm.com> Change-Id: Ib987ed79289ab0bcbd3130d54f5793408d9f1240 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/7510 Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Reviewed-by: Gunes Bayir <gunes.bayir@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2022-04-14Enable dynamic cl tuning for dynamically fused kernelsSiCong Li
* Add new tune_kernel_dynamic interface * Add generate_config_id Resolves: COMPMID-5154 Signed-off-by: SiCong Li <sicong.li@arm.com> Change-Id: I39870e59fceda875487970061ceb2048995c5a45 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/7400 Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Giorgio Arena <giorgio.arena@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2021-08-04Remove 21.08 deprecated codeFreddie Liardet
Also correct some deprecation macros Signed-off-by: Freddie Liardet <frederick.liardet@arm.com> Change-Id: Ib28d42571f55d7b01563a0877b92b8f3128831ba Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/6036 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2021-02-03Add WBSM tuning to CLTunerManuel Bottini
Add WBSM as possible parameter to be tuned Add helper functions to check WBSM support and setting the value in the kernel Update tuning parameter lists to use WBSM Update CLTuner to use WBSM The WBSM tuning is exposed as a parameter to be set at compile time by setting the CLTuningInfo CLTuningInfo contains information about the tuning mode and if wbsm tuning enabled Resolves: COMPMID-3936 Change-Id: Id53697c9c6d2cef41c049f368002f6197351b3ed Signed-off-by: Manuel Bottini <manuel.bottini@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4914 Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2021-01-26Remove iostream header file when not neededGian Marco Iodice
Change-Id: Ic6b3e94fc252c2d6fe17e6bc32900d11c68ae399 Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4916 Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2021-01-08Generalization of CLTunerManuel Bottini
Rename lws to tuning parameters in functions used externally Add new generalized objects for the OpenCL Tuner to accommodate further possible tuning parameters Resolves: COMPMID-3935 Change-Id: I0f2a0f89bca5dae4a4e4adce2f7c7cae32ecb84a Signed-off-by: Manuel Bottini <manuel.bottini@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4584 Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@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-08-05COMPMID-3392: Collapse TensorMaps into a single TensorPackGeorgios Pinitas
Collapse InputTensorMap and OutputTensorMap to a single TensorPack mechanism. Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com> Change-Id: Ie2fdfc6b07d84ad589169ec99ca64fcf45a00bec Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/c/VisualCompute/ComputeLibrary/+/253783 Tested-by: bsgcomp <bsgcomp@arm.com> Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com> Signed-off-by: Michalis Spyrou <michalis.spyrou@arm.com> Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3641 Tested-by: Arm Jenkins <bsgcomp@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: SiCong Li <sicong.li@arm.com> Reviewed-by: Sheri Zhang <sheri.zhang@arm.com>
2020-07-20COMPMID-3604: Graph failures during tuningGeorgios Pinitas
Update ICLTuner interface to account for the new memory injection interface. Redirect to appropriate kernel execution interface depending on if the kernel supports memory injection or not. Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com> Change-Id: I8ce29f5c22f1865c9e688d12b65e68ee4486f99c Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3588 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-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>
2019-10-16COMPMID-2486: Remove disabled compiler warningsMichalis Spyrou
Removed the following flags: -Wno-format-nonliteral: This had a side effect on Error.h that resulted in rewriting most of the macros. Since I was at it I removed all the va_args in order to comply with DCL50-CPP. -Wno-deprecated-increment-bool -Wno-vla-extension -Wno-mismatched-tags -Wno-redundant-move Change-Id: I7c593854ecc3b7d595b8edcbd6a86d3c2563c6bd Signed-off-by: Michalis Spyrou <michalis.spyrou@arm.com> Reviewed-on: https://review.mlplatform.org/c/2069 Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2019-07-30COMPMID-2467 Include IP information to the tuning configGiorgio Arena
Change-Id: I00e29f30eda0edc76ae393bef14f0092e57a5195 Signed-off-by: Giorgio Arena <giorgio.arena@arm.com> Reviewed-on: https://review.mlplatform.org/c/1634 Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2019-05-15COMPMID-2334: Extend search space for the EXHAUSTIVE OpenCL tunerGian Marco Iodice
Change-Id Ic71d875542f19d5a4f717f397e1e6ce9ce2536f Change-Id: I7f9cc86b149aa392acffed602b431451499704d7 Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Reviewed-on: https://review.mlplatform.org/c/1133 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2019-05-13COMPMID-1995: Change CLTuner default mode to NORMALMichalis Spyrou
Change-Id: I96a5fc460d5e403fcc4b64002ed4c71606c732f0 Signed-off-by: Michalis Spyrou <michalis.spyrou@arm.com> Reviewed-on: https://review.mlplatform.org/c/1118 Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2019-05-10COMPMID-2177 Fix clang warningsMichalis Spyrou
Change-Id: I78039db8c58d7b14a042c41e54c25fb9cb509bf7 Signed-off-by: Michalis Spyrou <michalis.spyrou@arm.com> Reviewed-on: https://review.mlplatform.org/c/1092 Reviewed-by: VidhyaSudhan Loganathan <vidhyasudhan.loganathan@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2019-04-25COMPMID-1974 : Extend CLTuner to support different of level of tuningVidhya Sudhan Loganathan
Change-Id: I52e4a00a25e7f7a17050038cee7c30e508553722 Signed-off-by: Vidhya Sudhan Loganathan <vidhyasudhan.loganathan@arm.com> Reviewed-on: https://review.mlplatform.org/c/977 Comments-Addressed: Pablo Marquez <pablo.tello@arm.com> Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
2019-02-12COMPMID-1919: Opencl - Tuning with Instrumentation does not workVidhya Sudhan Loganathan
Change-Id: Iee6a07d5bf6a35af04071865682bcc4a615c14f9 Signed-off-by: Vidhya Sudhan Loganathan <vidhyasudhan.loganathan@arm.com> Reviewed-on: https://review.mlplatform.org/638 Reviewed-by: Giuseppe Rossini <giuseppe.rossini@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Pablo Marquez <pablo.tello@arm.com>
2019-02-06COMPMID-1920: Failure exit order destruction when CLTuner enabled.Georgios Pinitas
Creates a profiling queue in every tuning iteration as it gives better control of the lifetime of the opencl objects. Change-Id: Id8629cc06086877b088c787cbb5f238c0979b6d9 Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com> Reviewed-on: https://review.mlplatform.org/631 Tested-by: Arm Jenkins <bsgcomp@arm.com> Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com>
2019-02-04COMPMID-1910: Improve CLTuner reducing the number of LWS to testGian Marco Iodice
Change-Id: I842120a2bcddc5bf8677ee4d0b1f9d379771b36b Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Reviewed-on: https://review.mlplatform.org/622 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-959 Refactor OpenCL interceptors to use lambda functionsAnthony Barbier
Change-Id: I29b73a311d7278255b77524f2a5eaaa4dccab711 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/128392 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-734: CLTuner reworkGeorgios Pinitas
Change-Id: I8f20d6ea8a09869d71003e7b08e0d33775282f6c Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/125802 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02COMPID-978: Throw exception on missing tuning fileAnthony Barbier
Change-Id: I09ad6f52cc32f5dcd38d2ba7c6143e7e9ab12b61 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/122767 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Matthew Bentham <matthew.bentham@arm.com> Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com>
2018-11-02COMPMID-978: Fixed bug in Load/Store tuning data and in the CLTuner interceptorAnthony Barbier
Change-Id: I6cdecef623ff5806cf82cb12a60aef8aefec32f5 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/122712 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-978 Load/Store tuning data from file (Part2)Anthony Barbier
Change-Id: I1819f42c0e456673543b267d51f730b6e80a0ad9 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/122629 Reviewed-by: Robert Hughes <robert.hughes@arm.com> Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-11-02COMPMID-978 Load/Store tuning data from fileAnthony Barbier
Change-Id: I1d1f402df3a58704c021b9866d489844fb5e7d7a Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/122395 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-11-02COMPMID-895 - Optimizing CLDepthwiseConvolution3x3KernelGian Marco
This patch brings the MACs utilisation up to 25 % when both stride_x and stride_y are equal to 1 Performance reported in the following confluence page: https://confluence.arm.com/display/MLENG/Depthwise+convolution+3x3+FP32+performance%3A+ACL+18.02 Change-Id: Ida1b64be9a88805902a3d90194559b58eb1224a3 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/119068 Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-891 - Use OpenCL timer in CLTunerGian Marco
Change-Id: I84a914c13b162c4f74321c9cafc30a18ad4ebbdb Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/118797 Tested-by: Jenkins <bsgcomp@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
2018-11-02COMPMID-765 - Added third dimension for CLTunerGian Marco
Change-Id: I0a7ea4cde1dbf8edd28908dfff80928ef7e996c4 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/117647 Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Tested-by: Jenkins <bsgcomp@arm.com>
2018-11-02COMPMID-417 - Fix for the config_id default stringGian Marco Iodice
The empty string used by default for config_id caused a memory corruption on Android. Change-Id: Ib9c764c8a8f7009936071f8c8fec1653bb000f90 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/89262 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
2018-11-02COMPMID-524 - Implemented CLTuner objectGian Marco
Change-Id: Idbdbecca1fc299ed042936119d90e2bed8db0938 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/87101 Reviewed-by: Anthony Barbier <anthony.barbier@arm.com> Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>