COMPMID-1478: Stop relying on static default OpenCL objects in cl2.hpp
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>
diff --git a/arm_compute/core/CL/CLHelpers.h b/arm_compute/core/CL/CLHelpers.h
index ca1345d..18d6bdf 100644
--- a/arm_compute/core/CL/CLHelpers.h
+++ b/arm_compute/core/CL/CLHelpers.h
@@ -69,7 +69,7 @@
*
* @return the GPU target
*/
-GPUTarget get_target_from_device(cl::Device &device);
+GPUTarget get_target_from_device(const cl::Device &device);
/** Helper function to get the highest OpenCL version supported
*
diff --git a/arm_compute/core/CL/CLKernelLibrary.h b/arm_compute/core/CL/CLKernelLibrary.h
index 18b6bb4..c1999b4 100644
--- a/arm_compute/core/CL/CLKernelLibrary.h
+++ b/arm_compute/core/CL/CLKernelLibrary.h
@@ -208,11 +208,11 @@
static CLKernelLibrary &get();
/** Initialises the kernel library.
*
- * @param[in] kernel_path (Optional) Path of the directory from which kernel sources are loaded.
- * @param[in] context (Optional) CL context used to create programs.
- * @param[in] device (Optional) CL device for which the programs are created.
+ * @param[in] kernel_path Path of the directory from which kernel sources are loaded.
+ * @param[in] context CL context used to create programs.
+ * @param[in] device CL device for which the programs are created.
*/
- void init(std::string kernel_path = ".", cl::Context context = cl::Context::getDefault(), cl::Device device = cl::Device::getDefault())
+ void init(std::string kernel_path, cl::Context context, cl::Device device)
{
_kernel_path = std::move(kernel_path);
_context = std::move(context);
diff --git a/arm_compute/core/CL/ICLKernel.h b/arm_compute/core/CL/ICLKernel.h
index 9e97064..150dd62 100644
--- a/arm_compute/core/CL/ICLKernel.h
+++ b/arm_compute/core/CL/ICLKernel.h
@@ -61,11 +61,17 @@
{
return 2 + 2 * dimension_size;
}
-
+ using IKernel::configure; //Prevent children from calling IKernel::configure() directly
public:
+ void configure_internal(const Window &window, cl::NDRange lws_hint = CLKernelLibrary::get().default_ndrange())
+ {
+ _lws_hint = lws_hint;
+ IKernel::configure(window);
+ }
+
/** Constructor */
ICLKernel()
- : _kernel(nullptr), _lws_hint(CLKernelLibrary::get().default_ndrange()), _target(GPUTarget::MIDGARD), _config_id(arm_compute::default_config_id), _max_workgroup_size(0)
+ : _kernel(nullptr), _target(GPUTarget::MIDGARD), _config_id(arm_compute::default_config_id), _max_workgroup_size(0), _lws_hint()
{
}
/** Returns a reference to the OpenCL kernel of this object.
@@ -196,6 +202,7 @@
*/
void set_lws_hint(const cl::NDRange &lws_hint)
{
+ ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // lws_hint will be overwritten by configure()
_lws_hint = lws_hint;
}
@@ -282,10 +289,11 @@
protected:
cl::Kernel _kernel; /**< OpenCL kernel to run */
- cl::NDRange _lws_hint; /**< Local workgroup size hint for the OpenCL kernel */
GPUTarget _target; /**< The targeted GPU */
std::string _config_id; /**< Configuration ID */
size_t _max_workgroup_size; /**< The maximum workgroup size for this kernel */
+private:
+ cl::NDRange _lws_hint; /**< Local workgroup size hint for the OpenCL kernel */
};
/** Add the kernel to the command queue with the given window.
diff --git a/arm_compute/core/CL/OpenCL.h b/arm_compute/core/CL/OpenCL.h
index 3305330..468e179 100644
--- a/arm_compute/core/CL/OpenCL.h
+++ b/arm_compute/core/CL/OpenCL.h
@@ -82,6 +82,7 @@
#define DECLARE_FUNCTION_PTR(func_name) \
std::function<decltype(func_name)> func_name##_ptr = nullptr
+ DECLARE_FUNCTION_PTR(clCreateContext);
DECLARE_FUNCTION_PTR(clCreateContextFromType);
DECLARE_FUNCTION_PTR(clCreateCommandQueue);
DECLARE_FUNCTION_PTR(clGetContextInfo);