diff --git a/MANIFEST.in b/MANIFEST.in
index 7a1dcbf027..2cb0991412 100644
--- a/MANIFEST.in
+++ b/MANIFEST.in
@@ -12,6 +12,8 @@ include dpctl/_sycl_event.h
 include dpctl/_sycl_event_api.h
 include dpctl/memory/_memory.h
 include dpctl/memory/_memory_api.h
+include dpctl/program/_program.h
+include dpctl/program/_program_api.h
 include dpctl/tensor/_usmarray.h
 include dpctl/tensor/_usmarray_api.h
 recursive-include dpctl/tensor/include *
diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd
index fdd19a6652..f4ce41f6a2 100644
--- a/dpctl/_backend.pxd
+++ b/dpctl/_backend.pxd
@@ -263,6 +263,7 @@ cdef extern from "syclinterface/dpctl_sycl_event_interface.h":
 cdef extern from "syclinterface/dpctl_sycl_kernel_interface.h":
     cdef size_t DPCTLKernel_GetNumArgs(const DPCTLSyclKernelRef KRef)
     cdef void DPCTLKernel_Delete(DPCTLSyclKernelRef KRef)
+    cdef DPCTLSyclKernelRef DPCTLKernel_Copy(const DPCTLSyclKernelRef KRef)
     cdef size_t DPCTLKernel_GetWorkGroupSize(const DPCTLSyclKernelRef KRef)
     cdef size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple(const DPCTLSyclKernelRef KRef)
     cdef size_t DPCTLKernel_GetPrivateMemSize(const DPCTLSyclKernelRef KRef)
@@ -341,6 +342,7 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h":
     cdef bool DPCTLKernelBundle_HasKernel(DPCTLSyclKernelBundleRef KBRef,
                                      const char *KernelName)
     cdef void DPCTLKernelBundle_Delete(DPCTLSyclKernelBundleRef KBRef)
+    cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy(const DPCTLSyclKernelBundleRef KBRef)
 
 
 cdef extern from "syclinterface/dpctl_sycl_queue_interface.h":
diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp
index 5b40f222ce..f2a5dc11e9 100644
--- a/dpctl/apis/include/dpctl4pybind11.hpp
+++ b/dpctl/apis/include/dpctl4pybind11.hpp
@@ -54,9 +54,9 @@ constexpr int platform_typeid_lookup(int I, Ints... Is)
                : platform_typeid_lookup<Concrete, Ts...>(Is...);
 }
 
-struct dpctl_capi
+class dpctl_capi
 {
-
+public:
     // dpctl type objects
     PyTypeObject *Py_SyclDeviceType_;
     PyTypeObject *PySyclDeviceType_;
@@ -71,6 +71,8 @@ struct dpctl_capi
     PyTypeObject *PyMemoryUSMSharedType_;
     PyTypeObject *PyMemoryUSMHostType_;
     PyTypeObject *PyUSMArrayType_;
+    PyTypeObject *PySyclProgramType_;
+    PyTypeObject *PySyclKernelType_;
 
     DPCTLSyclDeviceRef (*SyclDevice_GetDeviceRef_)(PySyclDeviceObject *);
     PySyclDeviceObject *(*SyclDevice_Make_)(DPCTLSyclDeviceRef);
@@ -94,6 +96,14 @@ struct dpctl_capi
                               DPCTLSyclQueueRef,
                               PyObject *);
 
+    // program
+    DPCTLSyclKernelRef (*SyclKernel_GetKernelRef_)(PySyclKernelObject *);
+    PySyclKernelObject *(*SyclKernel_Make_)(DPCTLSyclKernelRef, const char *);
+
+    DPCTLSyclKernelBundleRef (*SyclProgram_GetKernelBundleRef_)(
+        PySyclProgramObject *);
+    PySyclProgramObject *(*SyclProgram_Make_)(DPCTLSyclKernelBundleRef);
+
     // tensor
     char *(*UsmNDArray_GetData_)(PyUSMArrayObject *);
     int (*UsmNDArray_GetNDim_)(PyUSMArrayObject *);
@@ -131,6 +141,14 @@ struct dpctl_capi
     {
         return PyObject_TypeCheck(obj, PySyclQueueType_) != 0;
     }
+    bool PySyclKernel_Check_(PyObject *obj) const
+    {
+        return PyObject_TypeCheck(obj, PySyclKernelType_) != 0;
+    }
+    bool PySyclProgram_Check_(PyObject *obj) const
+    {
+        return PyObject_TypeCheck(obj, PySyclProgramType_) != 0;
+    }
 
     ~dpctl_capi(){};
 
@@ -142,19 +160,19 @@ struct dpctl_capi
 
     py::object default_sycl_queue_pyobj()
     {
-        return *default_sycl_queue;
+        return *default_sycl_queue_;
     }
     py::object default_usm_memory_pyobj()
     {
-        return *default_usm_memory;
+        return *default_usm_memory_;
     }
     py::object default_usm_ndarray_pyobj()
     {
-        return *default_usm_ndarray;
+        return *default_usm_ndarray_;
     }
     py::object as_usm_memory_pyobj()
     {
-        return *as_usm_memory;
+        return *as_usm_memory_;
     }
 
 private:
@@ -170,14 +188,14 @@ struct dpctl_capi
         }
     };
 
-    std::shared_ptr<py::object> default_sycl_queue;
-    std::shared_ptr<py::object> default_usm_memory;
-    std::shared_ptr<py::object> default_usm_ndarray;
-    std::shared_ptr<py::object> as_usm_memory;
+    std::shared_ptr<py::object> default_sycl_queue_;
+    std::shared_ptr<py::object> default_usm_memory_;
+    std::shared_ptr<py::object> default_usm_ndarray_;
+    std::shared_ptr<py::object> as_usm_memory_;
 
     dpctl_capi()
-        : default_sycl_queue{}, default_usm_memory{}, default_usm_ndarray{},
-          as_usm_memory{}
+        : default_sycl_queue_{}, default_usm_memory_{}, default_usm_ndarray_{},
+          as_usm_memory_{}
     {
         // Import Cython-generated C-API for dpctl
         // This imports python modules and initializes
@@ -201,6 +219,8 @@ struct dpctl_capi
         this->PyMemoryUSMSharedType_ = &PyMemoryUSMSharedType;
         this->PyMemoryUSMHostType_ = &PyMemoryUSMHostType;
         this->PyUSMArrayType_ = &PyUSMArrayType;
+        this->PySyclProgramType_ = &PySyclProgramType;
+        this->PySyclKernelType_ = &PySyclKernelType;
 
         // SyclDevice API
         this->SyclDevice_GetDeviceRef_ = SyclDevice_GetDeviceRef;
@@ -225,6 +245,12 @@ struct dpctl_capi
         this->Memory_GetNumBytes_ = Memory_GetNumBytes;
         this->Memory_Make_ = Memory_Make;
 
+        // dpctl.program API
+        this->SyclKernel_GetKernelRef_ = SyclKernel_GetKernelRef;
+        this->SyclKernel_Make_ = SyclKernel_Make;
+        this->SyclProgram_GetKernelBundleRef_ = SyclProgram_GetKernelBundleRef;
+        this->SyclProgram_Make_ = SyclProgram_Make;
+
         // dpctl.tensor.usm_ndarray API
         this->UsmNDArray_GetData_ = UsmNDArray_GetData;
         this->UsmNDArray_GetNDim_ = UsmNDArray_GetNDim;
@@ -284,18 +310,18 @@ struct dpctl_capi
         py::object py_sycl_queue = py::reinterpret_steal<py::object>(
             reinterpret_cast<PyObject *>(py_q_tmp));
 
-        default_sycl_queue = std::shared_ptr<py::object>(
+        default_sycl_queue_ = std::shared_ptr<py::object>(
             new py::object(py_sycl_queue), Deleter{});
 
         py::module_ mod_memory = py::module_::import("dpctl.memory");
         py::object py_as_usm_memory = mod_memory.attr("as_usm_memory");
-        as_usm_memory = std::shared_ptr<py::object>(
+        as_usm_memory_ = std::shared_ptr<py::object>(
             new py::object{py_as_usm_memory}, Deleter{});
 
         auto mem_kl = mod_memory.attr("MemoryUSMHost");
         py::object py_default_usm_memory =
             mem_kl(1, py::arg("queue") = py_sycl_queue);
-        default_usm_memory = std::shared_ptr<py::object>(
+        default_usm_memory_ = std::shared_ptr<py::object>(
             new py::object{py_default_usm_memory}, Deleter{});
 
         py::module_ mod_usmarray =
@@ -306,7 +332,7 @@ struct dpctl_capi
             tensor_kl(py::tuple(), py::arg("dtype") = py::str("u1"),
                       py::arg("buffer") = py_default_usm_memory);
 
-        default_usm_ndarray = std::shared_ptr<py::object>(
+        default_usm_ndarray_ = std::shared_ptr<py::object>(
             new py::object{py_default_usm_ndarray}, Deleter{});
     }
 
@@ -377,7 +403,7 @@ template <> struct type_caster<sycl::queue>
     bool load(handle src, bool)
     {
         PyObject *source = src.ptr();
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         if (api.PySyclQueue_Check_(source)) {
             DPCTLSyclQueueRef QRef = api.SyclQueue_GetQueueRef_(
                 reinterpret_cast<PySyclQueueObject *>(source));
@@ -393,7 +419,7 @@ template <> struct type_caster<sycl::queue>
 
     static handle cast(sycl::queue src, return_value_policy, handle)
     {
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         auto tmp =
             api.SyclQueue_Make_(reinterpret_cast<DPCTLSyclQueueRef>(&src));
         return handle(reinterpret_cast<PyObject *>(tmp));
@@ -412,7 +438,7 @@ template <> struct type_caster<sycl::device>
     bool load(handle src, bool)
     {
         PyObject *source = src.ptr();
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         if (api.PySyclDevice_Check_(source)) {
             DPCTLSyclDeviceRef DRef = api.SyclDevice_GetDeviceRef_(
                 reinterpret_cast<PySyclDeviceObject *>(source));
@@ -428,7 +454,7 @@ template <> struct type_caster<sycl::device>
 
     static handle cast(sycl::device src, return_value_policy, handle)
     {
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         auto tmp =
             api.SyclDevice_Make_(reinterpret_cast<DPCTLSyclDeviceRef>(&src));
         return handle(reinterpret_cast<PyObject *>(tmp));
@@ -447,7 +473,7 @@ template <> struct type_caster<sycl::context>
     bool load(handle src, bool)
     {
         PyObject *source = src.ptr();
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         if (api.PySyclContext_Check_(source)) {
             DPCTLSyclContextRef CRef = api.SyclContext_GetContextRef_(
                 reinterpret_cast<PySyclContextObject *>(source));
@@ -463,7 +489,7 @@ template <> struct type_caster<sycl::context>
 
     static handle cast(sycl::context src, return_value_policy, handle)
     {
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         auto tmp =
             api.SyclContext_Make_(reinterpret_cast<DPCTLSyclContextRef>(&src));
         return handle(reinterpret_cast<PyObject *>(tmp));
@@ -482,7 +508,7 @@ template <> struct type_caster<sycl::event>
     bool load(handle src, bool)
     {
         PyObject *source = src.ptr();
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         if (api.PySyclEvent_Check_(source)) {
             DPCTLSyclEventRef ERef = api.SyclEvent_GetEventRef_(
                 reinterpret_cast<PySyclEventObject *>(source));
@@ -498,7 +524,7 @@ template <> struct type_caster<sycl::event>
 
     static handle cast(sycl::event src, return_value_policy, handle)
     {
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         auto tmp =
             api.SyclEvent_Make_(reinterpret_cast<DPCTLSyclEventRef>(&src));
         return handle(reinterpret_cast<PyObject *>(tmp));
@@ -506,6 +532,86 @@ template <> struct type_caster<sycl::event>
 
     DPCTL_TYPE_CASTER(sycl::event, _("dpctl.SyclEvent"));
 };
+
+/* This type caster associates ``sycl::kernel`` C++ class with
+ * :class:`dpctl.program.SyclKernel` for the purposes of generation of
+ * Python bindings by pybind11.
+ */
+template <> struct type_caster<sycl::kernel>
+{
+public:
+    bool load(handle src, bool)
+    {
+        PyObject *source = src.ptr();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
+        if (api.PySyclKernel_Check_(source)) {
+            DPCTLSyclKernelRef KRef = api.SyclKernel_GetKernelRef_(
+                reinterpret_cast<PySyclKernelObject *>(source));
+            value = std::make_unique<sycl::kernel>(
+                *(reinterpret_cast<sycl::kernel *>(KRef)));
+            return true;
+        }
+        else {
+            throw py::type_error("Input is of unexpected type, expected "
+                                 "dpctl.program.SyclKernel");
+        }
+    }
+
+    static handle cast(sycl::kernel src, return_value_policy, handle)
+    {
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
+        auto tmp =
+            api.SyclKernel_Make_(reinterpret_cast<DPCTLSyclKernelRef>(&src),
+                                 "dpctl4pybind11_kernel");
+        return handle(reinterpret_cast<PyObject *>(tmp));
+    }
+
+    DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.program.SyclKernel"));
+};
+
+/* This type caster associates
+ * ``sycl::kernel_bundle<sycl::bundle_state::executable>`` C++ class with
+ * :class:`dpctl.program.SyclProgram` for the purposes of generation of
+ * Python bindings by pybind11.
+ */
+template <>
+struct type_caster<sycl::kernel_bundle<sycl::bundle_state::executable>>
+{
+public:
+    bool load(handle src, bool)
+    {
+        PyObject *source = src.ptr();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
+        if (api.PySyclProgram_Check_(source)) {
+            DPCTLSyclKernelBundleRef KBRef =
+                api.SyclProgram_GetKernelBundleRef_(
+                    reinterpret_cast<PySyclProgramObject *>(source));
+            value = std::make_unique<
+                sycl::kernel_bundle<sycl::bundle_state::executable>>(
+                *(reinterpret_cast<
+                    sycl::kernel_bundle<sycl::bundle_state::executable> *>(
+                    KBRef)));
+            return true;
+        }
+        else {
+            throw py::type_error("Input is of unexpected type, expected "
+                                 "dpctl.program.SyclProgram");
+        }
+    }
+
+    static handle cast(sycl::kernel_bundle<sycl::bundle_state::executable> src,
+                       return_value_policy,
+                       handle)
+    {
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
+        auto tmp = api.SyclProgram_Make_(
+            reinterpret_cast<DPCTLSyclKernelBundleRef>(&src));
+        return handle(reinterpret_cast<PyObject *>(tmp));
+    }
+
+    DPCTL_TYPE_CASTER(sycl::kernel_bundle<sycl::bundle_state::executable>,
+                      _("dpctl.program.SyclProgram"));
+};
 } // namespace detail
 } // namespace pybind11
 
@@ -544,7 +650,7 @@ class usm_memory : public py::object
     sycl::queue get_queue() const
     {
         Py_MemoryObject *mem_obj = reinterpret_cast<Py_MemoryObject *>(m_ptr);
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         DPCTLSyclQueueRef QRef = api.Memory_GetQueueRef_(mem_obj);
         sycl::queue *obj_q = reinterpret_cast<sycl::queue *>(QRef);
         return *obj_q;
@@ -553,14 +659,14 @@ class usm_memory : public py::object
     char *get_pointer() const
     {
         Py_MemoryObject *mem_obj = reinterpret_cast<Py_MemoryObject *>(m_ptr);
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         DPCTLSyclUSMRef MRef = api.Memory_GetUsmPointer_(mem_obj);
         return reinterpret_cast<char *>(MRef);
     }
 
     size_t get_nbytes() const
     {
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         Py_MemoryObject *mem_obj = reinterpret_cast<Py_MemoryObject *>(m_ptr);
         return api.Memory_GetNumBytes_(mem_obj);
     }
@@ -663,7 +769,7 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return api.UsmNDArray_GetData_(raw_ar);
     }
 
@@ -676,7 +782,7 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return api.UsmNDArray_GetNDim_(raw_ar);
     }
 
@@ -684,7 +790,7 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return api.UsmNDArray_GetShape_(raw_ar);
     }
 
@@ -698,7 +804,7 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return api.UsmNDArray_GetStrides_(raw_ar);
     }
 
@@ -706,7 +812,7 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         int ndim = api.UsmNDArray_GetNDim_(raw_ar);
         const py::ssize_t *shape = api.UsmNDArray_GetShape_(raw_ar);
 
@@ -723,7 +829,7 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         int nd = api.UsmNDArray_GetNDim_(raw_ar);
         const py::ssize_t *shape = api.UsmNDArray_GetShape_(raw_ar);
         const py::ssize_t *strides = api.UsmNDArray_GetStrides_(raw_ar);
@@ -757,7 +863,7 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         DPCTLSyclQueueRef QRef = api.UsmNDArray_GetQueueRef_(raw_ar);
         return *(reinterpret_cast<sycl::queue *>(QRef));
     }
@@ -766,7 +872,7 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return api.UsmNDArray_GetTypenum_(raw_ar);
     }
 
@@ -774,7 +880,7 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return api.UsmNDArray_GetFlags_(raw_ar);
     }
 
@@ -782,28 +888,28 @@ class usm_ndarray : public py::object
     {
         PyUSMArrayObject *raw_ar = this->usm_array_ptr();
 
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return api.UsmNDArray_GetElementSize_(raw_ar);
     }
 
     bool is_c_contiguous() const
     {
         int flags = this->get_flags();
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return static_cast<bool>(flags & api.USM_ARRAY_C_CONTIGUOUS_);
     }
 
     bool is_f_contiguous() const
     {
         int flags = this->get_flags();
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return static_cast<bool>(flags & api.USM_ARRAY_F_CONTIGUOUS_);
     }
 
     bool is_writable() const
     {
         int flags = this->get_flags();
-        auto &api = ::dpctl::detail::dpctl_capi::get();
+        auto const &api = ::dpctl::detail::dpctl_capi::get();
         return static_cast<bool>(flags & api.USM_ARRAY_WRITABLE_);
     }
 
diff --git a/dpctl/apis/include/dpctl_capi.h b/dpctl/apis/include/dpctl_capi.h
index d6c104581a..9715e42b38 100644
--- a/dpctl/apis/include/dpctl_capi.h
+++ b/dpctl/apis/include/dpctl_capi.h
@@ -40,6 +40,9 @@
 #include "../memory/_memory_api.h"
 #include "../tensor/_usmarray.h"
 #include "../tensor/_usmarray_api.h"
+#include "../program/_program.h"
+#include "../program/_program_api.h"
+
 // clang-format on
 
 /*
@@ -59,5 +62,6 @@ static inline void import_dpctl(void)
     import_dpctl___sycl_queue();
     import_dpctl__memory___memory();
     import_dpctl__tensor___usmarray();
+    import_dpctl__program___program();
     return;
 }
diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd
index 2133a7a667..86c338aff7 100644
--- a/dpctl/program/_program.pxd
+++ b/dpctl/program/_program.pxd
@@ -28,7 +28,7 @@ from .._sycl_device cimport SyclDevice
 from .._sycl_queue cimport SyclQueue
 
 
-cdef class SyclKernel:
+cdef api class SyclKernel [object PySyclKernelObject, type PySyclKernelType]:
     ''' Wraps a sycl::kernel object created from an OpenCL interoperability
         kernel.
     '''
@@ -40,7 +40,7 @@ cdef class SyclKernel:
     cdef SyclKernel _create (DPCTLSyclKernelRef kref, str name)
 
 
-cdef class SyclProgram:
+cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]:
     ''' Wraps a sycl::kernel_bundle<sycl::bundle_state::executable> object created by
         using SYCL interoperability layer for OpenCL and Level-Zero backends.
 
diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx
index 26a9c67a7e..82ff39de56 100644
--- a/dpctl/program/_program.pyx
+++ b/dpctl/program/_program.pyx
@@ -30,6 +30,7 @@ from libc.stdint cimport uint32_t
 
 from dpctl._backend cimport (  # noqa: E211, E402;
     DPCTLCString_Delete,
+    DPCTLKernel_Copy,
     DPCTLKernel_Delete,
     DPCTLKernel_GetCompileNumSubGroups,
     DPCTLKernel_GetCompileSubGroupSize,
@@ -38,6 +39,7 @@ from dpctl._backend cimport (  # noqa: E211, E402;
     DPCTLKernel_GetPreferredWorkGroupSizeMultiple,
     DPCTLKernel_GetPrivateMemSize,
     DPCTLKernel_GetWorkGroupSize,
+    DPCTLKernelBundle_Copy,
     DPCTLKernelBundle_CreateFromOCLSource,
     DPCTLKernelBundle_CreateFromSpirv,
     DPCTLKernelBundle_Delete,
@@ -165,6 +167,25 @@ cdef class SyclKernel:
         return n
 
 
+cdef api DPCTLSyclKernelRef SyclKernel_GetKernelRef(SyclKernel ker):
+    """ C-API function to access opaque kernel reference from
+    Python object of type :class:`dpctl.program.SyclKernel`.
+    """
+    return ker.get_kernel_ref()
+
+
+cdef api SyclKernel SyclKernel_Make(DPCTLSyclKernelRef KRef, const char *name):
+    """
+    C-API function to create :class:`dpctl.program.SyclKernel`
+    instance from opaque sycl kernel reference.
+    """
+    cdef DPCTLSyclKernelRef copied_KRef = DPCTLKernel_Copy(KRef)
+    if (name is NULL):
+        return SyclKernel._create(copied_KRef, "default_name")
+    else:
+        return SyclKernel._create(copied_KRef, name.decode("utf-8"))
+
+
 cdef class SyclProgram:
     """ Wraps a ``sycl::kernel_bundle<sycl::bundle_state::executable>`` object
     created using SYCL interoperability layer with underlying backends. Only the
@@ -290,3 +311,19 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,
         raise SyclProgramCompilationError()
 
     return SyclProgram._create(KBref)
+
+
+cdef api DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef(SyclProgram pro):
+    """ C-API function to access opaque kernel bundle reference from
+    Python object of type :class:`dpctl.program.SyclKernel`.
+    """
+    return pro.get_program_ref()
+
+
+cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef):
+    """
+    C-API function to create :class:`dpctl.program.SyclProgram`
+    instance from opaque sycl kernel bundle reference.
+    """
+    cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef)
+    return SyclProgram._create(copied_KBRef)
diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py
index ac55dce972..1cbd9114e4 100644
--- a/dpctl/tests/test_sycl_program.py
+++ b/dpctl/tests/test_sycl_program.py
@@ -31,6 +31,123 @@ def get_spirv_abspath(fn):
     return spirv_file
 
 
+def _check_cpython_api_SyclProgram_GetKernelBundleRef(sycl_prog):
+    """Checks Cython-generated C-API function
+    `SyclProgram_GetKernelBundleRef` defined in _program.pyx"""
+    import ctypes
+    import sys
+
+    assert type(sycl_prog) is dpctl_prog.SyclProgram
+    mod = sys.modules[sycl_prog.__class__.__module__]
+    # get capsule storing SyclProgram_GetKernelBundleRef function ptr
+    kb_ref_fn_cap = mod.__pyx_capi__["SyclProgram_GetKernelBundleRef"]
+    # construct Python callable to invoke "SyclProgram_GetKernelBundleRef"
+    cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer
+    cap_ptr_fn.restype = ctypes.c_void_p
+    cap_ptr_fn.argtypes = [ctypes.py_object, ctypes.c_char_p]
+    kb_ref_fn_ptr = cap_ptr_fn(
+        kb_ref_fn_cap,
+        b"DPCTLSyclKernelBundleRef (struct PySyclProgramObject *)",
+    )
+    # PYFUNCTYPE(result_type, *arg_types)
+    callable_maker = ctypes.PYFUNCTYPE(ctypes.c_void_p, ctypes.py_object)
+    get_kernel_bundle_ref_fn = callable_maker(kb_ref_fn_ptr)
+
+    r2 = sycl_prog.addressof_ref()
+    r1 = get_kernel_bundle_ref_fn(sycl_prog)
+    assert r1 == r2
+
+
+def _check_cpython_api_SyclProgram_Make(sycl_prog):
+    """Checks Cython-generated C-API function
+    `SyclProgram_Make` defined in _program.pyx"""
+    import ctypes
+    import sys
+
+    assert type(sycl_prog) is dpctl_prog.SyclProgram
+    mod = sys.modules[sycl_prog.__class__.__module__]
+    # get capsule storing SyclProgram_Make function ptr
+    make_prog_fn_cap = mod.__pyx_capi__["SyclProgram_Make"]
+    # construct Python callable to invoke "SyclProgram_Make"
+    cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer
+    cap_ptr_fn.restype = ctypes.c_void_p
+    cap_ptr_fn.argtypes = [ctypes.py_object, ctypes.c_char_p]
+    make_prog_fn_ptr = cap_ptr_fn(
+        make_prog_fn_cap,
+        b"struct PySyclProgramObject *(DPCTLSyclKernelBundleRef)",
+    )
+    # PYFUNCTYPE(result_type, *arg_types)
+    callable_maker = ctypes.PYFUNCTYPE(ctypes.py_object, ctypes.c_void_p)
+    make_prog_fn = callable_maker(make_prog_fn_ptr)
+
+    p2 = make_prog_fn(sycl_prog.addressof_ref())
+    assert p2.has_sycl_kernel("add")
+    assert p2.has_sycl_kernel("axpy")
+
+
+def _check_cpython_api_SyclKernel_GetKernelRef(krn):
+    """Checks Cython-generated C-API function
+    `SyclKernel_GetKernelRef` defined in _program.pyx"""
+    import ctypes
+    import sys
+
+    assert type(krn) is dpctl_prog.SyclKernel
+    mod = sys.modules[krn.__class__.__module__]
+    # get capsule storing SyclKernel_GetKernelRef function ptr
+    k_ref_fn_cap = mod.__pyx_capi__["SyclKernel_GetKernelRef"]
+    # construct Python callable to invoke "SyclKernel_GetKernelRef"
+    cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer
+    cap_ptr_fn.restype = ctypes.c_void_p
+    cap_ptr_fn.argtypes = [ctypes.py_object, ctypes.c_char_p]
+    k_ref_fn_ptr = cap_ptr_fn(
+        k_ref_fn_cap, b"DPCTLSyclKernelRef (struct PySyclKernelObject *)"
+    )
+    # PYFUNCTYPE(result_type, *arg_types)
+    callable_maker = ctypes.PYFUNCTYPE(ctypes.c_void_p, ctypes.py_object)
+    get_kernel_ref_fn = callable_maker(k_ref_fn_ptr)
+
+    r2 = krn.addressof_ref()
+    r1 = get_kernel_ref_fn(krn)
+    assert r1 == r2
+
+
+def _check_cpython_api_SyclKernel_Make(krn):
+    """Checks Cython-generated C-API function
+    `SyclKernel_Make` defined in _program.pyx"""
+    import ctypes
+    import sys
+
+    assert type(krn) is dpctl_prog.SyclKernel
+    mod = sys.modules[krn.__class__.__module__]
+    # get capsule storing SyclKernel_Make function ptr
+    k_make_fn_cap = mod.__pyx_capi__["SyclKernel_Make"]
+    # construct Python callable to invoke "SyclKernel_Make"
+    cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer
+    cap_ptr_fn.restype = ctypes.c_void_p
+    cap_ptr_fn.argtypes = [ctypes.py_object, ctypes.c_char_p]
+    k_make_fn_ptr = cap_ptr_fn(
+        k_make_fn_cap,
+        b"struct PySyclKernelObject *(DPCTLSyclKernelRef, char const *)",
+    )
+    # PYFUNCTYPE(result_type, *arg_types)
+    callable_maker = ctypes.PYFUNCTYPE(
+        ctypes.py_object, ctypes.c_void_p, ctypes.c_void_p
+    )
+    make_kernel_fn = callable_maker(k_make_fn_ptr)
+
+    k2 = make_kernel_fn(
+        krn.addressof_ref(), bytes(krn.get_function_name(), "utf-8")
+    )
+    assert krn.get_function_name() == k2.get_function_name()
+    assert krn.get_num_args() == k2.get_num_args()
+    assert krn.work_group_size == k2.work_group_size
+
+    k3 = make_kernel_fn(krn.addressof_ref(), ctypes.c_void_p(None))
+    assert k3.get_function_name() == "default_name"
+    assert krn.get_num_args() == k3.get_num_args()
+    assert krn.work_group_size == k3.work_group_size
+
+
 def _check_multi_kernel_program(prog):
     assert type(prog) is dpctl_prog.SyclProgram
 
@@ -49,6 +166,9 @@ def _check_multi_kernel_program(prog):
     assert type(axpyKernel.addressof_ref()) is int
 
     for krn in [addKernel, axpyKernel]:
+        _check_cpython_api_SyclKernel_GetKernelRef(krn)
+        _check_cpython_api_SyclKernel_Make(krn)
+
         na = krn.num_args
         assert na == krn.get_num_args()
         wgsz = krn.work_group_size
@@ -68,6 +188,9 @@ def _check_multi_kernel_program(prog):
         cmsgsz = krn.compile_sub_group_size
         assert type(cmsgsz) is int
 
+    _check_cpython_api_SyclProgram_GetKernelBundleRef(prog)
+    _check_cpython_api_SyclProgram_Make(prog)
+
 
 def test_create_program_from_source_ocl():
     oclSrc = "                                                             \
diff --git a/dpctl/tests/test_usm_ndarray_ctor.py b/dpctl/tests/test_usm_ndarray_ctor.py
index 0531eda402..e00747bfe9 100644
--- a/dpctl/tests/test_usm_ndarray_ctor.py
+++ b/dpctl/tests/test_usm_ndarray_ctor.py
@@ -825,8 +825,11 @@ def test_astype():
     X[:] = np.full((5, 5), 7, dtype="i4")
     Y = dpt.astype(X, "c8", order="C")
     assert np.allclose(dpt.to_numpy(Y), np.full((5, 5), 7, dtype="c8"))
-    Y = dpt.astype(X[::2, ::-1], "f2", order="K")
-    assert np.allclose(dpt.to_numpy(Y), np.full(Y.shape, 7, dtype="f2"))
+    if Y.sycl_device.has_aspect_fp16:
+        Y = dpt.astype(X[::2, ::-1], "f2", order="K")
+        assert np.allclose(dpt.to_numpy(Y), np.full(Y.shape, 7, dtype="f2"))
+    Y = dpt.astype(X[::2, ::-1], "f4", order="K")
+    assert np.allclose(dpt.to_numpy(Y), np.full(Y.shape, 7, dtype="f4"))
     Y = dpt.astype(X[::2, ::-1], "i4", order="K", copy=False)
     assert Y.usm_data is X.usm_data
 
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt b/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt
new file mode 100644
index 0000000000..441a0997c6
--- /dev/null
+++ b/examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt
@@ -0,0 +1,36 @@
+cmake_minimum_required(VERSION 3.21)
+
+project(use_queue_device LANGUAGES CXX)
+
+set(DPCTL_CMAKE_MODULES_PATH "${CMAKE_SOURCE_DIR}/../../../cmake")
+set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${DPCTL_CMAKE_MODULES_PATH})
+find_package(IntelDPCPP REQUIRED PATHS ${DPCTL_CMAKE_MODULES_PATH} NO_DEFAULT_PATH)
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED True)
+set(CMAKE_BUILD_TYPE Debug)
+
+# Fetch pybind11
+include(FetchContent)
+FetchContent_Declare(
+  pybind11
+  URL https://github.com/pybind/pybind11/archive/refs/tags/v2.10.0.tar.gz
+  URL_HASH SHA256=eacf582fa8f696227988d08cfc46121770823839fe9e301a20fbce67e7cd70ec
+)
+FetchContent_MakeAvailable(pybind11)
+
+find_package(PythonExtensions REQUIRED)
+find_package(Dpctl REQUIRED)
+find_package(NumPy REQUIRED)
+
+set(py_module_name _use_kernel)
+pybind11_add_module(${py_module_name}
+    MODULE
+    use_kernel/_example.cpp
+)
+target_include_directories(${py_module_name} PUBLIC ${Dpctl_INCLUDE_DIRS})
+install(TARGETS ${py_module_name}
+  DESTINATION use_kernel
+)
+
+set(ignoreMe "${SKBUILD}")
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/README.md b/examples/pybind11/use_dpctl_sycl_kernel/README.md
new file mode 100644
index 0000000000..77aa57bf6e
--- /dev/null
+++ b/examples/pybind11/use_dpctl_sycl_kernel/README.md
@@ -0,0 +1,25 @@
+# Usage of dpctl Entities in Pybind11
+
+## Description
+
+This extension demonstrates how you can use dpctl Python types,
+such as ``dpctl.SyclQueue`` and ``dpctl.program.SyclKernel``, in
+Pybind11 extensions.
+
+
+## Building
+
+To build the extension, run:
+```
+source /opt/intel/oneapi/compiler/latest/env/vars.sh
+CXX=icpx python setup.py build_ext --inplace
+python -m pytest tests
+python example.py
+```
+
+# Sample output
+
+```
+(dpctl) [17:25:27 ubuntu_vm use_dpctl_syclkernel]$ python example.py
+[ 0  2  4  6  8 10 12 14 16 18 20 22 24]
+```
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/example.py b/examples/pybind11/use_dpctl_sycl_kernel/example.py
new file mode 100644
index 0000000000..292cdf414d
--- /dev/null
+++ b/examples/pybind11/use_dpctl_sycl_kernel/example.py
@@ -0,0 +1,47 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2022 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+# coding: utf-8
+
+import use_kernel as eg
+
+import dpctl
+import dpctl.program as dppr
+import dpctl.tensor as dpt
+
+# create execution queue, targeting default selected device
+q = dpctl.SyclQueue()
+
+# read SPIR-V: a program in Khronos standardized intermediate form
+with open("resource/double_it.spv", "br") as fh:
+    il = fh.read()
+
+# Build the program for the selected device
+pr = dppr.create_program_from_spirv(q, il, "")
+assert pr.has_sycl_kernel("double_it")
+
+# Retrieve the kernel from the problem
+krn = pr.get_sycl_kernel("double_it")
+assert krn.num_args == 2
+
+# Construct the argument, and allocate memory for the result
+x = dpt.arange(0, stop=13, step=1, dtype="i4", sycl_queue=q)
+y = dpt.empty_like(x)
+
+eg.submit_custom_kernel(q, krn, src=x, dst=y)
+
+# output the result
+print(dpt.asnumpy(y))
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/resource/README.md b/examples/pybind11/use_dpctl_sycl_kernel/resource/README.md
new file mode 100644
index 0000000000..e83f8296a8
--- /dev/null
+++ b/examples/pybind11/use_dpctl_sycl_kernel/resource/README.md
@@ -0,0 +1,8 @@
+# Rebuilding SPIR-V file from source
+
+```bash
+export TOOLS_DIR=$(dirname $(dirname $(which icx)))/bin-llvm
+$TOOLS_DIR/clang -cc1 -triple spir double_it.cl -finclude-default-header -flto -emit -llvm-bc -o double_it.bc
+$TOOLS_DIR/llvm-spirv double_it.bc -o double_it.spv
+rm double_it.bc
+```
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.cl b/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.cl
new file mode 100644
index 0000000000..660e4b9b60
--- /dev/null
+++ b/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.cl
@@ -0,0 +1,30 @@
+/*
+                      Data Parallel Control (dpctl)
+
+   Copyright 2020-2022 Intel Corporation
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+      http://www.apache.org/licenses/LICENSE-2.0
+
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License.
+
+===----------------------------------------------------------------------===
+
+ \file
+ This file implements a sample OpenCL kernel to use in this example.
+
+===----------------------------------------------------------------------===
+*/
+
+__kernel void double_it(__global int *x, __global int *y) {
+    uint idx = get_global_id(0);
+
+    y[idx] = 2 * x[idx];
+}
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.spv b/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.spv
new file mode 100644
index 0000000000..c193a1f02b
Binary files /dev/null and b/examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.spv differ
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/setup.py b/examples/pybind11/use_dpctl_sycl_kernel/setup.py
new file mode 100644
index 0000000000..8888a9ff0d
--- /dev/null
+++ b/examples/pybind11/use_dpctl_sycl_kernel/setup.py
@@ -0,0 +1,26 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2022 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+from skbuild import setup
+
+setup(
+    name="use_kernel",
+    version="0.0.1",
+    description="an example of SYCL-powered Python package (with pybind11)",
+    author="Intel Scripting",
+    license="Apache 2.0",
+    packages=["use_kernel"],
+)
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py b/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py
new file mode 100644
index 0000000000..a8c96e09c5
--- /dev/null
+++ b/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py
@@ -0,0 +1,69 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2022 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+# coding: utf-8
+
+import os.path
+
+import numpy as np
+import pytest
+import use_kernel as uk
+
+import dpctl
+import dpctl.program as dpm
+import dpctl.tensor as dpt
+
+
+def _get_spv_path():
+    uk_dir = os.path.dirname(os.path.abspath(uk.__file__))
+    proj_dir = os.path.dirname(uk_dir)
+    return os.path.join(proj_dir, "resource", "double_it.spv")
+
+
+def test_spv_file_exists():
+    assert os.path.exists(_get_spv_path())
+
+
+def test_kernel_can_be_found():
+    fn = _get_spv_path()
+    with open(fn, "br") as f:
+        il = f.read()
+    try:
+        q = dpctl.SyclQueue()
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("Could not create default queue")
+    pr = dpm.create_program_from_spirv(q, il, "")
+    assert pr.has_sycl_kernel("double_it")
+
+
+def test_kernel_submit_through_extension():
+    fn = _get_spv_path()
+    with open(fn, "br") as f:
+        il = f.read()
+    try:
+        q = dpctl.SyclQueue()
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("Could not create default queue")
+    pr = dpm.create_program_from_spirv(q, il, "")
+    krn = pr.get_sycl_kernel("double_it")
+    assert krn.num_args == 2
+
+    x = dpt.arange(0, stop=13, step=1, dtype="i4", sycl_queue=q)
+    y = dpt.zeros_like(x)
+
+    uk.submit_custom_kernel(q, krn, x, y, [])
+
+    assert np.array_equal(dpt.asnumpy(y), np.arange(0, 26, step=2, dtype="i4"))
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py
new file mode 100644
index 0000000000..09ac9ffaf0
--- /dev/null
+++ b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py
@@ -0,0 +1,34 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2022 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+# coding: utf-8
+
+from ._use_kernel import submit_custom_kernel
+
+__all__ = [
+    "submit_custom_kernel",
+]
+
+__doc__ = """
+Example pybind11 extension demonstrating binding of dpctl entities to
+SYCL entities.
+
+dpctl provides type casters that bind ``sycl::kernel`` to
+`dpctl.program.SyclKernel`, ``sycl::device`` to `dpctl.SyclDevice`, etc.
+
+Use of these type casters simplifies writing of Python extensions and compile
+then using SYCL C++ compilers, such as Intel(R) oneAPI DPC++ compiler.
+"""
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp
new file mode 100644
index 0000000000..638db20320
--- /dev/null
+++ b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp
@@ -0,0 +1,92 @@
+//==- _example.cpp - Example of Pybind11 extension working with  =---------===//
+//  dpctl Python objects.
+//
+//                      Data Parallel Control (dpctl)
+//
+// Copyright 2020-2022 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file implements Pybind11-generated extension exposing functions that
+/// take dpctl Python objects, such as dpctl.SyclQueue, dpctl.SyclDevice as
+/// arguments.
+///
+//===----------------------------------------------------------------------===//
+
+#include "dpctl4pybind11.hpp"
+#include <CL/sycl.hpp>
+#include <cstdint>
+#include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+#include <vector>
+
+namespace py = pybind11;
+
+void submit_custom_kernel(sycl::queue q,
+                          sycl::kernel krn,
+                          dpctl::tensor::usm_ndarray x,
+                          dpctl::tensor::usm_ndarray y,
+                          const std::vector<sycl::event> &depends = {})
+{
+    if (x.get_ndim() != 1 || !x.is_c_contiguous() || y.get_ndim() != 1 ||
+        !y.is_c_contiguous())
+    {
+        throw py::value_error(
+            "src and dst arguments must be 1D and contiguous.");
+    }
+
+    auto const &api = dpctl::detail::dpctl_capi::get();
+    if (x.get_typenum() != api.UAR_INT32_ || y.get_typenum() != api.UAR_INT32_)
+    {
+        throw py::value_error(
+            "src and dst arguments must have int32 element data types.");
+    }
+
+    size_t n_x = x.get_size();
+    size_t n_y = y.get_size();
+
+    if (n_x != n_y) {
+        throw py::value_error("src and dst arguments must have equal size.");
+    }
+
+    if (!dpctl::utils::queues_are_compatible(q, {x.get_queue(), y.get_queue()}))
+    {
+        throw std::runtime_error(
+            "Execution queue is not compatible with allocation queues");
+    }
+
+    void *x_data = x.get_data<void>();
+    void *y_data = y.get_data<void>();
+
+    sycl::event e = q.submit([&](sycl::handler &cgh) {
+        cgh.depends_on(depends);
+        cgh.set_arg(0, x_data);
+        cgh.set_arg(1, y_data);
+        cgh.parallel_for(sycl::range<1>(n_x), krn);
+    });
+
+    e.wait();
+
+    return;
+}
+
+PYBIND11_MODULE(_use_kernel, m)
+{
+    m.def("submit_custom_kernel", &submit_custom_kernel,
+          "Submit given kernel with arguments (int *, int *) to queue",
+          py::arg("queue"), py::arg("kernel"), py::arg("src"), py::arg("dst"),
+          py::arg("depends") = py::list());
+}
diff --git a/libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h
index dae11bb3c8..97074a5bc6 100644
--- a/libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h
+++ b/libsyclinterface/include/dpctl_sycl_kernel_bundle_interface.h
@@ -117,4 +117,16 @@ bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef,
 DPCTL_API
 void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef);
 
+/*!
+ * @brief Returns a copy of the DPCTLSyclKernelBundleRef object.
+ *
+ * @param    KBRef           DPCTLSyclKernelBundleRef object to be copied.
+ * @return   A new DPCTLSyclKernelBundleRef created by copying the passed in
+ * DPCTLSyclKernelBundleRef object.
+ * @ingroup KernelBundleInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclKernelBundleRef
+DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef);
+
 DPCTL_C_EXTERN_C_END
diff --git a/libsyclinterface/include/dpctl_sycl_kernel_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_interface.h
index 43c5dd7a8c..ad46ecb3f1 100644
--- a/libsyclinterface/include/dpctl_sycl_kernel_interface.h
+++ b/libsyclinterface/include/dpctl_sycl_kernel_interface.h
@@ -26,6 +26,7 @@
 
 #pragma once
 
+#include "Config/dpctl_config.h"
 #include "Support/DllExport.h"
 #include "Support/ExternC.h"
 #include "Support/MemOwnershipAttrs.h"
@@ -62,6 +63,18 @@ size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef);
 DPCTL_API
 void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef);
 
+/*!
+ * @brief Returns a copy of the DPCTLSyclKernelRef object.
+ *
+ * @param    KRef           DPCTLSyclKernelRef object to be copied.
+ * @return   A new DPCTLSyclKernelRef created by copying the passed in
+ * DPCTLSyclKernelRef object.
+ * @ingroup KernelInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclKernelRef
+DPCTLKernel_Copy(__dpctl_keep const DPCTLSyclKernelRef KRef);
+
 /*!
  * !brief Wrapper around
  * `kernel::get_info<info::kernel_device_specific::work_group_size>()`.
@@ -117,7 +130,7 @@ DPCTL_API
 uint32_t
 DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef);
 
-#if 0
+#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_2023_SWITCHOVER
 /*!
  * !brief Wrapper around
  * `kernel::get_info<info::kernel_device_specific::max_sub_group_size>()`.
diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp
index a6e0492421..b196264350 100644
--- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp
+++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp
@@ -739,3 +739,23 @@ void DPCTLKernelBundle_Delete(__dpctl_take DPCTLSyclKernelBundleRef KBRef)
 {
     delete unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
 }
+
+__dpctl_give DPCTLSyclKernelBundleRef
+DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef)
+{
+    auto Bundle = unwrap<kernel_bundle<bundle_state::executable>>(KBRef);
+    if (!Bundle) {
+        error_handler(
+            "Cannot copy DPCTLSyclKernelBundleRef as input is a nullptr",
+            __FILE__, __func__, __LINE__);
+        return nullptr;
+    }
+    try {
+        auto CopiedBundle =
+            new kernel_bundle<bundle_state::executable>(*Bundle);
+        return wrap<kernel_bundle<bundle_state::executable>>(CopiedBundle);
+    } catch (std::exception const &e) {
+        error_handler(e, __FILE__, __func__, __LINE__);
+        return nullptr;
+    }
+}
diff --git a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp
index cdf009e5b4..95a800a34e 100644
--- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp
+++ b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp
@@ -25,6 +25,7 @@
 //===----------------------------------------------------------------------===//
 
 #include "dpctl_sycl_kernel_interface.h"
+#include "Config/dpctl_config.h"
 #include "dpctl_error_handlers.h"
 #include "dpctl_string_utils.hpp"
 #include "dpctl_sycl_type_casters.hpp"
@@ -57,6 +58,24 @@ void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef)
     delete unwrap<kernel>(KRef);
 }
 
+__dpctl_give DPCTLSyclKernelRef
+DPCTLKernel_Copy(__dpctl_keep const DPCTLSyclKernelRef KRef)
+{
+    auto Kernel = unwrap<kernel>(KRef);
+    if (!Kernel) {
+        error_handler("Cannot copy DPCTLSyclKernelRef as input is a nullptr",
+                      __FILE__, __func__, __LINE__);
+        return nullptr;
+    }
+    try {
+        auto CopiedKernel = new kernel(*Kernel);
+        return wrap<kernel>(CopiedKernel);
+    } catch (std::exception const &e) {
+        error_handler(e, __FILE__, __func__, __LINE__);
+        return nullptr;
+    }
+}
+
 size_t DPCTLKernel_GetWorkGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
 {
     if (!KRef) {
@@ -142,7 +161,7 @@ DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef)
     return static_cast<uint32_t>(v);
 }
 
-#if 0
+#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_2023_SWITCHOVER
 // commented out due to bug in DPC++ runtime, get_info for max_sub_group_size
 // exported by libsycl has different, not SPEC-compliant signature
 uint32_t
@@ -161,8 +180,9 @@ DPCTLKernel_GetMaxSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
                       __FILE__, __func__, __LINE__);
         return 0;
     }
-    auto v = sycl_kern
-      ->get_info<info::kernel_device_specific::max_sub_group_size>(devs[0]);
+    auto v =
+        sycl_kern->get_info<info::kernel_device_specific::max_sub_group_size>(
+            devs[0]);
     return v;
 }
 #endif
diff --git a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp
index 3560100be6..6383b730a0 100644
--- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp
+++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp
@@ -102,6 +102,27 @@ TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirv)
     ASSERT_FALSE(DPCTLKernelBundle_HasKernel(KBRef, nullptr));
 }
 
+TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCopy)
+{
+    DPCTLSyclKernelBundleRef Copied_KBRef = nullptr;
+    ASSERT_TRUE(KBRef != nullptr);
+
+    EXPECT_NO_FATAL_FAILURE(Copied_KBRef = DPCTLKernelBundle_Copy(KBRef));
+    ASSERT_TRUE(DPCTLKernelBundle_HasKernel(Copied_KBRef, "add"));
+    ASSERT_TRUE(DPCTLKernelBundle_HasKernel(Copied_KBRef, "axpy"));
+
+    EXPECT_NO_FATAL_FAILURE(DPCTLKernelBundle_Delete(Copied_KBRef));
+}
+
+TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCopyNullArgument)
+{
+    DPCTLSyclKernelBundleRef Null_KBRef = nullptr;
+    DPCTLSyclKernelBundleRef Copied_KBRef = nullptr;
+
+    EXPECT_NO_FATAL_FAILURE(Copied_KBRef = DPCTLKernelBundle_Copy(Null_KBRef));
+    ASSERT_TRUE(Copied_KBRef == nullptr);
+}
+
 TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirvNull)
 {
     DPCTLSyclContextRef Null_CRef = nullptr;
diff --git a/libsyclinterface/tests/test_sycl_kernel_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_interface.cpp
index a45eae034f..b679fe09d5 100644
--- a/libsyclinterface/tests/test_sycl_kernel_interface.cpp
+++ b/libsyclinterface/tests/test_sycl_kernel_interface.cpp
@@ -24,6 +24,7 @@
 ///
 //===----------------------------------------------------------------------===//
 
+#include "Config/dpctl_config.h"
 #include "dpctl_sycl_context_interface.h"
 #include "dpctl_sycl_device_interface.h"
 #include "dpctl_sycl_device_selector_interface.h"
@@ -99,11 +100,18 @@ struct TestDPCTLSyclKernelInterface
 
 TEST_P(TestDPCTLSyclKernelInterface, CheckGetNumArgs)
 {
-
     ASSERT_EQ(DPCTLKernel_GetNumArgs(AddKRef), 3ul);
     ASSERT_EQ(DPCTLKernel_GetNumArgs(AxpyKRef), 4ul);
 }
 
+TEST_P(TestDPCTLSyclKernelInterface, CheckCopy)
+{
+    DPCTLSyclKernelRef Copied_KRef = nullptr;
+    EXPECT_NO_FATAL_FAILURE(Copied_KRef = DPCTLKernel_Copy(AddKRef));
+    ASSERT_EQ(DPCTLKernel_GetNumArgs(Copied_KRef), 3ul);
+    EXPECT_NO_FATAL_FAILURE(DPCTLKernel_Delete(Copied_KRef));
+}
+
 TEST_P(TestDPCTLSyclKernelInterface, CheckGetWorkGroupSize)
 {
 
@@ -159,7 +167,7 @@ TEST_P(TestDPCTLSyclKernelInterface, CheckGetMaxNumSubGroups)
     ASSERT_TRUE(axpy_mnsg != 0);
 }
 
-/*
+#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_2023_SWITCHOVER
 TEST_P(TestDPCTLSyclKernelInterface, CheckGetMaxSubGroupSize)
 {
 
@@ -172,7 +180,7 @@ TEST_P(TestDPCTLSyclKernelInterface, CheckGetMaxSubGroupSize)
     ASSERT_TRUE(add_msg_sz != 0);
     ASSERT_TRUE(axpy_msg_sz != 0);
 }
-*/
+#endif
 
 TEST_P(TestDPCTLSyclKernelInterface, CheckGetCompileNumSubGroups)
 {
@@ -215,6 +223,11 @@ TEST_F(TestDPCTLSyclKernelNullArgs, CheckNumArgsNullKRef)
     ASSERT_EQ(DPCTLKernel_GetNumArgs(Null_KRef), -1);
 }
 
+TEST_F(TestDPCTLSyclKernelNullArgs, CheckCopyNullKRef)
+{
+    ASSERT_TRUE(DPCTLKernel_Copy(Null_KRef) == nullptr);
+}
+
 TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetWorkGroupsSizeNullKRef)
 {
     DPCTLSyclKernelRef NullKRef = nullptr;
@@ -244,14 +257,14 @@ TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetMaxNumSubGroupsNullKRef)
     ASSERT_EQ(DPCTLKernel_GetMaxNumSubGroups(NullKRef), 0);
 }
 
-/*
+#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_2023_SWITCHOVER
 TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetMaxSubGroupSizeNullKRef)
 {
     DPCTLSyclKernelRef NullKRef = nullptr;
 
     ASSERT_EQ(DPCTLKernel_GetMaxSubGroupSize(NullKRef), 0);
 }
-*/
+#endif
 
 TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetCompileNumSubGroupsNullKRef)
 {