8423792464

Committed 25 Mar 2024 04:47PM UTC coverage: 81.694% (+2.6%) from 79.08%

Build # 8423792464

Build Type

Pull #1404

github

Committed by

web-flow

Commit Message

Merge 8785696b1 into 7149d3728

Pull Request Pull Request #1404: Remove old files not used anymore.

Run Details

1402 of 1949 branches covered (71.93%)

Branch coverage included in aggregate %.

8 of 8 new or added lines in 3 files covered. (100.0%)

10 existing lines in 4 files now uncovered.

5908 of 6999 relevant lines covered (84.41%)

0.84 hits per line

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

90.91

/numba_dpex/core/decorators.py

# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
#
# SPDX-License-Identifier: Apache-2.0

import inspect
from warnings import warn

from numba.core import decorators, sigutils, typeinfer
from numba.core.target_extension import (
    jit_registry,
    resolve_dispatcher_from_str,
    target_registry,
)

from numba_dpex.core import config
from numba_dpex.core.pipelines.dpjit_compiler import get_compiler
from numba_dpex.core.targets.dpjit_target import DPEX_TARGET_NAME
from numba_dpex.experimental.target import DPEX_KERNEL_EXP_TARGET_NAME
from numba_dpex.kernel_api_impl.spirv.dispatcher import SPIRVKernelDispatcher
from numba_dpex.kernel_api_impl.spirv.target import CompilationMode


def _parse_func_or_sig(signature_or_function):
    # Handle signature (borrowed from numba). swapped signature and list check
    if signature_or_function is None:
        # No signature, no function
        pyfunc = None
        sigs = []
    elif sigutils.is_signature(signature_or_function):
        # A single signature is passed
        pyfunc = None
        sigs = [signature_or_function]
    elif isinstance(signature_or_function, list):
        # A list of signatures is passed
        pyfunc = None
        sigs = signature_or_function
    else:
        # A function is passed
        pyfunc = signature_or_function
        sigs = []

    return pyfunc, sigs


def kernel(function_or_signature=None, **options):
    """A decorator to compile a function written using :py:mod:`numba_dpex.kernel_api`.

    The ``kernel`` decorator triggers the compilation of a function written
    using the data-parallel kernel programming API exposed by
    :py:mod:`numba_dpex.kernel_api`. Such a function is conceptually
    equivalent to a kernel function written in the C++ SYCL eDSL. The
    decorator will compile the function based on the types of the arguments
    to a SPIR-V binary that can be executed either on OpenCL CPU, GPU
    devices or Intel Level Zero GPU devices.

    Any function to be compilable using the kernel decorator should
    adhere to the following semantic rules:

    - The first argument to the function should be either an instance of the
      :class:`numba_dpex.kernel_api.Item` class or an instance of the
      :class:`numba_dpex.kernel_api.NdItem`.

    - The function should not return any value.

    - The function should have at least one array type argument that can
      either be an instance of ``dpnp.ndarray`` or an instance of
      ``dpctl.tensor.usm_ndarray``.


    Args:
        signature_or_function (optional): An optional signature or list of
            signatures for which a function is to be compiled. Passing in a
            signature "specializes" the decorated function and no other versions
            of the function will be compiled. A function can also be
            directly passed instead of a signature and the signature will get
            inferred from the function. The actual compilation happens on every
            invocation of the :func:`numba_dpex.experimental.call_kernel`
            function where the decorated function is passed in as an argument
            along with the argument values for the decorated function.
        options (optional):
            - **debug** (bool): Whether the compilation should happen in debug
              mode. *(Default = False)*
            - **inline_threshold** (int): Specifies the level of inlining that
              the compiler should attempt. *(Default = 2)*
    Returns:
        An instance of
        :class:`numba_dpex.kernel_api_impl.spirv.dispatcher.KernelDispatcher`.
        The ``KernelDispatcher`` object compiles the decorated function when
        passed in to :func:`numba_dpex.experimental.call_kernel`.

    Examples:

    1. Decorate a function and pass it to ``call_kernel`` for compilation and
       execution.

    .. code-block:: python

        import dpnp
        import numba_dpex as dpex
        from numba_dpex import kernel_api as kapi


        # Data parallel kernel implementing vector sum
        @dpex.kernel
        def vecadd(item: kapi.Item, a, b, c):
            i = item.get_id(0)
            c[i] = a[i] + b[i]


        N = 1024
        a = dpnp.ones(N)
        b = dpnp.ones_like(a)
        c = dpnp.zeros_like(a)
        dpex.call_kernel(vecadd, kapi.Range(N), a, b, c)

    2. Specializes a kernel and then compiles it directly before executing it
       via ``call_kernel``. The kernel is specialized to expect a 1-D
       ``dpnp.ndarray`` with either ``float32`` type elements or ``int64`` type
       elements.

    .. code-block:: python

        import dpnp
        import numba_dpex as dpex
        from numba_dpex import kernel_api as kapi
        from numba_dpex import DpnpNdArray, float32, int64
        from numba_dpex.core.types.kernel_api.index_space_ids import ItemType

        i64arrty = DpnpNdArray(ndim=1, dtype=int64, layout="C")
        f32arrty = DpnpNdArray(ndim=1, dtype=float32, layout="C")
        item_ty = ItemType(ndim=1)

        specialized_kernel = dpex.kernel(
            [
                (item_ty, i64arrty, i64arrty, i64arrty),
                (item_ty, f32arrty, f32arrty, f32arrty),
            ]
        )


        def vecadd(item: kapi.Item, a, b, c):
            i = item.get_id(0)
            c[i] = a[i] + b[i]


        # Compile all specializations for vecadd
        precompiled_kernels = specialized_kernel(vecadd)
        N = 1024
        a = dpnp.ones(N, dtype=dpnp.int64)
        b = dpnp.ones_like(a)
        c = dpnp.zeros_like(a)
        # Call a specific pre-compiled version of vecadd
        dpex.call_kernel(precompiled_kernels, kapi.Range(N), a, b, c)

    """

    # dispatcher is a type:
    # <class 'numba_dpex.experimental.kernel_dispatcher.KernelDispatcher'>
    dispatcher = resolve_dispatcher_from_str(DPEX_KERNEL_EXP_TARGET_NAME)
    if "_compilation_mode" in options:
        user_compilation_mode = options["_compilation_mode"]
        warn(
            "_compilation_mode is an internal flag that should not be set "
            "in the decorator. The decorator defined option "
            f"{user_compilation_mode} is going to be ignored."
        )
    options["_compilation_mode"] = CompilationMode.KERNEL

    # TODO: The options need to be evaluated and checked here like it is
    # done in numba.core.decorators.jit

    func, sigs = _parse_func_or_sig(function_or_signature)
    for sig in sigs:
        if isinstance(sig, str):
            raise NotImplementedError(
                "Specifying signatures as string is not yet supported"
            )

    def _kernel_dispatcher(pyfunc):
        disp: SPIRVKernelDispatcher = dispatcher(
            pyfunc=pyfunc,
            targetoptions=options,
        )

        if len(sigs) > 0:
            with typeinfer.register_dispatcher(disp):
                for sig in sigs:
                    disp.compile(sig)
                disp.disable_compile()

        return disp

    if func is None:
        return _kernel_dispatcher

    if not inspect.isfunction(func):
        raise ValueError(
            "Argument passed to the kernel decorator is neither a "
            "function object, nor a signature. If you are trying to "
            "specialize the kernel that takes a single argument, specify "
            "the return type as None explicitly."
        )
    return _kernel_dispatcher(func)


def device_func(function_or_signature=None, **options):
    """Compiles a device-callable function that can be only invoked from a kernel.

    The decorator is used to  express auxiliary device-only functions that can
    be called from a kernel or another device function, but are not callable
    from the host. This decorator :func:`numba_dpex.experimental.device_func`
    has no direct analogue in SYCL and primarily is provided to help programmers
    make their kapi applications modular.

    A ``device_func`` decorated function does not require the first argument to
    be a :class:`numba_dpex.kernel_api.Item` object or a
    :class:`numba_dpex.kernel_api.NdItem` object, and unlike a ``kernel``
    decorated function is allowed to return any value.
    All :py:mod:`numba_dpex.kernel_api` functionality can be used in a
    ``device_func`` decorated function.

    The decorator is also used to compile overloads in the ``DpexKernelTarget``.

    A ``device_func`` decorated function is not compiled down to device binary
    and instead is compiled down to LLVM IR. Final compilation to binary happens
    when the function is invoked from a ``kernel`` decorated function. The
    compilation happens this was to allow a ``device_func`` decorated function
    to be internally linked into the kernel module at the LLVM level, leading to
    more optimization opportunities.

    Args:
        signature_or_function (optional): An optional signature or list of
            signatures for which a function is to be compiled. Passing in a
            signature "specializes" the decorated function and no other versions
            of the function will be compiled. A function can also be
            directly passed instead of a signature and the signature will get
            inferred from the function. The actual compilation happens on every
            invocation of the decorated function from another ``device_func`` or
            ``kernel`` decorated function.
        options (optional):
            - **debug** (bool): Whether the compilation should happen in debug
              mode. *(Default = False)*
            - **inline_threshold** (int): Specifies the level of inlining that
              the compiler should attempt. *(Default = 2)*

    Returns:
        An instance of
        :class:`numba_dpex.kernel_api_impl.spirv.dispatcher.KernelDispatcher`.
        The ``KernelDispatcher`` object compiles the decorated function when
        it is called from another function.


    Example:

    .. code-block:: python

        import dpnp

        from numba_dpex import experimental as dpex_exp
        from numba_dpex import kernel_api as kapi


        @dpex_exp.device_func
        def increment_value(nd_item: NdItem, a):
            i = nd_item.get_global_id(0)

            a[i] += 1
            group_barrier(nd_item.get_group(), MemoryScope.DEVICE)

            if i == 0:
                for idx in range(1, a.size):
                    a[0] += a[idx]


        @dpex_exp.kernel
        def another_kernel(nd_item: NdItem, a):
            increment_value(nd_item, a)


        N = 16
        b = dpnp.ones(N, dtype=dpnp.int32)

        dpex_exp.call_kernel(another_kernel, dpex.NdRange((N,), (N,)), b)
    """
    dispatcher = resolve_dispatcher_from_str(DPEX_KERNEL_EXP_TARGET_NAME)

    if "_compilation_mode" in options:
        user_compilation_mode = options["_compilation_mode"]
        warn(
            "_compilation_mode is an internal flag that should not be set "
            "in the decorator. The decorator defined option "
            f"{user_compilation_mode} is going to be ignored."
        )
    options["_compilation_mode"] = CompilationMode.DEVICE_FUNC

    func, sigs = _parse_func_or_sig(function_or_signature)
    for sig in sigs:
        if isinstance(sig, str):
            raise NotImplementedError(
                "Specifying signatures as string is not yet supported"
            )

    def _kernel_dispatcher(pyfunc):
        disp: SPIRVKernelDispatcher = dispatcher(
            pyfunc=pyfunc,
            targetoptions=options,
        )

        if len(sigs) > 0:
            with typeinfer.register_dispatcher(disp):
                for sig in sigs:
                    disp.compile(sig)
                disp.disable_compile()

        return disp

    if func is None:
        return _kernel_dispatcher

    return _kernel_dispatcher(function_or_signature)


# ----------------- Experimental dpjit decorator ------------------------------#


def dpjit(*args, **kws):
    if "nopython" in kws and kws["nopython"] is not True:
        warn("nopython is set for dpjit and is ignored", RuntimeWarning)
    if "forceobj" in kws:
        warn("forceobj is set for dpjit and is ignored", RuntimeWarning)
        del kws["forceobj"]
    if "pipeline_class" in kws:
        warn("pipeline class is set for dpjit and is ignored", RuntimeWarning)
        del kws["pipeline_class"]

    use_mlir = kws.pop("use_mlir", bool(config.USE_MLIR))

    kws.update({"nopython": True})
    kws.update({"parallel": True})
    kws.update({"pipeline_class": get_compiler(use_mlir)})

    kws.update({"_target": DPEX_TARGET_NAME})

    return decorators.jit(*args, **kws)


# add it to the decorator registry, this is so e.g. @overload can look up a
# JIT function to do the compilation work.
jit_registry[target_registry[DPEX_TARGET_NAME]] = dpjit
jit_registry[target_registry[DPEX_KERNEL_EXP_TARGET_NAME]] = device_func

1	# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
2	#
3	# SPDX-License-Identifier: Apache-2.0
4
5	import inspect	1✔
6	from warnings import warn	1✔
7
8	from numba.core import decorators, sigutils, typeinfer	1✔
9	from numba.core.target_extension import (	1✔
10	jit_registry,
11	resolve_dispatcher_from_str,
12	target_registry,
13	)
14
15	from numba_dpex.core import config	1✔
16	from numba_dpex.core.pipelines.dpjit_compiler import get_compiler	1✔
17	from numba_dpex.core.targets.dpjit_target import DPEX_TARGET_NAME	1✔
18	from numba_dpex.experimental.target import DPEX_KERNEL_EXP_TARGET_NAME	1✔
19	from numba_dpex.kernel_api_impl.spirv.dispatcher import SPIRVKernelDispatcher	1✔
20	from numba_dpex.kernel_api_impl.spirv.target import CompilationMode	1✔
21
22
23	def _parse_func_or_sig(signature_or_function):	1✔
24	# Handle signature (borrowed from numba). swapped signature and list check
25	if signature_or_function is None:	1✔
26	# No signature, no function
27	pyfunc = None	1✔
28	sigs = []	1✔
29	elif sigutils.is_signature(signature_or_function):	1✔
30	# A single signature is passed
31	pyfunc = None	1✔
32	sigs = [signature_or_function]	1✔
33	elif isinstance(signature_or_function, list):	1✔
34	# A list of signatures is passed
35	pyfunc = None	1✔
36	sigs = signature_or_function	1✔
37	else:
38	# A function is passed
39	pyfunc = signature_or_function	1✔
40	sigs = []	1✔
41
42	return pyfunc, sigs	1✔
43
44
45	def kernel(function_or_signature=None, **options):	1✔
46	"""A decorator to compile a function written using :py:mod:`numba_dpex.kernel_api`.
47
48	The ``kernel`` decorator triggers the compilation of a function written
49	using the data-parallel kernel programming API exposed by
50	:py:mod:`numba_dpex.kernel_api`. Such a function is conceptually
51	equivalent to a kernel function written in the C++ SYCL eDSL. The
52	decorator will compile the function based on the types of the arguments
53	to a SPIR-V binary that can be executed either on OpenCL CPU, GPU
54	devices or Intel Level Zero GPU devices.
55
56	Any function to be compilable using the kernel decorator should
57	adhere to the following semantic rules:
58
59	- The first argument to the function should be either an instance of the
60	:class:`numba_dpex.kernel_api.Item` class or an instance of the
61	:class:`numba_dpex.kernel_api.NdItem`.
62
63	- The function should not return any value.
64
65	- The function should have at least one array type argument that can
66	either be an instance of ``dpnp.ndarray`` or an instance of
67	``dpctl.tensor.usm_ndarray``.
68
69
70	Args:
71	signature_or_function (optional): An optional signature or list of
72	signatures for which a function is to be compiled. Passing in a
73	signature "specializes" the decorated function and no other versions
74	of the function will be compiled. A function can also be
75	directly passed instead of a signature and the signature will get
76	inferred from the function. The actual compilation happens on every
77	invocation of the :func:`numba_dpex.experimental.call_kernel`
78	function where the decorated function is passed in as an argument
79	along with the argument values for the decorated function.
80	options (optional):
81	- debug (bool): Whether the compilation should happen in debug
82	mode. (Default = False)
83	- inline_threshold (int): Specifies the level of inlining that
84	the compiler should attempt. (Default = 2)
85	Returns:
86	An instance of
87	:class:`numba_dpex.kernel_api_impl.spirv.dispatcher.KernelDispatcher`.
88	The ``KernelDispatcher`` object compiles the decorated function when
89	passed in to :func:`numba_dpex.experimental.call_kernel`.
90
91	Examples:
92
93	1. Decorate a function and pass it to ``call_kernel`` for compilation and
94	execution.
95
96	.. code-block:: python
97
98	import dpnp
99	import numba_dpex as dpex
100	from numba_dpex import kernel_api as kapi
101
102
103	# Data parallel kernel implementing vector sum
104	@dpex.kernel
105	def vecadd(item: kapi.Item, a, b, c):
106	i = item.get_id(0)
107	c[i] = a[i] + b[i]
108
109
110	N = 1024
111	a = dpnp.ones(N)
112	b = dpnp.ones_like(a)
113	c = dpnp.zeros_like(a)
114	dpex.call_kernel(vecadd, kapi.Range(N), a, b, c)
115
116	2. Specializes a kernel and then compiles it directly before executing it
117	via ``call_kernel``. The kernel is specialized to expect a 1-D
118	``dpnp.ndarray`` with either ``float32`` type elements or ``int64`` type
119	elements.
120
121	.. code-block:: python
122
123	import dpnp
124	import numba_dpex as dpex
125	from numba_dpex import kernel_api as kapi
126	from numba_dpex import DpnpNdArray, float32, int64
127	from numba_dpex.core.types.kernel_api.index_space_ids import ItemType
128
129	i64arrty = DpnpNdArray(ndim=1, dtype=int64, layout="C")
130	f32arrty = DpnpNdArray(ndim=1, dtype=float32, layout="C")
131	item_ty = ItemType(ndim=1)
132
133	specialized_kernel = dpex.kernel(
134	[
135	(item_ty, i64arrty, i64arrty, i64arrty),
136	(item_ty, f32arrty, f32arrty, f32arrty),
137	]
138	)
139
140
141	def vecadd(item: kapi.Item, a, b, c):
142	i = item.get_id(0)
143	c[i] = a[i] + b[i]
144
145
146	# Compile all specializations for vecadd
147	precompiled_kernels = specialized_kernel(vecadd)
148	N = 1024
149	a = dpnp.ones(N, dtype=dpnp.int64)
150	b = dpnp.ones_like(a)
151	c = dpnp.zeros_like(a)
152	# Call a specific pre-compiled version of vecadd
153	dpex.call_kernel(precompiled_kernels, kapi.Range(N), a, b, c)
154
155	"""
156
157	# dispatcher is a type:
158	# <class 'numba_dpex.experimental.kernel_dispatcher.KernelDispatcher'>
159	dispatcher = resolve_dispatcher_from_str(DPEX_KERNEL_EXP_TARGET_NAME)	1✔
160	if "_compilation_mode" in options:	1✔
161	user_compilation_mode = options["_compilation_mode"]	1✔
162	warn(	1✔
163	"_compilation_mode is an internal flag that should not be set "
164	"in the decorator. The decorator defined option "
165	f"{user_compilation_mode} is going to be ignored."
166	)
167	options["_compilation_mode"] = CompilationMode.KERNEL	1✔
168
169	# TODO: The options need to be evaluated and checked here like it is
170	# done in numba.core.decorators.jit
171
172	func, sigs = _parse_func_or_sig(function_or_signature)	1✔
173	for sig in sigs:	1✔
174	if isinstance(sig, str):	1✔
175	raise NotImplementedError(
176	"Specifying signatures as string is not yet supported"
177	)
178
179	def _kernel_dispatcher(pyfunc):	1✔
180	disp: SPIRVKernelDispatcher = dispatcher(	1✔
181	pyfunc=pyfunc,
182	targetoptions=options,
183	)
184
185	if len(sigs) > 0:	1✔
186	with typeinfer.register_dispatcher(disp):	1✔
187	for sig in sigs:	1✔
188	disp.compile(sig)	1✔
189	disp.disable_compile()	1✔
190
191	return disp	1✔
192
193	if func is None:	1✔
194	return _kernel_dispatcher	1✔
195
196	if not inspect.isfunction(func):	1✔
197	raise ValueError(	1✔
198	"Argument passed to the kernel decorator is neither a "
199	"function object, nor a signature. If you are trying to "
200	"specialize the kernel that takes a single argument, specify "
201	"the return type as None explicitly."
202	)
203	return _kernel_dispatcher(func)	1✔
204
205
206	def device_func(function_or_signature=None, **options):	1✔
207	"""Compiles a device-callable function that can be only invoked from a kernel.
208
209	The decorator is used to express auxiliary device-only functions that can
210	be called from a kernel or another device function, but are not callable
211	from the host. This decorator :func:`numba_dpex.experimental.device_func`
212	has no direct analogue in SYCL and primarily is provided to help programmers
213	make their kapi applications modular.
214
215	A ``device_func`` decorated function does not require the first argument to
216	be a :class:`numba_dpex.kernel_api.Item` object or a
217	:class:`numba_dpex.kernel_api.NdItem` object, and unlike a ``kernel``
218	decorated function is allowed to return any value.
219	All :py:mod:`numba_dpex.kernel_api` functionality can be used in a
220	``device_func`` decorated function.
221
222	The decorator is also used to compile overloads in the ``DpexKernelTarget``.
223
224	A ``device_func`` decorated function is not compiled down to device binary
225	and instead is compiled down to LLVM IR. Final compilation to binary happens
226	when the function is invoked from a ``kernel`` decorated function. The
227	compilation happens this was to allow a ``device_func`` decorated function
228	to be internally linked into the kernel module at the LLVM level, leading to
229	more optimization opportunities.
230
231	Args:
232	signature_or_function (optional): An optional signature or list of
233	signatures for which a function is to be compiled. Passing in a
234	signature "specializes" the decorated function and no other versions
235	of the function will be compiled. A function can also be
236	directly passed instead of a signature and the signature will get
237	inferred from the function. The actual compilation happens on every
238	invocation of the decorated function from another ``device_func`` or
239	``kernel`` decorated function.
240	options (optional):
241	- debug (bool): Whether the compilation should happen in debug
242	mode. (Default = False)
243	- inline_threshold (int): Specifies the level of inlining that
244	the compiler should attempt. (Default = 2)
245
246	Returns:
247	An instance of
248	:class:`numba_dpex.kernel_api_impl.spirv.dispatcher.KernelDispatcher`.
249	The ``KernelDispatcher`` object compiles the decorated function when
250	it is called from another function.
251
252
253	Example:
254
255	.. code-block:: python
256
257	import dpnp
258
259	from numba_dpex import experimental as dpex_exp
260	from numba_dpex import kernel_api as kapi
261
262
263	@dpex_exp.device_func
264	def increment_value(nd_item: NdItem, a):
265	i = nd_item.get_global_id(0)
266
267	a[i] += 1
268	group_barrier(nd_item.get_group(), MemoryScope.DEVICE)
269
270	if i == 0:
271	for idx in range(1, a.size):
272	a[0] += a[idx]
273
274
275	@dpex_exp.kernel
276	def another_kernel(nd_item: NdItem, a):
277	increment_value(nd_item, a)
278
279
280	N = 16
281	b = dpnp.ones(N, dtype=dpnp.int32)
282
283	dpex_exp.call_kernel(another_kernel, dpex.NdRange((N,), (N,)), b)
284	"""
285	dispatcher = resolve_dispatcher_from_str(DPEX_KERNEL_EXP_TARGET_NAME)	1✔
286
287	if "_compilation_mode" in options:	1!
UNCOV 288	user_compilation_mode = options["_compilation_mode"]	×
UNCOV 289	warn(	×
290	"_compilation_mode is an internal flag that should not be set "
291	"in the decorator. The decorator defined option "
292	f"{user_compilation_mode} is going to be ignored."
293	)
294	options["_compilation_mode"] = CompilationMode.DEVICE_FUNC	1✔
295
296	func, sigs = _parse_func_or_sig(function_or_signature)	1✔
297	for sig in sigs:	1✔
298	if isinstance(sig, str):	1✔
299	raise NotImplementedError(
300	"Specifying signatures as string is not yet supported"
301	)
302
303	def _kernel_dispatcher(pyfunc):	1✔
304	disp: SPIRVKernelDispatcher = dispatcher(	1✔
305	pyfunc=pyfunc,
306	targetoptions=options,
307	)
308
309	if len(sigs) > 0:	1✔
310	with typeinfer.register_dispatcher(disp):	1✔
311	for sig in sigs:	1✔
312	disp.compile(sig)	1✔
313	disp.disable_compile()	1✔
314
315	return disp	1✔
316
317	if func is None:	1✔
318	return _kernel_dispatcher	1✔
319
320	return _kernel_dispatcher(function_or_signature)	1✔
321
322
323	# ----------------- Experimental dpjit decorator ------------------------------#
324
325
326	def dpjit(args, *kws):	1✔
327	if "nopython" in kws and kws["nopython"] is not True:	1!
328	warn("nopython is set for dpjit and is ignored", RuntimeWarning)	×
329	if "forceobj" in kws:	1!
UNCOV 330	warn("forceobj is set for dpjit and is ignored", RuntimeWarning)	×
331	del kws["forceobj"]	×
332	if "pipeline_class" in kws:	1!
UNCOV 333	warn("pipeline class is set for dpjit and is ignored", RuntimeWarning)	×
UNCOV 334	del kws["pipeline_class"]	×
335
336	use_mlir = kws.pop("use_mlir", bool(config.USE_MLIR))	1✔
337
338	kws.update({"nopython": True})	1✔
339	kws.update({"parallel": True})	1✔
340	kws.update({"pipeline_class": get_compiler(use_mlir)})	1✔
341
342	kws.update({"_target": DPEX_TARGET_NAME})	1✔
343
344	return decorators.jit(args, *kws)	1✔
345
346
347	# add it to the decorator registry, this is so e.g. @overload can look up a
348	# JIT function to do the compilation work.
349	jit_registry[target_registry[DPEX_TARGET_NAME]] = dpjit	1✔
350	jit_registry[target_registry[DPEX_KERNEL_EXP_TARGET_NAME]] = device_func	1✔

IntelPython / numba-dpex / 8423792464

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