Feature/refactor program

.github/workflows/generate-coverage.yaml

@@ -85,7 +85,7 @@ jobs:
         shell: bash -l {0}
         run: |
           source /opt/intel/oneapi/setvars.sh
-          python scripts/gen_coverage.py
+          SYCL_ENABLE_HOST_DEVICE=1 python scripts/gen_coverage.py
 
       - name: Install coverall dependencies
         shell: bash -l {0}

dpctl/_backend.pxd

@@ -120,7 +120,7 @@ cdef extern from "syclinterface/dpctl_sycl_types.h":
     cdef struct DPCTLOpaqueSyclEvent
     cdef struct DPCTLOpaqueSyclKernel
     cdef struct DPCTLOpaqueSyclPlatform
-    cdef struct DPCTLOpaqueSyclProgram
+    cdef struct DPCTLOpaqueSyclKernelBundle
     cdef struct DPCTLOpaqueSyclQueue
     cdef struct DPCTLOpaqueSyclUSM
 
@@ -130,7 +130,7 @@ cdef extern from "syclinterface/dpctl_sycl_types.h":
     ctypedef DPCTLOpaqueSyclEvent          *DPCTLSyclEventRef
     ctypedef DPCTLOpaqueSyclKernel         *DPCTLSyclKernelRef
     ctypedef DPCTLOpaqueSyclPlatform       *DPCTLSyclPlatformRef
-    ctypedef DPCTLOpaqueSyclProgram        *DPCTLSyclProgramRef
+    ctypedef DPCTLOpaqueSyclKernelBundle   *DPCTLSyclKernelBundleRef
     ctypedef DPCTLOpaqueSyclQueue          *DPCTLSyclQueueRef
     ctypedef DPCTLOpaqueSyclUSM            *DPCTLSyclUSMRef
 
@@ -250,7 +250,6 @@ cdef extern from "syclinterface/dpctl_sycl_event_interface.h":
 
 
 cdef extern from "syclinterface/dpctl_sycl_kernel_interface.h":
-    cdef const char* DPCTLKernel_GetFunctionName(const DPCTLSyclKernelRef KRef)
     cdef size_t DPCTLKernel_GetNumArgs(const DPCTLSyclKernelRef KRef)
     cdef void DPCTLKernel_Delete(DPCTLSyclKernelRef KRef)
 
@@ -305,22 +304,24 @@ cdef extern from "syclinterface/dpctl_sycl_context_interface.h":
     cdef void DPCTLContext_Delete(DPCTLSyclContextRef CtxRef)
 
 
-cdef extern from "syclinterface/dpctl_sycl_program_interface.h":
-    cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromSpirv(
+cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h":
+    cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSpirv(
         const DPCTLSyclContextRef Ctx,
+        const DPCTLSyclDeviceRef Dev,
         const void *IL,
         size_t Length,
         const char *CompileOpts)
-    cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSource(
+    cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource(
         const DPCTLSyclContextRef Ctx,
+        const DPCTLSyclDeviceRef Dev,
         const char *Source,
         const char *CompileOpts)
-    cdef DPCTLSyclKernelRef DPCTLProgram_GetKernel(
-        DPCTLSyclProgramRef PRef,
+    cdef DPCTLSyclKernelRef DPCTLKernelBundle_GetKernel(
+        DPCTLSyclKernelBundleRef KBRef,
         const char *KernelName)
-    cdef bool DPCTLProgram_HasKernel(DPCTLSyclProgramRef PRef,
+    cdef bool DPCTLKernelBundle_HasKernel(DPCTLSyclKernelBundleRef KBRef,
                                      const char *KernelName)
-    cdef void DPCTLProgram_Delete(DPCTLSyclProgramRef PRef)
+    cdef void DPCTLKernelBundle_Delete(DPCTLSyclKernelBundleRef KBRef)
 
 
 cdef extern from "syclinterface/dpctl_sycl_queue_interface.h":

dpctl/program/_program.pxd

@@ -22,7 +22,7 @@
 """
 
 
-from .._backend cimport DPCTLSyclKernelRef, DPCTLSyclProgramRef
+from .._backend cimport DPCTLSyclKernelBundleRef, DPCTLSyclKernelRef
 from .._sycl_context cimport SyclContext
 from .._sycl_device cimport SyclDevice
 from .._sycl_queue cimport SyclQueue
@@ -33,26 +33,26 @@ cdef class SyclKernel:
         kernel.
     '''
     cdef DPCTLSyclKernelRef _kernel_ref
-    cdef const char *_function_name
+    cdef str _function_name
     cdef DPCTLSyclKernelRef get_kernel_ref (self)
 
     @staticmethod
-    cdef SyclKernel _create (DPCTLSyclKernelRef kref)
+    cdef SyclKernel _create (DPCTLSyclKernelRef kref, str name)
 
 
 cdef class SyclProgram:
-    ''' Wraps a sycl::program object created from an OpenCL interoperability
-        program.
+    ''' Wraps a sycl::kernel_bundle<sycl::bundle_state::executable> object created by
+        using SYCL interoperability layer for OpenCL and Level-Zero backends.
 
-        SyclProgram exposes the C API from dpctl_sycl_program_interface.h. A
+        SyclProgram exposes the C API from dpctl_sycl_kernel_bundle_interface.h. A
         SyclProgram can be created from either a source string or a SPIR-V
         binary file.
     '''
-    cdef DPCTLSyclProgramRef _program_ref
+    cdef DPCTLSyclKernelBundleRef _program_ref
 
     @staticmethod
-    cdef  SyclProgram _create (DPCTLSyclProgramRef pref)
-    cdef  DPCTLSyclProgramRef get_program_ref (self)
+    cdef  SyclProgram _create (DPCTLSyclKernelBundleRef pref)
+    cdef  DPCTLSyclKernelBundleRef get_program_ref (self)
     cpdef SyclKernel get_sycl_kernel(self, str kernel_name)
 
 

dpctl/program/_program.pyx

@@ -30,16 +30,16 @@ cimport cython.array
 from dpctl._backend cimport (  # noqa: E211, E402
     DPCTLCString_Delete,
     DPCTLKernel_Delete,
-    DPCTLKernel_GetFunctionName,
     DPCTLKernel_GetNumArgs,
-    DPCTLProgram_CreateFromOCLSource,
-    DPCTLProgram_CreateFromSpirv,
-    DPCTLProgram_Delete,
-    DPCTLProgram_GetKernel,
-    DPCTLProgram_HasKernel,
+    DPCTLKernelBundle_CreateFromOCLSource,
+    DPCTLKernelBundle_CreateFromSpirv,
+    DPCTLKernelBundle_Delete,
+    DPCTLKernelBundle_GetKernel,
+    DPCTLKernelBundle_HasKernel,
     DPCTLSyclContextRef,
+    DPCTLSyclDeviceRef,
+    DPCTLSyclKernelBundleRef,
     DPCTLSyclKernelRef,
-    DPCTLSyclProgramRef,
 )
 
 __all__ = [
@@ -51,8 +51,8 @@ __all__ = [
 ]
 
 cdef class SyclProgramCompilationError(Exception):
-    """This exception is raised when a ``sycl::program`` could not be built from
-       either a SPIR-V binary file or a string source.
+    """This exception is raised when a ``sycl::kernel_bundle`` could not be
+       built from either a SPIR-V binary file or a string source.
     """
     pass
 
@@ -61,20 +61,19 @@ cdef class SyclKernel:
     """
     """
     @staticmethod
-    cdef SyclKernel _create(DPCTLSyclKernelRef kref):
+    cdef SyclKernel _create(DPCTLSyclKernelRef kref, str name):
         cdef SyclKernel ret = SyclKernel.__new__(SyclKernel)
         ret._kernel_ref = kref
-        ret._function_name = DPCTLKernel_GetFunctionName(kref)
+        ret._function_name = name
         return ret
 
     def __dealloc__(self):
         DPCTLKernel_Delete(self._kernel_ref)
-        DPCTLCString_Delete(self._function_name)
 
     def get_function_name(self):
         """ Returns the name of the ``sycl::kernel`` function.
         """
-        return self._function_name.decode()
+        return self._function_name
 
     def get_num_args(self):
         """ Returns the number of arguments for this kernel function.
@@ -98,42 +97,45 @@ cdef class SyclKernel:
 
 
 cdef class SyclProgram:
-    """ Wraps a ``sycl::program`` object created from an OpenCL interoperability
-        program.
+    """ Wraps a ``sycl::kernel_bundle<sycl::bundle_state::executable>`` object
+    created using SYCL interoperability layer with underlying backends. Only the
+    OpenCL and Level-Zero backends are currently supported.
 
-        SyclProgram exposes the C API from ``dpctl_sycl_program_interface.h``. A
-        SyclProgram can be created from either a source string or a SPIR-V
-        binary file.
+    SyclProgram exposes the C API from ``dpctl_sycl_kernel_bundle_interface.h``.
+    A SyclProgram can be created from either a source string or a SPIR-V
+    binary file.
     """
 
     @staticmethod
-    cdef SyclProgram _create(DPCTLSyclProgramRef pref):
+    cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef):
         cdef SyclProgram ret = SyclProgram.__new__(SyclProgram)
-        ret._program_ref = pref
+        ret._program_ref = KBRef
         return ret
 
     def __dealloc__(self):
-        DPCTLProgram_Delete(self._program_ref)
+        DPCTLKernelBundle_Delete(self._program_ref)
 
-    cdef DPCTLSyclProgramRef get_program_ref(self):
+    cdef DPCTLSyclKernelBundleRef get_program_ref(self):
         return self._program_ref
 
     cpdef SyclKernel get_sycl_kernel(self, str kernel_name):
         name = kernel_name.encode('utf8')
-        return SyclKernel._create(DPCTLProgram_GetKernel(self._program_ref,
-                                                         name))
+        return SyclKernel._create(
+            DPCTLKernelBundle_GetKernel(self._program_ref, name),
+            kernel_name
+        )
 
     def has_sycl_kernel(self, str kernel_name):
         name = kernel_name.encode('utf8')
-        return DPCTLProgram_HasKernel(self._program_ref, name)
+        return DPCTLKernelBundle_HasKernel(self._program_ref, name)
 
     def addressof_ref(self):
-        """Returns the address of the C API DPCTLSyclProgramRef pointer
+        """Returns the address of the C API DPCTLSyclKernelBundleRef pointer
         as a long.
 
         Returns:
-            The address of the ``DPCTLSyclProgramRef`` pointer used to create
-            this :class:`dpctl.SyclProgram` object cast to a ``size_t``.
+            The address of the ``DPCTLSyclKernelBundleRef`` pointer used to
+            create this :class:`dpctl.SyclProgram` object cast to a ``size_t``.
         """
         return int(<size_t>self._program_ref)
 
@@ -142,9 +144,10 @@ cpdef create_program_from_source(SyclQueue q, unicode src, unicode copts=""):
     """
         Creates a Sycl interoperability program from an OpenCL source string.
 
-        We use the ``DPCTLProgram_CreateFromOCLSource()`` C API function to
-        create a ``sycl::program`` from an OpenCL source program that can
-        contain multiple kernels. Note currently only supported for OpenCL.
+        We use the ``DPCTLKernelBundle_CreateFromOCLSource()`` C API function
+        to create a ``sycl::kernel_bundle<sycl::bundle_state::executable>``
+        from an OpenCL source program that can contain multiple kernels.
+        Note: This function is currently only supported for the OpenCL backend.
 
         Parameters:
             q (SyclQueue)   : The :class:`SyclQueue` for which the
@@ -155,33 +158,37 @@ cpdef create_program_from_source(SyclQueue q, unicode src, unicode copts=""):
 
         Returns:
             program (SyclProgram): A :class:`SyclProgram` object wrapping the
-            ``sycl::program`` returned by the C API.
+            ``sycl::kernel_bundle<sycl::bundle_state::executable>`` returned
+            by the C API.
 
         Raises:
-            SyclProgramCompilationError: If a SYCL program could not be created.
+            SyclProgramCompilationError: If a SYCL kernel bundle could not be
+            created.
     """
 
-    cdef DPCTLSyclProgramRef Pref
+    cdef DPCTLSyclKernelBundleRef KBref
     cdef bytes bSrc = src.encode('utf8')
     cdef bytes bCOpts = copts.encode('utf8')
     cdef const char *Src = <const char*>bSrc
     cdef const char *COpts = <const char*>bCOpts
     cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
-    Pref = DPCTLProgram_CreateFromOCLSource(CRef, Src, COpts)
+    cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()
+    KBref = DPCTLKernelBundle_CreateFromOCLSource(CRef, DRef, Src, COpts)
 
-    if Pref is NULL:
+    if KBref is NULL:
         raise SyclProgramCompilationError()
 
-    return SyclProgram._create(Pref)
+    return SyclProgram._create(KBref)
 
 
 cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,
                                 unicode copts=""):
     """
         Creates a Sycl interoperability program from an SPIR-V binary.
 
-        We use the ``DPCTLProgram_CreateFromOCLSpirv()`` C API function to
-        create a ``sycl::program`` object from an compiled SPIR-V binary file.
+        We use the ``DPCTLKernelBundle_CreateFromOCLSpirv()`` C API function to
+        create a ``sycl::kernel_bundle<sycl::bundle_state::executable>`` object
+        from an compiled SPIR-V binary file.
 
         Parameters:
             q (SyclQueue): The :class:`SyclQueue` for which the
@@ -192,20 +199,25 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,
 
         Returns:
             program (SyclProgram): A :class:`SyclProgram` object wrapping the
-            ``sycl::program`` returned by the C API.
+            ``sycl::kernel_bundle<sycl::bundle_state::executable>`` returned by
+            the C API.
 
         Raises:
-            SyclProgramCompilationError: If a SYCL program could not be created.
+            SyclProgramCompilationError: If a SYCL kernel bundle could not be
+            created.
     """
 
-    cdef DPCTLSyclProgramRef Pref
+    cdef DPCTLSyclKernelBundleRef KBref
     cdef const unsigned char *dIL = &IL[0]
     cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
+    cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()
     cdef size_t length = IL.shape[0]
     cdef bytes bCOpts = copts.encode('utf8')
     cdef const char *COpts = <const char*>bCOpts
-    Pref = DPCTLProgram_CreateFromSpirv(CRef, <const void*>dIL, length, COpts)
-    if Pref is NULL:
+    KBref = DPCTLKernelBundle_CreateFromSpirv(
+        CRef, DRef, <const void*>dIL, length, COpts
+    )
+    if KBref is NULL:
         raise SyclProgramCompilationError()
 
-    return SyclProgram._create(Pref)
+    return SyclProgram._create(KBref)