commit be9f9f9139b759d314f4f2a6d2ee747079666504
author Manuel Bottini <manuel.bottini@arm.com> Mon Jan 25 15:07:17 2021 +0000
committer Georgios Pinitas <georgios.pinitas@arm.com> Wed Feb 03 17:35:00 2021 +0000
tree 461690abb95caeaeca40261fd85816a906c8446c
parent 7061eb283969f9a020c08349454447564e4dd5b3

Add WBSM tuning to CLTuner

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>

src/core/CL/CLCompileContext.cpp

diff --git a/src/core/CL/CLCompileContext.cpp b/src/core/CL/CLCompileContext.cpp
index 0afb7e5..3db0fe5 100644
--- a/src/core/CL/CLCompileContext.cpp
+++ b/src/core/CL/CLCompileContext.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020 Arm Limited.
+ * Copyright (c) 2020-2021 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -137,15 +137,16 @@
 {
 }
 CLCompileContext::CLCompileContext()
-    : _context(), _device(), _programs_map(), _built_programs_map()
+    : _context(), _device(), _programs_map(), _built_programs_map(), _is_wbsm_supported()
 {
 }
 
 CLCompileContext::CLCompileContext(cl::Context context, const cl::Device &device)
-    : _context(), _device(), _programs_map(), _built_programs_map()
+    : _context(), _device(), _programs_map(), _built_programs_map(), _is_wbsm_supported()
 {
-    _context = std::move(context);
-    _device  = CLDevice(device);
+    _context           = std::move(context);
+    _device            = CLDevice(device);
+    _is_wbsm_supported = get_wbsm_support_info(device);
 }
 
 Kernel CLCompileContext::create_kernel(const std::string &kernel_name, const std::string &program_name, const std::string &program_source,
@@ -318,7 +319,8 @@
 
 void CLCompileContext::set_device(cl::Device device)
 {
-    _device = std::move(device);
+    _device            = std::move(device);
+    _is_wbsm_supported = get_wbsm_support_info(device);
 }
 
 cl::NDRange CLCompileContext::default_ndrange() const
@@ -346,6 +348,11 @@
     return _device.supported("cl_khr_int64_base_atomics");
 }
 
+bool CLCompileContext::is_wbsm_supported() const
+{
+    return _is_wbsm_supported;
+}
+
 size_t CLCompileContext::max_local_workgroup_size(const cl::Kernel &kernel) const
 {
     size_t result;