14929317467

Committed 09 May 2025 12:49PM UTC coverage: 85.335% (-1.0%) from 86.372%

Build # 14929317467

Build Type

Pull #2049

github

Committed by

web-flow

Commit Message

Merge 153ec69eb into 3aff6ca34

Pull Request Pull Request #2049: Support compilation from SYCL source code

Run Details

2965 of 3732 branches covered (79.45%)

Branch coverage included in aggregate %.

122 of 199 new or added lines in 4 files covered. (61.31%)

94 existing lines in 17 files now uncovered.

12298 of 14154 relevant lines covered (86.89%)

6867.74 hits per line

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

64.39

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

IntelPython / dpctl / 14929317467

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