7124424362

Committed 07 Dec 2023 06:05AM UTC coverage: 79.24% (+0.1%) from 79.106%

Build # 7124424362

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #1230 from IntelPython/experimental/enable_overloads

Enables adding overload to DpexExpKernelTarget and fully inline them into the final module.

Run Details

1187 of 1763 branches covered (0.0%)

Branch coverage included in aggregate %.

57 of 68 new or added lines in 6 files covered. (83.82%)

1 existing line in 1 file now uncovered.

5233 of 6339 relevant lines covered (82.55%)

0.83 hits per line

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

78.87

/numba_dpex/experimental/kernel_dispatcher.py

# SPDX-FileCopyrightText: 2023 Intel Corporation
#
# SPDX-License-Identifier: Apache-2.0

"""Implements a new numba dispatcher class and a compiler class to compile and
call numba_dpex.kernel decorated function.
"""
from collections import namedtuple
from contextlib import ExitStack
from typing import Tuple

import numba.core.event as ev
from numba.core import errors, sigutils, types
from numba.core.compiler import CompileResult, Flags
from numba.core.compiler_lock import global_compiler_lock
from numba.core.dispatcher import Dispatcher, _FunctionCompiler
from numba.core.target_extension import dispatcher_registry, target_registry
from numba.core.types import void
from numba.core.typing.typeof import Purpose, typeof

from numba_dpex import config, spirv_generator
from numba_dpex.core.exceptions import (
    ExecutionQueueInferenceError,
    KernelHasReturnValueError,
    UnsupportedKernelArgumentError,
)
from numba_dpex.core.pipelines import kernel_compiler
from numba_dpex.core.targets.kernel_target import CompilationMode
from numba_dpex.core.types import DpnpNdArray

from .target import DPEX_KERNEL_EXP_TARGET_NAME, dpex_exp_kernel_target

_KernelModule = namedtuple("_KernelModule", ["kernel_name", "kernel_bitcode"])

_KernelCompileResult = namedtuple(
    "_KernelCompileResult", CompileResult._fields + ("kernel_device_ir_module",)
)


class _KernelCompiler(_FunctionCompiler):
    """A special compiler class used to compile numba_dpex.kernel decorated
    functions.
    """

    def check_queue_equivalence_of_args(
        self, py_func_name: str, args: [types.Type, ...]
    ):
        """Evaluates if all DpnpNdArray arguments passed to a kernel function
        has the same DpctlSyclQueue type.

        Args:
            py_func_name (str): Name of the kernel that is being evaluated
            args (types.Type, ...]): List of numba inferred types for each
            argument passed to the kernel

        Raises:
            ExecutionQueueInferenceError: If all DpnpNdArray were not allocated
            on the same dpctl.SyclQueue
            ExecutionQueueInferenceError: If there were not DpnpNdArray
            arguments passed to the kernel.
        """
        common_queue = None

        for arg in args:
            if isinstance(arg, DpnpNdArray):
                if common_queue is None:
                    common_queue = arg.queue
                elif common_queue != arg.queue:
                    raise ExecutionQueueInferenceError(
                        kernel_name=py_func_name, usmarray_argnum_list=[]
                    )

        if common_queue is None:
            raise ExecutionQueueInferenceError(
                kernel_name=py_func_name, usmarray_argnum_list=None
            )

    def _compile_to_spirv(
        self, kernel_library, kernel_fndesc, kernel_targetctx
    ):
        kernel_func = kernel_library.get_function(kernel_fndesc.llvm_func_name)

        # Create a spir_kernel wrapper function
        kernel_fn = kernel_targetctx.prepare_spir_kernel(
            kernel_func, kernel_fndesc.argtypes
        )
        # Get the compiler flags that were passed through the target descriptor
        flags = Flags()
        self.targetdescr.options.parse_as_flags(flags, self.targetoptions)

        # If the inline_threshold option was set then set the property in the
        # kernel_library to force inlining ``overload`` calls into a kernel.
        inline_threshold = flags.inline_threshold  # pylint: disable=E1101
        kernel_library.inline_threshold = inline_threshold

        # Call finalize on the LLVM module. Finalization will result in
        # all linking libraries getting linked together and final optimization
        # including inlining of functions if an inlining level is specified.
        kernel_library.finalize()
        # Compiled the LLVM IR to SPIR-V
        kernel_spirv_module = spirv_generator.llvm_to_spirv(
            kernel_targetctx,
            kernel_library.final_module,
            kernel_library.final_module.as_bitcode(),
        )
        return _KernelModule(
            kernel_name=kernel_fn.name, kernel_bitcode=kernel_spirv_module
        )

    def compile(self, args, return_type):
        status, kcres = self._compile_cached(args, return_type)
        if status:
            return kcres

        raise kcres

    def _compile_cached(
        self, args, return_type: types.Type
    ) -> Tuple[bool, _KernelCompileResult]:
        """Compiles the kernel function to bitcode and generates a host-callable
        wrapper to submit the kernel to a SYCL queue.

        The LLVM IR generated for the kernel function is available in the
        CompileResult objected returned by
        numba_dpex.core.pipeline.kernel_compiler.KernelCompiler.

        Once the kernel decorated function is compiled down to LLVM IR, the
        following steps are performed:

            a) compile the IR into SPIR-V kernel
            b) generate a host callable wrapper function that will create a
               sycl::kernel_bundle from the SPIR-V and then submits the
               kernel_bundle to a sycl::queue
            c) create a cpython_wrapper_function for the host callable wrapper
               function.
            d) create a cfunc_wrapper_function to make the host callable wrapper
               function callable inside another JIT-compiled function.

        Args:
            args (tuple(types.Type)): A tuple of numba.core.Type instances each
            representing the numba-inferred type of a kernel argument.

            return_type (types.Type): The numba-inferred type of the returned
            value from the kernel. Should always be types.NoneType.

        Returns:
            CompileResult: A CompileResult object storing the LLVM library for
            the host-callable wrapper function.
        """
        key = tuple(args), return_type
        try:
            return False, self._failed_cache[key]
        except KeyError:
            pass

        try:
            cres: CompileResult = self._compile_core(args, return_type)

            if (
                self.targetoptions["_compilation_mode"]
                == CompilationMode.KERNEL
            ):
                kernel_device_ir_module: _KernelModule = self._compile_to_spirv(
                    cres.library, cres.fndesc, cres.target_context
                )
            else:
                kernel_device_ir_module = None

            kcres_attrs = []

            for cres_field in cres._fields:
                cres_attr = getattr(cres, cres_field)
                if cres_field == "entry_point":
                    if cres_attr is not None:
                        raise AssertionError(
                            "Compiled kernel and device_func should be "
                            "compiled with compile_cfunc option turned off"
                        )
                    cres_attr = cres.fndesc.qualname
                kcres_attrs.append(cres_attr)

            kcres_attrs.append(kernel_device_ir_module)

            if config.DUMP_KERNEL_LLVM:
                with open(
                    cres.fndesc.llvm_func_name + ".ll",
                    "w",
                    encoding="UTF-8",
                ) as f:
                    f.write(cres.library.final_module)

        except errors.TypingError as e:
            self._failed_cache[key] = e
            return False, e

        return True, _KernelCompileResult(*kcres_attrs)


class KernelDispatcher(Dispatcher):
    """Dispatcher class designed to compile kernel decorated functions. The
    dispatcher inherits the Numba Dispatcher class, but has a different
    compilation strategy. Instead of compiling a kernel decorated function to
    an executable binary, the dispatcher compiles it to SPIR-V and then caches
    that SPIR-V bitcode.

    """

    targetdescr = dpex_exp_kernel_target
    _fold_args = False

    Dispatcher._impl_kinds["kernel"] = _KernelCompiler

    def __init__(
        self,
        pyfunc,
        local_vars_to_numba_types=None,
        targetoptions=None,
        pipeline_class=kernel_compiler.KernelCompiler,
    ):
        if targetoptions is None:
            targetoptions = {}

        if local_vars_to_numba_types is None:
            local_vars_to_numba_types = {}

        targetoptions["nopython"] = True
        targetoptions["experimental"] = True

        self._kernel_name = pyfunc.__name__

        super().__init__(
            py_func=pyfunc,
            locals=local_vars_to_numba_types,
            impl_kind="kernel",
            targetoptions=targetoptions,
            pipeline_class=pipeline_class,
        )

    def typeof_pyval(self, val):
        """
        Resolve the Numba type of Python value *val*.
        This is called from numba._dispatcher as a fallback if the native code
        cannot decide the type.
        """
        # Not going through the resolve_argument_type() indirection
        # can save a couple µs.
        try:
            tp = typeof(val, Purpose.argument)
            if isinstance(tp, types.Array) and not isinstance(tp, DpnpNdArray):
                raise UnsupportedKernelArgumentError(
                    type=str(type(val)), value=val
                )
        except ValueError:
            tp = types.pyobject
        else:
            if tp is None:
                tp = types.pyobject
        self._types_active_call.append(tp)
        return tp

    def add_overload(self, cres):
        args = tuple(cres.signature.args)
        self.overloads[args] = cres

    def get_overload_device_ir(self, sig):
        """
        Return the compiled device bitcode for the given signature.
        """
        args, _ = sigutils.normalize_signature(sig)
        return self.overloads[tuple(args)].kernel_device_ir_module

    def compile(self, sig) -> any:
        disp = self._get_dispatcher_for_current_target()
        if disp is not self:
            return disp.compile(sig)

        with ExitStack() as scope:
            cres = None

            def cb_compiler(dur):
                if cres is not None:
                    self._callback_add_compiler_timer(dur, cres)

            def cb_llvm(dur):
                if cres is not None:
                    self._callback_add_llvm_timer(dur, cres)

            scope.enter_context(
                ev.install_timer("numba:compiler_lock", cb_compiler)
            )
            scope.enter_context(ev.install_timer("numba:llvm_lock", cb_llvm))
            scope.enter_context(global_compiler_lock)

            if not self._can_compile:
                raise RuntimeError("compilation disabled")
            # Use counter to track recursion compilation depth
            with self._compiling_counter:
                args, return_type = sigutils.normalize_signature(sig)

                if (
                    self.targetoptions["_compilation_mode"]
                    == CompilationMode.KERNEL
                ):
                    # Compute follows data based queue equivalence is only
                    # evaluated for kernel functions whose arguments are
                    # supposed to be arrays. For device_func decorated
                    # functions, the arguments can be scalar and we skip queue
                    # equivalence check.
                    try:
                        self._compiler.check_queue_equivalence_of_args(
                            self._kernel_name, args
                        )
                    except ExecutionQueueInferenceError as eqie:
                        raise eqie

                    # A function being compiled in the KERNEL compilation mode
                    # cannot have a non-void return value
                    if return_type and return_type != void:
                        raise KernelHasReturnValueError(
                            kernel_name=None, return_type=return_type, sig=sig
                        )

                # Don't recompile if signature already exists
                existing = self.overloads.get(tuple(args))
                if existing is not None:
                    return existing.entry_point

                # TODO: Enable caching
                # Add code to enable on disk caching of a binary spirv kernel.
                # Refer: https://github.com/IntelPython/numba-dpex/issues/1197
                self._cache_misses[sig] += 1
                ev_details = {
                    "dispatcher": self,
                    "args": args,
                    "return_type": return_type,
                }
                with ev.trigger_event("numba_dpex:compile", data=ev_details):
                    try:
                        kcres: _KernelCompileResult = self._compiler.compile(
                            args, return_type
                        )
                    except errors.ForceLiteralArg as e:

                        def folded(args, kws):
                            return self._compiler.fold_argument_types(
                                args, kws
                            )[1]

                        raise e.bind_fold_arguments(folded)
                    self.add_overload(kcres)

                    kcres.target_context.insert_user_function(
                        kcres.entry_point, kcres.fndesc, [kcres.library]
                    )

                # TODO: enable caching of kernel_module
                # https://github.com/IntelPython/numba-dpex/issues/1197

                return kcres.entry_point

    def __getitem__(self, args):
        """Square-bracket notation for configuring launch arguments is not
        supported.
        """

        raise NotImplementedError

    def __call__(self, *args, **kw_args):
        """Functor to launch a kernel."""

        raise NotImplementedError


_dpex_target = target_registry[DPEX_KERNEL_EXP_TARGET_NAME]
dispatcher_registry[_dpex_target] = KernelDispatcher

1	# SPDX-FileCopyrightText: 2023 Intel Corporation
2	#
3	# SPDX-License-Identifier: Apache-2.0
4
5	"""Implements a new numba dispatcher class and a compiler class to compile and	1✔
6	call numba_dpex.kernel decorated function.
7	"""
8	from collections import namedtuple	1✔
9	from contextlib import ExitStack	1✔
10	from typing import Tuple	1✔
11
12	import numba.core.event as ev	1✔
13	from numba.core import errors, sigutils, types	1✔
14	from numba.core.compiler import CompileResult, Flags	1✔
15	from numba.core.compiler_lock import global_compiler_lock	1✔
16	from numba.core.dispatcher import Dispatcher, _FunctionCompiler	1✔
17	from numba.core.target_extension import dispatcher_registry, target_registry	1✔
18	from numba.core.types import void	1✔
19	from numba.core.typing.typeof import Purpose, typeof	1✔
20
21	from numba_dpex import config, spirv_generator	1✔
22	from numba_dpex.core.exceptions import (	1✔
23	ExecutionQueueInferenceError,
24	KernelHasReturnValueError,
25	UnsupportedKernelArgumentError,
26	)
27	from numba_dpex.core.pipelines import kernel_compiler	1✔
28	from numba_dpex.core.targets.kernel_target import CompilationMode	1✔
29	from numba_dpex.core.types import DpnpNdArray	1✔
30
31	from .target import DPEX_KERNEL_EXP_TARGET_NAME, dpex_exp_kernel_target	1✔
32
33	_KernelModule = namedtuple("_KernelModule", ["kernel_name", "kernel_bitcode"])	1✔
34
35	_KernelCompileResult = namedtuple(	1✔
36	"_KernelCompileResult", CompileResult._fields + ("kernel_device_ir_module",)
37	)
38
39
40	class _KernelCompiler(_FunctionCompiler):	1✔
41	"""A special compiler class used to compile numba_dpex.kernel decorated
42	functions.
43	"""
44
45	def check_queue_equivalence_of_args(	1✔
46	self, py_func_name: str, args: [types.Type, ...]
47	):
48	"""Evaluates if all DpnpNdArray arguments passed to a kernel function
49	has the same DpctlSyclQueue type.
50
51	Args:
52	py_func_name (str): Name of the kernel that is being evaluated
53	args (types.Type, ...]): List of numba inferred types for each
54	argument passed to the kernel
55
56	Raises:
57	ExecutionQueueInferenceError: If all DpnpNdArray were not allocated
58	on the same dpctl.SyclQueue
59	ExecutionQueueInferenceError: If there were not DpnpNdArray
60	arguments passed to the kernel.
61	"""
62	common_queue = None	1✔
63
64	for arg in args:	1✔
65	if isinstance(arg, DpnpNdArray):	1✔
66	if common_queue is None:	1✔
67	common_queue = arg.queue	1✔
68	elif common_queue != arg.queue:	1✔
69	raise ExecutionQueueInferenceError(	1✔
70	kernel_name=py_func_name, usmarray_argnum_list=[]
71	)
72
73	if common_queue is None:	1✔
74	raise ExecutionQueueInferenceError(	1✔
75	kernel_name=py_func_name, usmarray_argnum_list=None
76	)
77
78	def _compile_to_spirv(	1✔
79	self, kernel_library, kernel_fndesc, kernel_targetctx
80	):
81	kernel_func = kernel_library.get_function(kernel_fndesc.llvm_func_name)	1✔
82
83	# Create a spir_kernel wrapper function
84	kernel_fn = kernel_targetctx.prepare_spir_kernel(	1✔
85	kernel_func, kernel_fndesc.argtypes
86	)
87	# Get the compiler flags that were passed through the target descriptor
88	flags = Flags()	1✔
89	self.targetdescr.options.parse_as_flags(flags, self.targetoptions)	1✔
90
91	# If the inline_threshold option was set then set the property in the
92	# kernel_library to force inlining ``overload`` calls into a kernel.
93	inline_threshold = flags.inline_threshold # pylint: disable=E1101	1✔
94	kernel_library.inline_threshold = inline_threshold	1✔
95
96	# Call finalize on the LLVM module. Finalization will result in
97	# all linking libraries getting linked together and final optimization
98	# including inlining of functions if an inlining level is specified.
99	kernel_library.finalize()	1✔
100	# Compiled the LLVM IR to SPIR-V
101	kernel_spirv_module = spirv_generator.llvm_to_spirv(	1✔
102	kernel_targetctx,
103	kernel_library.final_module,
104	kernel_library.final_module.as_bitcode(),
105	)
106	return _KernelModule(	1✔
107	kernel_name=kernel_fn.name, kernel_bitcode=kernel_spirv_module
108	)
109
110	def compile(self, args, return_type):	1✔
111	status, kcres = self._compile_cached(args, return_type)	1✔
112	if status:	1!
113	return kcres	1✔
114
115	raise kcres	×
116
117	def _compile_cached(	1✔
118	self, args, return_type: types.Type
119	) -> Tuple[bool, _KernelCompileResult]:
120	"""Compiles the kernel function to bitcode and generates a host-callable
121	wrapper to submit the kernel to a SYCL queue.
122
123	The LLVM IR generated for the kernel function is available in the
124	CompileResult objected returned by
125	numba_dpex.core.pipeline.kernel_compiler.KernelCompiler.
126
127	Once the kernel decorated function is compiled down to LLVM IR, the
128	following steps are performed:
129
130	a) compile the IR into SPIR-V kernel
131	b) generate a host callable wrapper function that will create a
132	sycl::kernel_bundle from the SPIR-V and then submits the
133	kernel_bundle to a sycl::queue
134	c) create a cpython_wrapper_function for the host callable wrapper
135	function.
136	d) create a cfunc_wrapper_function to make the host callable wrapper
137	function callable inside another JIT-compiled function.
138
139	Args:
140	args (tuple(types.Type)): A tuple of numba.core.Type instances each
141	representing the numba-inferred type of a kernel argument.
142
143	return_type (types.Type): The numba-inferred type of the returned
144	value from the kernel. Should always be types.NoneType.
145
146	Returns:
147	CompileResult: A CompileResult object storing the LLVM library for
148	the host-callable wrapper function.
149	"""
150	key = tuple(args), return_type	1✔
151	try:	1✔
152	return False, self._failed_cache[key]	1✔
153	except KeyError:	1✔
154	pass	1✔
155
156	try:	1✔
157	cres: CompileResult = self._compile_core(args, return_type)	1✔
158
159	if (	1✔
160	self.targetoptions["_compilation_mode"]
161	== CompilationMode.KERNEL
162	):
163	kernel_device_ir_module: _KernelModule = self._compile_to_spirv(	1✔
164	cres.library, cres.fndesc, cres.target_context
165	)
166	else:
167	kernel_device_ir_module = None	1✔
168
169	kcres_attrs = []	1✔
170
171	for cres_field in cres._fields:	1✔
172	cres_attr = getattr(cres, cres_field)	1✔
173	if cres_field == "entry_point":	1✔
174	if cres_attr is not None:	1!
175	raise AssertionError(	×
176	"Compiled kernel and device_func should be "
177	"compiled with compile_cfunc option turned off"
178	)
179	cres_attr = cres.fndesc.qualname	1✔
180	kcres_attrs.append(cres_attr)	1✔
181
182	kcres_attrs.append(kernel_device_ir_module)	1✔
183
184	if config.DUMP_KERNEL_LLVM:	1!
185	with open(	×
186	cres.fndesc.llvm_func_name + ".ll",
187	"w",
188	encoding="UTF-8",
189	) as f:
190	f.write(cres.library.final_module)	×
191
192	except errors.TypingError as e:	×
193	self._failed_cache[key] = e	×
194	return False, e	×
195
196	return True, _KernelCompileResult(*kcres_attrs)	1✔
197
198
199	class KernelDispatcher(Dispatcher):	1✔
200	"""Dispatcher class designed to compile kernel decorated functions. The
201	dispatcher inherits the Numba Dispatcher class, but has a different
202	compilation strategy. Instead of compiling a kernel decorated function to
203	an executable binary, the dispatcher compiles it to SPIR-V and then caches
204	that SPIR-V bitcode.
205
206	"""
207
208	targetdescr = dpex_exp_kernel_target	1✔
209	_fold_args = False	1✔
210
211	Dispatcher._impl_kinds["kernel"] = _KernelCompiler	1✔
212
213	def __init__(	1✔
214	self,
215	pyfunc,
216	local_vars_to_numba_types=None,
217	targetoptions=None,
218	pipeline_class=kernel_compiler.KernelCompiler,
219	):
220	if targetoptions is None:	1!
221	targetoptions = {}	×
222
223	if local_vars_to_numba_types is None:	1!
224	local_vars_to_numba_types = {}	1✔
225
226	targetoptions["nopython"] = True	1✔
227	targetoptions["experimental"] = True	1✔
228
229	self._kernel_name = pyfunc.__name__	1✔
230
231	super().__init__(	1✔
232	py_func=pyfunc,
233	locals=local_vars_to_numba_types,
234	impl_kind="kernel",
235	targetoptions=targetoptions,
236	pipeline_class=pipeline_class,
237	)
238
239	def typeof_pyval(self, val):	1✔
240	"""
241	Resolve the Numba type of Python value val.
242	This is called from numba._dispatcher as a fallback if the native code
243	cannot decide the type.
244	"""
245	# Not going through the resolve_argument_type() indirection
246	# can save a couple µs.
247	try:	×
248	tp = typeof(val, Purpose.argument)	×
249	if isinstance(tp, types.Array) and not isinstance(tp, DpnpNdArray):	×
250	raise UnsupportedKernelArgumentError(	×
251	type=str(type(val)), value=val
252	)
253	except ValueError:	×
254	tp = types.pyobject	×
255	else:
256	if tp is None:	×
257	tp = types.pyobject	×
258	self._types_active_call.append(tp)	×
259	return tp	×
260
261	def add_overload(self, cres):	1✔
262	args = tuple(cres.signature.args)	1✔
263	self.overloads[args] = cres	1✔
264
265	def get_overload_device_ir(self, sig):	1✔
266	"""
267	Return the compiled device bitcode for the given signature.
268	"""
269	args, _ = sigutils.normalize_signature(sig)	1✔
270	return self.overloads[tuple(args)].kernel_device_ir_module	1✔
271
272	def compile(self, sig) -> any:	1✔
273	disp = self._get_dispatcher_for_current_target()	1✔
274	if disp is not self:	1!
275	return disp.compile(sig)	×
276
277	with ExitStack() as scope:	1✔
278	cres = None	1✔
279
280	def cb_compiler(dur):	1✔
281	if cres is not None:	1!
282	self._callback_add_compiler_timer(dur, cres)	×
283
284	def cb_llvm(dur):	1✔
285	if cres is not None:	1!
286	self._callback_add_llvm_timer(dur, cres)	×
287
288	scope.enter_context(	1✔
289	ev.install_timer("numba:compiler_lock", cb_compiler)
290	)
291	scope.enter_context(ev.install_timer("numba:llvm_lock", cb_llvm))	1✔
292	scope.enter_context(global_compiler_lock)	1✔
293
294	if not self._can_compile:	1!
295	raise RuntimeError("compilation disabled")	×
296	# Use counter to track recursion compilation depth
297	with self._compiling_counter:	1✔
298	args, return_type = sigutils.normalize_signature(sig)	1✔
299
300	if (	1✔
301	self.targetoptions["_compilation_mode"]
302	== CompilationMode.KERNEL
303	):
304	# Compute follows data based queue equivalence is only
305	# evaluated for kernel functions whose arguments are
306	# supposed to be arrays. For device_func decorated
307	# functions, the arguments can be scalar and we skip queue
308	# equivalence check.
309	try:	1✔
310	self._compiler.check_queue_equivalence_of_args(	1✔
311	self._kernel_name, args
312	)
313	except ExecutionQueueInferenceError as eqie:	1✔
314	raise eqie	1✔
315
316	# A function being compiled in the KERNEL compilation mode
317	# cannot have a non-void return value
318	if return_type and return_type != void:	1!
NEW 319	raise KernelHasReturnValueError(	×
320	kernel_name=None, return_type=return_type, sig=sig
321	)
322
323	# Don't recompile if signature already exists
324	existing = self.overloads.get(tuple(args))	1✔
325	if existing is not None:	1✔
326	return existing.entry_point	1✔
327
328	# TODO: Enable caching
329	# Add code to enable on disk caching of a binary spirv kernel.
330	# Refer: https://github.com/IntelPython/numba-dpex/issues/1197
331	self._cache_misses[sig] += 1	1✔
332	ev_details = {	1✔
333	"dispatcher": self,
334	"args": args,
335	"return_type": return_type,
336	}
337	with ev.trigger_event("numba_dpex:compile", data=ev_details):	1✔
338	try:	1✔
339	kcres: _KernelCompileResult = self._compiler.compile(	1✔
340	args, return_type
341	)
342	except errors.ForceLiteralArg as e:	×
343
344	def folded(args, kws):	×
345	return self._compiler.fold_argument_types(	×
346	args, kws
347	)[1]
348
349	raise e.bind_fold_arguments(folded)	×
350	self.add_overload(kcres)	1✔
351
352	kcres.target_context.insert_user_function(	1✔
353	kcres.entry_point, kcres.fndesc, [kcres.library]
354	)
355
356	# TODO: enable caching of kernel_module
357	# https://github.com/IntelPython/numba-dpex/issues/1197
358
359	return kcres.entry_point	1✔
360
361	def __getitem__(self, args):	1✔
362	"""Square-bracket notation for configuring launch arguments is not
363	supported.
364	"""
365
366	raise NotImplementedError
367
368	def __call__(self, args, *kw_args):	1✔
369	"""Functor to launch a kernel."""
370
371	raise NotImplementedError
372
373
374	_dpex_target = target_registry[DPEX_KERNEL_EXP_TARGET_NAME]	1✔
375	dispatcher_registry[_dpex_target] = KernelDispatcher	1✔

IntelPython / numba-dpex / 7124424362

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