Merge pull request #720 from jszuppe/pr_kernel_clone

Add clCloneKernel() wrapper

Author: jszuppe

include/boost/compute/device.hpp

@@ -407,7 +407,8 @@ class device
         return m_id != other.m_id;
     }
 
-    /// \internal_
+    /// Returns \c true if the device OpenCL version is major.minor
+    /// or newer; otherwise returns \c false.
     bool check_version(int major, int minor) const
     {
         std::stringstream stream;

include/boost/compute/kernel.hpp

@@ -128,6 +128,21 @@ class kernel
         }
     }
 
+    #if defined(BOOST_COMPUTE_CL_VERSION_2_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
+    /// Creates a new kernel object based on a shallow copy of
+    /// the undelying OpenCL kernel object.
+    ///
+    /// \opencl_version_warning{2,1}
+    ///
+    /// \see_opencl_ref{clCloneKernel}
+    kernel clone()
+    {
+        cl_int ret = 0;
+        cl_kernel k = clCloneKernel(m_kernel, &ret);
+        return kernel(k, false);
+    }
+    #endif // BOOST_COMPUTE_CL_VERSION_2_1
+
     /// Returns a reference to the underlying OpenCL kernel object.
     cl_kernel& get() const
     {

include/boost/compute/platform.hpp

@@ -211,6 +211,23 @@ class platform
         return m_platform != other.m_platform;
     }
 
+    /// Returns \c true if the platform OpenCL version is major.minor
+    /// or newer; otherwise returns \c false.
+    bool check_version(int major, int minor) const
+    {
+        std::stringstream stream;
+        stream << version();
+
+        int actual_major, actual_minor;
+        stream.ignore(7); // 'OpenCL '
+        stream >> actual_major;
+        stream.ignore(1); // '.'
+        stream >> actual_minor;
+
+        return actual_major > major ||
+               (actual_major == major && actual_minor >= minor);
+    }
+
 private:
     cl_platform_id m_platform;
 };

test/opencl_version_check.hpp

@@ -14,4 +14,7 @@
 #define REQUIRES_OPENCL_VERSION(major, minor) \
     if (!device.check_version(major, minor)) return
 
+#define REQUIRES_OPENCL_PLATFORM_VERSION(major, minor) \
+    if (!device.platform().check_version(major, minor)) return
+
 #endif

test/test_kernel.cpp

@@ -17,6 +17,7 @@
 #include <boost/compute/utility/source.hpp>
 
 #include "context_setup.hpp"
+#include "check_macros.hpp"
 
 namespace compute = boost::compute;
 
@@ -288,4 +289,25 @@ BOOST_AUTO_TEST_CASE(get_sub_group_info_core)
 }
 #endif // BOOST_COMPUTE_CL_VERSION_2_1
 
+#ifdef BOOST_COMPUTE_CL_VERSION_2_1
+BOOST_AUTO_TEST_CASE(clone_kernel)
+{
+    REQUIRES_OPENCL_PLATFORM_VERSION(2, 1);
+
+    compute::kernel k1 = compute::kernel::create_with_source(
+        "__kernel void test(__global int * x) { x[get_global_id(0)] = get_global_id(0); }",
+        "test", context
+    );
+
+    compute::buffer x(context, 5 * sizeof(compute::int_));
+    k1.set_arg(0, x);
+
+    // Clone k1 kernel
+    compute::kernel k2 = k1.clone();
+    // After clone k2 0th argument (__global float * x) should be set,
+    // so we should be able to enqueue k2 kernel without problems
+    queue.enqueue_1d_range_kernel(k2, 0, x.size() / sizeof(compute::int_), 0).wait();
+}
+#endif // BOOST_COMPUTE_CL_VERSION_2_1
+
 BOOST_AUTO_TEST_SUITE_END()