24128732776

Committed 08 Apr 2026 09:39AM UTC coverage: 77.558%. First build

Build # 24128732776

Build Type

Pull #2206

github

Committed by

web-flow

Commit Message

Merge 4656febb8 into 7d8bbc823

Pull Request Pull Request #2206: Support SYCL source compilation

Coverage Stats

831 of 1156 branches covered (71.89%)

Branch coverage included in aggregate %.

194 of 224 new or added lines in 4 files covered. (86.61%)

5445 of 6936 relevant lines covered (78.5%)

164.29 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

92.81

/dpctl/program/_program.pyx

#                      Data Parallel Control (dpctl)
#
# Copyright 2020-2025 Intel Corporation
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
#    http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.

# distutils: language = c++
# cython: language_level=3
# cython: linetrace=True

"""Implements a Python interface for SYCL's program and kernel runtime classes.

The module also provides functions to create a SYCL program from either
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,
    DPCTLKernel_GetCompileSubGroupSize,
    DPCTLKernel_GetMaxNumSubGroups,
    DPCTLKernel_GetMaxSubGroupSize,
    DPCTLKernel_GetNumArgs,
    DPCTLKernel_GetPreferredWorkGroupSizeMultiple,
    DPCTLKernel_GetPrivateMemSize,
    DPCTLKernel_GetWorkGroupSize,
    DPCTLKernelBuildLog_Create,
    DPCTLKernelBuildLog_Delete,
    DPCTLKernelBuildLog_Get,
    DPCTLKernelBuildLogRef,
    DPCTLKernelBundle_Copy,
    DPCTLKernelBundle_CreateFromOCLSource,
    DPCTLKernelBundle_CreateFromSpirv,
    DPCTLKernelBundle_CreateFromSYCLSource,
    DPCTLKernelBundle_CreateFromSYCLSource_Available,
    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__ = [
    "create_program_from_source",
    "create_program_from_spirv",
    "is_sycl_source_compilation_available",
    "SyclKernel",
    "SyclProgram",
    "SyclProgramCompilationError",
]

cdef class SyclProgramCompilationError(Exception):
    """This exception is raised when a ``sycl::kernel_bundle`` could not be
       built from either a SPIR-V binary file or a string source.
    """
    pass


cpdef bint is_sycl_source_compilation_available():
    """Returns True if dpctl was built with compiler that supports the DPC++
       `kernel_compiler` extension API used by
       :func:`create_program_from_sycl_source`.

       Device support is separate; callers should also check
       ``q.sycl_device.can_compile('sycl')`` (or similar) for specific devices.
    """
    return DPCTLKernelBundle_CreateFromSYCLSource_Available()


cdef class SyclKernel:
    """
    """
    @staticmethod
    cdef SyclKernel _create(DPCTLSyclKernelRef kref, str name):
        cdef SyclKernel ret = SyclKernel.__new__(SyclKernel)
        ret._kernel_ref = kref
        ret._function_name = name
        return ret

    def __dealloc__(self):
        DPCTLKernel_Delete(self._kernel_ref)

    def get_function_name(self):
        """ Returns the name of the ``sycl::kernel`` function.
        """
        return self._function_name

    def get_num_args(self):
        """ Returns the number of arguments for this kernel function.
        """
        return DPCTLKernel_GetNumArgs(self._kernel_ref)

    cdef DPCTLSyclKernelRef get_kernel_ref(self):
        """ Returns the ``DPCTLSyclKernelRef`` pointer for this SyclKernel.
        """
        return self._kernel_ref

    def addressof_ref(self):
        """ Returns the address of the C API ``DPCTLSyclKernelRef`` pointer
        as a ``size_t``.

        Returns:
            The address of the ``DPCTLSyclKernelRef`` pointer used to create
            this :class:`dpctl.SyclKernel` object cast to a ``size_t``.
        """
        return int(<size_t>self._kernel_ref)

    @property
    def num_args(self):
        """ Property equivalent to method call `SyclKernel.get_num_args()`
        """
        return self.get_num_args()

    @property
    def work_group_size(self):
        """ Returns the maximum number of work-items in a work-group that can
        be used to execute the kernel on device it was built for.
        """
        cdef size_t v = DPCTLKernel_GetWorkGroupSize(self._kernel_ref)
        return v

    @property
    def preferred_work_group_size_multiple(self):
        """ Returns a value, of which work-group size is preferred to be
        a multiple, for executing the kernel on the device it was built for.
        """
        cdef size_t v = DPCTLKernel_GetPreferredWorkGroupSizeMultiple(
            self._kernel_ref
        )
        return v

    @property
    def private_mem_size(self):
        """ Returns the minimum amount of private memory, in bytes, used by each
        work-item in the kernel.
        """
        cdef size_t v = DPCTLKernel_GetPrivateMemSize(self._kernel_ref)
        return v

    @property
    def max_num_sub_groups(self):
        """ Returns the maximum number of sub-groups for this kernel.
        """
        cdef uint32_t n = DPCTLKernel_GetMaxNumSubGroups(self._kernel_ref)
        return n

    @property
    def max_sub_group_size(self):
        """ Returns the maximum sub-groups size for this kernel.
        """
        cdef uint32_t sz = DPCTLKernel_GetMaxSubGroupSize(self._kernel_ref)
        return sz

    @property
    def compile_num_sub_groups(self):
        """ Returns the number of sub-groups specified by this kernel,
        or 0 (if not specified).
        """
        cdef size_t n = DPCTLKernel_GetCompileNumSubGroups(self._kernel_ref)
        return n

    @property
    def compile_sub_group_size(self):
        """ Returns the required sub-group size specified by this kernel,
        or 0 (if not specified).
        """
        cdef size_t n = DPCTLKernel_GetCompileSubGroupSize(self._kernel_ref)
        return n


cdef api DPCTLSyclKernelRef SyclKernel_GetKernelRef(SyclKernel ker):
    """ C-API function to access opaque kernel reference from
    Python object of type :class:`dpctl.program.SyclKernel`.
    """
    return ker.get_kernel_ref()


cdef api SyclKernel SyclKernel_Make(DPCTLSyclKernelRef KRef, const char *name):
    """
    C-API function to create :class:`dpctl.program.SyclKernel`
    instance from opaque sycl kernel reference.
    """
    cdef DPCTLSyclKernelRef copied_KRef = DPCTLKernel_Copy(KRef)
    if (name is NULL):
        return SyclKernel._create(copied_KRef, "default_name")
    else:
        return SyclKernel._create(copied_KRef, name.decode("utf-8"))


cdef class SyclProgram:
    """ Wraps a ``sycl::kernel_bundle<sycl::bundle_state::executable>`` object
    created using SYCL interoperability layer with underlying backends. Only the
    OpenCL and Level-Zero backends are currently supported.

    SyclProgram exposes the C API from ``dpctl_sycl_kernel_bundle_interface.h``.
    A SyclProgram can be created from either a source string or a SPIR-V
    binary file.
    """

    @staticmethod
    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):
        DPCTLKernelBundle_Delete(self._program_ref)

    cdef DPCTLSyclKernelBundleRef get_program_ref(self):
        return self._program_ref

    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):
        """Returns the address of the C API DPCTLSyclKernelBundleRef pointer
        as a long.

        Returns:
            The address of the ``DPCTLSyclKernelBundleRef`` pointer used to
            create this :class:`dpctl.SyclProgram` object cast to a ``size_t``.
        """
        return int(<size_t>self._program_ref)


cpdef create_program_from_source(SyclQueue q, str src, str copts=""):
    """
        Creates a Sycl interoperability program from an OpenCL source string.

        We use the ``DPCTLKernelBundle_CreateFromOCLSource()`` C API function
        to create a ``sycl::kernel_bundle<sycl::bundle_state::executable>``
        from an OpenCL source program that can contain multiple kernels.
        Note: This function is currently only supported for the OpenCL backend.

        Parameters:
            q (:class:`dpctl.SyclQueue`)
                The :class:`dpctl.SyclQueue` for which the
                :class:`.SyclProgram` is going to be built.
            src (str)
                Source string for an OpenCL program.
            copts (str, optional)
                Optional 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 bytes bSrc = src.encode("utf8")
    cdef bytes bCOpts = copts.encode("utf8")
    cdef const char *Src = <const char*>bSrc
    cdef const char *COpts = <const char*>bCOpts
    cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
    cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()
    KBref = DPCTLKernelBundle_CreateFromOCLSource(CRef, DRef, Src, COpts)

    if KBref is NULL:
        raise SyclProgramCompilationError()

    return SyclProgram._create(KBref, False)


cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,
                                str copts=""):
    """
        Creates a Sycl interoperability program from an SPIR-V binary.

        We use the :c:func:`DPCTLKernelBundle_CreateFromOCLSpirv` C API function
        to create a ``sycl::kernel_bundle<sycl::bundle_state::executable>``
        object from an compiled SPIR-V binary file.

        Parameters:
            q (:class:`dpctl.SyclQueue`)
                The :class:`dpctl.SyclQueue` for which the
                :class:`.SyclProgram` is going to be built.
            IL (bytes)
                SPIR-V binary IL file for an OpenCL program.
            copts (str, optional)
                Optional 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 const unsigned char *dIL = &IL[0]
    cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
    cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()
    cdef size_t length = IL.shape[0]
    cdef bytes bCOpts = copts.encode("utf8")
    cdef const char *COpts = <const char*>bCOpts
    KBref = DPCTLKernelBundle_CreateFromSpirv(
        CRef, DRef, <const void*>dIL, length, COpts
    )
    if KBref is NULL:
        raise SyclProgramCompilationError()

    return SyclProgram._create(KBref, False)


cpdef create_program_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:`.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. 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 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
    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)

    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 SyclProgramCompilationError(buildLogStr)

    DPCTLBuildOptionList_Delete(BuildOpts)
    DPCTLKernelNameList_Delete(KernelNames)
    DPCTLVirtualHeaderList_Delete(VirtualHeaders)
    DPCTLKernelBuildLog_Delete(BuildLog)

    return SyclProgram._create(KBref, True)


cdef api DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef(
    SyclProgram pro
):
    """ C-API function to access opaque kernel bundle reference from
    Python object of type :class:`dpctl.program.SyclKernel`.
    """
    return pro.get_program_ref()


cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef):
    """
    C-API function to create :class:`dpctl.program.SyclProgram`
    instance from opaque ``sycl::kernel_bundle<sycl::bundle_state::executable>``
    reference.
    """
    cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef)
    return SyclProgram._create(copied_KBRef, False)

1	# Data Parallel Control (dpctl)	1✔
2	#
3	# Copyright 2020-2025 Intel Corporation
4	#
5	# Licensed under the Apache License, Version 2.0 (the "License");
6	# you may not use this file except in compliance with the License.
7	# You may obtain a copy of the License at
8	#
9	# http://www.apache.org/licenses/LICENSE-2.0
10	#
11	# Unless required by applicable law or agreed to in writing, software
12	# distributed under the License is distributed on an "AS IS" BASIS,
13	# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
14	# See the License for the specific language governing permissions and
15	# limitations under the License.
16
17	# distutils: language = c++
18	# cython: language_level=3
19	# cython: linetrace=True
20
21	"""Implements a Python interface for SYCL's program and kernel runtime classes.
22
23	The module also provides functions to create a SYCL program from either
24	a OpenCL source string or a SPIR-V binary file.
25
26	"""
27
28	from libc.stdint cimport uint32_t
29
30	from dpctl._backend cimport ( # noqa: E211, E402;
31	DPCTLBuildOptionList_Append,
32	DPCTLBuildOptionList_Create,
33	DPCTLBuildOptionList_Delete,
34	DPCTLBuildOptionListRef,
35	DPCTLKernel_Copy,
36	DPCTLKernel_Delete,
37	DPCTLKernel_GetCompileNumSubGroups,
38	DPCTLKernel_GetCompileSubGroupSize,
39	DPCTLKernel_GetMaxNumSubGroups,
40	DPCTLKernel_GetMaxSubGroupSize,
41	DPCTLKernel_GetNumArgs,
42	DPCTLKernel_GetPreferredWorkGroupSizeMultiple,
43	DPCTLKernel_GetPrivateMemSize,
44	DPCTLKernel_GetWorkGroupSize,
45	DPCTLKernelBuildLog_Create,
46	DPCTLKernelBuildLog_Delete,
47	DPCTLKernelBuildLog_Get,
48	DPCTLKernelBuildLogRef,
49	DPCTLKernelBundle_Copy,
50	DPCTLKernelBundle_CreateFromOCLSource,
51	DPCTLKernelBundle_CreateFromSpirv,
52	DPCTLKernelBundle_CreateFromSYCLSource,
53	DPCTLKernelBundle_CreateFromSYCLSource_Available,
54	DPCTLKernelBundle_Delete,
55	DPCTLKernelBundle_GetKernel,
56	DPCTLKernelBundle_GetSyclKernel,
57	DPCTLKernelBundle_HasKernel,
58	DPCTLKernelBundle_HasSyclKernel,
59	DPCTLKernelNameList_Append,
60	DPCTLKernelNameList_Create,
61	DPCTLKernelNameList_Delete,
62	DPCTLKernelNameListRef,
63	DPCTLSyclContextRef,
64	DPCTLSyclDeviceRef,
65	DPCTLSyclKernelBundleRef,
66	DPCTLSyclKernelRef,
67	DPCTLVirtualHeaderList_Append,
68	DPCTLVirtualHeaderList_Create,
69	DPCTLVirtualHeaderList_Delete,
70	DPCTLVirtualHeaderListRef,
71	)
72
73	__all__ = [	1✔
74	"create_program_from_source",
75	"create_program_from_spirv",
76	"is_sycl_source_compilation_available",
77	"SyclKernel",
78	"SyclProgram",
79	"SyclProgramCompilationError",
80	]
81
82	cdef class SyclProgramCompilationError(Exception):
83	"""This exception is raised when a ``sycl::kernel_bundle`` could not be
84	built from either a SPIR-V binary file or a string source.
85	"""
86	pass
87
88
89	cpdef bint is_sycl_source_compilation_available():	1✔
90	"""Returns True if dpctl was built with compiler that supports the DPC++
91	`kernel_compiler` extension API used by
92	:func:`create_program_from_sycl_source`.
93
94	Device support is separate; callers should also check
95	``q.sycl_device.can_compile('sycl')`` (or similar) for specific devices.
96	"""
97	return DPCTLKernelBundle_CreateFromSYCLSource_Available()	1✔
98
99
100	cdef class SyclKernel:
101	"""
102	"""
103	@staticmethod	1✔
104	cdef SyclKernel _create(DPCTLSyclKernelRef kref, str name):
105	cdef SyclKernel ret = SyclKernel.__new__(SyclKernel)	1✔
106	ret._kernel_ref = kref	1✔
107	ret._function_name = name	1✔
108	return ret	1✔
109
110	def __dealloc__(self):
111	DPCTLKernel_Delete(self._kernel_ref)	1✔
112
113	def get_function_name(self):	1✔
114	""" Returns the name of the ``sycl::kernel`` function.
115	"""
116	return self._function_name	1✔
117
118	def get_num_args(self):	1✔
119	""" Returns the number of arguments for this kernel function.
120	"""
121	return DPCTLKernel_GetNumArgs(self._kernel_ref)	1✔
122
123	cdef DPCTLSyclKernelRef get_kernel_ref(self):	1✔
124	""" Returns the ``DPCTLSyclKernelRef`` pointer for this SyclKernel.
125	"""
126	return self._kernel_ref	1✔
127
128	def addressof_ref(self):	1✔
129	""" Returns the address of the C API ``DPCTLSyclKernelRef`` pointer
130	as a ``size_t``.
131
132	Returns:
133	The address of the ``DPCTLSyclKernelRef`` pointer used to create
134	this :class:`dpctl.SyclKernel` object cast to a ``size_t``.
135	"""
136	return int(<size_t>self._kernel_ref)	1✔
137
138	@property
139	def num_args(self):
140	""" Property equivalent to method call `SyclKernel.get_num_args()`
141	"""
142	return self.get_num_args()	1✔
143
144	@property
145	def work_group_size(self):
146	""" Returns the maximum number of work-items in a work-group that can
147	be used to execute the kernel on device it was built for.
148	"""
149	cdef size_t v = DPCTLKernel_GetWorkGroupSize(self._kernel_ref)	1✔
150	return v	1✔
151
152	@property
153	def preferred_work_group_size_multiple(self):
154	""" Returns a value, of which work-group size is preferred to be
155	a multiple, for executing the kernel on the device it was built for.
156	"""
157	cdef size_t v = DPCTLKernel_GetPreferredWorkGroupSizeMultiple(	1✔
158	self._kernel_ref
159	)
160	return v	1✔
161
162	@property
163	def private_mem_size(self):
164	""" Returns the minimum amount of private memory, in bytes, used by each
165	work-item in the kernel.
166	"""
167	cdef size_t v = DPCTLKernel_GetPrivateMemSize(self._kernel_ref)	1✔
168	return v	1✔
169
170	@property
171	def max_num_sub_groups(self):
172	""" Returns the maximum number of sub-groups for this kernel.
173	"""
174	cdef uint32_t n = DPCTLKernel_GetMaxNumSubGroups(self._kernel_ref)	1✔
175	return n	1✔
176
177	@property
178	def max_sub_group_size(self):
179	""" Returns the maximum sub-groups size for this kernel.
180	"""
181	cdef uint32_t sz = DPCTLKernel_GetMaxSubGroupSize(self._kernel_ref)	1✔
182	return sz	1✔
183
184	@property
185	def compile_num_sub_groups(self):
186	""" Returns the number of sub-groups specified by this kernel,
187	or 0 (if not specified).
188	"""
189	cdef size_t n = DPCTLKernel_GetCompileNumSubGroups(self._kernel_ref)	1✔
190	return n	1✔
191
192	@property
193	def compile_sub_group_size(self):
194	""" Returns the required sub-group size specified by this kernel,
195	or 0 (if not specified).
196	"""
197	cdef size_t n = DPCTLKernel_GetCompileSubGroupSize(self._kernel_ref)	1✔
198	return n	1✔
199
200
201	cdef api DPCTLSyclKernelRef SyclKernel_GetKernelRef(SyclKernel ker):	1✔
202	""" C-API function to access opaque kernel reference from
203	Python object of type :class:`dpctl.program.SyclKernel`.
204	"""
205	return ker.get_kernel_ref()	1✔
206
207
208	cdef api SyclKernel SyclKernel_Make(DPCTLSyclKernelRef KRef, const char *name):	1✔
209	"""
210	C-API function to create :class:`dpctl.program.SyclKernel`
211	instance from opaque sycl kernel reference.
212	"""
213	cdef DPCTLSyclKernelRef copied_KRef = DPCTLKernel_Copy(KRef)	1✔
214	if (name is NULL):	1✔
215	return SyclKernel._create(copied_KRef, "default_name")	1✔
216	else:
217	return SyclKernel._create(copied_KRef, name.decode("utf-8"))	1✔
218
219
220	cdef class SyclProgram:
221	""" Wraps a ``sycl::kernel_bundle<sycl::bundle_state::executable>`` object
222	created using SYCL interoperability layer with underlying backends. Only the
223	OpenCL and Level-Zero backends are currently supported.
224
225	SyclProgram exposes the C API from ``dpctl_sycl_kernel_bundle_interface.h``.
226	A SyclProgram can be created from either a source string or a SPIR-V
227	binary file.
228	"""
229
230	@staticmethod	1✔
231	cdef SyclProgram _create(DPCTLSyclKernelBundleRef KBRef,
232	bint is_sycl_source):
233	cdef SyclProgram ret = SyclProgram.__new__(SyclProgram)	1✔
234	ret._program_ref = KBRef	1✔
235	ret._is_sycl_source = is_sycl_source	1✔
236	return ret	1✔
237
238	def __dealloc__(self):
239	DPCTLKernelBundle_Delete(self._program_ref)	1✔
240
241	cdef DPCTLSyclKernelBundleRef get_program_ref(self):	1✔
242	return self._program_ref	1✔
243
244	cpdef SyclKernel get_sycl_kernel(self, str kernel_name):	1✔
245	name = kernel_name.encode("utf8")	1✔
246	if self._is_sycl_source:	1✔
247	return SyclKernel._create(	1✔
248	DPCTLKernelBundle_GetSyclKernel(self._program_ref, name),	1✔
249	kernel_name)
250	return SyclKernel._create(	1✔
251	DPCTLKernelBundle_GetKernel(self._program_ref, name),	1✔
252	kernel_name
253	)
254
255	def has_sycl_kernel(self, str kernel_name):	1✔
256	name = kernel_name.encode("utf8")	1✔
257	if self._is_sycl_source:	1✔
258	return DPCTLKernelBundle_HasSyclKernel(self._program_ref, name)	1✔
259	return DPCTLKernelBundle_HasKernel(self._program_ref, name)	1✔
260
261	def addressof_ref(self):	1✔
262	"""Returns the address of the C API DPCTLSyclKernelBundleRef pointer
263	as a long.
264
265	Returns:
266	The address of the ``DPCTLSyclKernelBundleRef`` pointer used to
267	create this :class:`dpctl.SyclProgram` object cast to a ``size_t``.
268	"""
269	return int(<size_t>self._program_ref)	1✔
270
271
272	cpdef create_program_from_source(SyclQueue q, str src, str copts=""):	1✔
273	"""
274	Creates a Sycl interoperability program from an OpenCL source string.
275
276	We use the ``DPCTLKernelBundle_CreateFromOCLSource()`` C API function
277	to create a ``sycl::kernel_bundle<sycl::bundle_state::executable>``
278	from an OpenCL source program that can contain multiple kernels.
279	Note: This function is currently only supported for the OpenCL backend.
280
281	Parameters:
282	q (:class:`dpctl.SyclQueue`)
283	The :class:`dpctl.SyclQueue` for which the
284	:class:`.SyclProgram` is going to be built.
285	src (str)
286	Source string for an OpenCL program.
287	copts (str, optional)
288	Optional compilation flags that will be used
289	when compiling the program. Default: ``""``.
290
291	Returns:
292	program (:class:`.SyclProgram`)
293	A :class:`.SyclProgram` object wrapping the
294	``sycl::kernel_bundle<sycl::bundle_state::executable>``
295	returned by the C API.
296
297	Raises:
298	SyclProgramCompilationError
299	If a SYCL kernel bundle could not be created.
300	"""
301
302	cdef DPCTLSyclKernelBundleRef KBref
303	cdef bytes bSrc = src.encode("utf8")	1✔
304	cdef bytes bCOpts = copts.encode("utf8")	1✔
305	cdef const char Src = <const char>bSrc	1✔
306	cdef const char COpts = <const char>bCOpts	1✔
307	cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()	1✔
308	cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()	1✔
309	KBref = DPCTLKernelBundle_CreateFromOCLSource(CRef, DRef, Src, COpts)	1✔
310
311	if KBref is NULL:	1✔
312	raise SyclProgramCompilationError()	1✔
313
314	return SyclProgram._create(KBref, False)	1✔
315
316
317	cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL,	1✔
318	str copts=""):
319	"""
320	Creates a Sycl interoperability program from an SPIR-V binary.
321
322	We use the :c:func:`DPCTLKernelBundle_CreateFromOCLSpirv` C API function
323	to create a ``sycl::kernel_bundle<sycl::bundle_state::executable>``
324	object from an compiled SPIR-V binary file.
325
326	Parameters:
327	q (:class:`dpctl.SyclQueue`)
328	The :class:`dpctl.SyclQueue` for which the
329	:class:`.SyclProgram` is going to be built.
330	IL (bytes)
331	SPIR-V binary IL file for an OpenCL program.
332	copts (str, optional)
333	Optional compilation flags that will be used
334	when compiling the program. Default: ``""``.
335
336	Returns:
337	program (:class:`.SyclProgram`)
338	A :class:`.SyclProgram` object wrapping the
339	``sycl::kernel_bundle<sycl::bundle_state::executable>``
340	returned by the C API.
341
342	Raises:
343	SyclProgramCompilationError
344	If a SYCL kernel bundle could not be created.
345	"""
346
347	cdef DPCTLSyclKernelBundleRef KBref
348	cdef const unsigned char *dIL = &IL[0]	1✔
349	cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()	1✔
350	cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()	1✔
351	cdef size_t length = IL.shape[0]	1✔
352	cdef bytes bCOpts = copts.encode("utf8")	1✔
353	cdef const char COpts = <const char>bCOpts	1✔
354	KBref = DPCTLKernelBundle_CreateFromSpirv(	1✔
355	CRef, DRef, <const void*>dIL, length, COpts
356	)
357	if KBref is NULL:	1✔
358	raise SyclProgramCompilationError()	×
359
360	return SyclProgram._create(KBref, False)	1✔
361
362
363	cpdef create_program_from_sycl_source(SyclQueue q, unicode source,	1✔
364	list headers=None,
365	list registered_names=None,
366	list copts=None):	1✔
367	"""
368	Creates an executable SYCL kernel_bundle from SYCL source code.
369
370	This uses the DPC++ ``kernel_compiler`` extension to create a
371	``sycl::kernel_bundle<sycl::bundle_state::executable>`` object from
372	SYCL source code.
373
374	Parameters:
375	q (:class:`dpctl.SyclQueue`)
376	The :class:`dpctl.SyclQueue` for which the
377	:class:`.SyclProgram` is going to be built.
378	source (unicode)
379	SYCL source code string.
380	headers (list)
381	Optional list of virtual headers, where each entry in the list
382	needs to be a tuple of header name and header content. See the
383	documentation of the ``include_files`` property in the DPC++
384	``kernel_compiler`` extension for more information.
385	Default: []
386	registered_names (list, optional)
387	Optional list of kernel names to register. See the
388	documentation of the ``registered_names`` property in the DPC++
389	``kernel_compiler`` extension for more information.
390	Default: []
391	copts (list)
392	Optional list of compilation flags that will be used
393	when compiling the program. Default: ``""``.
394
395	Returns:
396	program (:class:`.SyclProgram`)
397	A :class:`.SyclProgram` object wrapping the
398	``sycl::kernel_bundle<sycl::bundle_state::executable>``
399	returned by the C API.
400
401	Raises:
402	SyclProgramCompilationError
403	If a SYCL kernel bundle could not be created. The exception
404	message contains the build log for more details.
405	"""
406	cdef DPCTLSyclKernelBundleRef KBref
407	cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()	1✔
408	cdef DPCTLSyclDeviceRef DRef = q.get_sycl_device().get_device_ref()	1✔
409	cdef bytes bSrc = source.encode("utf8")	1✔
410	cdef const char Src = <const char>bSrc	1✔
411	cdef DPCTLBuildOptionListRef BuildOpts = DPCTLBuildOptionList_Create()	1✔
412	cdef bytes bOpt
413	cdef const char* sOpt
414	cdef bytes bName
415	cdef const char* sName
416	cdef bytes bContent
417	cdef const char* sContent
418	cdef const char* buildLogContent
419	for opt in copts:	1✔
420	if not isinstance(opt, unicode):	1✔
NEW 421	DPCTLBuildOptionList_Delete(BuildOpts)	×
NEW 422	raise SyclProgramCompilationError()	×
423	bOpt = opt.encode("utf8")	1✔
424	sOpt = <const char*>bOpt	1✔
425	DPCTLBuildOptionList_Append(BuildOpts, sOpt)	1✔
426
427	cdef DPCTLKernelNameListRef KernelNames = DPCTLKernelNameList_Create()	1✔
428	for name in registered_names:	1✔
429	if not isinstance(name, unicode):	1✔
NEW 430	DPCTLBuildOptionList_Delete(BuildOpts)	×
NEW 431	DPCTLKernelNameList_Delete(KernelNames)	×
NEW 432	raise SyclProgramCompilationError()	×
433	bName = name.encode("utf8")	1✔
434	sName = <const char*>bName	1✔
435	DPCTLKernelNameList_Append(KernelNames, sName)	1✔
436
437	cdef DPCTLVirtualHeaderListRef VirtualHeaders
438	VirtualHeaders = DPCTLVirtualHeaderList_Create()	1✔
439
440	for name, content in headers:	1✔
441	if not isinstance(name, unicode) or not isinstance(content, unicode):	1✔
NEW 442	DPCTLBuildOptionList_Delete(BuildOpts)	×
NEW 443	DPCTLKernelNameList_Delete(KernelNames)	×
NEW 444	DPCTLVirtualHeaderList_Delete(VirtualHeaders)	×
NEW 445	raise SyclProgramCompilationError()	×
446	bName = name.encode("utf8")	1✔
447	sName = <const char*>bName	1✔
448	bContent = content.encode("utf8")	1✔
449	sContent = <const char*>bContent	1✔
450	DPCTLVirtualHeaderList_Append(VirtualHeaders, sName, sContent)	1✔
451
452	cdef DPCTLKernelBuildLogRef BuildLog
453	BuildLog = DPCTLKernelBuildLog_Create()	1✔
454
455	KBref = DPCTLKernelBundle_CreateFromSYCLSource(CRef, DRef, Src,	1✔
456	VirtualHeaders, KernelNames,
457	BuildOpts, BuildLog)
458
459	if KBref is NULL:	1✔
460	buildLogContent = DPCTLKernelBuildLog_Get(BuildLog)	1✔
461	buildLogStr = str(buildLogContent, "utf-8")	1✔
462	DPCTLBuildOptionList_Delete(BuildOpts)	1✔
463	DPCTLKernelNameList_Delete(KernelNames)	1✔
464	DPCTLVirtualHeaderList_Delete(VirtualHeaders)	1✔
465	DPCTLKernelBuildLog_Delete(BuildLog)	1✔
466	raise SyclProgramCompilationError(buildLogStr)	1✔
467
468	DPCTLBuildOptionList_Delete(BuildOpts)	1✔
469	DPCTLKernelNameList_Delete(KernelNames)	1✔
470	DPCTLVirtualHeaderList_Delete(VirtualHeaders)	1✔
471	DPCTLKernelBuildLog_Delete(BuildLog)	1✔
472
473	return SyclProgram._create(KBref, True)	1✔
474
475
476	cdef api DPCTLSyclKernelBundleRef SyclProgram_GetKernelBundleRef(	1✔
477	SyclProgram pro
478	):
479	""" C-API function to access opaque kernel bundle reference from
480	Python object of type :class:`dpctl.program.SyclKernel`.
481	"""
482	return pro.get_program_ref()	1✔
483
484
485	cdef api SyclProgram SyclProgram_Make(DPCTLSyclKernelBundleRef KBRef):	1✔
486	"""
487	C-API function to create :class:`dpctl.program.SyclProgram`
488	instance from opaque ``sycl::kernel_bundle<sycl::bundle_state::executable>``
489	reference.
490	"""
491	cdef DPCTLSyclKernelBundleRef copied_KBRef = DPCTLKernelBundle_Copy(KBRef)	1✔
492	return SyclProgram._create(copied_KBRef, False)	1✔

IntelPython / dpctl / 24128732776

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous