Support SYCL source compilation

CHANGELOG.md

@@ -9,6 +9,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 ### Added
 * `dpctl.SyclQueue.copy` and `dpctl.SyclQueue.copy_async` methods [gh-2273](https://github.com/IntelPython/dpctl/pull/2273)
 * Added a number of `sycl::device` info queries to `dpctl.SyclDevice` [gh-2324](https://github.com/IntelPython/dpctl/pull/2324)
+* Added `create_kernel_bundle_from_sycl_source`, `is_sycl_source_compilation_available`, and `dpctl.SyclDevice.can_compile` for supporting the creation of `dpctl.SyclKernelBundle`s from SYCL source strings via DPC++ extension, as well as corresponding C-API functions to support it [gh-2206](https://github.com/IntelPython/dpctl/pull/2206)
 
 ### Changed
 * Bump minimum NumPy version to 1.26 [gh-2192](https://github.com/IntelPython/dpctl/pull/2192)

docs/doc_sources/api_reference/dpctl/program.rst

@@ -22,8 +22,10 @@ execution via :py:meth:`dpctl.SyclQueue.submit`.
 
     create_kernel_bundle_from_source
     create_kernel_bundle_from_spirv
+    create_kernel_bundle_from_sycl_source
     create_program_from_source
     create_program_from_spirv
+    is_sycl_source_compilation_available
 
 .. autosummary::
     :toctree: generated

dpctl/_backend.pxd

@@ -331,7 +331,6 @@ 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 uint32_t DPCTLDevice_GetVendorId(const DPCTLSyclDeviceRef DRef)
@@ -374,6 +373,10 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h":
         const DPCTLSyclDeviceRef DRef, size_t *res_len)
     cdef int *DPCTLDevice_GetPartitionAffinityDomains(
         const DPCTLSyclDeviceRef DRef, size_t *res_len)
+    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(
@@ -542,6 +545,53 @@ 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 bool DPCTLKernelBundle_CreateFromSYCLSource_Available()
+
+    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

@@ -27,6 +27,9 @@ from ._backend cimport (  # noqa: E211
     DPCTLDefaultSelector_Create,
     DPCTLDevice_AreEq,
     DPCTLDevice_CanAccessPeer,
+    DPCTLDevice_CanCompileOpenCL,
+    DPCTLDevice_CanCompileSPIRV,
+    DPCTLDevice_CanCompileSYCL,
     DPCTLDevice_Copy,
     DPCTLDevice_CreateFromSelector,
     DPCTLDevice_CreateSubDevicesByAffinity,
@@ -2869,6 +2872,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,15 +29,19 @@
     SyclKernelBundleCompilationError,
     create_kernel_bundle_from_source,
     create_kernel_bundle_from_spirv,
+    create_kernel_bundle_from_sycl_source,
     create_program_from_source,
     create_program_from_spirv,
+    is_sycl_source_compilation_available,
 )
 
 __all__ = [
     "create_kernel_bundle_from_source",
     "create_kernel_bundle_from_spirv",
     "create_program_from_source",
     "create_program_from_spirv",
+    "create_kernel_bundle_from_sycl_source",
+    "is_sycl_source_compilation_available",
     "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)
 
@@ -72,3 +74,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=*)