diff options
Diffstat (limited to 'arm_compute')
-rw-r--r-- | arm_compute/core/CL/CLHelpers.h | 2 | ||||
-rw-r--r-- | arm_compute/core/CL/CLKernelLibrary.h | 8 | ||||
-rw-r--r-- | arm_compute/core/CL/ICLKernel.h | 14 | ||||
-rw-r--r-- | arm_compute/core/CL/OpenCL.h | 1 | ||||
-rw-r--r-- | arm_compute/graph/GraphContext.h | 1 | ||||
-rw-r--r-- | arm_compute/graph/IDeviceBackend.h | 7 | ||||
-rw-r--r-- | arm_compute/graph/Utils.h | 7 | ||||
-rw-r--r-- | arm_compute/graph/backends/CL/CLDeviceBackend.h | 9 | ||||
-rw-r--r-- | arm_compute/graph/backends/GLES/GCDeviceBackend.h | 1 | ||||
-rw-r--r-- | arm_compute/graph/backends/NEON/NEDeviceBackend.h | 1 | ||||
-rw-r--r-- | arm_compute/graph/frontend/Stream.h | 6 | ||||
-rw-r--r-- | arm_compute/runtime/CL/CLScheduler.h | 4 |
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); |