Merge pull request #10231 from alalek:ocl_refactor_program_api

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;

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>
+#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

samples/tapi/opencl_custom_kernel.cpp

+#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;
+}