Support compilation from SYCL source code

Enable SYCL source compilation, but only for DPC++ versions that actually
support the compilation, based on the __SYCL_COMPILER_VERSION reported.

Uses the correct naming for the property based on DPC++ version,
detected through C++ type traits to check which property actually refers
to a fully defined type.

This commit also works around a bug in DPC++ version 2025.1. The constructor
with no parameter of class `include_files` was only declared, but never
defined. Calling it when creating a SYCL source kernel bundle therefore
leads to references to undefined symbols with DPC++ version 2025.1. This
change works around this issue by calling an alternative constructor,
which is defined in the release.

Signed-off-by: Lukas Sommer <lukas.sommer@codeplay.com>

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,51 @@ 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
+    cdef struct DPCTLKernelBuildLog
+    ctypedef DPCTLBuildOptionList* DPCTLBuildOptionListRef
+    ctypedef DPCTLKernelNameList* DPCTLKernelNameListRef
+    ctypedef DPCTLVirtualHeaderList* DPCTLVirtualHeaderListRef
+    ctypedef DPCTLKernelBuildLog* DPCTLKernelBuildLogRef
+
+    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 DPCTLKernelBuildLogRef DPCTLKernelBuildLog_Create()
+    cdef void DPCTLKernelBuildLog_Delete(DPCTLKernelBuildLogRef Ref)
+    cdef const char *DPCTLKernelBuildLog_Get(DPCTLKernelBuildLogRef)
+
+    cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource(
+        const DPCTLSyclContextRef Ctx,
+        const DPCTLSyclDeviceRef Dev,
+        const char *Source,
+        DPCTLVirtualHeaderListRef Headers,
+        DPCTLKernelNameListRef Names,
+        DPCTLBuildOptionListRef BuildOptions,
+        DPCTLKernelBuildLogRef BuildLog)
+
+    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,
@@ -2363,6 +2366,34 @@ 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_CanCompileSPIRV(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

@@ -29,13 +29,15 @@
     create_kernel_bundle_from_spirv,
     create_program_from_source,
     create_program_from_spirv,
+    create_program_from_sycl_source,
 )
 
 __all__ = [
     "create_kernel_bundle_from_source",
     "create_kernel_bundle_from_spirv",
     "create_program_from_source",
     "create_program_from_spirv",
+    "create_program_from_sycl_source",
     "SyclKernel",
     "SyclKernelBundle",
     "SyclKernelBundleCompilationError",

dpctl/program/_program.pxd

@@ -52,9 +52,11 @@ cdef api class SyclKernelBundle [
     binary file.
     """
     cdef DPCTLSyclKernelBundleRef _kernel_bundle_ref
+    cdef bint _is_sycl_source
 
     @staticmethod
-    cdef  SyclKernelBundle _create (DPCTLSyclKernelBundleRef kbref)
+    cdef  SyclKernelBundle _create (DPCTLSyclKernelBundleRef kbref,
+                                    bint _is_sycl_source)
     cdef  DPCTLSyclKernelBundleRef get_kernel_bundle_ref (self)
     cpdef SyclKernel get_sycl_kernel(self, str kernel_name)
 
@@ -69,3 +71,7 @@ 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_kernel_bundle_from_sycl_source(SyclQueue q, unicode source,
+                                            list headers=*,
+                                            list registered_names=*,
+                                            list copts=*)

dpctl/program/_program.pyx

@@ -31,6 +31,10 @@ from libc.stdint cimport uint32_t
 import warnings
 
 from dpctl._backend cimport (  # noqa: E211, E402;
+    DPCTLBuildOptionList_Append,
+    DPCTLBuildOptionList_Create,
+    DPCTLBuildOptionList_Delete,
+    DPCTLBuildOptionListRef,
     DPCTLKernel_Copy,
     DPCTLKernel_Delete,
     DPCTLKernel_GetCompileNumSubGroups,
@@ -41,16 +45,31 @@ from dpctl._backend cimport (  # noqa: E211, E402;
     DPCTLKernel_GetPreferredWorkGroupSizeMultiple,
     DPCTLKernel_GetPrivateMemSize,
     DPCTLKernel_GetWorkGroupSize,
+    DPCTLKernelBuildLog_Create,
+    DPCTLKernelBuildLog_Delete,
+    DPCTLKernelBuildLog_Get,
+    DPCTLKernelBuildLogRef,
     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__ = [
@@ -199,9 +218,11 @@ cdef class SyclKernelBundle:
     """
 
     @staticmethod
-    cdef SyclKernelBundle _create(DPCTLSyclKernelBundleRef KBRef):
+    cdef SyclKernelBundle _create(DPCTLSyclKernelBundleRef KBRef,
+                                  bint is_sycl_source):
         cdef SyclKernelBundle ret = SyclKernelBundle.__new__(SyclKernelBundle)
         ret._kernel_bundle_ref = KBRef
+        ret._is_sycl_source = is_sycl_source
         return ret
 
     def __dealloc__(self):
@@ -212,13 +233,24 @@ cdef class SyclKernelBundle:
 
     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._kernel_bundle_ref, name
+                    ),
+                    kernel_name
+                )
         return SyclKernel._create(
             DPCTLKernelBundle_GetKernel(self._kernel_bundle_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._kernel_bundle_ref, name
+            )
         return DPCTLKernelBundle_HasKernel(self._kernel_bundle_ref, name)
 
     def addressof_ref(self):
@@ -249,7 +281,7 @@ cdef api SyclKernelBundle SyclKernelBundle_Make(DPCTLSyclKernelBundleRef KBRef):
     reference.
     """
     cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef)
-    return SyclKernelBundle._create(copied_KBRef)
+    return SyclKernelBundle._create(copied_KBRef, False)
 
 
 cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""):
@@ -295,7 +327,7 @@ cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""):
     if KBref is NULL:
         raise SyclKernelBundleCompilationError()
 
-    return SyclKernelBundle._create(KBref)
+    return SyclKernelBundle._create(KBref, False)
 
 
 cpdef create_kernel_bundle_from_spirv(
@@ -342,7 +374,121 @@ cpdef create_kernel_bundle_from_spirv(
     if KBref is NULL:
         raise SyclKernelBundleCompilationError()
 
-    return SyclKernelBundle._create(KBref)
+    return SyclKernelBundle._create(KBref, False)
+
+
+cpdef create_kernel_bundle_from_sycl_source(SyclQueue q,
+                                            unicode source,
+                                            list headers=None,
+                                            list registered_names=None,
+                                            list copts=None):
+    """
+        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:`.SyclKernelBundle` 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:
+            kernel_bundle (:class:`.SyclKernelBundle`)
+                A :class:`.SyclKernelBundle` object wrapping the
+                ``sycl::kernel_bundle<sycl::bundle_state::executable>``
+                returned by the C API.
+
+        Raises:
+            SyclKernelBundleCompilationError
+                If a SYCL kernel bundle could not be created. The exception
+                message contains the build log for more details.
+    """
+    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
+    cdef const char* buildLogContent
+    for opt in copts:
+        if not isinstance(opt, unicode):
+            DPCTLBuildOptionList_Delete(BuildOpts)
+            raise SyclKernelBundleCompilationError()
+        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 SyclKernelBundleCompilationError()
+        bName = name.encode("utf8")
+        sName = <const char*>bName
+        DPCTLKernelNameList_Append(KernelNames, sName)
+
+    cdef DPCTLVirtualHeaderListRef VirtualHeaders
+    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 SyclKernelBundleCompilationError()
+        bName = name.encode("utf8")
+        sName = <const char*>bName
+        bContent = content.encode("utf8")
+        sContent = <const char*>bContent
+        DPCTLVirtualHeaderList_Append(VirtualHeaders, sName, sContent)
+
+    cdef DPCTLKernelBuildLogRef BuildLog
+    BuildLog = DPCTLKernelBuildLog_Create()
+
+    KBref = DPCTLKernelBundle_CreateFromSYCLSource(CRef, DRef, Src,
+                                                   VirtualHeaders, KernelNames,
+                                                   BuildOpts, BuildLog)
+
+    if KBref is NULL:
+        buildLogContent = DPCTLKernelBuildLog_Get(BuildLog)
+        buildLogStr = str(buildLogContent, "utf-8")
+        DPCTLBuildOptionList_Delete(BuildOpts)
+        DPCTLKernelNameList_Delete(KernelNames)
+        DPCTLVirtualHeaderList_Delete(VirtualHeaders)
+        DPCTLKernelBuildLog_Delete(BuildLog)
+        raise SyclKernelBundleCompilationError(buildLogStr)
+
+    DPCTLBuildOptionList_Delete(BuildOpts)
+    DPCTLKernelNameList_Delete(KernelNames)
+    DPCTLVirtualHeaderList_Delete(VirtualHeaders)
+    DPCTLKernelBuildLog_Delete(BuildLog)
+
+    return SyclKernelBundle._create(KBref, True)
 
 
 cpdef create_program_from_source(SyclQueue q, str src, str copts=""):