Merge pull request #2206 from IntelPython/feature/sycl-source-compilation

Support SYCL source compilation

Author: ndgrigorian

CHANGELOG.md

@@ -10,6 +10,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 * `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 `sycl::info::context` queries `sycl_platform`, `atomic_memory_order_capabilities`, `atomic_fence_order_capabilities`, `atomic_memory_scope_capabilities`, and `atomic_fence_scope_capabilities` to `dpctl.SyclContext` [gh-2354](https://github.com/IntelPython/dpctl/pull/2354)
+* 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

@@ -4,8 +4,8 @@
 =======================
 
 :py:mod:`dpctl.program` provides a way to create a SYCL kernel
-from either an OpenCL* program source code represented as a string
-or a SPIR-V binary file.
+from an OpenCL* program source code represented as a string, SYCL
+source code represented as a string, or a SPIR-V binary file.
 
 It implements creation of interoperability
 ``sycl::kernel_bundle<sycl::bundle_state_executable>`` (a collection of kernels),
@@ -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

@@ -328,7 +328,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)
@@ -371,6 +370,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(
@@ -553,6 +556,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,
@@ -2854,6 +2857,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" or "spv" for
+                SPIR-V binary files, "opencl" or "ocl" 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/_sycl_queue.pxd

@@ -68,16 +68,16 @@ cdef public api class SyclQueue (_SyclQueue) [
     cdef int _populate_range(self, size_t Range[3], list gS, size_t nGS)
 
     @staticmethod
-    cdef  SyclQueue _create(DPCTLSyclQueueRef qref)
+    cdef SyclQueue _create(DPCTLSyclQueueRef qref)
 
     @staticmethod
-    cdef  SyclQueue _create_from_context_and_device(
+    cdef SyclQueue _create_from_context_and_device(
         SyclContext ctx, SyclDevice dev, int props=*
     )
     cdef cpp_bool equals(self, SyclQueue q)
     cpdef SyclContext get_sycl_context(self)
     cpdef SyclDevice get_sycl_device(self)
-    cdef  DPCTLSyclQueueRef get_queue_ref(self)
+    cdef DPCTLSyclQueueRef get_queue_ref(self)
     cpdef SyclEvent _submit_keep_args_alive(
         self,
         object args,

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,10 +52,12 @@ cdef api class SyclKernelBundle [
     binary file.
     """
     cdef DPCTLSyclKernelBundleRef _kernel_bundle_ref
+    cdef bint _is_sycl_source
 
     @staticmethod
-    cdef  SyclKernelBundle _create (DPCTLSyclKernelBundleRef kbref)
-    cdef  DPCTLSyclKernelBundleRef get_kernel_bundle_ref (self)
+    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=*)