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>

12 files changed, 43 insertions, 18 deletions

diff --git a/arm_compute/core/CL/CLHelpers.h b/arm_compute/core/CL/CLHelpers.h
index ca1345d807..18d6bdf49f 100644
--- a/arm_compute/core/CL/CLHelpers.h
+++ b/arm_compute/core/CL/CLHelpers.h
@@ -69,7 +69,7 @@ std::string get_underlying_cl_type_from_data_type(const DataType &dt);
  *
  * @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 18b6bb4434..c1999b45e1 100644
--- a/arm_compute/core/CL/CLKernelLibrary.h
+++ b/arm_compute/core/CL/CLKernelLibrary.h
@@ -208,11 +208,11 @@ public:
     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 9e970641c1..150dd62a89 100644
--- a/arm_compute/core/CL/ICLKernel.h
+++ b/arm_compute/core/CL/ICLKernel.h
@@ -61,11 +61,17 @@ private:
     {
         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 @@ public:
      */
     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 @@ private:
 
 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 33053308ec..468e1792f0 100644
--- a/arm_compute/core/CL/OpenCL.h
+++ b/arm_compute/core/CL/OpenCL.h
@@ -82,6 +82,7 @@ public:
 #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);
diff --git a/arm_compute/graph/GraphContext.h b/arm_compute/graph/GraphContext.h
index 1831cc2c8b..ce6f86f611 100644
--- a/arm_compute/graph/GraphContext.h
+++ b/arm_compute/graph/GraphContext.h
@@ -50,6 +50,7 @@ class GraphContext final
 public:
     /** Constructor */
     GraphContext();
+    ~GraphContext();
     /** Prevent instances of this class from being copied (As this class contains pointers) */
     GraphContext(const GraphContext &) = delete;
     /** Default move constructor */
diff --git a/arm_compute/graph/IDeviceBackend.h b/arm_compute/graph/IDeviceBackend.h
index f28cb1ab42..358d26af81 100644
--- a/arm_compute/graph/IDeviceBackend.h
+++ b/arm_compute/graph/IDeviceBackend.h
@@ -53,9 +53,14 @@ public:
     virtual void initialize_backend() = 0;
     /** Setups the given graph context
      *
-     * @param[in] ctx Graph context
+     * @param[in,out] ctx Graph context
      */
     virtual void setup_backend_context(GraphContext &ctx) = 0;
+    /** Release the backend specific resources associated to a given graph context
+     *
+     * @param[in,out] ctx Graph context
+     */
+    virtual void release_backend_context(GraphContext &ctx) = 0;
     /** Checks if an instantiated backend is actually supported
      *
      * @return True if the backend is supported else false
diff --git a/arm_compute/graph/Utils.h b/arm_compute/graph/Utils.h
index 582d47e406..3604bad4af 100644
--- a/arm_compute/graph/Utils.h
+++ b/arm_compute/graph/Utils.h
@@ -91,9 +91,14 @@ void force_target_to_graph(Graph &g, Target target);
 PassManager create_default_pass_manager(Target target);
 /** Default setups the graph context if not done manually
  *
- * @param[in] ctx Graph Context
+ * @param[in,out] ctx Graph Context
  */
 void setup_default_graph_context(GraphContext &ctx);
+/** Default releases the graph context if not done manually
+ *
+ * @param[in,out] ctx Graph Context
+ */
+void release_default_graph_context(GraphContext &ctx);
 /** Get size of a tensor's given dimension depending on its layout
  *
  * @param[in] descriptor            Descriptor
diff --git a/arm_compute/graph/backends/CL/CLDeviceBackend.h b/arm_compute/graph/backends/CL/CLDeviceBackend.h
index c1a6a28e6c..cc8d55239e 100644
--- a/arm_compute/graph/backends/CL/CLDeviceBackend.h
+++ b/arm_compute/graph/backends/CL/CLDeviceBackend.h
@@ -54,6 +54,7 @@ public:
     // Inherited overridden methods
     void initialize_backend() override;
     void setup_backend_context(GraphContext &ctx) override;
+    void release_backend_context(GraphContext &ctx) override;
     bool                           is_backend_supported() override;
     IAllocator                    *backend_allocator() override;
     std::unique_ptr<ITensorHandle> create_tensor(const Tensor &tensor) override;
@@ -63,10 +64,10 @@ public:
     std::shared_ptr<arm_compute::IMemoryManager> create_memory_manager(MemoryManagerAffinity affinity) override;
 
 private:
-    bool                               _initialized; /**< Flag that specifies if the backend has been default initialized */
-    CLTuner                            _tuner;       /**< CL kernel tuner */
-    std::unique_ptr<CLBufferAllocator> _allocator;   /**< CL buffer affinity allocator */
-    std::string                        _tuner_file;  /** Filename to load/store the tuner's values from */
+    int                                _context_count; /**< Counts how many contexts are currently using the backend */
+    CLTuner                            _tuner;         /**< CL kernel tuner */
+    std::unique_ptr<CLBufferAllocator> _allocator;     /**< CL buffer affinity allocator */
+    std::string                        _tuner_file;    /** Filename to load/store the tuner's values from */
 };
 } // namespace backends
 } // namespace graph
diff --git a/arm_compute/graph/backends/GLES/GCDeviceBackend.h b/arm_compute/graph/backends/GLES/GCDeviceBackend.h
index ba789221e3..ca2d3734eb 100644
--- a/arm_compute/graph/backends/GLES/GCDeviceBackend.h
+++ b/arm_compute/graph/backends/GLES/GCDeviceBackend.h
@@ -44,6 +44,7 @@ public:
     // Inherited overridden methods
     void initialize_backend() override;
     void setup_backend_context(GraphContext &ctx) override;
+    void release_backend_context(GraphContext &ctx) override;
     bool                           is_backend_supported() override;
     IAllocator                    *backend_allocator() override;
     std::unique_ptr<ITensorHandle> create_tensor(const Tensor &tensor) override;
diff --git a/arm_compute/graph/backends/NEON/NEDeviceBackend.h b/arm_compute/graph/backends/NEON/NEDeviceBackend.h
index c1e2e0c078..abc17d9e83 100644
--- a/arm_compute/graph/backends/NEON/NEDeviceBackend.h
+++ b/arm_compute/graph/backends/NEON/NEDeviceBackend.h
@@ -43,6 +43,7 @@ public:
     // Inherited overridden methods
     void initialize_backend() override;
     void setup_backend_context(GraphContext &ctx) override;
+    void release_backend_context(GraphContext &ctx) override;
     bool                           is_backend_supported() override;
     IAllocator                    *backend_allocator() override;
     std::unique_ptr<ITensorHandle> create_tensor(const Tensor &tensor) override;
diff --git a/arm_compute/graph/frontend/Stream.h b/arm_compute/graph/frontend/Stream.h
index 244d18e753..c8e24eeae2 100644
--- a/arm_compute/graph/frontend/Stream.h
+++ b/arm_compute/graph/frontend/Stream.h
@@ -74,11 +74,13 @@ public:
     const Graph &graph() const override;
 
 private:
-    GraphManager _manager; /**< Graph manager */
+    //Important: GraphContext must be declared *before* the GraphManager because the GraphManager
+    //allocates resources from the context and therefore needs to be destroyed before the context during clean up).
     GraphContext _ctx;     /**< Graph context to use */
+    GraphManager _manager; /**< Graph manager */
     Graph        _g;       /**< Internal graph representation of the stream */
 };
 } // namespace frontend
 } // namespace graph
 } // namespace arm_compute
-#endif /* __ARM_COMPUTE_GRAPH_STREAM_H__ */
\ No newline at end of file
+#endif /* __ARM_COMPUTE_GRAPH_STREAM_H__ */
diff --git a/arm_compute/runtime/CL/CLScheduler.h b/arm_compute/runtime/CL/CLScheduler.h
index 8eb287c942..807d8bb448 100644
--- a/arm_compute/runtime/CL/CLScheduler.h
+++ b/arm_compute/runtime/CL/CLScheduler.h
@@ -74,8 +74,8 @@ public:
      * @param[in] cl_tuner (Optional) Pointer to OpenCL tuner (default=nullptr)
      *                     Note: It is caller's responsibility to release the allocated memory for CLTuner
      */
-    void init(cl::Context context = cl::Context::getDefault(), cl::CommandQueue queue = cl::CommandQueue::getDefault(),
-              cl::Device device = cl::Device::getDefault(), ICLTuner *cl_tuner = nullptr)
+    void init(cl::Context context, cl::CommandQueue queue,
+              cl::Device device, ICLTuner *cl_tuner = nullptr)
     {
         set_context(context);
         _queue          = std::move(queue);