20317447685

Committed 17 Dec 2025 09:13PM UTC coverage: 86.229% (-0.02%) from 86.25%

Build # 20317447685

Build Type

Pull #2206

github

Committed by

web-flow

Commit Message

Merge 22475d963 into e041fca3f

Pull Request Pull Request #2206: Support SYCL source compilation

Coverage Stats

3128 of 3668 branches covered (85.28%)

Branch coverage included in aggregate %.

188 of 218 new or added lines in 4 files covered. (86.24%)

1 existing line in 1 file now uncovered.

12464 of 14414 relevant lines covered (86.47%)

5804.09 hits per line

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

92.7

/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_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",
    "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


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

IntelPython / dpctl / 20317447685

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