COMPMID-3279: Create CLCompiler interface

Change-Id: Ic9dd5288d72a690651aa03d474f2bfd6e1ebe8b2
Signed-off-by: Michalis Spyrou <michalis.spyrou@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2957
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>

17 files changed, 1123 insertions, 487 deletions

diff --git a/Android.bp b/Android.bp
index 6e34cb5cca..f31dabbe07 100644
--- a/Android.bp
+++ b/Android.bp
@@ -51,6 +51,7 @@ cc_library_static {
         "src/core/AccessWindowAutoPadding.cpp",
         "src/core/AccessWindowStatic.cpp",
         "src/core/AccessWindowTranspose.cpp",
+        "src/core/CL/CLCompileContext.cpp",
         "src/core/CL/CLCoreRuntimeContext.cpp",
         "src/core/CL/CLHelpers.cpp",
         "src/core/CL/CLKernelLibrary.cpp",
diff --git a/arm_compute/core/CL/CLCompileContext.h b/arm_compute/core/CL/CLCompileContext.h
new file mode 100644
index 0000000000..2b6d8cd2cb
--- /dev/null
+++ b/arm_compute/core/CL/CLCompileContext.h
@@ -0,0 +1,324 @@
+/*
+ * Copyright (c) 2020 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_CLCOMPILECONTEXT_H
+#define ARM_COMPUTE_CLCOMPILECONTEXT_H
+
+#include "arm_compute/core/CL/CLDevice.h"
+#include "arm_compute/core/CL/OpenCL.h"
+
+#include <map>
+#include <set>
+#include <string>
+#include <utility>
+
+namespace arm_compute
+{
+/** Build options */
+class CLBuildOptions final
+{
+    using StringSet = std::set<std::string>;
+
+public:
+    /** Default constructor. */
+    CLBuildOptions();
+    /** Adds option to the existing build option list
+     *
+     * @param[in] option Option to add
+     */
+    void add_option(std::string option);
+    /** Adds option if a given condition is true;
+     *
+     * @param[in] cond   Condition to check
+     * @param[in] option Option to add if condition is true
+     */
+    void add_option_if(bool cond, std::string option);
+    /** Adds first option if condition is true else the second one
+     *
+     * @param[in] cond         Condition to check
+     * @param[in] option_true  Option to add if condition is true
+     * @param[in] option_false Option to add if condition is false
+     */
+    void add_option_if_else(bool cond, std::string option_true, std::string option_false);
+    /** Appends given build options to the current's objects options.
+     *
+     * @param[in] options Build options to append
+     */
+    void add_options(const StringSet &options);
+    /** Appends given build options to the current's objects options if a given condition is true.
+     *
+     * @param[in] cond    Condition to check
+     * @param[in] options Option to add if condition is true
+     */
+    void add_options_if(bool cond, const StringSet &options);
+    /** Gets the current options list set
+     *
+     * @return Build options set
+     */
+    const StringSet &options() const;
+
+private:
+    StringSet _build_opts; /**< Build options set */
+};
+
+/** Program class */
+class Program final
+{
+public:
+    /** Default constructor. */
+    Program();
+    /** Construct program from source file.
+     *
+     * @param[in] context CL context used to create the program.
+     * @param[in] name    Program name.
+     * @param[in] source  Program source.
+     */
+    Program(cl::Context context, std::string name, std::string source);
+    /** Construct program from binary file.
+     *
+     * @param[in] context CL context used to create the program.
+     * @param[in] device  CL device for which the programs are created.
+     * @param[in] name    Program name.
+     * @param[in] binary  Program binary.
+     */
+    Program(cl::Context context, cl::Device device, std::string name, std::vector<unsigned char> binary);
+    /** Default Copy Constructor. */
+    Program(const Program &) = default;
+    /** Default Move Constructor. */
+    Program(Program &&) = default;
+    /** Default copy assignment operator */
+    Program &operator=(const Program &) = default;
+    /** Default move assignment operator */
+    Program &operator=(Program &&) = default;
+    /** Returns program name.
+     *
+     * @return Program's name.
+     */
+    std::string name() const
+    {
+        return _name;
+    }
+    /** User-defined conversion to the underlying CL program.
+     *
+     * @return The CL program object.
+     */
+    explicit operator cl::Program() const;
+    /** Build the given CL program.
+     *
+     * @param[in] program       The CL program to build.
+     * @param[in] build_options Options to build the CL program.
+     *
+     * @return True if the CL program builds successfully.
+     */
+    static bool build(const cl::Program &program, const std::string &build_options = "");
+    /** Build the underlying CL program.
+     *
+     * @param[in] build_options Options used to build the CL program.
+     *
+     * @return A reference to itself.
+     */
+    cl::Program build(const std::string &build_options = "") const;
+
+private:
+    cl::Context                _context;   /**< Underlying CL context. */
+    cl::Device                 _device;    /**< CL device for which the programs are created. */
+    bool                       _is_binary; /**< Create program from binary? */
+    std::string                _name;      /**< Program name. */
+    std::string                _source;    /**< Source code for the program. */
+    std::vector<unsigned char> _binary;    /**< Binary from which to create the program. */
+};
+
+/** Kernel class */
+class Kernel final
+{
+public:
+    /** Default Constructor. */
+    Kernel();
+    /** Default Copy Constructor. */
+    Kernel(const Kernel &) = default;
+    /** Default Move Constructor. */
+    Kernel(Kernel &&) = default;
+    /** Default copy assignment operator */
+    Kernel &operator=(const Kernel &) = default;
+    /** Default move assignment operator */
+    Kernel &operator=(Kernel &&) = default;
+    /** Constructor.
+     *
+     * @param[in] name    Kernel name.
+     * @param[in] program Built program.
+     */
+    Kernel(std::string name, const cl::Program &program);
+    /** Returns kernel name.
+     *
+     * @return Kernel's name.
+     */
+    std::string name() const
+    {
+        return _name;
+    }
+    /** Returns OpenCL kernel.
+     *
+     * @return OpenCL Kernel.
+     */
+    explicit operator cl::Kernel() const
+    {
+        return _kernel;
+    }
+
+private:
+    std::string _name;   /**< Kernel name */
+    cl::Kernel  _kernel; /**< OpenCL Kernel */
+};
+
+/** CLCompileContext class */
+class CLCompileContext final
+{
+    using StringSet = std::set<std::string>;
+
+public:
+    /** Constructor */
+    CLCompileContext();
+    /** Constructor
+     *
+     * @param[in] context A CL context.
+     * @param[in] device  A CL device.
+     * */
+    CLCompileContext(cl::Context context, const cl::Device &device);
+
+    /** Accessor for the associated CL context.
+     *
+     * @return A CL context.
+     */
+    cl::Context &context();
+
+    /** Sets the CL context used to create programs.
+     *
+     * @note Setting the context also resets the device to the
+     *       first one available in the new context.
+     *
+     * @param[in] context A CL context.
+     */
+    void set_context(cl::Context context);
+
+    /** Gets the CL device for which the programs are created. */
+    const cl::Device &get_device() const;
+
+    /** Sets the CL device for which the programs are created.
+     *
+     * @param[in] device A CL device.
+     */
+    void set_device(cl::Device device);
+
+    /** Creates an OpenCL kernel.
+     *
+     * @param[in] kernel_name       Kernel name.
+     * @param[in] program_name      Program name.
+     * @param[in] program_source    Program source.
+     * @param[in] kernel_path       CL kernel path.
+     * @param[in] build_options_set Kernel build options as a set.
+     * @param[in] is_binary         Flag to indicate if the program source is binary.
+     *
+     * @return The created kernel.
+     */
+    Kernel create_kernel(const std::string &kernel_name, const std::string &program_name, const std::string &program_source,
+                         const std::string &kernel_path, const StringSet &build_options_set, bool is_binary) const;
+
+    /** Clear the library's cache of binary programs
+     */
+    void clear_programs_cache();
+
+    /** Access the cache of built OpenCL programs */
+    const std::map<std::string, cl::Program> &get_built_programs() const;
+
+    /** Add a new built program to the cache
+     *
+     * @param[in] built_program_name Name of the program
+     * @param[in] program            Built program to add to the cache
+     */
+    void add_built_program(const std::string &built_program_name, const cl::Program &program) const;
+
+    /** Returns true if FP16 is supported by the CL device
+     *
+     * @return true if the CL device supports FP16
+     */
+    bool fp16_supported() const;
+
+    /** Return the maximum number of compute units in the device
+     *
+     * @return The content of CL_DEVICE_MAX_COMPUTE_UNITS
+     */
+    cl_uint get_num_compute_units() const;
+    /** Find the maximum number of local work items in a workgroup can be supported for the kernel.
+     *
+     */
+    size_t max_local_workgroup_size(const cl::Kernel &kernel) const;
+    /** Return the default NDRange for the device.
+     *
+     */
+    cl::NDRange default_ndrange() const;
+    /** Return the device version
+     *
+     * @return The content of CL_DEVICE_VERSION
+     */
+    std::string get_device_version() const;
+
+    /** Returns true if int64_base_atomics extension is supported by the CL device
+     *
+     * @return true if the CL device supports int64_base_atomics extension
+     */
+    bool int64_base_atomics_supported() const;
+
+private:
+    /** Load program and its dependencies.
+     *
+     * @param[in] program_name   Name of the program to load.
+     * @param[in] program_source Source of the program.
+     * @param[in] is_binary      Flag to indicate if the program source is binary.
+     */
+    const Program &load_program(const std::string &program_name, const std::string &program_source, bool is_binary) const;
+
+    /** Generates the build options given a string of user defined ones
+     *
+     * @param[in] build_options User defined build options
+     * @param[in] kernel_path   Path of the CL kernels
+     *
+     * @return Generated build options
+     */
+    std::string generate_build_options(const StringSet &build_options, const std::string &kernel_path) const;
+
+    /** Concatenates contents of a set into a single string.
+     *
+     * @param[in] s           Input set to concatenate.
+     * @param[in] kernel_path Path of the CL kernels
+     *
+     * @return Concatenated string.
+     */
+    std::string stringify_set(const StringSet &s, const std::string &kernel_path) const;
+
+    cl::Context _context;                                             /**< Underlying CL context. */
+    CLDevice    _device;                                              /**< Underlying CL device. */
+    mutable std::map<std::string, const Program> _programs_map;       /**< Map with all already loaded program data. */
+    mutable std::map<std::string, cl::Program>   _built_programs_map; /**< Map with all already built program data. */
+};
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_CLCOMPILECONTEXT_H */
diff --git a/arm_compute/core/CL/CLDevice.h b/arm_compute/core/CL/CLDevice.h
new file mode 100644
index 0000000000..812834743d
--- /dev/null
+++ b/arm_compute/core/CL/CLDevice.h
@@ -0,0 +1,152 @@
+/*
+ * Copyright (c) 2020 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_CLDEVICE_H
+#define ARM_COMPUTE_CLDEVICE_H
+
+#include "arm_compute/core/CL/CLHelpers.h"
+#include "arm_compute/core/CL/CLTypes.h"
+#include "arm_compute/core/GPUTarget.h"
+#include "arm_compute/core/IDevice.h"
+
+#include <set>
+#include <string>
+
+namespace arm_compute
+{
+/**  OpenCL device type class
+ *
+ *   Initializes and stores all the information about a cl device,
+ *   working mainly as a cache mechanism.
+ * */
+class CLDevice : public IDevice
+{
+public:
+    /** Default Constructor */
+    CLDevice()
+        : _device(cl::Device()), _options()
+    {
+    }
+
+    /** Constructor
+     *
+     * @param[in] cl_device OpenCL device
+     */
+    CLDevice(const cl::Device &cl_device)
+        : _device(), _options()
+    {
+        _device = cl_device;
+
+        // Get device target
+        std::string device_name = _device.getInfo<CL_DEVICE_NAME>();
+        _options.gpu_target     = get_target_from_name(device_name);
+
+        // Fill extensions
+        std::string extensions = _device.getInfo<CL_DEVICE_EXTENSIONS>();
+
+        std::istringstream iss(extensions);
+        for(std::string s; iss >> s;)
+        {
+            _options.extensions.insert(s);
+        }
+
+        // SW workaround for G76
+        if(_options.gpu_target == GPUTarget::G76)
+        {
+            _options.extensions.insert("cl_arm_integer_dot_product_int8");
+        }
+
+        // Get device version
+        _options.version = get_cl_version(_device);
+
+        // Get compute units
+        _options.compute_units = _device.getInfo<CL_DEVICE_MAX_COMPUTE_UNITS>();
+
+        // Get device version
+        _options.device_version = _device.getInfo<CL_DEVICE_VERSION>();
+    }
+
+    /** Returns the GPU target of the cl device
+     *
+     * @return The GPU target
+     */
+    const GPUTarget &target() const
+    {
+        return _options.gpu_target;
+    }
+
+    /** Returns the number of compute units available
+     *
+     * @return Number of compute units
+     */
+    size_t compute_units() const
+    {
+        return _options.compute_units;
+    }
+
+    /** Returns the underlying cl device object
+     *
+     * @return A cl device
+     */
+    const cl::Device &cl_device() const
+    {
+        return _device;
+    }
+
+    /** Returns the device's CL version
+     *
+     * @return CLVersion of the device
+     */
+    CLVersion version() const
+    {
+        return _options.version;
+    }
+
+    /** Returns the device version as a string
+     *
+     * @return CLVersion of the device
+     */
+    std::string device_version() const
+    {
+        return _options.device_version;
+    }
+
+    // Inherrited methods
+    DeviceType type() const override
+    {
+        return DeviceType::CL;
+    }
+
+    bool supported(const std::string &extension) const override
+    {
+        return _options.extensions.count(extension) != 0;
+    }
+
+private:
+    cl::Device             _device;  /**< OpenCL device. */
+    struct CLDeviceOptions _options; /**< OpenCL device options */
+};
+
+} // namespace arm_compute
+
+#endif /* ARM_COMPUTE_CLDEVICE_H */
diff --git a/arm_compute/core/CL/CLHelpers.h b/arm_compute/core/CL/CLHelpers.h
index 453384e14e..ee6397af7a 100644
--- a/arm_compute/core/CL/CLHelpers.h
+++ b/arm_compute/core/CL/CLHelpers.h
@@ -27,11 +27,13 @@
 #include "arm_compute/core/CL/CLTypes.h"
 #include "arm_compute/core/CL/OpenCL.h"
 
+#include <set>
 #include <string>
 
 namespace arm_compute
 {
 class CLCoreRuntimeContext;
+class CLCompileContext;
 class CLBuildOptions;
 
 enum class DataType;
@@ -196,6 +198,16 @@ bool preferred_dummy_work_items_support(const cl::Device &device);
  */
 cl::Kernel create_opencl_kernel(CLCoreRuntimeContext *ctx, const std::string &kernel_name, const CLBuildOptions &build_opts);
 
+/** Creates an opencl kernel using a compile context
+ *
+ * @param[in] ctx         A compile context to be used to create the opencl kernel.
+ * @param[in] kernel_name The kernel name.
+ * @param[in] build_opts  The build options to be used for the opencl kernel compilation.
+ *
+ * @return An opencl kernel
+ */
+cl::Kernel create_kernel(CLCompileContext &ctx, const std::string &kernel_name, const std::set<std::string> &build_opts);
+
 /** Creates a suitable LWS hint object for parallel implementations. Sets the number of WG based on the input size.
  *  If input width is smaller than 128 we can use fewer threads than 8.
  *
diff --git a/arm_compute/core/CL/CLKernelLibrary.h b/arm_compute/core/CL/CLKernelLibrary.h
index 2d55351c95..6c5df6cb08 100644
--- a/arm_compute/core/CL/CLKernelLibrary.h
+++ b/arm_compute/core/CL/CLKernelLibrary.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016-2019 ARM Limited.
+ * Copyright (c) 2016-2020 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -24,6 +24,7 @@
 #ifndef ARM_COMPUTE_CLKERNELLIBRARY_H
 #define ARM_COMPUTE_CLKERNELLIBRARY_H
 
+#include "arm_compute/core/CL/CLCompileContext.h"
 #include "arm_compute/core/CL/OpenCL.h"
 
 #include <map>
@@ -33,173 +34,18 @@
 
 namespace arm_compute
 {
-/** Build options */
-class CLBuildOptions final
-{
-    using StringSet = std::set<std::string>;
-
-public:
-    /** Default constructor. */
-    CLBuildOptions();
-    /** Adds option to the existing build option list
-     *
-     * @param[in] option Option to add
-     */
-    void add_option(std::string option);
-    /** Adds option if a given condition is true;
-     *
-     * @param[in] cond   Condition to check
-     * @param[in] option Option to add if condition is true
-     */
-    void add_option_if(bool cond, std::string option);
-    /** Adds first option if condition is true else the second one
-     *
-     * @param[in] cond         Condition to check
-     * @param[in] option_true  Option to add if condition is true
-     * @param[in] option_false Option to add if condition is false
-     */
-    void add_option_if_else(bool cond, std::string option_true, std::string option_false);
-    /** Appends given build options to the current's objects options.
-     *
-     * @param[in] options Build options to append
-     */
-    void add_options(const StringSet &options);
-    /** Appends given build options to the current's objects options if a given condition is true.
-     *
-     * @param[in] cond    Condition to check
-     * @param[in] options Option to add if condition is true
-     */
-    void add_options_if(bool cond, const StringSet &options);
-    /** Gets the current options list set
-     *
-     * @return Build options set
-     */
-    const StringSet &options() const;
-
-private:
-    StringSet _build_opts; /**< Build options set */
-};
-/** Program class */
-class Program final
-{
-public:
-    /** Default constructor. */
-    Program();
-    /** Construct program from source file.
-     *
-     * @param[in] context CL context used to create the program.
-     * @param[in] name    Program name.
-     * @param[in] source  Program source.
-     */
-    Program(cl::Context context, std::string name, std::string source);
-    /** Construct program from binary file.
-     *
-     * @param[in] context CL context used to create the program.
-     * @param[in] device  CL device for which the programs are created.
-     * @param[in] name    Program name.
-     * @param[in] binary  Program binary.
-     */
-    Program(cl::Context context, cl::Device device, std::string name, std::vector<unsigned char> binary);
-    /** Default Copy Constructor. */
-    Program(const Program &) = default;
-    /** Default Move Constructor. */
-    Program(Program &&) = default;
-    /** Default copy assignment operator */
-    Program &operator=(const Program &) = default;
-    /** Default move assignment operator */
-    Program &operator=(Program &&) = default;
-    /** Returns program name.
-     *
-     * @return Program's name.
-     */
-    std::string name() const
-    {
-        return _name;
-    }
-    /** User-defined conversion to the underlying CL program.
-     *
-     * @return The CL program object.
-     */
-    explicit operator cl::Program() const;
-    /** Build the given CL program.
-     *
-     * @param[in] program       The CL program to build.
-     * @param[in] build_options Options to build the CL program.
-     *
-     * @return True if the CL program builds successfully.
-     */
-    static bool build(const cl::Program &program, const std::string &build_options = "");
-    /** Build the underlying CL program.
-     *
-     * @param[in] build_options Options used to build the CL program.
-     *
-     * @return A reference to itself.
-     */
-    cl::Program build(const std::string &build_options = "") const;
-
-private:
-    cl::Context                _context;   /**< Underlying CL context. */
-    cl::Device                 _device;    /**< CL device for which the programs are created. */
-    bool                       _is_binary; /**< Create program from binary? */
-    std::string                _name;      /**< Program name. */
-    std::string                _source;    /**< Source code for the program. */
-    std::vector<unsigned char> _binary;    /**< Binary from which to create the program. */
-};
-
-/** Kernel class */
-class Kernel final
-{
-public:
-    /** Default Constructor. */
-    Kernel();
-    /** Default Copy Constructor. */
-    Kernel(const Kernel &) = default;
-    /** Default Move Constructor. */
-    Kernel(Kernel &&) = default;
-    /** Default copy assignment operator */
-    Kernel &operator=(const Kernel &) = default;
-    /** Default move assignment operator */
-    Kernel &operator=(Kernel &&) = default;
-    /** Constructor.
-     *
-     * @param[in] name    Kernel name.
-     * @param[in] program Built program.
-     */
-    Kernel(std::string name, const cl::Program &program);
-    /** Returns kernel name.
-     *
-     * @return Kernel's name.
-     */
-    std::string name() const
-    {
-        return _name;
-    }
-    /** Returns OpenCL kernel.
-     *
-     * @return OpenCL Kernel.
-     */
-    explicit operator cl::Kernel() const
-    {
-        return _kernel;
-    }
-
-private:
-    std::string _name;   /**< Kernel name */
-    cl::Kernel  _kernel; /**< OpenCL Kernel */
-};
-
 /** CLKernelLibrary class */
 class CLKernelLibrary final
 {
-    using StringSet = std::set<std::string>;
-
-public:
+private:
     /** Default Constructor. */
     CLKernelLibrary();
     /** Prevent instances of this class from being copied */
     CLKernelLibrary(const CLKernelLibrary &) = delete;
     /** Prevent instances of this class from being copied */
     const CLKernelLibrary &operator=(const CLKernelLibrary &) = delete;
+
+public:
     /** Access the KernelLibrary singleton.
      * This method has been deprecated and will be removed in the next release.
      * @return The KernelLibrary instance.
@@ -224,17 +70,9 @@ public:
      *
      * @param[in] program_name Program name.
      *
-     * @return Source of the selected program.
+     * @return A pair with the source (false) or the binary (true), of the selected program.
      */
-    std::string get_program_source(const std::string &program_name);
-    /** Sets the CL context used to create programs.
-     *
-     * @note Setting the context also resets the device to the
-     *       first one available in the new context.
-     *
-     * @param[in] context A CL context.
-     */
-    void set_context(cl::Context context);
+    std::pair<std::string, bool> get_program(const std::string &program_name) const;
 
     /** Accessor for the associated CL context.
      *
@@ -243,7 +81,7 @@ public:
     cl::Context &context();
 
     /** Gets the CL device for which the programs are created. */
-    cl::Device &get_device();
+    const cl::Device &get_device();
 
     /** Sets the CL device for which the programs are created.
      *
@@ -268,7 +106,7 @@ public:
      *
      * @return The created kernel.
      */
-    Kernel create_kernel(const std::string &kernel_name, const StringSet &build_options_set = {}) const;
+    Kernel create_kernel(const std::string &kernel_name, const std::set<std::string> &build_options_set = {}) const;
     /** Find the maximum number of local work items in a workgroup can be supported for the kernel.
      *
      */
@@ -304,28 +142,33 @@ public:
      */
     bool int64_base_atomics_supported() const;
 
-private:
-    /** Load program and its dependencies.
+    /** Returns the program name given a kernel name
      *
-     * @param[in] program_name Name of the program to load.
+     * @return Program name
      */
-    const Program &load_program(const std::string &program_name) const;
-    /** Concatenates contents of a set into a single string.
+    std::string get_program_name(const std::string &kernel_name) const;
+
+    /** Sets the CL context used to create programs.
      *
-     * @param[in] s Input set to concatenate.
+     * @note Setting the context also resets the device to the
+     *       first one available in the new context.
      *
-     * @return Concatenated string.
+     * @param[in] context A CL context.
      */
-    std::string stringify_set(const StringSet &s) const;
+    void set_context(cl::Context context);
 
-    cl::Context _context;                                                /**< Underlying CL context. */
-    cl::Device  _device;                                                 /**< Underlying CL device. */
-    std::string _kernel_path;                                            /**< Path to the kernels folder. */
-    mutable std::map<std::string, const Program>    _programs_map;       /**< Map with all already loaded program data. */
-    mutable std::map<std::string, cl::Program>      _built_programs_map; /**< Map with all already built program data. */
+    /** Gets the compile context used
+     *
+     * @return The used compile context
+     */
+    CLCompileContext &get_compile_context();
+
+private:
+    CLCompileContext _compile_context;                                   /**< Compile Context. */
+    std::string      _kernel_path;                                       /**< Path to the kernels folder. */
     static const std::map<std::string, std::string> _kernel_program_map; /**< Map that associates kernel names with programs. */
     static const std::map<std::string, std::string> _program_source_map; /**< Contains sources for all programs.
-                                                                              Used for compile-time kernel inclusion. >*/
+                                                                            Used for compile-time kernel inclusion. >*/
 };
 } // namespace arm_compute
 #endif /* ARM_COMPUTE_CLKERNELLIBRARY_H */
diff --git a/arm_compute/core/CL/CLTypes.h b/arm_compute/core/CL/CLTypes.h
index 9f6ff6a82b..3643b178d3 100644
--- a/arm_compute/core/CL/CLTypes.h
+++ b/arm_compute/core/CL/CLTypes.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017-2019 ARM Limited.
+ * Copyright (c) 2017-2020 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -27,6 +27,7 @@
 #include "arm_compute/core/CL/ICLArray.h"
 #include "arm_compute/core/GPUTarget.h"
 
+#include <set>
 #include <string>
 
 namespace arm_compute
@@ -47,12 +48,14 @@ enum class CLVersion
 /** OpenCL device options */
 struct CLDeviceOptions
 {
-    std::string name;        /**< Device name */
-    std::string extensions;  /**< List of supported extensions */
-    std::string ddk_version; /**< DDK version */
-    GPUTarget   gpu_target;  /**< GPU target architecture/instance */
-    size_t      num_cores;   /**< Number of cores */
-    size_t      cache_size;  /**< Cache size */
+    std::string           name{};           /**< Device name */
+    std::string           device_version{}; /**< Device version string */
+    std::set<std::string> extensions{};     /**< List of supported extensions */
+    std::string           ddk_version{};    /**< DDK version */
+    GPUTarget             gpu_target{};     /**< GPU target architecture/instance */
+    CLVersion             version{};        /**< Device OpenCL version */
+    size_t                compute_units{};  /**< Number of compute units */
+    size_t                cache_size{};     /**< Cache size */
 };
 
 /** OpenCL quantization data */
diff --git a/arm_compute/core/CL/kernels/CLFloorKernel.h b/arm_compute/core/CL/kernels/CLFloorKernel.h
index 00f77dcd6b..a3ccb96c61 100644
--- a/arm_compute/core/CL/kernels/CLFloorKernel.h
+++ b/arm_compute/core/CL/kernels/CLFloorKernel.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017-2019 ARM Limited.
+ * Copyright (c) 2017-2020 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -52,6 +52,15 @@ public:
      * @param[out] output Destination tensor. Same as @p input
      */
     void configure(const ICLTensor *input, ICLTensor *output);
+
+    /** Set the source, destination of the kernel
+     *
+     * @param[in]  compile_context The compile context to be used.
+     * @param[in]  input           Source tensor. Data type supported: F16/F32.
+     * @param[out] output          Destination tensor. Same as @p input
+     */
+    void configure(CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output);
+
     /** Static function to check if given info will lead to a valid configuration of @ref CLFloorKernel
      *
      * @param[in] input  Source tensor info. Data type supported: F16/F32.
diff --git a/arm_compute/core/IDevice.h b/arm_compute/core/IDevice.h
new file mode 100644
index 0000000000..5cffe646d4
--- /dev/null
+++ b/arm_compute/core/IDevice.h
@@ -0,0 +1,60 @@
+/*
+ * Copyright (c) 2020 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+
+#ifndef ARM_COMPUTE_IDEVICE_H
+#define ARM_COMPUTE_IDEVICE_H
+
+#include <string>
+
+namespace arm_compute
+{
+/** Device types */
+enum class DeviceType
+{
+    NEON,
+    CL,
+    GLES
+};
+
+/** Interface for device object */
+class IDevice
+{
+public:
+    /** Virtual Destructor */
+    virtual ~IDevice() = default;
+
+    /** Device type accessor */
+    virtual DeviceType type() const = 0;
+
+    /** Check if extensions on a device are supported
+     *
+     * @param[in] extension An extension to check if it's supported.
+     *
+     * @return True if the extension is supported else false
+     */
+    virtual bool supported(const std::string &extension) const = 0;
+};
+} // namespace arm_compute
+
+#endif /* ARM_COMPUTE_IDEVICE_H */
diff --git a/arm_compute/runtime/CL/CLRuntimeContext.h b/arm_compute/runtime/CL/CLRuntimeContext.h
index 791d1deaa7..54c7d3cd23 100644
--- a/arm_compute/runtime/CL/CLRuntimeContext.h
+++ b/arm_compute/runtime/CL/CLRuntimeContext.h
@@ -58,7 +58,6 @@ private:
     std::unique_ptr<CLScheduler> _gpu_owned_scheduler{ nullptr };
     CLScheduler                 *_gpu_scheduler{ nullptr };
     CLTuner                      _tuner{ false };
-    CLKernelLibrary              _kernel_lib{};
     CLSymbols                    _symbols{};
     CLCoreRuntimeContext         _core_context{};
 };
diff --git a/arm_compute/runtime/CL/functions/CLFloor.h b/arm_compute/runtime/CL/functions/CLFloor.h
index 33d03217e7..c4a893fdeb 100644
--- a/arm_compute/runtime/CL/functions/CLFloor.h
+++ b/arm_compute/runtime/CL/functions/CLFloor.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017-2019 ARM Limited.
+ * Copyright (c) 2017-2020 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -42,6 +42,13 @@ public:
      * @param[out] output Destination tensor. Same as @p input
      */
     void configure(const ICLTensor *input, ICLTensor *output);
+    /** Set the source, destination of the kernel
+     *
+     * @param[in]  compile_context The compile context to be used.
+     * @param[in]  input           Source tensor. Data type supported: F16/F32.
+     * @param[out] output          Destination tensor. Same as @p input
+     */
+    void configure(CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output);
     /** Static function to check if given info will lead to a valid configuration of @ref CLFloor
      *
      * @param[in] input  Source tensor info. Data type supported: F16/F32.
diff --git a/src/core/CL/CLCompileContext.cpp b/src/core/CL/CLCompileContext.cpp
new file mode 100644
index 0000000000..48cc64c387
--- /dev/null
+++ b/src/core/CL/CLCompileContext.cpp
@@ -0,0 +1,369 @@
+/*
+ * Copyright (c) 2020 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/CL/CLCompileContext.h"
+#include "arm_compute/core/CL/OpenCL.h"
+
+#include "arm_compute/core/CL/CLHelpers.h"
+#include "arm_compute/core/Error.h"
+#include "arm_compute/core/Utils.h"
+#include "support/StringSupport.h"
+
+namespace arm_compute
+{
+CLBuildOptions::CLBuildOptions()
+    : _build_opts()
+{
+}
+
+void CLBuildOptions::add_option(std::string option)
+{
+    _build_opts.emplace(std::move(option));
+}
+
+void CLBuildOptions::add_option_if(bool cond, std::string option)
+{
+    if(cond)
+    {
+        add_option(std::move(option));
+    }
+}
+
+void CLBuildOptions::add_option_if_else(bool cond, std::string option_true, std::string option_false)
+{
+    (cond) ? add_option(std::move(option_true)) : add_option(std::move(option_false));
+}
+
+void CLBuildOptions::add_options(const StringSet &options)
+{
+    _build_opts.insert(options.begin(), options.end());
+}
+
+void CLBuildOptions::add_options_if(bool cond, const StringSet &options)
+{
+    if(cond)
+    {
+        add_options(options);
+    }
+}
+
+const CLBuildOptions::StringSet &CLBuildOptions::options() const
+{
+    return _build_opts;
+}
+
+Program::Program()
+    : _context(), _device(), _is_binary(false), _name(), _source(), _binary()
+{
+}
+
+Program::Program(cl::Context context, std::string name, std::string source)
+    : _context(std::move(context)), _device(), _is_binary(false), _name(std::move(name)), _source(std::move(source)), _binary()
+{
+}
+
+Program::Program(cl::Context context, cl::Device device, std::string name, std::vector<unsigned char> binary)
+    : _context(std::move(context)), _device(std::move(device)), _is_binary(true), _name(std::move(name)), _source(), _binary(std::move(binary))
+{
+}
+
+Program::operator cl::Program() const
+{
+    if(_is_binary)
+    {
+        return cl::Program(_context, { _device }, { _binary });
+    }
+    else
+    {
+        return cl::Program(_context, _source, false);
+    }
+}
+
+bool Program::build(const cl::Program &program, const std::string &build_options)
+{
+    try
+    {
+        return program.build(build_options.c_str()) == CL_SUCCESS;
+    }
+    catch(const cl::Error &e)
+    {
+        cl_int     err        = CL_SUCCESS;
+        const auto build_info = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(&err);
+
+        for(auto &pair : build_info)
+        {
+            std::cerr << pair.second << std::endl;
+        }
+
+        return false;
+    }
+}
+
+cl::Program Program::build(const std::string &build_options) const
+{
+    cl::Program cl_program = static_cast<cl::Program>(*this);
+    build(cl_program, build_options);
+    return cl_program;
+}
+
+Kernel::Kernel()
+    : _name(), _kernel()
+{
+}
+
+Kernel::Kernel(std::string name, const cl::Program &program)
+    : _name(std::move(name)),
+      _kernel(cl::Kernel(program, _name.c_str()))
+{
+}
+CLCompileContext::CLCompileContext()
+    : _context(), _device(), _programs_map(), _built_programs_map()
+{
+}
+
+CLCompileContext::CLCompileContext(cl::Context context, const cl::Device &device)
+    : _context(), _device(), _programs_map(), _built_programs_map()
+{
+    _context = std::move(context);
+    _device  = CLDevice(device);
+}
+
+Kernel CLCompileContext::create_kernel(const std::string &kernel_name, const std::string &program_name, const std::string &program_source,
+                                       const std::string &kernel_path, const StringSet &build_options_set, bool is_binary) const
+{
+    const std::string build_options      = generate_build_options(build_options_set, kernel_path);
+    const std::string built_program_name = program_name + "_" + build_options;
+    auto              built_program_it   = _built_programs_map.find(built_program_name);
+    cl::Program       cl_program;
+
+    if(_built_programs_map.end() != built_program_it)
+    {
+        // If program has been built, retrieve to create kernel from it
+        cl_program = built_program_it->second;
+    }
+    else
+    {
+        Program program = load_program(program_name, program_source, is_binary);
+
+        // Build program
+        cl_program = program.build(build_options);
+
+        // Add built program to internal map
+        _built_programs_map.emplace(program_name, cl_program);
+    }
+
+    // Create and return kernel
+    return Kernel(kernel_name, cl_program);
+}
+
+const Program &CLCompileContext::load_program(const std::string &program_name, const std::string &program_source, bool is_binary) const
+{
+    const auto program_it = _programs_map.find(program_name);
+
+    if(program_it != _programs_map.end())
+    {
+        return program_it->second;
+    }
+
+    Program program;
+
+#ifdef EMBEDDED_KERNELS
+    ARM_COMPUTE_UNUSED(is_binary);
+    program = Program(_context, program_name, program_source);
+#else  /* EMBEDDED_KERNELS */
+    if(is_binary)
+    {
+        program = Program(_context, _device.cl_device(), program_name, std::vector<unsigned char>(program_source.begin(), program_source.end()));
+    }
+    else
+    {
+        program = Program(_context, program_name, program_source);
+    }
+#endif /* EMBEDDED_KERNELS */
+
+    // Insert program to program map
+    const auto new_program = _programs_map.emplace(program_name, std::move(program));
+
+    return new_program.first->second;
+}
+
+void CLCompileContext::set_context(cl::Context context)
+{
+    _context = std::move(context);
+    if(_context.get() != nullptr)
+    {
+        const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
+
+        if(!cl_devices.empty())
+        {
+            _device = CLDevice(cl_devices[0]);
+        }
+    }
+}
+
+std::string CLCompileContext::generate_build_options(const StringSet &build_options_set, const std::string &kernel_path) const
+{
+    std::string concat_str;
+
+#if defined(ARM_COMPUTE_DEBUG_ENABLED)
+    // Enable debug properties in CL kernels
+    concat_str += " -DARM_COMPUTE_DEBUG_ENABLED";
+#endif // defined(ARM_COMPUTE_DEBUG_ENABLED)
+
+    GPUTarget gpu_arch = get_arch_from_target(_device.target());
+    concat_str += " -DGPU_ARCH=" + support::cpp11::to_string(
+                      static_cast<std::underlying_type<GPUTarget>::type>(gpu_arch));
+
+    if(_device.supported("cl_khr_fp16"))
+    {
+        concat_str += " -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 ";
+    }
+
+    if(_device.supported("cl_arm_integer_dot_product_int8"))
+    {
+        concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ENABLED=1 ";
+    }
+
+    if(_device.supported("cl_arm_integer_dot_product_accumulate_int8"))
+    {
+        concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED=1 ";
+    }
+
+    if(_device.version() == CLVersion::CL20)
+    {
+        concat_str += " -cl-std=CL2.0 ";
+    }
+    else if(_device.supported("cl_arm_non_uniform_work_group_size"))
+    {
+        concat_str += " -cl-arm-non-uniform-work-group-size ";
+    }
+    else
+    {
+        ARM_COMPUTE_ERROR("Non uniform workgroup size is not supported!!");
+    }
+
+    std::string build_options = stringify_set(build_options_set, kernel_path) + concat_str;
+
+    return build_options;
+}
+
+bool CLCompileContext::fp16_supported() const
+{
+    return _device.supported("cl_khr_fp16");
+}
+
+std::string CLCompileContext::stringify_set(const StringSet &s, const std::string &kernel_path) const
+{
+    std::string concat_set;
+#ifndef EMBEDDED_KERNELS
+    concat_set += "-I" + kernel_path + " ";
+#else  /* EMBEDDED_KERNELS */
+    ARM_COMPUTE_UNUSED(kernel_path);
+#endif /* EMBEDDED_KERNELS */
+
+    // Concatenate set
+    for(const auto &el : s)
+    {
+        concat_set += " " + el;
+    }
+
+    return concat_set;
+}
+
+void CLCompileContext::add_built_program(const std::string &built_program_name, const cl::Program &program) const
+{
+    _built_programs_map.emplace(built_program_name, program);
+}
+
+void CLCompileContext::clear_programs_cache()
+{
+    _programs_map.clear();
+    _built_programs_map.clear();
+}
+
+const std::map<std::string, cl::Program> &CLCompileContext::get_built_programs() const
+{
+    return _built_programs_map;
+}
+
+cl::Context &CLCompileContext::context()
+{
+    return _context;
+}
+
+const cl::Device &CLCompileContext::get_device() const
+{
+    return _device.cl_device();
+}
+
+void CLCompileContext::set_device(cl::Device device)
+{
+    _device = std::move(device);
+}
+
+cl::NDRange CLCompileContext::default_ndrange() const
+{
+    GPUTarget   _target = get_target_from_device(_device.cl_device());
+    cl::NDRange default_range;
+
+    switch(_target)
+    {
+        case GPUTarget::MIDGARD:
+        case GPUTarget::T600:
+        case GPUTarget::T700:
+        case GPUTarget::T800:
+            default_range = cl::NDRange(128u, 1);
+            break;
+        default:
+            default_range = cl::NullRange;
+    }
+
+    return default_range;
+}
+
+bool CLCompileContext::int64_base_atomics_supported() const
+{
+    return _device.supported("cl_khr_int64_base_atomics");
+}
+
+size_t CLCompileContext::max_local_workgroup_size(const cl::Kernel &kernel) const
+{
+    size_t result;
+
+    size_t err = kernel.getWorkGroupInfo(_device.cl_device(), CL_KERNEL_WORK_GROUP_SIZE, &result);
+    ARM_COMPUTE_ERROR_ON_MSG(err != 0, "clGetKernelWorkGroupInfo failed to return the maximum workgroup size for the kernel");
+    ARM_COMPUTE_UNUSED(err);
+
+    return result;
+}
+
+std::string CLCompileContext::get_device_version() const
+{
+    return _device.device_version();
+}
+
+cl_uint CLCompileContext::get_num_compute_units() const
+{
+    return _device.compute_units();
+}
+} // namespace arm_compute
diff --git a/src/core/CL/CLHelpers.cpp b/src/core/CL/CLHelpers.cpp
index 84de380cc9..7d1b0ea6c7 100644
--- a/src/core/CL/CLHelpers.cpp
+++ b/src/core/CL/CLHelpers.cpp
@@ -384,6 +384,14 @@ cl::Kernel create_opencl_kernel(CLCoreRuntimeContext *ctx, const std::string &ke
     }
 }
 
+cl::Kernel create_kernel(CLCompileContext &ctx, const std::string &kernel_name, const std::set<std::string> &build_opts)
+{
+    const std::string program_name = CLKernelLibrary::get().get_program_name(kernel_name);
+    std::pair<std::string, bool> kernel_src = CLKernelLibrary::get().get_program(program_name);
+    const std::string kernel_path = CLKernelLibrary::get().get_kernel_path();
+    return static_cast<cl::Kernel>(ctx.create_kernel(kernel_name, program_name, kernel_src.first, kernel_path, build_opts, kernel_src.second));
+}
+
 cl::NDRange create_lws_hint_parallel_implementations(unsigned int input_dimension, unsigned int vector_size)
 {
     const unsigned int width_leftover = input_dimension % vector_size;
diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index c6c88569ce..7437f1bf22 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -35,113 +35,6 @@
 #include <vector>
 
 using namespace arm_compute;
-
-CLBuildOptions::CLBuildOptions()
-    : _build_opts()
-{
-}
-
-void CLBuildOptions::add_option(std::string option)
-{
-    _build_opts.emplace(std::move(option));
-}
-
-void CLBuildOptions::add_option_if(bool cond, std::string option)
-{
-    if(cond)
-    {
-        add_option(std::move(option));
-    }
-}
-
-void CLBuildOptions::add_option_if_else(bool cond, std::string option_true, std::string option_false)
-{
-    (cond) ? add_option(std::move(option_true)) : add_option(std::move(option_false));
-}
-
-void CLBuildOptions::add_options(const StringSet &options)
-{
-    _build_opts.insert(options.begin(), options.end());
-}
-
-void CLBuildOptions::add_options_if(bool cond, const StringSet &options)
-{
-    if(cond)
-    {
-        add_options(options);
-    }
-}
-
-const CLBuildOptions::StringSet &CLBuildOptions::options() const
-{
-    return _build_opts;
-}
-
-Program::Program()
-    : _context(), _device(), _is_binary(false), _name(), _source(), _binary()
-{
-}
-
-Program::Program(cl::Context context, std::string name, std::string source)
-    : _context(std::move(context)), _device(), _is_binary(false), _name(std::move(name)), _source(std::move(source)), _binary()
-{
-}
-
-Program::Program(cl::Context context, cl::Device device, std::string name, std::vector<unsigned char> binary)
-    : _context(std::move(context)), _device(std::move(device)), _is_binary(true), _name(std::move(name)), _source(), _binary(std::move(binary))
-{
-}
-
-Program::operator cl::Program() const
-{
-    if(_is_binary)
-    {
-        return cl::Program(_context, { _device }, { _binary });
-    }
-    else
-    {
-        return cl::Program(_context, _source, false);
-    }
-}
-
-bool Program::build(const cl::Program &program, const std::string &build_options)
-{
-    try
-    {
-        return program.build(build_options.c_str()) == CL_SUCCESS;
-    }
-    catch(const cl::Error &e)
-    {
-        cl_int     err        = CL_SUCCESS;
-        const auto build_info = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(&err);
-
-        for(auto &pair : build_info)
-        {
-            std::cerr << pair.second << std::endl;
-        }
-
-        return false;
-    }
-}
-
-cl::Program Program::build(const std::string &build_options) const
-{
-    cl::Program cl_program = static_cast<cl::Program>(*this);
-    build(cl_program, build_options);
-    return cl_program;
-}
-
-Kernel::Kernel()
-    : _name(), _kernel()
-{
-}
-
-Kernel::Kernel(std::string name, const cl::Program &program)
-    : _name(std::move(name)),
-      _kernel(cl::Kernel(program, _name.c_str()))
-{
-}
-
 const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map =
 {
     { "absdiff", "absdiff.cl" },
@@ -1066,7 +959,7 @@ const std::map<std::string, std::string> CLKernelLibrary::_program_source_map =
 };
 
 CLKernelLibrary::CLKernelLibrary()
-    : _context(), _device(), _kernel_path("."), _programs_map(), _built_programs_map()
+    : _compile_context(), _kernel_path()
 {
     opencl_is_available(); // Make sure the OpenCL symbols are initialised *before* the CLKernelLibrary is built
 }
@@ -1077,7 +970,15 @@ CLKernelLibrary &CLKernelLibrary::get()
     return _kernel_library;
 }
 
-Kernel CLKernelLibrary::create_kernel(const std::string &kernel_name, const StringSet &build_options_set) const
+Kernel CLKernelLibrary::create_kernel(const std::string &kernel_name, const std::set<std::string> &build_options_set) const
+{
+    const std::string program_name = get_program_name(kernel_name);
+    auto              program      = get_program(program_name);
+
+    return _compile_context.create_kernel(kernel_name, program_name, program.first, _kernel_path, build_options_set, program.second);
+}
+
+std::string CLKernelLibrary::get_program_name(const std::string &kernel_name) const
 {
     // Find which program contains the kernel
     auto kernel_program_it = _kernel_program_map.find(kernel_name);
@@ -1086,99 +987,41 @@ Kernel CLKernelLibrary::create_kernel(const std::string &kernel_name, const Stri
     {
         ARM_COMPUTE_ERROR_VAR("Kernel %s not found in the CLKernelLibrary", kernel_name.c_str());
     }
-    std::string concat_str;
-
-#if defined(ARM_COMPUTE_DEBUG_ENABLED)
-    // Enable debug properties in CL kernels
-    concat_str += " -DARM_COMPUTE_DEBUG_ENABLED";
-#endif // defined(ARM_COMPUTE_DEBUG_ENABLED)
 
-    GPUTarget gpu_arch = get_arch_from_target(get_target_from_device(_device));
-    concat_str += " -DGPU_ARCH=" + support::cpp11::to_string(
-                      static_cast<std::underlying_type<GPUTarget>::type>(gpu_arch));
-    if(fp16_supported())
-    {
-        concat_str += " -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 ";
-    }
+    const std::string program_name = kernel_program_it->second;
 
-    if(dot8_supported(_device))
-    {
-        concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ENABLED=1 ";
-    }
-
-    if(dot8_acc_supported(_device))
-    {
-        concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED=1 ";
-    }
-
-    if(get_cl_version(_device) == CLVersion::CL20)
-    {
-        concat_str += " -cl-std=CL2.0 ";
-    }
-    else if(arm_non_uniform_workgroup_supported(_device))
-    {
-        concat_str += " -cl-arm-non-uniform-work-group-size ";
-    }
-    else
-    {
-        ARM_COMPUTE_ERROR("Non uniform workgroup size is not supported!!");
-    }
-
-    // Check if the program has been built before with same build options.
-    const std::string program_name  = kernel_program_it->second;
-    const std::string build_options = stringify_set(build_options_set) + concat_str;
-
-    const std::string built_program_name = program_name + "_" + build_options;
-    auto              built_program_it   = _built_programs_map.find(built_program_name);
-
-    cl::Program cl_program;
-
-    if(_built_programs_map.end() != built_program_it)
-    {
-        // If program has been built, retrieve to create kernel from it
-        cl_program = built_program_it->second;
-    }
-    else
-    {
-        // Get program
-        Program program = load_program(program_name);
-
-        // Build program
-        cl_program = program.build(build_options);
-
-        // Add built program to internal map
-        _built_programs_map.emplace(built_program_name, cl_program);
-    }
-
-    // Create and return kernel
-    return Kernel(kernel_name, cl_program);
+    return program_name;
 }
 
 void CLKernelLibrary::init(std::string kernel_path, cl::Context context, cl::Device device)
 {
-    _kernel_path = std::move(kernel_path);
-    _context     = std::move(context);
-    _device      = std::move(device);
+    _compile_context = CLCompileContext(context, device);
+    _kernel_path     = kernel_path;
 }
 
 void CLKernelLibrary::set_kernel_path(const std::string &kernel_path)
 {
-    _kernel_path = kernel_path;
+    _kernel_path = std::move(kernel_path);
 }
 
 cl::Context &CLKernelLibrary::context()
 {
-    return _context;
+    return _compile_context.context();
 }
 
-cl::Device &CLKernelLibrary::get_device()
+const cl::Device &CLKernelLibrary::get_device()
 {
-    return _device;
+    return _compile_context.get_device();
 }
 
 void CLKernelLibrary::set_device(cl::Device device)
 {
-    _device = std::move(device);
+    _compile_context.set_device(device);
+}
+
+void CLKernelLibrary::set_context(cl::Context context)
+{
+    _compile_context.set_context(context);
 }
 
 std::string CLKernelLibrary::get_kernel_path()
@@ -1188,164 +1031,86 @@ std::string CLKernelLibrary::get_kernel_path()
 
 void CLKernelLibrary::clear_programs_cache()
 {
-    _programs_map.clear();
-    _built_programs_map.clear();
+    _compile_context.clear_programs_cache();
 }
 
 const std::map<std::string, cl::Program> &CLKernelLibrary::get_built_programs() const
 {
-    return _built_programs_map;
+    return _compile_context.get_built_programs();
 }
 
 void CLKernelLibrary::add_built_program(const std::string &built_program_name, const cl::Program &program)
 {
-    _built_programs_map.emplace(built_program_name, program);
+    _compile_context.add_built_program(built_program_name, program);
 }
 
 bool CLKernelLibrary::fp16_supported() const
 {
-    return ::fp16_supported(_device);
+    return _compile_context.fp16_supported();
 }
 
 bool CLKernelLibrary::int64_base_atomics_supported() const
 {
-    return device_supports_extension(_device, "cl_khr_int64_base_atomics");
+    return _compile_context.int64_base_atomics_supported();
 }
 
-const Program &CLKernelLibrary::load_program(const std::string &program_name) const
+std::pair<std::string, bool> CLKernelLibrary::get_program(const std::string &program_name) const
 {
-    const auto program_it = _programs_map.find(program_name);
-
-    if(program_it != _programs_map.end())
-    {
-        return program_it->second;
-    }
-
-    Program program;
-
 #ifdef EMBEDDED_KERNELS
     const auto program_source_it = _program_source_map.find(program_name);
 
-    if(_program_source_map.end() == program_source_it)
+    if(program_source_it == _program_source_map.end())
     {
         ARM_COMPUTE_ERROR_VAR("Embedded program for %s does not exist.", program_name.c_str());
     }
 
-    program = Program(_context, program_name, program_source_it->second);
+    return std::make_pair(program_source_it->second, false);
 #else  /* EMBEDDED_KERNELS */
     // Check for binary
     std::string source_name = _kernel_path + program_name;
     std::string binary_name = source_name + "bin";
+    std::string program_source{};
+    bool        is_binary = false;
 
     if(std::ifstream(binary_name).is_open())
     {
-        const std::string program_binary = read_file(binary_name, true);
-        program                          = Program(_context, _device, program_name, std::vector<unsigned char>(program_binary.begin(), program_binary.end()));
+        program_source = read_file(binary_name, true);
+        is_binary      = true;
     }
     else if(std::ifstream(source_name).is_open())
     {
-        program = Program(_context, program_name, read_file(source_name, false));
+        program_source = read_file(source_name, false);
     }
     else
     {
         ARM_COMPUTE_ERROR_VAR("Kernel file %s does not exist.", source_name.c_str());
     }
-#endif /* EMBEDDED_KERNELS */
 
-    // Insert program to program map
-    const auto new_program = _programs_map.emplace(program_name, std::move(program));
-
-    return new_program.first->second;
-}
-
-void CLKernelLibrary::set_context(cl::Context context)
-{
-    _context = std::move(context);
-    if(_context.get() == nullptr)
-    {
-        _device = cl::Device();
-    }
-    else
-    {
-        const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
-
-        if(cl_devices.empty())
-        {
-            _device = cl::Device();
-        }
-        else
-        {
-            _device = cl_devices[0];
-        }
-    }
-}
-
-std::string CLKernelLibrary::stringify_set(const StringSet &s) const
-{
-    std::string concat_set;
-
-#ifndef EMBEDDED_KERNELS
-    concat_set += "-I" + _kernel_path + " ";
+    return std::make_pair(program_source, is_binary);
 #endif /* EMBEDDED_KERNELS */
-
-    // Concatenate set
-    for(const auto &el : s)
-    {
-        concat_set += " " + el;
-    }
-
-    return concat_set;
-}
-
-std::string CLKernelLibrary::get_program_source(const std::string &program_name)
-{
-    const auto program_source_it = _program_source_map.find(program_name);
-
-    if(program_source_it == _program_source_map.end())
-    {
-        ARM_COMPUTE_ERROR_VAR("Embedded program for %s does not exist.", program_name.c_str());
-    }
-
-    return program_source_it->second;
 }
 
 size_t CLKernelLibrary::max_local_workgroup_size(const cl::Kernel &kernel) const
 {
-    size_t result;
-
-    size_t err = kernel.getWorkGroupInfo(_device, CL_KERNEL_WORK_GROUP_SIZE, &result);
-    ARM_COMPUTE_ERROR_ON_MSG(err != 0, "clGetKernelWorkGroupInfo failed to return the maximum workgroup size for the kernel");
-    ARM_COMPUTE_UNUSED(err);
-
-    return result;
+    return _compile_context.max_local_workgroup_size(kernel);
 }
 
 cl::NDRange CLKernelLibrary::default_ndrange() const
 {
-    GPUTarget   _target = get_target_from_device(_device);
-    cl::NDRange default_range;
-
-    switch(_target)
-    {
-        case GPUTarget::MIDGARD:
-        case GPUTarget::T600:
-        case GPUTarget::T700:
-        case GPUTarget::T800:
-            default_range = cl::NDRange(128u, 1);
-            break;
-        default:
-            default_range = cl::NullRange;
-    }
-
-    return default_range;
+    return _compile_context.default_ndrange();
 }
 
 std::string CLKernelLibrary::get_device_version()
 {
-    return _device.getInfo<CL_DEVICE_VERSION>();
+    return _compile_context.get_device_version();
 }
 
 cl_uint CLKernelLibrary::get_num_compute_units()
 {
-    return _device.getInfo<CL_DEVICE_MAX_COMPUTE_UNITS>();
+    return _compile_context.get_num_compute_units();
+}
+
+CLCompileContext &CLKernelLibrary::get_compile_context()
+{
+    return _compile_context;
 }
diff --git a/src/core/CL/kernels/CLFloorKernel.cpp b/src/core/CL/kernels/CLFloorKernel.cpp
index 8f0043f08a..abfed8d18e 100644
--- a/src/core/CL/kernels/CLFloorKernel.cpp
+++ b/src/core/CL/kernels/CLFloorKernel.cpp
@@ -77,7 +77,7 @@ CLFloorKernel::CLFloorKernel()
 {
 }
 
-void CLFloorKernel::configure(const ICLTensor *input, ICLTensor *output)
+void CLFloorKernel::configure(CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output)
 {
     ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
 
@@ -90,13 +90,13 @@ void CLFloorKernel::configure(const ICLTensor *input, ICLTensor *output)
     _input  = input;
     _output = output;
 
-    const unsigned int num_elems_processed_per_iteration = 16 / input->info()->element_size();
-
-    // Create kernel
+    const unsigned int    num_elems_processed_per_iteration = 16 / input->info()->element_size();
     std::set<std::string> build_opts;
     build_opts.emplace(("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())));
     build_opts.emplace(("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration)));
-    _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("floor_layer", build_opts));
+
+    // Create kernel
+    _kernel = create_kernel(compile_context, "floor_layer", build_opts);
 
     // Configure kernel window
     auto win_config = validate_and_configure_window(input->info(), output->info());
@@ -104,6 +104,11 @@ void CLFloorKernel::configure(const ICLTensor *input, ICLTensor *output)
     ICLKernel::configure_internal(win_config.second);
 }
 
+void CLFloorKernel::configure(const ICLTensor *input, ICLTensor *output)
+{
+    configure(CLKernelLibrary::get().get_compile_context(), input, output);
+}
+
 Status CLFloorKernel::validate(const ITensorInfo *input, const ITensorInfo *output)
 {
     ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output));
diff --git a/src/runtime/CL/CLRuntimeContext.cpp b/src/runtime/CL/CLRuntimeContext.cpp
index 49e4c10c84..4d70edac2f 100644
--- a/src/runtime/CL/CLRuntimeContext.cpp
+++ b/src/runtime/CL/CLRuntimeContext.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019 ARM Limited.
+ * Copyright (c) 2019-2020 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -39,13 +39,13 @@ CLRuntimeContext::CLRuntimeContext()
     cl::CommandQueue queue = cl::CommandQueue(ctx, dev);
     _gpu_owned_scheduler->init(ctx, queue, dev, &_tuner);
     const std::string cl_kernels_folder("./cl_kernels");
-    _kernel_lib.init(cl_kernels_folder, ctx, dev);
-    _core_context = CLCoreRuntimeContext(&_kernel_lib, _gpu_owned_scheduler->context(), _gpu_owned_scheduler->queue());
+    CLKernelLibrary::get().init(cl_kernels_folder, ctx, dev);
+    _core_context = CLCoreRuntimeContext(&CLKernelLibrary::get(), _gpu_owned_scheduler->context(), _gpu_owned_scheduler->queue());
 }
 
 CLKernelLibrary &CLRuntimeContext::kernel_library()
 {
-    return _kernel_lib;
+    return CLKernelLibrary::get();
 }
 
 CLCoreRuntimeContext *CLRuntimeContext::core_runtime_context()
diff --git a/src/runtime/CL/functions/CLFloor.cpp b/src/runtime/CL/functions/CLFloor.cpp
index 525810d400..204ca7400c 100644
--- a/src/runtime/CL/functions/CLFloor.cpp
+++ b/src/runtime/CL/functions/CLFloor.cpp
@@ -30,8 +30,13 @@ namespace arm_compute
 {
 void CLFloor::configure(const ICLTensor *input, ICLTensor *output)
 {
+    configure(CLKernelLibrary::get().get_compile_context(), input, output);
+}
+
+void CLFloor::configure(CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output)
+{
     auto k = arm_compute::support::cpp14::make_unique<CLFloorKernel>();
-    k->configure(input, output);
+    k->configure(compile_context, input, output);
     _kernel = std::move(k);
 }
 
diff --git a/tests/validation/CL/UNIT/CompileContext.cpp b/tests/validation/CL/UNIT/CompileContext.cpp
new file mode 100644
index 0000000000..5245044323
--- /dev/null
+++ b/tests/validation/CL/UNIT/CompileContext.cpp
@@ -0,0 +1,74 @@
+/*
+ * Copyright (c) 2020 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/CL/CLCompileContext.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/Utils.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/validation/Validation.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CL)
+TEST_SUITE(UNIT)
+TEST_SUITE(CompileContext)
+
+TEST_CASE(CompileContextCache, framework::DatasetMode::ALL)
+{
+    // Create compile context
+    CLCompileContext compile_context(CLKernelLibrary::get().context(), CLKernelLibrary::get().get_device());
+
+    // Check if the program cache is empty
+    ARM_COMPUTE_EXPECT(compile_context.get_built_programs().size() == 0, framework::LogLevel::ERRORS);
+
+    // Create a kernel using the compile context
+    const std::string kernel_name  = "floor_layer";
+    const std::string program_name = CLKernelLibrary::get().get_program_name(kernel_name);
+    std::pair<std::string, bool> kernel_src = CLKernelLibrary::get().get_program(program_name);
+    const std::string kernel_path = CLKernelLibrary::get().get_kernel_path();
+
+    std::set<std::string> build_opts;
+    build_opts.emplace("-DDATA_TYPE=float");
+    build_opts.emplace("-DVEC_SIZE=16");
+    compile_context.create_kernel(kernel_name, program_name, kernel_src.first, kernel_path, build_opts, kernel_src.second);
+
+    // Check if the program is stored in the cache
+    ARM_COMPUTE_EXPECT(compile_context.get_built_programs().size() == 1, framework::LogLevel::ERRORS);
+
+    // Try to build the same program and check if the program cache stayed the same
+    compile_context.create_kernel(kernel_name, program_name, kernel_src.first, kernel_path, build_opts, kernel_src.second);
+    ARM_COMPUTE_EXPECT(compile_context.get_built_programs().size() == 1, framework::LogLevel::ERRORS);
+}
+
+TEST_SUITE_END() // CompileContext
+TEST_SUITE_END() // UNIT
+TEST_SUITE_END() // CL
+} // namespace validation
+} // namespace test
+} // namespace arm_compute