Support compilation from SYCL source code

Author: sommerlukas

dpctl/_backend.pxd

@@ -287,9 +287,12 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h":
                                         _peer_access PT)
     cdef void DPCTLDevice_EnablePeerAccess(const DPCTLSyclDeviceRef DRef,
                                            const DPCTLSyclDeviceRef PDRef)
-
     cdef void DPCTLDevice_DisablePeerAccess(const DPCTLSyclDeviceRef DRef,
                                             const DPCTLSyclDeviceRef PDRef)
+    cdef bool DPCTLDevice_CanCompileSPIRV(const DPCTLSyclDeviceRef DRef)
+    cdef bool DPCTLDevice_CanCompileOpenCL(const DPCTLSyclDeviceRef DRef)
+    cdef bool DPCTLDevice_CanCompileSYCL(const DPCTLSyclDeviceRef DRef)
+
 
 cdef extern from "syclinterface/dpctl_sycl_device_manager.h":
     cdef DPCTLDeviceVectorRef DPCTLDeviceVector_CreateFromArray(
@@ -452,6 +455,43 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h":
     cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_Copy(
         const DPCTLSyclKernelBundleRef KBRef)
 
+    cdef struct DPCTLBuildOptionList
+    cdef struct DPCTLKernelNameList
+    cdef struct DPCTLVirtualHeaderList
+    ctypedef DPCTLBuildOptionList* DPCTLBuildOptionListRef
+    ctypedef DPCTLKernelNameList* DPCTLKernelNameListRef
+    ctypedef DPCTLVirtualHeaderList* DPCTLVirtualHeaderListRef
+
+    cdef DPCTLBuildOptionListRef DPCTLBuildOptionList_Create()
+    cdef void DPCTLBuildOptionList_Delete(DPCTLBuildOptionListRef Ref)
+    cdef void DPCTLBuildOptionList_Append(DPCTLBuildOptionListRef Ref,
+                                 const char *Option)
+
+    cdef DPCTLKernelNameListRef DPCTLKernelNameList_Create()
+    cdef void DPCTLKernelNameList_Delete(DPCTLKernelNameListRef Ref)
+    cdef void DPCTLKernelNameList_Append(DPCTLKernelNameListRef Ref,
+                                const char *Option)
+
+    cdef DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create()
+    cdef void DPCTLVirtualHeaderList_Delete(DPCTLVirtualHeaderListRef Ref)
+    cdef void DPCTLVirtualHeaderList_Append(DPCTLVirtualHeaderListRef Ref,
+                                   const char *Name,
+                                   const char *Content)
+
+    cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource(
+        const DPCTLSyclContextRef Ctx,
+        const DPCTLSyclDeviceRef Dev,
+        const char *Source,
+        DPCTLVirtualHeaderListRef Headers,
+        DPCTLKernelNameListRef Names,
+        DPCTLBuildOptionListRef BuildOptions)
+
+    cdef DPCTLSyclKernelRef DPCTLKernelBundle_GetSyclKernel(DPCTLSyclKernelBundleRef KBRef,
+                                const char *KernelName)
+
+    cdef bool DPCTLKernelBundle_HasSyclKernel(DPCTLSyclKernelBundleRef KBRef,
+                                     const char *KernelName);
+
 
 cdef extern from "syclinterface/dpctl_sycl_queue_interface.h":
     ctypedef struct _md_local_accessor "MDLocalAccessor":

dpctl/_sycl_device.pxd

@@ -61,3 +61,4 @@ cdef public api class SyclDevice(_SyclDevice) [
     cdef int get_overall_ordinal(self)
     cdef int get_backend_ordinal(self)
     cdef int get_backend_and_device_type_ordinal(self)
+    cpdef bint can_compile(self, str language)

dpctl/_sycl_device.pyx

@@ -26,6 +26,9 @@ from ._backend cimport (  # noqa: E211
     DPCTLDefaultSelector_Create,
     DPCTLDevice_AreEq,
     DPCTLDevice_CanAccessPeer,
+    DPCTLDevice_CanCompileOpenCL,
+    DPCTLDevice_CanCompileSPIRV,
+    DPCTLDevice_CanCompileSYCL,
     DPCTLDevice_Copy,
     DPCTLDevice_CreateFromSelector,
     DPCTLDevice_CreateSubDevicesByAffinity,
@@ -2367,6 +2370,35 @@ cdef class SyclDevice(_SyclDevice):
             raise ValueError("device could not be found")
         return dev_id
 
+    cpdef bint can_compile(self, str language):
+        """
+        Check whether it is possible to create an executable kernel_bundle
+        for this device from the given source language.
+
+        Parameters:
+            language
+                Input language. Possible values are "spirv" for SPIR-V binary
+                files, "opencl" for OpenCL C device code and "sycl" for SYCL
+                device code.
+
+        Returns:
+            bool:
+                True if compilation is supported, False otherwise.
+
+        Raises:
+            ValueError:
+                If an unknown source language is used.
+        """
+        if language == "spirv" or language == "spv":
+            return DPCTLDevice_CanCompileSYCL(self._device_ref)
+        if language == "opencl" or language == "ocl":
+            return DPCTLDevice_CanCompileOpenCL(self._device_ref)
+        if language == "sycl":
+            return DPCTLDevice_CanCompileSYCL(self._device_ref)
+
+        raise ValueError(f"Unknown source language {language}")
+
+
 
 cdef api DPCTLSyclDeviceRef SyclDevice_GetDeviceRef(SyclDevice dev):
     """

dpctl/program/__init__.py

@@ -26,11 +26,13 @@
     SyclProgramCompilationError,
     create_program_from_source,
     create_program_from_spirv,
+    create_program_from_sycl_source,
 )
 
 __all__ = [
     "create_program_from_source",
     "create_program_from_spirv",
+    "create_program_from_sycl_source",
     "SyclKernel",
     "SyclProgram",
     "SyclProgramCompilationError",

dpctl/program/_program.pxd

@@ -49,13 +49,17 @@ cdef api class SyclProgram [object PySyclProgramObject, type PySyclProgramType]:
     binary file.
     """
     cdef DPCTLSyclKernelBundleRef _program_ref
+    cdef bint _is_sycl_source
 
     @staticmethod
-    cdef  SyclProgram _create (DPCTLSyclKernelBundleRef pref)
+    cdef  SyclProgram _create (DPCTLSyclKernelBundleRef pref, bint _is_sycl_source)
     cdef  DPCTLSyclKernelBundleRef get_program_ref (self)
     cpdef SyclKernel get_sycl_kernel(self, str kernel_name)
 
 
 cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*)
 cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL,
                                  unicode copts=*)
+cpdef create_program_from_sycl_source(SyclQueue q, unicode source,
+                                      list headers=*, list registered_names=*,
+                                      list copts=*)

dpctl/program/_program.pyx

@@ -28,6 +28,10 @@ a OpenCL source string or a SPIR-V binary file.
 from libc.stdint cimport uint32_t
 
 from dpctl._backend cimport (  # noqa: E211, E402;
+    DPCTLBuildOptionList_Append,
+    DPCTLBuildOptionList_Create,
+    DPCTLBuildOptionList_Delete,
+    DPCTLBuildOptionListRef,
     DPCTLKernel_Copy,
     DPCTLKernel_Delete,
     DPCTLKernel_GetCompileNumSubGroups,
@@ -41,13 +45,24 @@ from dpctl._backend cimport (  # noqa: E211, E402;
     DPCTLKernelBundle_Copy,
     DPCTLKernelBundle_CreateFromOCLSource,
     DPCTLKernelBundle_CreateFromSpirv,
+    DPCTLKernelBundle_CreateFromSYCLSource,
     DPCTLKernelBundle_Delete,
     DPCTLKernelBundle_GetKernel,
+    DPCTLKernelBundle_GetSyclKernel,
     DPCTLKernelBundle_HasKernel,
+    DPCTLKernelBundle_HasSyclKernel,
+    DPCTLKernelNameList_Append,
+    DPCTLKernelNameList_Create,
+    DPCTLKernelNameList_Delete,
+    DPCTLKernelNameListRef,
     DPCTLSyclContextRef,
     DPCTLSyclDeviceRef,
     DPCTLSyclKernelBundleRef,
     DPCTLSyclKernelRef,
+    DPCTLVirtualHeaderList_Append,
+    DPCTLVirtualHeaderList_Create,
+    DPCTLVirtualHeaderList_Delete,
+    DPCTLVirtualHeaderListRef,
 )
 
 __all__ = [
@@ -196,9 +211,10 @@ cdef class SyclProgram:
     """
 
     @staticmethod
-    cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef):
+    cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef, bint is_sycl_source):
         cdef SyclProgram ret = SyclProgram.__new__(SyclProgram)
         ret._program_ref = KBRef
+        ret._is_sycl_source = is_sycl_source
         return ret
 
     def __dealloc__(self):
@@ -209,13 +225,19 @@ cdef class SyclProgram:
 
     cpdef SyclKernel get_sycl_kernel(self, str kernel_name):
         name = kernel_name.encode("utf8")
+        if self._is_sycl_source:
+            return SyclKernel._create(
+                    DPCTLKernelBundle_GetSyclKernel(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")
+        if self._is_sycl_source:
+            return DPCTLKernelBundle_HasSyclKernel(self._program_ref, name)
         return DPCTLKernelBundle_HasKernel(self._program_ref, name)
 
     def addressof_ref(self):
@@ -271,7 +293,7 @@ cpdef create_program_from_source(SyclQueue q, str src, str copts=""):
     if KBref is NULL:
         raise SyclProgramCompilationError()
 
-    return SyclProgram._create(KBref)
+    return SyclProgram._create(KBref, False)
 
 
 cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,
@@ -317,7 +339,107 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,
     if KBref is NULL:
         raise SyclProgramCompilationError()
 
-    return SyclProgram._create(KBref)
+    return SyclProgram._create(KBref, False)
+
+
+cpdef create_program_from_sycl_source(SyclQueue q, unicode source, list headers=[], list registered_names=[], list copts=[]):
+    """
+        Creates an executable SYCL kernel_bundle from SYCL source code.
+
+        This uses the DPC++ ``kernel_compiler`` extension to create a
+        ``sycl::kernel_bundle<sycl::bundle_state::executable>`` object from
+        SYCL source code.
+
+        Parameters:
+            q (:class:`dpctl.SyclQueue`)
+                The :class:`dpctl.SyclQueue` for which the
+                :class:`.SyclProgram` is going to be built.
+            source (unicode)
+                SYCL source code string.
+            headers (list)
+                Optional list of virtual headers, where each entry in the list
+                needs to be a tuple of header name and header content. See the
+                documentation of the ``include_files`` property in the DPC++
+                ``kernel_compiler`` extension for more information.
+                Default: []
+            registered_names (list, optional)
+                Optional list of kernel names to register. See the
+                documentation of the ``registered_names`` property in the DPC++
+                ``kernel_compiler`` extension for more information.
+                Default: []
+            copts (list)
+                Optional list of compilation flags that will be used
+                when compiling the program. Default: ``""``.
+
+        Returns:
+            program (:class:`.SyclProgram`)
+                A :class:`.SyclProgram` object wrapping the
+                ``sycl::kernel_bundle<sycl::bundle_state::executable>``
+                returned by the C API.
+
+        Raises:
+            SyclProgramCompilationError
+                If a SYCL kernel bundle could not be created.
+    """
+    cdef DPCTLSyclKernelBundleRef KBref
+    cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
+    cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()
+    cdef bytes bSrc = source.encode('utf8')
+    cdef const char *Src = <const char*>bSrc
+    cdef DPCTLBuildOptionListRef BuildOpts = DPCTLBuildOptionList_Create()
+    cdef bytes bOpt
+    cdef const char* sOpt
+    cdef bytes bName
+    cdef const char* sName
+    cdef bytes bContent
+    cdef const char* sContent
+    for opt in copts:
+        if not isinstance(opt, unicode):
+            DPCTLBuildOptionList_Delete(BuildOpts)
+            raise SyclProgramCompilationError()
+        bOpt = opt.encode('utf8')
+        sOpt = <const char*>bOpt
+        DPCTLBuildOptionList_Append(BuildOpts, sOpt)
+
+    cdef DPCTLKernelNameListRef KernelNames = DPCTLKernelNameList_Create()
+    for name in registered_names:
+        if not isinstance(name, unicode):
+            DPCTLBuildOptionList_Delete(BuildOpts)
+            DPCTLKernelNameList_Delete(KernelNames)
+            raise SyclProgramCompilationError()
+        bName = name.encode('utf8')
+        sName = <const char*>bName
+        DPCTLKernelNameList_Append(KernelNames, sName)
+
+
+    cdef DPCTLVirtualHeaderListRef VirtualHeaders = DPCTLVirtualHeaderList_Create()
+    for name, content in headers:
+        if not isinstance(name, unicode) or not isinstance(content, unicode):
+            DPCTLBuildOptionList_Delete(BuildOpts)
+            DPCTLKernelNameList_Delete(KernelNames)
+            DPCTLVirtualHeaderList_Delete(VirtualHeaders)
+            raise SyclProgramCompilationError()
+        bName = name.encode('utf8')
+        sName = <const char*>bName
+        bContent = content.encode('utf8')
+        sContent = <const char*>bContent
+        DPCTLVirtualHeaderList_Append(VirtualHeaders, sName, sContent)
+
+    KBref = DPCTLKernelBundle_CreateFromSYCLSource(CRef, DRef, Src,
+                                                   VirtualHeaders, KernelNames,
+                                                   BuildOpts)
+
+    if KBref is NULL:
+        DPCTLBuildOptionList_Delete(BuildOpts)
+        DPCTLKernelNameList_Delete(KernelNames)
+        DPCTLVirtualHeaderList_Delete(VirtualHeaders)
+        raise SyclProgramCompilationError()
+
+    DPCTLBuildOptionList_Delete(BuildOpts)
+    DPCTLKernelNameList_Delete(KernelNames)
+    DPCTLVirtualHeaderList_Delete(VirtualHeaders)
+
+    return SyclProgram._create(KBref, True)
 
 
 cdef api DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef(
@@ -336,4 +458,4 @@ cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef):
     reference.
     """
     cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef)
-    return SyclProgram._create(copied_KBRef)
+    return SyclProgram._create(copied_KBRef, False)