Skip to content

Support compilation from SYCL source code #2049

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: master
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
40 changes: 40 additions & 0 deletions dpctl/_backend.pxd
Original file line number Diff line number Diff line change
Expand Up @@ -221,6 +221,9 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h":
cdef uint64_t DPCTLDevice_GetMaxMemAllocSize(const DPCTLSyclDeviceRef DRef)
cdef DPCTLSyclDeviceRef DPCTLDevice_GetCompositeDevice(const DPCTLSyclDeviceRef DRef)
cdef DPCTLDeviceVectorRef DPCTLDevice_GetComponentDevices(const DPCTLSyclDeviceRef DRef)
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":
Expand Down Expand Up @@ -367,6 +370,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':
Expand Down
1 change: 1 addition & 0 deletions dpctl/_sycl_device.pxd
Original file line number Diff line number Diff line change
Expand Up @@ -59,3 +59,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)
32 changes: 32 additions & 0 deletions dpctl/_sycl_device.pyx
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,9 @@ from ._backend cimport ( # noqa: E211
DPCTLCString_Delete,
DPCTLDefaultSelector_Create,
DPCTLDevice_AreEq,
DPCTLDevice_CanCompileOpenCL,
DPCTLDevice_CanCompileSPIRV,
DPCTLDevice_CanCompileSYCL,
DPCTLDevice_Copy,
DPCTLDevice_CreateFromSelector,
DPCTLDevice_CreateSubDevicesByAffinity,
Expand Down Expand Up @@ -2164,6 +2167,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):
"""
Expand Down
2 changes: 2 additions & 0 deletions dpctl/program/__init__.py
Original file line number Diff line number Diff line change
Expand Up @@ -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",
Expand Down
6 changes: 5 additions & 1 deletion dpctl/program/_program.pxd
Original file line number Diff line number Diff line change
Expand Up @@ -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=*)
130 changes: 126 additions & 4 deletions dpctl/program/_program.pyx
Original file line number Diff line number Diff line change
Expand Up @@ -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,
DPCTLCString_Delete,
DPCTLKernel_Copy,
DPCTLKernel_Delete,
Expand All @@ -42,13 +46,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__ = [
Expand Down Expand Up @@ -197,9 +212,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):
Expand All @@ -210,13 +226,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):
Expand Down Expand Up @@ -272,7 +294,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,
Expand Down Expand Up @@ -318,7 +340,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(SyclProgram pro):
Expand All @@ -335,4 +457,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)
Loading
Loading