Add clone() to kernel class

This adds wrapper for clCloneKernel() OpenCL 2.1 API
function to the kernel class.

Author: jszuppe

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
     {

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()