ocl: low-level API to support OpenCL binary programs

modules/core/include/opencv2/core/ocl.hpp

@@ -606,17 +606,26 @@ public:
 
     bool create(const ProgramSource& src,
                 const String& buildflags, String& errmsg);
-    bool read(const String& buf, const String& buildflags);
-    bool write(String& buf) const;
+    bool read(const String& buf, const String& buildflags); // deprecated
+    bool write(String& buf) const; // deprecated
 
-    const ProgramSource& source() const;
+    const ProgramSource& source() const; // deprecated
     void* ptr() const;
 
-    String getPrefix() const;
-    static String getPrefix(const String& buildflags);
+    String getPrefix() const; // deprecated
+    static String getPrefix(const String& buildflags); // deprecated
 
 
-    struct Impl;
+    /**
+     * @brief Query device-specific program binary.
+     *
+     * @sa ProgramSource::fromBinary
+     *
+     * @param[out] binary output buffer
+     */
+    void getBinary(std::vector<char>& binary) const;
+
+    struct Impl; friend struct Impl;
     inline Impl* getImpl() const { return (Impl*)p; }
 protected:
     Impl* p;
@@ -636,10 +645,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;

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,

modules/core/test/ocl/test_opencl.cpp

+// 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>
+
+namespace opencv_test { namespace {
+
+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);
+    }
+
+    { // Run kernel
+        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));
+    }
+}
+
+}} // namespace