diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp
index 540e2f80ea..888477e154 100644
--- a/modules/core/include/opencv2/core/ocl.hpp
+++ b/modules/core/include/opencv2/core/ocl.hpp
@@ -606,20 +606,33 @@ public:
 
     bool create(const ProgramSource& src,
                 const String& buildflags, String& errmsg);
-    bool read(const String& buf, const String& buildflags);
-    bool write(String& buf) const;
 
-    const ProgramSource& source() const;
     void* ptr() const;
 
-    String getPrefix() const;
-    static String getPrefix(const String& buildflags);
-
+    /**
+     * @brief Query device-specific program binary.
+     *
+     * Returns RAW OpenCL executable binary without additional attachments.
+     *
+     * @sa ProgramSource::fromBinary
+     *
+     * @param[out] binary output buffer
+     */
+    void getBinary(std::vector<char>& binary) const;
 
-    struct Impl;
+    struct Impl; friend struct Impl;
     inline Impl* getImpl() const { return (Impl*)p; }
 protected:
     Impl* p;
+public:
+#ifndef OPENCV_REMOVE_DEPRECATED_API
+    // TODO Remove this
+    CV_DEPRECATED bool read(const String& buf, const String& buildflags); // removed, use ProgramSource instead
+    CV_DEPRECATED bool write(String& buf) const; // removed, use getBinary() method instead (RAW OpenCL binary)
+    CV_DEPRECATED const ProgramSource& source() const; // implementation removed
+    CV_DEPRECATED String getPrefix() const; // deprecated, implementation replaced
+    CV_DEPRECATED static String getPrefix(const String& buildflags); // deprecated, implementation replaced
+#endif
 };
 
 
@@ -636,10 +649,59 @@ public:
     ProgramSource(const ProgramSource& prog);
     ProgramSource& operator = (const ProgramSource& prog);
 
-    const String& source() const;
+    const String& source() const; // deprecated
     hash_t hash() const; // deprecated
 
-    struct Impl;
+
+    /** @brief Describe OpenCL program binary.
+     * Do not call clCreateProgramWithBinary() and/or clBuildProgram().
+     *
+     * Caller should guarantee binary buffer lifetime greater than ProgramSource object (and any of its copies).
+     *
+     * This kind of binary is not portable between platforms in general - it is specific to OpenCL vendor / device / driver version.
+     *
+     * @param module name of program owner module
+     * @param name unique name of program (module+name is used as key for OpenCL program caching)
+     * @param binary buffer address. See buffer lifetime requirement in description.
+     * @param size buffer size
+     * @param buildOptions additional program-related build options passed to clBuildProgram()
+     * @return created ProgramSource object
+     */
+    static ProgramSource fromBinary(const String& module, const String& name,
+            const unsigned char* binary, const size_t size,
+            const cv::String& buildOptions = cv::String());
+
+    /** @brief Describe OpenCL program in SPIR format.
+     * Do not call clCreateProgramWithBinary() and/or clBuildProgram().
+     *
+     * Supports SPIR 1.2 by default (pass '-spir-std=X.Y' in buildOptions to override this behavior)
+     *
+     * Caller should guarantee binary buffer lifetime greater than ProgramSource object (and any of its copies).
+     *
+     * Programs in this format are portable between OpenCL implementations with 'khr_spir' extension:
+     * https://www.khronos.org/registry/OpenCL/sdk/2.0/docs/man/xhtml/cl_khr_spir.html
+     * (but they are not portable between different platforms: 32-bit / 64-bit)
+     *
+     * Note: these programs can't support vendor specific extensions, like 'cl_intel_subgroups'.
+     *
+     * @param module name of program owner module
+     * @param name unique name of program (module+name is used as key for OpenCL program caching)
+     * @param binary buffer address. See buffer lifetime requirement in description.
+     * @param size buffer size
+     * @param buildOptions additional program-related build options passed to clBuildProgram()
+     *        (these options are added automatically: '-x spir' and '-spir-std=1.2')
+     * @return created ProgramSource object.
+     */
+    static ProgramSource fromSPIR(const String& module, const String& name,
+            const unsigned char* binary, const size_t size,
+            const cv::String& buildOptions = cv::String());
+
+    //OpenCL 2.1+ only
+    //static Program fromSPIRV(const String& module, const String& name,
+    //        const unsigned char* binary, const size_t size,
+    //        const cv::String& buildOptions = cv::String());
+
+    struct Impl; friend struct Impl;
     inline Impl* getImpl() const { return (Impl*)p; }
 protected:
     Impl* p;
diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
index eac630e3be..0d670aa253 100644
--- a/modules/core/src/ocl.cpp
+++ b/modules/core/src/ocl.cpp
@@ -102,6 +102,17 @@
 #ifdef HAVE_OPENCL
 #include "opencv2/core/opencl/runtime/opencl_core.hpp"
 #else
+#if defined(_MSC_VER)
+    #pragma warning(push)
+    #pragma warning(disable : 4100)
+    #pragma warning(disable : 4702)
+#elif defined(__clang__)
+    #pragma clang diagnostic push
+    #pragma clang diagnostic ignored "-Wunused-parameter"
+#elif defined(__GNUC__)
+    #pragma GCC diagnostic push
+    #pragma GCC diagnostic ignored "-Wunused-parameter"
+#endif
 // TODO FIXIT: This file can't be build without OPENCL
 #include "ocl_deprecated.hpp"
 #endif // HAVE_OPENCL
@@ -114,6 +125,34 @@
 
 namespace cv { namespace ocl {
 
+#define IMPLEMENT_REFCOUNTABLE() \
+    void addref() { CV_XADD(&refcount, 1); } \
+    void release() { if( CV_XADD(&refcount, -1) == 1 && !cv::__termination) delete this; } \
+    int refcount
+
+#ifndef HAVE_OPENCL
+#define CV_OPENCL_NO_SUPPORT() CV_ErrorNoReturn(cv::Error::OpenCLApiCallError, "OpenCV build without OpenCL support")
+namespace {
+struct DummyImpl
+{
+    DummyImpl() { CV_OPENCL_NO_SUPPORT(); }
+    ~DummyImpl() { /* do not throw in desctructors */ }
+    IMPLEMENT_REFCOUNTABLE();
+};
+} // namespace
+
+// TODO Replace to empty body (without HAVE_OPENCL)
+#define CV_OCL_TRACE_CHECK_RESULT(status, message) /* nothing */
+#define CV_OCL_API_ERROR_MSG(check_result, msg) cv::String()
+#define CV_OCL_CHECK_RESULT(check_result, msg) (void)check_result
+#define CV_OCL_CHECK_(expr, check_result) expr; (void)check_result
+#define CV_OCL_CHECK(expr) do { cl_int __cl_result = (expr); CV_OCL_CHECK_RESULT(__cl_result, #expr); } while (0)
+#define CV_OCL_DBG_CHECK_RESULT(check_result, msg) (void)check_result
+#define CV_OCL_DBG_CHECK_(expr, check_result) expr; (void)check_result
+#define CV_OCL_DBG_CHECK(expr) do { cl_int __cl_result = (expr); CV_OCL_CHECK_RESULT(__cl_result, #expr); } while (0)
+
+#else // HAVE_OPENCL
+
 #ifndef _DEBUG
 static bool isRaiseError()
 {
@@ -186,6 +225,7 @@ static const bool CV_OPENCL_CACHE_CLEANUP = utils::getConfigurationParameterBool
 static const bool CV_OPENCL_VALIDATE_BINARY_PROGRAMS_VALUE = utils::getConfigurationParameterBool("OPENCV_OPENCL_VALIDATE_BINARY_PROGRAMS", false);
 #endif
 
+#endif // HAVE_OPENCL
 
 struct UMat2D
 {
@@ -246,7 +286,7 @@ static uint64 crc64( const uchar* data, size_t size, uint64 crc0=0 )
     return ~crc;
 }
 
-#if OPENCV_HAVE_FILESYSTEM_SUPPORT
+#if defined HAVE_OPENCL && OPENCV_HAVE_FILESYSTEM_SUPPORT
 struct OpenCLBinaryCacheConfigurator
 {
     cv::String cache_path_;
@@ -1032,11 +1072,6 @@ void finish()
     Queue::getDefault().finish();
 }
 
-#define IMPLEMENT_REFCOUNTABLE() \
-    void addref() { CV_XADD(&refcount, 1); } \
-    void release() { if( CV_XADD(&refcount, -1) == 1 && !cv::__termination) delete this; } \
-    int refcount
-
 /////////////////////////////////////////// Platform /////////////////////////////////////////////
 
 struct Platform::Impl
@@ -1194,6 +1229,17 @@ struct Device::Impl
             vendorID_ = VENDOR_NVIDIA;
         else
             vendorID_ = UNKNOWN_VENDOR;
+
+#if 0
+        if (isExtensionSupported("cl_khr_spir"))
+        {
+#ifndef CL_DEVICE_SPIR_VERSIONS
+#define CL_DEVICE_SPIR_VERSIONS                     0x40E0
+#endif
+            cv::String spir_versions = getStrProp(CL_DEVICE_SPIR_VERSIONS);
+            std::cout << spir_versions << std::endl;
+        }
+#endif
     }
 
     template<typename _TpCL, typename _TpOut>
@@ -1217,7 +1263,7 @@ struct Device::Impl
 
     String getStrProp(cl_device_info prop) const
     {
-        char buf[1024];
+        char buf[4096];
         size_t sz=0;
         return clGetDeviceInfo(handle, prop, sizeof(buf)-16, buf, &sz) == CL_SUCCESS &&
             sz < sizeof(buf) ? String(buf) : String();
@@ -1859,6 +1905,7 @@ static unsigned int getSVMCapabilitiesMask()
 } // namespace
 #endif
 
+#ifdef HAVE_OPENCL
 static size_t getProgramCountLimit()
 {
     static bool initialized = false;
@@ -1870,6 +1917,7 @@ static size_t getProgramCountLimit()
     }
     return count;
 }
+#endif
 
 struct Context::Impl
 {
@@ -1989,56 +2037,7 @@ struct Context::Impl
         devices.clear();
     }
 
-    Program getProg(const ProgramSource& src,
-                    const String& buildflags, String& errmsg)
-    {
-        size_t limit = getProgramCountLimit();
-        String key = cv::format("codehash=%08llx ", src.hash()) + Program::getPrefix(buildflags);
-        {
-            cv::AutoLock lock(program_cache_mutex);
-            phash_t::iterator it = phash.find(key);
-            if (it != phash.end())
-            {
-                // TODO LRU cache
-                CacheList::iterator i = std::find(cacheList.begin(), cacheList.end(), key);
-                if (i != cacheList.end() && i != cacheList.begin())
-                {
-                    cacheList.erase(i);
-                    cacheList.push_front(key);
-                }
-                return it->second;
-            }
-            { // cleanup program cache
-                size_t sz = phash.size();
-                if (limit > 0 && sz >= limit)
-                {
-                    static bool warningFlag = false;
-                    if (!warningFlag)
-                    {
-                        printf("\nWARNING: OpenCV-OpenCL:\n"
-                            "    In-memory cache for OpenCL programs is full, older programs will be unloaded.\n"
-                            "    You can change cache size via OPENCV_OPENCL_PROGRAM_CACHE environment variable\n\n");
-                        warningFlag = true;
-                    }
-                    while (!cacheList.empty())
-                    {
-                        size_t c = phash.erase(cacheList.back());
-                        cacheList.pop_back();
-                        if (c != 0)
-                            break;
-                    }
-                }
-            }
-        }
-        Program prog(src, buildflags, errmsg);
-        // Cache result of build failures too (to prevent unnecessary compiler invocations)
-        {
-            cv::AutoLock lock(program_cache_mutex);
-            phash.insert(std::pair<std::string, Program>(key, prog));
-            cacheList.push_front(key);
-        }
-        return prog;
-    }
+    Program getProg(const ProgramSource& src, const String& buildflags, String& errmsg);
 
     void unloadProg(Program& prog)
     {
@@ -2253,6 +2252,8 @@ struct Context::Impl
         return;
     }
 #endif
+
+    friend class Program;
 };
 
 
@@ -2887,7 +2888,7 @@ bool Kernel::create(const char* kname, const ProgramSource& src,
     }
     String tempmsg;
     if( !errmsg ) errmsg = &tempmsg;
-    const Program& prog = Context::getDefault().getProg(src, buildopts, *errmsg);
+    const Program prog = Context::getDefault().getProg(src, buildopts, *errmsg);
     return create(kname, prog);
 }
 
@@ -3207,46 +3208,147 @@ size_t Kernel::localMemSize() const
 
 struct ProgramSource::Impl
 {
+    IMPLEMENT_REFCOUNTABLE();
+
+    enum KIND {
+        PROGRAM_SOURCE_CODE = 0,
+        PROGRAM_BINARIES,
+        PROGRAM_SPIR,
+        PROGRAM_SPIRV
+    } kind_;
+
     Impl(const String& src)
     {
-        init(cv::String(), cv::String(), src, cv::String());
+        init(PROGRAM_SOURCE_CODE, cv::String(), cv::String());
+        initFromSource(src, cv::String());
     }
     Impl(const String& module, const String& name, const String& codeStr, const String& codeHash)
     {
-        init(module, name, codeStr, codeHash);
+        init(PROGRAM_SOURCE_CODE, module, name);
+        initFromSource(codeStr, codeHash);
     }
-    void init(const String& module, const String& name, const String& codeStr, const String& codeHash)
+
+    /// reset fields
+    void init(enum KIND kind, const String& module, const String& name)
     {
         refcount = 1;
+        kind_ = kind;
         module_ = module;
         name_ = name;
-        codeStr_ = codeStr;
-        codeHash_ = codeHash;
 
+        sourceAddr_ = NULL;
+        sourceSize_ = 0;
         isHashUpdated = false;
-        if (codeHash_.empty())
+    }
+
+    void initFromSource(const String& codeStr, const String& codeHash)
+    {
+        codeStr_ = codeStr;
+        sourceHash_ = codeHash;
+        if (sourceHash_.empty())
         {
             updateHash();
-            codeHash_ = cv::format("%08llx", hash_);
+        }
+        else
+        {
+            isHashUpdated = true;
         }
     }
 
-    void updateHash()
+    void updateHash(const char* hashStr = NULL)
     {
-        hash_ = crc64((uchar*)codeStr_.c_str(), codeStr_.size());
+        if (hashStr)
+        {
+            sourceHash_ = cv::String(hashStr);
+            isHashUpdated = true;
+            return;
+        }
+        uint64 hash = 0;
+        switch (kind_)
+        {
+        case PROGRAM_SOURCE_CODE:
+            if (sourceAddr_)
+            {
+                CV_Assert(codeStr_.empty());
+                hash = crc64(sourceAddr_, sourceSize_); // static storage
+            }
+            else
+            {
+                CV_Assert(!codeStr_.empty());
+                hash = crc64((uchar*)codeStr_.c_str(), codeStr_.size());
+            }
+            break;
+        case PROGRAM_BINARIES:
+        case PROGRAM_SPIR:
+        case PROGRAM_SPIRV:
+            hash = crc64(sourceAddr_, sourceSize_);
+            break;
+        default:
+            CV_ErrorNoReturn(Error::StsInternal, "Internal error");
+        }
+        sourceHash_ = cv::format("%08llx", hash);
         isHashUpdated = true;
     }
 
-    IMPLEMENT_REFCOUNTABLE();
+    Impl(enum KIND kind,
+            const String& module, const String& name,
+            const unsigned char* binary, const size_t size,
+            const cv::String& buildOptions = cv::String())
+    {
+        init(kind, module, name);
+
+        sourceAddr_ = binary;
+        sourceSize_ = size;
+
+        buildOptions_ = buildOptions;
+    }
+
+    static ProgramSource fromSourceWithStaticLifetime(const String& module, const String& name,
+            const char* sourceCodeStaticStr, const char* hashStaticStr,
+            const cv::String& buildOptions)
+    {
+        ProgramSource result;
+        result.p = new Impl(PROGRAM_SOURCE_CODE, module, name,
+                (const unsigned char*)sourceCodeStaticStr, strlen(sourceCodeStaticStr), buildOptions);
+        result.p->updateHash(hashStaticStr);
+        return result;
+    }
+
+    static ProgramSource fromBinary(const String& module, const String& name,
+            const unsigned char* binary, const size_t size,
+            const cv::String& buildOptions)
+    {
+        ProgramSource result;
+        result.p = new Impl(PROGRAM_BINARIES, module, name, binary, size, buildOptions);
+        return result;
+    }
+
+    static ProgramSource fromSPIR(const String& module, const String& name,
+            const unsigned char* binary, const size_t size,
+            const cv::String& buildOptions)
+    {
+        ProgramSource result;
+        result.p = new Impl(PROGRAM_SPIR, module, name, binary, size, buildOptions);
+        return result;
+    }
 
     String module_;
     String name_;
-    String codeStr_;
-    String codeHash_;
+
     // TODO std::vector<ProgramSource> includes_;
+    String codeStr_; // PROGRAM_SOURCE_CODE only
+
+    const unsigned char* sourceAddr_;
+    size_t sourceSize_;
 
+    cv::String buildOptions_;
+
+    String sourceHash_;
     bool isHashUpdated;
-    ProgramSource::hash_t hash_;
+
+    friend struct Program::Impl;
+    friend struct internal::ProgramEntry;
+    friend struct Context::Impl;
 };
 
 
@@ -3297,15 +3399,32 @@ ProgramSource& ProgramSource::operator = (const ProgramSource& prog)
 const String& ProgramSource::source() const
 {
     CV_Assert(p);
+    CV_Assert(p->kind_ == Impl::PROGRAM_SOURCE_CODE);
+    CV_Assert(p->sourceAddr_ == NULL); // method returns reference - can't construct temporary object
     return p->codeStr_;
 }
 
 ProgramSource::hash_t ProgramSource::hash() const
 {
-    CV_Assert(p);
-    if (!p->isHashUpdated)
-        p->updateHash();
-    return p->hash_;
+    CV_ErrorNoReturn(Error::StsNotImplemented, "Removed method: ProgramSource::hash()");
+}
+
+ProgramSource ProgramSource::fromBinary(const String& module, const String& name,
+        const unsigned char* binary, const size_t size,
+        const cv::String& buildOptions)
+{
+    CV_Assert(binary);
+    CV_Assert(size > 0);
+    return Impl::fromBinary(module, name, binary, size, buildOptions);
+}
+
+ProgramSource ProgramSource::fromSPIR(const String& module, const String& name,
+        const unsigned char* binary, const size_t size,
+        const cv::String& buildOptions)
+{
+    CV_Assert(binary);
+    CV_Assert(size > 0);
+    return Impl::fromBinary(module, name, binary, size, buildOptions);
 }
 
 
@@ -3316,8 +3435,9 @@ internal::ProgramEntry::operator ProgramSource&() const
         cv::AutoLock lock(cv::getInitializationMutex());
         if (this->pProgramSource == NULL)
         {
-            ProgramSource* ps = new ProgramSource(this->module, this->name, this->programCode, this->programHash);
-            const_cast<ProgramEntry*>(this)->pProgramSource = ps;
+            ProgramSource ps = ProgramSource::Impl::fromSourceWithStaticLifetime(this->module, this->name, this->programCode, this->programHash, cv::String());
+            ProgramSource* ptr = new ProgramSource(ps);
+            const_cast<ProgramEntry*>(this)->pProgramSource = ptr;
         }
     }
     return *this->pProgramSource;
@@ -3327,39 +3447,84 @@ internal::ProgramEntry::operator ProgramSource&() const
 
 /////////////////////////////////////////// Program /////////////////////////////////////////////
 
+#ifdef HAVE_OPENCL
+
+static
+cv::String joinBuildOptions(const cv::String& a, const cv::String& b)
+{
+    if (b.empty())
+        return a;
+    if (a.empty())
+        return b;
+    if (b[0] == ' ')
+        return a + b;
+    return a + (cv::String(" ") + b);
+}
+
 struct Program::Impl
 {
-    Impl(const ProgramSource& _src,
+    IMPLEMENT_REFCOUNTABLE();
+
+    Impl(const ProgramSource& src,
          const String& _buildflags, String& errmsg) :
-         src(_src),
-         buildflags(_buildflags),
-         handle(NULL)
-    {
-        refcount = 1;
+         refcount(1),
+         handle(NULL),
+         buildflags(_buildflags)
+    {
+        const ProgramSource::Impl* src_ = src.getImpl();
+        CV_Assert(src_);
+        sourceModule_ = src_->module_;
+        sourceName_ = src_->name_;
         const Context ctx = Context::getDefault();
         Device device = ctx.device(0);
         if (ctx.ptr() == NULL || device.ptr() == NULL)
             return;
-        if (device.isAMD())
-            buildflags += " -D AMD_DEVICE";
-        else if (device.isIntel())
-            buildflags += " -D INTEL_DEVICE";
-        compile(ctx, errmsg);
+        buildflags = joinBuildOptions(buildflags, src_->buildOptions_);
+        if (src.getImpl()->kind_ == ProgramSource::Impl::PROGRAM_SOURCE_CODE)
+        {
+            if (device.isAMD())
+                buildflags = joinBuildOptions(buildflags, " -D AMD_DEVICE");
+            else if (device.isIntel())
+                buildflags = joinBuildOptions(buildflags, " -D INTEL_DEVICE");
+        }
+        compile(ctx, src_, errmsg);
     }
 
-    bool compile(const Context& ctx, String& errmsg)
+    bool compile(const Context& ctx, const ProgramSource::Impl* src_, String& errmsg)
     {
-#if OPENCV_HAVE_FILESYSTEM_SUPPORT
         CV_Assert(ctx.getImpl());
+        CV_Assert(src_);
+
+        // We don't cache OpenCL binaries
+        if (src_->kind_ == ProgramSource::Impl::PROGRAM_BINARIES)
+        {
+            bool isLoaded = createFromBinary(ctx, src_->sourceAddr_, src_->sourceSize_, errmsg);
+            return isLoaded;
+        }
+        return compileWithCache(ctx, src_, errmsg);
+    }
+
+    bool compileWithCache(const Context& ctx, const ProgramSource::Impl* src_, String& errmsg)
+    {
+        CV_Assert(ctx.getImpl());
+        CV_Assert(src_);
+        CV_Assert(src_->kind_ != ProgramSource::Impl::PROGRAM_BINARIES);
+
+#if OPENCV_HAVE_FILESYSTEM_SUPPORT
         OpenCLBinaryCacheConfigurator& config = OpenCLBinaryCacheConfigurator::getSingletonInstance();
         const std::string base_dir = config.prepareCacheDirectoryForContext(
                 ctx.getImpl()->getPrefixString(),
                 ctx.getImpl()->getPrefixBase()
         );
-        const std::string fname = base_dir.empty() ? std::string() :
-                std::string(base_dir + src.getImpl()->module_.c_str() + "--" + src.getImpl()->name_ + "_" + src.getImpl()->codeHash_ + ".bin");
+        const String& hash_str = src_->sourceHash_;
+        cv::String fname;
+        if (!base_dir.empty() && !src_->module_.empty() && !src_->name_.empty())
+        {
+            CV_Assert(!hash_str.empty());
+            fname = src_->module_ + "--" + src_->name_ + "_" + hash_str + ".bin";
+            fname = utils::fs::join(base_dir, fname);
+        }
         const cv::Ptr<utils::fs::FileLock> fileLock = config.cache_lock_; // can be empty
-        const String& hash_str = src.getImpl()->codeHash_;
         if (!fname.empty() && CV_OPENCL_CACHE_ENABLE)
         {
             try
@@ -3391,9 +3556,31 @@ struct Program::Impl
         }
 #endif // OPENCV_HAVE_FILESYSTEM_SUPPORT
         CV_Assert(handle == NULL);
-        if (!buildFromSources(ctx, errmsg))
+        if (src_->kind_ == ProgramSource::Impl::PROGRAM_SOURCE_CODE)
         {
-            return true;
+            if (!buildFromSources(ctx, src_, errmsg))
+            {
+                return false;
+            }
+        }
+        else if (src_->kind_ == ProgramSource::Impl::PROGRAM_SPIR)
+        {
+            buildflags = joinBuildOptions(buildflags, " -x spir");
+            if ((cv::String(" ") + buildflags).find(" -spir-std=") == cv::String::npos)
+            {
+                buildflags = joinBuildOptions(buildflags, " -spir-std=1.2");
+            }
+            bool isLoaded = createFromBinary(ctx, src_->sourceAddr_, src_->sourceSize_, errmsg);
+            if (!isLoaded)
+                return false;
+        }
+        else if (src_->kind_ == ProgramSource::Impl::PROGRAM_SPIRV)
+        {
+            CV_ErrorNoReturn(Error::StsNotImplemented, "OpenCL: SPIR-V is not supported");
+        }
+        else
+        {
+            CV_ErrorNoReturn(Error::StsInternal, "Internal error");
         }
         CV_Assert(handle != NULL);
 #if OPENCV_HAVE_FILESYSTEM_SUPPORT
@@ -3462,24 +3649,28 @@ struct Program::Impl
 
         errmsg = String(buffer);
         printf("OpenCL program build log: %s/%s\nStatus %d: %s\n%s\n%s\n",
-                src.getImpl()->module_.c_str(), src.getImpl()->name_.c_str(),
+                sourceModule_.c_str(), sourceName_.c_str(),
                 result, getOpenCLErrorString(result),
                 buildflags.c_str(), errmsg.c_str());
         fflush(stdout);
     }
 
-    bool buildFromSources(const Context& ctx, String& errmsg)
+    bool buildFromSources(const Context& ctx, const ProgramSource::Impl* src_, String& errmsg)
     {
+        CV_Assert(src_);
+        CV_Assert(src_->kind_ == ProgramSource::Impl::PROGRAM_SOURCE_CODE);
         CV_Assert(handle == NULL);
         CV_INSTRUMENT_REGION_OPENCL_COMPILE(cv::format("Build OpenCL program: %s/%s %" PRIx64 " options: %s",
-                src.getImpl()->module_.c_str(), src.getImpl()->name_.c_str(),
+                sourceModule_.c_str(), sourceName_.c_str(),
                 src.hash(), buildflags.c_str()).c_str());
 
-        CV_LOG_VERBOSE(NULL, 0, "Compile... " << src.getImpl()->module_.c_str() << "/" << src.getImpl()->name_.c_str());
+        CV_LOG_VERBOSE(NULL, 0, "Compile... " << sourceModule_.c_str() << "/" << sourceName_.c_str());
+
+        const char* srcptr = src_->sourceAddr_ ? ((const char*)src_->sourceAddr_) : src_->codeStr_.c_str();
+        size_t srclen = src_->sourceAddr_ ? src_->sourceSize_ : src_->codeStr_.size();
+        CV_Assert(srcptr != NULL);
+        CV_Assert(srclen > 0);
 
-        const String& srcstr = src.source();
-        const char* srcptr = srcstr.c_str();
-        size_t srclen = srcstr.size();
         cl_int retval = 0;
 
         handle = clCreateProgramWithSource((cl_context)ctx.ptr(), 1, &srcptr, &srclen, &retval);
@@ -3496,6 +3687,7 @@ struct Program::Impl
             }
 
             retval = clBuildProgram(handle, (cl_uint)n, deviceList, buildflags.c_str(), 0, 0);
+            CV_OCL_TRACE_CHECK_RESULT(/*don't throw: retval*/CL_SUCCESS, cv::format("clBuildProgram(source: %s)", buildflags.c_str()).c_str());
 #if !CV_OPENCL_ALWAYS_SHOW_BUILD_LOG
             if (retval != CL_SUCCESS)
 #endif
@@ -3510,63 +3702,25 @@ struct Program::Impl
                     handle = NULL;
                 }
             }
+#if CV_OPENCL_VALIDATE_BINARY_PROGRAMS
+            if (handle && CV_OPENCL_VALIDATE_BINARY_PROGRAMS_VALUE)
+            {
+                CV_LOG_INFO(NULL, "OpenCL: query kernel names (build from sources)...");
+                size_t retsz = 0;
+                char kernels_buffer[4096] = {0};
+                cl_int result = clGetProgramInfo(handle, CL_PROGRAM_KERNEL_NAMES, sizeof(kernels_buffer), &kernels_buffer[0], &retsz);
+                if (retsz < sizeof(kernels_buffer))
+                    kernels_buffer[retsz] = 0;
+                else
+                    kernels_buffer[0] = 0;
+                CV_LOG_INFO(NULL, result << ": Kernels='" << kernels_buffer << "'");
+            }
+#endif
 
         }
         return handle != NULL;
     }
 
-    Impl(const String& _buf, const String& _buildflags)
-    {
-        refcount = 1;
-        handle = 0;
-        buildflags = _buildflags;
-        if(_buf.empty())
-            return;
-        String prefix0 = Program::getPrefix(buildflags);
-        const Context& ctx = Context::getDefault();
-        const Device& dev = Device::getDefault();
-        const char* pos0 = _buf.c_str();
-        const char* pos1 = strchr(pos0, '\n');
-        if(!pos1)
-            return;
-        const char* pos2 = strchr(pos1+1, '\n');
-        if(!pos2)
-            return;
-        const char* pos3 = strchr(pos2+1, '\n');
-        if(!pos3)
-            return;
-        size_t prefixlen = (pos3 - pos0)+1;
-        String prefix(pos0, prefixlen);
-        if( prefix != prefix0 )
-            return;
-        const uchar* bin = (uchar*)(pos3+1);
-        void* devid = dev.ptr();
-        size_t codelen = _buf.length() - prefixlen;
-        cl_int binstatus = 0, retval = 0;
-        handle = clCreateProgramWithBinary((cl_context)ctx.ptr(), 1, (cl_device_id*)&devid,
-                                           &codelen, &bin, &binstatus, &retval);
-        CV_OCL_DBG_CHECK_RESULT(retval, "clCreateProgramWithBinary");
-    }
-
-    String store()
-    {
-        if(!handle)
-            return String();
-        size_t progsz = 0, retsz = 0;
-        String prefix = Program::getPrefix(buildflags);
-        size_t prefixlen = prefix.length();
-        if(clGetProgramInfo(handle, CL_PROGRAM_BINARY_SIZES, sizeof(progsz), &progsz, &retsz) != CL_SUCCESS)
-            return String();
-        AutoBuffer<uchar> bufbuf(prefixlen + progsz + 16);
-        uchar* buf = bufbuf;
-        memcpy(buf, prefix.c_str(), prefixlen);
-        buf += prefixlen;
-        if(clGetProgramInfo(handle, CL_PROGRAM_BINARIES, sizeof(buf), &buf, &retsz) != CL_SUCCESS)
-            return String();
-        buf[progsz] = (uchar)'\0';
-        return String((const char*)(uchar*)bufbuf, prefixlen + progsz);
-    }
-
     void getProgramBinary(std::vector<char>& buf)
     {
         CV_Assert(handle);
@@ -3575,30 +3729,19 @@ struct Program::Impl
         buf.resize(sz);
         uchar* ptr = (uchar*)&buf[0];
         CV_OCL_CHECK(clGetProgramInfo(handle, CL_PROGRAM_BINARIES, sizeof(ptr), &ptr, NULL));
-#if CV_OPENCL_VALIDATE_BINARY_PROGRAMS
-        if (CV_OPENCL_VALIDATE_BINARY_PROGRAMS_VALUE)
-        {
-            CV_LOG_INFO(NULL, "OpenCL: query kernel names (compiled)...");
-            size_t retsz = 0;
-            char kernels_buffer[4096] = {0};
-            cl_int result = clGetProgramInfo(handle, CL_PROGRAM_KERNEL_NAMES, sizeof(kernels_buffer), &kernels_buffer[0], &retsz);
-            if (retsz < sizeof(kernels_buffer))
-                kernels_buffer[retsz] = 0;
-            else
-                kernels_buffer[0] = 0;
-            CV_LOG_INFO(NULL, result << ": Kernels='" << kernels_buffer << "'");
-        }
-#endif
     }
 
     bool createFromBinary(const Context& ctx, const std::vector<char>& buf, String& errmsg)
+    {
+        return createFromBinary(ctx, (const unsigned char*)&buf[0], buf.size(), errmsg);
+    }
+
+    bool createFromBinary(const Context& ctx, const unsigned char* binaryAddr, const size_t binarySize, String& errmsg)
     {
         CV_Assert(handle == NULL);
         CV_INSTRUMENT_REGION_OPENCL_COMPILE("Load OpenCL program");
         CV_LOG_VERBOSE(NULL, 0, "Load from binary... " << src.getImpl()->module_.c_str() << "/" << src.getImpl()->name_.c_str());
 
-        const uchar* binaryPtr = (uchar*)&buf[0];
-        size_t binarySize = buf.size();
         CV_Assert(binarySize > 0);
 
         size_t ndevices = (int)ctx.ndevices();
@@ -3612,7 +3755,7 @@ struct Program::Impl
         for (size_t i = 0; i < ndevices; i++)
         {
             devices[i] = (cl_device_id)ctx.device(i).ptr();
-            binaryPtrs[i] = binaryPtr;
+            binaryPtrs[i] = binaryAddr;
             binarySizes[i] = binarySize;
         }
 
@@ -3641,7 +3784,7 @@ struct Program::Impl
         else
         {
             result = clBuildProgram(handle, (cl_uint)ndevices, (cl_device_id*)devices_, buildflags.c_str(), 0, 0);
-            CV_OCL_DBG_CHECK_RESULT(result, cv::format("clBuildProgram(binary: %s/%s)", src.getImpl()->module_.c_str(), src.getImpl()->name_.c_str()).c_str());
+            CV_OCL_DBG_CHECK_RESULT(result, cv::format("clBuildProgram(binary: %s/%s)", sourceModule_.c_str(), sourceName_.c_str()).c_str());
             if (result != CL_SUCCESS)
             {
                 dumpBuildLog_(result, devices, errmsg);
@@ -3710,13 +3853,17 @@ struct Program::Impl
         }
     }
 
-    IMPLEMENT_REFCOUNTABLE();
+    cl_program handle;
 
-    ProgramSource src;
     String buildflags;
-    cl_program handle;
+    String sourceModule_;
+    String sourceName_;
 };
 
+#else // HAVE_OPENCL
+struct Program::Impl : public DummyImpl {};
+#endif // HAVE_OPENCL
+
 
 Program::Program() { p = 0; }
 
@@ -3755,7 +3902,11 @@ bool Program::create(const ProgramSource& src,
             const String& buildflags, String& errmsg)
 {
     if(p)
+    {
         p->release();
+        p = NULL;
+    }
+#ifdef HAVE_OPENCL
     p = new Impl(src, buildflags, errmsg);
     if(!p->handle)
     {
@@ -3763,50 +3914,133 @@ bool Program::create(const ProgramSource& src,
         p = 0;
     }
     return p != 0;
+#else
+    CV_OPENCL_NO_SUPPORT();
+#endif
 }
 
-const ProgramSource& Program::source() const
+void* Program::ptr() const
 {
-    static ProgramSource dummy;
-    return p ? p->src : dummy;
+#ifdef HAVE_OPENCL
+    return p ? p->handle : 0;
+#else
+    CV_OPENCL_NO_SUPPORT();
+#endif
 }
 
-void* Program::ptr() const
+#ifndef OPENCV_REMOVE_DEPRECATED_API
+const ProgramSource& Program::source() const
 {
-    return p ? p->handle : 0;
+    CV_ErrorNoReturn(Error::StsNotImplemented, "Removed API");
 }
 
 bool Program::read(const String& bin, const String& buildflags)
 {
-    if(p)
-        p->release();
-    p = new Impl(bin, buildflags);
-    return p->handle != 0;
+    CV_UNUSED(bin); CV_UNUSED(buildflags);
+    CV_ErrorNoReturn(Error::StsNotImplemented, "Removed API");
 }
 
 bool Program::write(String& bin) const
 {
-    if(!p)
-        return false;
-    bin = p->store();
-    return !bin.empty();
+    CV_UNUSED(bin);
+    CV_ErrorNoReturn(Error::StsNotImplemented, "Removed API");
 }
 
 String Program::getPrefix() const
 {
+#ifdef HAVE_OPENCL
     if(!p)
         return String();
-    return getPrefix(p->buildflags);
+    Context::Impl* ctx_ = Context::getDefault().getImpl();
+    CV_Assert(ctx_);
+    return cv::format("opencl=%s\nbuildflags=%s", ctx_->getPrefixString().c_str(), p->buildflags.c_str());
+#else
+    CV_OPENCL_NO_SUPPORT();
+#endif
 }
 
 String Program::getPrefix(const String& buildflags)
 {
-    const Context& ctx = Context::getDefault();
-    const Device& dev = ctx.device(0);
-    return format("name=%s\ndriver=%s\nbuildflags=%s\n",
-                  dev.name().c_str(), dev.driverVersion().c_str(), buildflags.c_str());
+#ifdef HAVE_OPENCL
+        Context::Impl* ctx_ = Context::getDefault().getImpl();
+        CV_Assert(ctx_);
+        return cv::format("opencl=%s\nbuildflags=%s", ctx_->getPrefixString().c_str(), buildflags.c_str());
+#else
+    CV_OPENCL_NO_SUPPORT();
+#endif
 }
+#endif
 
+void Program::getBinary(std::vector<char>& binary) const
+{
+#ifdef HAVE_OPENCL
+    CV_Assert(p && "Empty program");
+    p->getProgramBinary(binary);
+#else
+    binary.clear();
+    CV_OPENCL_NO_SUPPORT();
+#endif
+}
+
+Program Context::Impl::getProg(const ProgramSource& src,
+                               const String& buildflags, String& errmsg)
+{
+#ifdef HAVE_OPENCL
+    size_t limit = getProgramCountLimit();
+    const ProgramSource::Impl* src_ = src.getImpl();
+    CV_Assert(src_);
+    String key = cv::format("module=%s name=%s codehash=%s\nopencl=%s\nbuildflags=%s",
+            src_->module_.c_str(), src_->name_.c_str(), src_->sourceHash_.c_str(),
+            getPrefixString().c_str(),
+            buildflags.c_str());
+    {
+        cv::AutoLock lock(program_cache_mutex);
+        phash_t::iterator it = phash.find(key);
+        if (it != phash.end())
+        {
+            // TODO LRU cache
+            CacheList::iterator i = std::find(cacheList.begin(), cacheList.end(), key);
+            if (i != cacheList.end() && i != cacheList.begin())
+            {
+                cacheList.erase(i);
+                cacheList.push_front(key);
+            }
+            return it->second;
+        }
+        { // cleanup program cache
+            size_t sz = phash.size();
+            if (limit > 0 && sz >= limit)
+            {
+                static bool warningFlag = false;
+                if (!warningFlag)
+                {
+                    printf("\nWARNING: OpenCV-OpenCL:\n"
+                        "    In-memory cache for OpenCL programs is full, older programs will be unloaded.\n"
+                        "    You can change cache size via OPENCV_OPENCL_PROGRAM_CACHE environment variable\n\n");
+                    warningFlag = true;
+                }
+                while (!cacheList.empty())
+                {
+                    size_t c = phash.erase(cacheList.back());
+                    cacheList.pop_back();
+                    if (c != 0)
+                        break;
+                }
+            }
+        }
+    }
+    Program prog(src, buildflags, errmsg);
+    // Cache result of build failures too (to prevent unnecessary compiler invocations)
+    {
+        cv::AutoLock lock(program_cache_mutex);
+        phash.insert(std::pair<std::string, Program>(key, prog));
+        cacheList.push_front(key);
+    }
+    return prog;
+#else
+    CV_OPENCL_NO_SUPPORT();
+#endif
+}
 
 
 //////////////////////////////////////////// OpenCLAllocator //////////////////////////////////////////////////
@@ -6351,4 +6585,13 @@ uint64 Timer::durationNS() const
     return p->durationNS();
 }
 
+#ifndef HAVE_OPENCL
+#if defined(_MSC_VER)
+    #pragma warning(pop)
+#elif defined(__clang__)
+    #pragma clang diagnostic pop
+#elif defined(__GNUC__)
+    #pragma GCC diagnostic pop
+#endif
+#endif
 }} // namespace
diff --git a/modules/core/src/ocl_deprecated.hpp b/modules/core/src/ocl_deprecated.hpp
index 3cf261b8e7..753e8c312b 100644
--- a/modules/core/src/ocl_deprecated.hpp
+++ b/modules/core/src/ocl_deprecated.hpp
@@ -968,7 +968,7 @@ OCL_FUNC(cl_int, clRetainMemObject, (cl_mem memobj), (memobj))
 
 OCL_FUNC(cl_int, clReleaseMemObject, (cl_mem memobj), (memobj))
 
-
+/*
 OCL_FUNC_P(cl_program, clCreateProgramWithSource,
     (cl_context context,
     cl_uint count,
@@ -1014,7 +1014,7 @@ OCL_FUNC(cl_int, clGetProgramBuildInfo,
     void * param_value,
     size_t * param_value_size_ret),
     (program, device, param_name, param_value_size, param_value, param_value_size_ret))
-
+*/
 OCL_FUNC_P(cl_kernel, clCreateKernel,
     (cl_program program,
     const char * kernel_name,
diff --git a/modules/core/test/ocl/test_opencl.cpp b/modules/core/test/ocl/test_opencl.cpp
new file mode 100644
index 0000000000..b8ff3a131c
--- /dev/null
+++ b/modules/core/test/ocl/test_opencl.cpp
@@ -0,0 +1,131 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+#include "../test_precomp.hpp"
+
+#include <opencv2/core/ocl.hpp>
+#include <fstream>
+
+namespace opencv_test { namespace {
+
+static void testOpenCLKernel(cv::ocl::Kernel& k)
+{
+    ASSERT_FALSE(k.empty());
+    cv::UMat src(cv::Size(4096, 2048), CV_8UC1, cv::Scalar::all(100));
+    cv::UMat dst(src.size(), CV_8UC1);
+    size_t globalSize[2] = {(size_t)src.cols, (size_t)src.rows};
+    size_t localSize[2] = {8, 8};
+    int64 kernel_time = k.args(
+            cv::ocl::KernelArg::ReadOnlyNoSize(src), // size is not used (similar to 'dst' size)
+            cv::ocl::KernelArg::WriteOnly(dst),
+            (int)5
+        ).runProfiling(2, globalSize, localSize);
+    ASSERT_GE(kernel_time, (int64)0);
+    std::cout << "Kernel time: " << (kernel_time * 1e-6) << " ms" << std::endl;
+    cv::Mat res, reference(src.size(), CV_8UC1, cv::Scalar::all(105));
+    dst.copyTo(res);
+    EXPECT_EQ(0, cvtest::norm(reference, res, cv::NORM_INF));
+}
+
+TEST(OpenCL, support_binary_programs)
+{
+    cv::ocl::Context ctx = cv::ocl::Context::getDefault();
+    if (!ctx.ptr())
+    {
+        throw cvtest::SkipTestException("OpenCL is not available");
+    }
+    cv::ocl::Device device = cv::ocl::Device::getDefault();
+    if (!device.compilerAvailable())
+    {
+        throw cvtest::SkipTestException("OpenCL compiler is not available");
+    }
+    std::vector<char> program_binary_code;
+
+    cv::String module_name; // empty to disable OpenCL cache
+
+    { // Generate program binary from OpenCL C source
+        static const char* opencl_kernel_src =
+"__kernel void test_kernel(__global const uchar* src, int src_step, int src_offset,\n"
+"                          __global uchar* dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,\n"
+"                          int c)\n"
+"{\n"
+"   int x = get_global_id(0);\n"
+"   int y = get_global_id(1);\n"
+"   if (x < dst_cols && y < dst_rows)\n"
+"   {\n"
+"       int src_idx = y * src_step + x + src_offset;\n"
+"       int dst_idx = y * dst_step + x + dst_offset;\n"
+"       dst[dst_idx] = src[src_idx] + c;\n"
+"   }\n"
+"}\n";
+        cv::ocl::ProgramSource src(module_name, "simple", opencl_kernel_src, "");
+        cv::String errmsg;
+        cv::ocl::Program program(src, "", errmsg);
+        ASSERT_TRUE(program.ptr() != NULL);
+        cv::ocl::Kernel k("test_kernel", program);
+        EXPECT_FALSE(k.empty());
+        program.getBinary(program_binary_code);
+        std::cout << "Program binary size: " << program_binary_code.size() << " bytes" << std::endl;
+    }
+
+    cv::ocl::Kernel k;
+
+    { // Load program from binary (without sources)
+        ASSERT_FALSE(program_binary_code.empty());
+        cv::ocl::ProgramSource src = cv::ocl::ProgramSource::fromBinary(module_name, "simple_binary", (uchar*)&program_binary_code[0], program_binary_code.size(), "");
+        cv::String errmsg;
+        cv::ocl::Program program(src, "", errmsg);
+        ASSERT_TRUE(program.ptr() != NULL);
+        k.create("test_kernel", program);
+    }
+
+    testOpenCLKernel(k);
+}
+
+
+TEST(OpenCL, support_SPIR_programs)
+{
+    cv::ocl::Context ctx = cv::ocl::Context::getDefault();
+    if (!ctx.ptr())
+    {
+        throw cvtest::SkipTestException("OpenCL is not available");
+    }
+    cv::ocl::Device device = cv::ocl::Device::getDefault();
+    if (!device.isExtensionSupported("cl_khr_spir"))
+    {
+        throw cvtest::SkipTestException("'cl_khr_spir' extension is not supported by OpenCL device");
+    }
+    std::vector<char> program_binary_code;
+    cv::String fname = cv::format("test_kernel.spir%d", device.addressBits());
+    std::string full_path = cvtest::findDataFile(std::string("opencl/") + fname);
+
+    {
+        std::fstream f(full_path.c_str(), std::ios::in|std::ios::binary);
+        ASSERT_TRUE(f.is_open());
+        size_t pos = (size_t)f.tellg();
+        f.seekg(0, std::fstream::end);
+        size_t fileSize = (size_t)f.tellg();
+        std::cout << "Program SPIR size: " << fileSize << " bytes" << std::endl;
+        f.seekg(pos, std::fstream::beg);
+        program_binary_code.resize(fileSize);
+        f.read(&program_binary_code[0], fileSize);
+        ASSERT_FALSE(f.fail());
+    }
+
+    cv::String module_name; // empty to disable OpenCL cache
+
+    cv::ocl::Kernel k;
+
+    { // Load program from SPIR format
+        ASSERT_FALSE(program_binary_code.empty());
+        cv::ocl::ProgramSource src = cv::ocl::ProgramSource::fromSPIR(module_name, "simple_spir", (uchar*)&program_binary_code[0], program_binary_code.size(), "");
+        cv::String errmsg;
+        cv::ocl::Program program(src, "", errmsg);
+        ASSERT_TRUE(program.ptr() != NULL);
+        k.create("test_kernel", program);
+    }
+
+    testOpenCLKernel(k);
+}
+
+}} // namespace
diff --git a/samples/tapi/opencl_custom_kernel.cpp b/samples/tapi/opencl_custom_kernel.cpp
new file mode 100644
index 0000000000..87f5b9a24a
--- /dev/null
+++ b/samples/tapi/opencl_custom_kernel.cpp
@@ -0,0 +1,160 @@
+#include "opencv2/core.hpp"
+#include "opencv2/core/ocl.hpp"
+#include "opencv2/highgui.hpp"
+#include "opencv2/imgcodecs.hpp"
+#include "opencv2/imgproc.hpp"
+
+#include <iostream>
+
+using namespace std;
+using namespace cv;
+
+static const char* opencl_kernel_src =
+"__kernel void magnutude_filter_8u(\n"
+"       __global const uchar* src, int src_step, int src_offset,\n"
+"       __global uchar* dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,\n"
+"       float scale)\n"
+"{\n"
+"   int x = get_global_id(0);\n"
+"   int y = get_global_id(1);\n"
+"   if (x < dst_cols && y < dst_rows)\n"
+"   {\n"
+"       int dst_idx = y * dst_step + x + dst_offset;\n"
+"       if (x > 0 && x < dst_cols - 1 && y > 0 && y < dst_rows - 2)\n"
+"       {\n"
+"           int src_idx = y * src_step + x + src_offset;\n"
+"           int dx = (int)src[src_idx]*2 - src[src_idx - 1]          - src[src_idx + 1];\n"
+"           int dy = (int)src[src_idx]*2 - src[src_idx - 1*src_step] - src[src_idx + 1*src_step];\n"
+"           dst[dst_idx] = convert_uchar_sat(sqrt((float)(dx*dx + dy*dy)) * scale);\n"
+"       }\n"
+"       else\n"
+"       {\n"
+"           dst[dst_idx] = 0;\n"
+"       }\n"
+"   }\n"
+"}\n";
+
+int main(int argc, char** argv)
+{
+    const char* keys =
+        "{ i input    | | specify input image }"
+        "{ h help     | | print help message }";
+
+    cv::CommandLineParser args(argc, argv, keys);
+    if (args.has("help"))
+    {
+        cout << "Usage : " << argv[0] << " [options]" << endl;
+        cout << "Available options:" << endl;
+        args.printMessage();
+        return EXIT_SUCCESS;
+    }
+
+    cv::ocl::Context ctx = cv::ocl::Context::getDefault();
+    if (!ctx.ptr())
+    {
+        cerr << "OpenCL is not available" << endl;
+        return 1;
+    }
+    cv::ocl::Device device = cv::ocl::Device::getDefault();
+    if (!device.compilerAvailable())
+    {
+        cerr << "OpenCL compiler is not available" << endl;
+        return 1;
+    }
+
+
+    UMat src;
+    {
+        string image_file = args.get<string>("i");
+        if (!image_file.empty())
+        {
+            Mat image = imread(image_file);
+            if (image.empty())
+            {
+                cout << "error read image: " << image_file << endl;
+                return 1;
+            }
+            cvtColor(image, src, COLOR_BGR2GRAY);
+        }
+        else
+        {
+            Mat frame(cv::Size(640, 480), CV_8U, Scalar::all(128));
+            Point p(frame.cols / 2, frame.rows / 2);
+            line(frame, Point(0, frame.rows / 2), Point(frame.cols, frame.rows / 2), 1);
+            circle(frame, p, 200, Scalar(32, 32, 32), 8, LINE_AA);
+            string str = "OpenCL";
+            int baseLine = 0;
+            Size box = getTextSize(str, FONT_HERSHEY_COMPLEX, 2, 5, &baseLine);
+            putText(frame, str, Point((frame.cols - box.width) / 2, (frame.rows - box.height) / 2 + baseLine),
+                    FONT_HERSHEY_COMPLEX, 2, Scalar(255, 255, 255), 5, LINE_AA);
+            frame.copyTo(src);
+        }
+    }
+
+
+    cv::String module_name; // empty to disable OpenCL cache
+
+    {
+        cout << "OpenCL program source: " << endl;
+        cout << "======================================================================================================" << endl;
+        cout << opencl_kernel_src << endl;
+        cout << "======================================================================================================" << endl;
+        //! [Define OpenCL program source]
+        cv::ocl::ProgramSource source(module_name, "simple", opencl_kernel_src, "");
+        //! [Define OpenCL program source]
+
+        //! [Compile/build OpenCL for current OpenCL device]
+        cv::String errmsg;
+        cv::ocl::Program program(source, "", errmsg);
+        if (program.ptr() == NULL)
+        {
+            cerr << "Can't compile OpenCL program:" << endl << errmsg << endl;
+            return 1;
+        }
+        //! [Compile/build OpenCL for current OpenCL device]
+
+        if (!errmsg.empty())
+        {
+            cout << "OpenCL program build log:" << endl << errmsg << endl;
+        }
+
+        //! [Get OpenCL kernel by name]
+        cv::ocl::Kernel k("magnutude_filter_8u", program);
+        if (k.empty())
+        {
+            cerr << "Can't get OpenCL kernel" << endl;
+            return 1;
+        }
+        //! [Get OpenCL kernel by name]
+
+        UMat result(src.size(), CV_8UC1);
+
+        //! [Define kernel parameters and run]
+        size_t globalSize[2] = {(size_t)src.cols, (size_t)src.rows};
+        size_t localSize[2] = {8, 8};
+        bool executionResult = k
+            .args(
+                cv::ocl::KernelArg::ReadOnlyNoSize(src), // size is not used (similar to 'dst' size)
+                cv::ocl::KernelArg::WriteOnly(result),
+                (float)2.0
+            )
+            .run(2, globalSize, localSize, true);
+        if (!executionResult)
+        {
+            cerr << "OpenCL kernel launch failed" << endl;
+            return 1;
+        }
+        //! [Define kernel parameters and run]
+
+        imshow("Source", src);
+        imshow("Result", result);
+
+        for (;;)
+        {
+            int key = waitKey();
+            if (key == 27/*ESC*/ || key == 'q' || key == 'Q')
+                break;
+        }
+    }
+    return 0;
+}