Refactoring of program support in dpctl and SyclInterface

Removed uses of `sycl::program` throughout the code.

In SyclInterface replace DPCTLSyclProgramRef with DPCTLSyclKernelBundleRef
which is reference to `sycl::kernel_bundle<sycl::bundle_state::executable>`.

Functions `DPCTLProgram_*` were replaced with `DCPTLKernelBundle_*`.
Functions to create program now take both context and device.

Tests were modified.

dpctl.SyclProgram stays with this name, but it now encapsulates
DPCTLSyclKernelBundleRef instead of removed DPCTLSyclProgramRef.

OpenCL functions are no longer directly used (instead of used via loader,
like in the case of level-zero backend), hence removed linkage to OpenCL
library.

Author: oleksandr-pavlyk

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
 
@@ -305,21 +305,23 @@ cdef extern from "syclinterface/dpctl_sycl_context_interface.h":
 
 
 cdef extern from "syclinterface/dpctl_sycl_program_interface.h":
-    cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromSpirv(
+    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
@@ -41,18 +41,18 @@ 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 from
+        using SYCL interoperability layer for OpenCL and Level-Zero backends.
 
         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.
     '''
-    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

@@ -31,14 +31,15 @@ from dpctl._backend cimport (  # noqa: E211, E402
     DPCTLCString_Delete,
     DPCTLKernel_Delete,
     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__ = [
@@ -50,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
 
@@ -105,33 +106,35 @@ cdef class SyclProgram:
     """
 
     @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), kernel_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)
 
@@ -140,9 +143,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
@@ -153,33 +157,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
@@ -190,20 +198,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)

libsyclinterface/CMakeLists.txt

@@ -71,6 +71,23 @@ if(DPCTL_ENABLE_L0_PROGRAM_CREATION)
     endif()
 endif()
 
+if (UNIX)
+  find_library(PI_OPENCL_LIB
+      NAMES pi_opencl
+      HINTS ${IntelSycl_LIBRARY_DIR}
+  )
+  find_program(READELF_PROG readelf)
+  find_program(GREP_PROG grep)
+  execute_process(
+    COMMAND ${READELF_PROG} -d ${PI_OPENCL_LIB}
+    COMMAND ${GREP_PROG} OpenCL
+    COMMAND ${GREP_PROG} -Po "libOpenCL[^\]]*"
+    OUTPUT_VARIABLE LIBCL_LOADER_FILENAME
+    OUTPUT_STRIP_TRAILING_WHITESPACE
+    ERROR_STRIP_TRAILING_WHITESPACE
+  )
+endif()
+
 configure_file(
     ${CMAKE_CURRENT_SOURCE_DIR}/include/Config/dpctl_config.h.in
     ${CMAKE_CURRENT_SOURCE_DIR}/include/Config/dpctl_config.h
@@ -193,7 +210,6 @@ target_include_directories(DPCTLSyclInterface
 )
 target_link_libraries(DPCTLSyclInterface
   PRIVATE ${IntelSycl_SYCL_LIBRARY}
-  PRIVATE ${IntelSycl_OPENCL_LIBRARY}
 )
 
 if(DPCTL_ENABLE_GLOG)

libsyclinterface/helper/include/dpctl_dynamic_lib_helper.h

@@ -75,13 +75,13 @@ class DynamicLibHelper final
         void *sym = dlsym(_handle, symName);
         char *error = dlerror();
 
-        if (NULL != error) {
+        if (nullptr != error) {
             return nullptr;
         }
 #elif defined(_WIN32) || defined(_WIN64)
         void *sym = (void *)GetProcAddress((HMODULE)_handle, symName);
 
-        if (NULL == sym) {
+        if (nullptr == sym) {
             return nullptr;
         }
 #endif

libsyclinterface/include/Config/dpctl_config.h.in

@@ -26,9 +26,11 @@
 #pragma once
 
 /* Defined when dpctl was built with level zero program creation enabled. */
-#cmakedefine DPCTL_ENABLE_L0_PROGRAM_CREATION @DPCTL_ENABLE_L0_PROGRAM_CREATION@
+#cmakedefine DPCTL_ENABLE_L0_PROGRAM_CREATION                                  \
+    @DPCTL_ENABLE_L0_PROGRAM_CREATION @
 
 /* The DPCPP version used to build dpctl */
 #define DPCTL_DPCPP_VERSION "@IntelSycl_VERSION@"
 
 #define DPCTL_LIBZE_LOADER_FILENAME "@LIBZE_LOADER_FILENAME@"
+#define DPCTL_LIBCL_LOADER_FILENAME "@LIBCL_LOADER_FILENAME@"