4981627980

Build Type

github

Committed by GitHub

Commit Message

Merge 969f6c414 into 87206facf

Pull Request Pull Request #999: Enable dpnp parfor reduction

Run Details

2153 of 3737 branches covered (57.61%)

Branch coverage included in aggregate %.

538 of 594 new or added lines in 9 files covered. (90.57%)

35 existing lines in 1 file now uncovered.

7381 of 10676 relevant lines covered (69.14%)

0.69 hits per line

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

86.21

/numba_dpex/core/utils/kernel_templates/partial_reduction_template.py

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

import operator
import sys

import dpnp
from numba.core import compiler

import numba_dpex as dpex


class TreeReduceIntermediateKernelTemplate:
    def __init__(
        self,
        kernel_name,
        kernel_params,  # parfor_params
        ivar_names,
        sentinel_name,
        loop_ranges,
        param_dict,
        parfor_dim,
        redvars,
        parfor_args,
        parfor_reddict,
        redvars_dict,
        typemap,
        work_group_size,
    ) -> None:
        self._kernel_name = kernel_name
        self._kernel_params = kernel_params
        self._ivar_names = ivar_names
        self._sentinel_name = sentinel_name
        self._loop_ranges = loop_ranges
        self._param_dict = param_dict
        self._parfor_dim = parfor_dim
        self._redvars = redvars

        # added
        self._parfor_args = parfor_args
        self._parfor_reddict = parfor_reddict
        self._redvars_dict = redvars_dict
        self._typemap = typemap
        self._work_group_size = work_group_size

        self._kernel_txt = self._generate_main_kernel_as_string()
        self._kernel_ir = self._generate_kernel_ir()

    def _generate_main_kernel_as_string(self):
        """Generates the indexing intrinsic calls inside a dpex kernel.

        Args:
            parfor_dim (_type_): _description_
            ivar_names (_type_): _description_
            loop_ranges (_type_): _description_
            param_dict (_type_): _description_

        Returns:
            _type_: _description_
        """
        gufunc_txt = ""
        gufunc_txt += "def " + self._kernel_name
        gufunc_txt += "(" + (", ".join(self._kernel_params)) + "):\n"
        global_id_dim = 0
        for_loop_dim = self._parfor_dim

        if self._parfor_dim > 3:
            raise NotImplementedError
        else:
            global_id_dim = self._parfor_dim

        for dim in range(global_id_dim):
            dstr = str(dim)
            gufunc_txt += (
                f"    {self._ivar_names[dim]} = dpex.get_global_id({dstr})\n"
            )
            gufunc_txt += f"    local_id{dim} = dpex.get_local_id({dstr})\n"
            gufunc_txt += f"    local_size{dim} = dpex.get_local_size({dstr})\n"
            gufunc_txt += f"    group_id{dim} = dpex.get_group_id({dstr})\n"

        # Allocate local_sums arrays for each reduction variable.
        for redvar in self._redvars:
            rtyp = str(self._typemap[redvar])
            redvar = self._redvars_dict[redvar]
            gufunc_txt += f"    local_sums_{redvar} = dpex.local.array({self._work_group_size}, dpnp.{rtyp})\n"

        for dim in range(global_id_dim, for_loop_dim):
            for indent in range(1 + (dim - global_id_dim)):
                gufunc_txt += "    "

            start, stop, step = self._loop_ranges[dim]
            st = str(self._param_dict.get(str(start), start))
            en = str(self._param_dict.get(str(stop), stop))
            gufunc_txt += (
                f"for {self._ivar_names[dim]} in range({st}, {en} + 1):\n"
            )

        for dim in range(global_id_dim, for_loop_dim):
            for indent in range(1 + (dim - global_id_dim)):
                gufunc_txt += "    "
        # Add the sentinel assignment so that we can find the loop body position
        # in the IR.
        for redvar in self._redvars:
            legal_redvar = self._redvars_dict[redvar]
            gufunc_txt += "    "
            gufunc_txt += legal_redvar + " = 0\n"
        gufunc_txt += "    "
        gufunc_txt += self._sentinel_name + " = 0\n"

        # Generate local_sum[local_id0] = redvar, for each reduction variable
        for redvar in self._redvars:
            legal_redvar = self._redvars_dict[redvar]
            gufunc_txt += (
                "    "
                + f"local_sums_{legal_redvar}[local_id0] = {legal_redvar}\n"
            )

        gufunc_txt += (
            "    stride0 = local_size0 // 2\n"
            + "    while stride0 > 0:\n"
            + "        dpex.barrier(dpex.LOCAL_MEM_FENCE)\n"
            + "        if local_id0 < stride0:\n"
        )

        for redvar in self._redvars:
            redop = self._parfor_reddict[redvar].redop
            redvar_legal = self._redvars_dict[redvar]
            if redop == operator.iadd:
                gufunc_txt += (
                    "            "
                    f"local_sums_{redvar_legal}[local_id0] "
                    f"+= local_sums_{redvar_legal}[local_id0 + stride0]\n"
                )
            elif redop == operator.imul:
                gufunc_txt += (
                    "            "
                    f"local_sums_{redvar_legal}[local_id0] "
                    f"*= local_sums_{redvar_legal}[local_id0 + stride0]\n"
                )
            else:
                raise NotImplementedError

        gufunc_txt += "        stride0 >>= 1\n"
        gufunc_txt += "    if local_id0 == 0:\n"
        for redvar in self._redvars:
            for i, arg in enumerate(self._parfor_args):
                if arg == redvar:
                    partial_sum_var = self._kernel_params[i]
                    redvar_legal = self._redvars_dict[redvar]
                    gufunc_txt += (
                        "        "
                        f"{partial_sum_var}[group_id0] = "
                        f"local_sums_{redvar_legal}[0]\n"
                    )

        gufunc_txt += "    return None\n"

        return gufunc_txt

    def _generate_kernel_ir(self):
        """Exec the kernel_txt string into a Python function object and then
        compile it using Numba's compiler front end.

        Returns: The Numba functionIR object for the compiled kernel_txt string.

        """
        globls = {"dpnp": dpnp, "dpex": dpex}
        locls = {}
        exec(self._kernel_txt, globls, locls)
        kernel_fn = locls[self._kernel_name]

        return compiler.run_frontend(kernel_fn)

    @property
    def kernel_ir(self):
        """Returns the Numba IR generated for a RangeKernelTemplate.

        Returns: The Numba functionIR object for the compiled kernel_txt string.
        """
        return self._kernel_ir

    @property
    def kernel_string(self):
        """Returns the function string generated for a RangeKernelTemplate.

        Returns:
            str: A string representing a stub kernel function for the parfor.
        """
        return self._kernel_txt

    def dump_kernel_string(self):
        """Helper to print the kernel function string."""
        print(self._kernel_txt)
        sys.stdout.flush()

    def dump_kernel_ir(self):
        """Helper to dump the Numba IR for the RangeKernelTemplate."""
        self._kernel_ir.dump()


class RemainderReduceIntermediateKernelTemplate:
    def __init__(
        self,
        kernel_name,
        kernel_params,
        sentinel_name,
        redvars,
        parfor_reddict,
        redvars_dict,
        typemap,
        legal_loop_indices,
        global_size_var_name,
        global_size_mod_var_name,
        partial_sum_size_var_name,
        partial_sum_var_name,
        final_sum_var_name,
        reductionKernelVar,
    ) -> None:
        self._kernel_name = kernel_name
        self._kernel_params = kernel_params
        self._sentinel_name = sentinel_name
        self._redvars = redvars
        self._parfor_reddict = parfor_reddict
        self._redvars_dict = redvars_dict
        self._typemap = typemap
        self._legal_loop_indices = legal_loop_indices
        self._global_size_var_name = global_size_var_name
        self._global_size_mod_var_name = global_size_mod_var_name
        self._partial_sum_size_var_name = partial_sum_size_var_name
        self._partial_sum_var_name = partial_sum_var_name
        self._final_sum_var_name = final_sum_var_name
        self._reductionKernelVar = reductionKernelVar

        self._kernel_txt = self._generate_remainder_kernel_as_string()
        self._kernel_ir = self._generate_kernel_ir()

    def _generate_remainder_kernel_as_string(self):
        gufunc_txt = ""
        gufunc_txt += "def " + self._kernel_name
        gufunc_txt += "(" + (", ".join(self._kernel_params))

        for i in range(len(self._redvars)):
            gufunc_txt += (
                ", "
                + f"{self._global_size_var_name[i]}, "
                + f"{self._global_size_mod_var_name[i]}, "
            )
            gufunc_txt += (
                f"{self._partial_sum_size_var_name[i]}, "
                + f"{self._final_sum_var_name[i]}"
            )

        gufunc_txt += "):\n"

        gufunc_txt += "    "

        gufunc_txt += (
            "for j" + f" in range({self._partial_sum_size_var_name[0]}):\n"
        )

        for i, redvar in enumerate(self._redvars):
            gufunc_txt += f"        {self._final_sum_var_name[i]}[0] += {self._partial_sum_var_name[i]}[j]\n"

        gufunc_txt += (
            f"    for j in range ({self._global_size_mod_var_name[0]}) :\n"
        )

        for redvar in self._redvars:
            legal_redvar = self._redvars_dict[redvar]
            gufunc_txt += "        "
            gufunc_txt += legal_redvar + " = 0\n"

        gufunc_txt += (
            "        "
            + self._legal_loop_indices[0]
            + " = "
            + f"{self._global_size_var_name[0]} + j\n"
        )

        for redvar in self._redvars:
            rtyp = str(self._typemap[redvar])
            redvar = self._redvars_dict[redvar]
            gufunc_txt += (
                "        "
                + f"local_sums_{redvar} = "
                + f"dpex.local.array(1, dpnp.{rtyp})\n"
            )

        gufunc_txt += "        " + self._sentinel_name + " = 0\n"

        for i, redvar in enumerate(self._redvars):
            legal_redvar = self._redvars_dict[redvar]
            gufunc_txt += (
                "        " + f"local_sums_{legal_redvar}[0] = {legal_redvar}\n"
            )

        for i, redvar in enumerate(self._redvars):
            legal_redvar = self._redvars_dict[redvar]
            redop = self._parfor_reddict[redvar].redop
            if redop == operator.iadd:
                gufunc_txt += f"        {self._final_sum_var_name[i]}[0] +=  local_sums_{legal_redvar}[0]\n"
            elif redop == operator.imul:
                gufunc_txt += f"        {self._final_sum_var_name[i]}[0] *=  local_sums_{legal_redvar}[0]\n"
            else:
                raise NotImplementedError

        return gufunc_txt

    def _generate_kernel_ir(self):
        """Exec the kernel_txt string into a Python function object and then
        compile it using Numba's compiler front end.

        Returns: The Numba functionIR object for the compiled kernel_txt string.

        """
        globls = {"dpnp": dpnp, "dpex": dpex}
        locls = {}
        exec(self._kernel_txt, globls, locls)
        kernel_fn = locls[self._kernel_name]

        return compiler.run_frontend(kernel_fn)

    @property
    def kernel_ir(self):
        """Returns the Numba IR generated for a RangeKernelTemplate.

        Returns: The Numba functionIR object for the compiled kernel_txt string.
        """
        return self._kernel_ir

    @property
    def kernel_string(self):
        """Returns the function string generated for a RangeKernelTemplate.

        Returns:
            str: A string representing a stub kernel function for the parfor.
        """
        return self._kernel_txt

    def dump_kernel_string(self):
        """Helper to print the kernel function string."""
        print(self._kernel_txt)
        sys.stdout.flush()

    def dump_kernel_ir(self):
        """Helper to dump the Numba IR for the RangeKernelTemplate."""
        self._kernel_ir.dump()

1	# SPDX-FileCopyrightText: 2023 Intel Corporation
2	#
3	# SPDX-License-Identifier: Apache-2.0
4
5	import operator	1✔
6	import sys	1✔
7
8	import dpnp	1✔
9	from numba.core import compiler	1✔
10
11	import numba_dpex as dpex	1✔
12
13
14	class TreeReduceIntermediateKernelTemplate:	1✔
15	def __init__(	1✔
16	self,
17	kernel_name,
18	kernel_params, # parfor_params
19	ivar_names,
20	sentinel_name,
21	loop_ranges,
22	param_dict,
23	parfor_dim,
24	redvars,
25	parfor_args,
26	parfor_reddict,
27	redvars_dict,
28	typemap,
29	work_group_size,
30	) -> None:
31	self._kernel_name = kernel_name	1✔
32	self._kernel_params = kernel_params	1✔
33	self._ivar_names = ivar_names	1✔
34	self._sentinel_name = sentinel_name	1✔
35	self._loop_ranges = loop_ranges	1✔
36	self._param_dict = param_dict	1✔
37	self._parfor_dim = parfor_dim	1✔
38	self._redvars = redvars	1✔
39
40	# added
41	self._parfor_args = parfor_args	1✔
42	self._parfor_reddict = parfor_reddict	1✔
43	self._redvars_dict = redvars_dict	1✔
44	self._typemap = typemap	1✔
45	self._work_group_size = work_group_size	1✔
46
47	self._kernel_txt = self._generate_main_kernel_as_string()	1✔
48	self._kernel_ir = self._generate_kernel_ir()	1✔
49
50	def _generate_main_kernel_as_string(self):	1✔
51	"""Generates the indexing intrinsic calls inside a dpex kernel.
52
53	Args:
54	parfor_dim (_type_): _description_
55	ivar_names (_type_): _description_
56	loop_ranges (_type_): _description_
57	param_dict (_type_): _description_
58
59	Returns:
60	_type_: _description_
61	"""
62	gufunc_txt = ""	1✔
63	gufunc_txt += "def " + self._kernel_name	1✔
64	gufunc_txt += "(" + (", ".join(self._kernel_params)) + "):\n"	1✔
65	global_id_dim = 0	1✔
66	for_loop_dim = self._parfor_dim	1✔
67
68	if self._parfor_dim > 3:	1✔
69	raise NotImplementedError
70	else:
71	global_id_dim = self._parfor_dim	1✔
72
73	for dim in range(global_id_dim):	1✔
74	dstr = str(dim)	1✔
75	gufunc_txt += (	1✔
76	f" {self._ivar_names[dim]} = dpex.get_global_id({dstr})\n"
77	)
78	gufunc_txt += f" local_id{dim} = dpex.get_local_id({dstr})\n"	1✔
79	gufunc_txt += f" local_size{dim} = dpex.get_local_size({dstr})\n"	1✔
80	gufunc_txt += f" group_id{dim} = dpex.get_group_id({dstr})\n"	1✔
81
82	# Allocate local_sums arrays for each reduction variable.
83	for redvar in self._redvars:	1✔
84	rtyp = str(self._typemap[redvar])	1✔
85	redvar = self._redvars_dict[redvar]	1✔
86	gufunc_txt += f" local_sums_{redvar} = dpex.local.array({self._work_group_size}, dpnp.{rtyp})\n"	1✔
87
88	for dim in range(global_id_dim, for_loop_dim):	1!
NEW 89	for indent in range(1 + (dim - global_id_dim)):	×
NEW 90	gufunc_txt += " "	×
91
NEW 92	start, stop, step = self._loop_ranges[dim]	×
NEW 93	st = str(self._param_dict.get(str(start), start))	×
NEW 94	en = str(self._param_dict.get(str(stop), stop))	×
NEW 95	gufunc_txt += (	×
96	f"for {self._ivar_names[dim]} in range({st}, {en} + 1):\n"
97	)
98
99	for dim in range(global_id_dim, for_loop_dim):	1!
NEW 100	for indent in range(1 + (dim - global_id_dim)):	×
NEW 101	gufunc_txt += " "	×
102	# Add the sentinel assignment so that we can find the loop body position
103	# in the IR.
104	for redvar in self._redvars:	1✔
105	legal_redvar = self._redvars_dict[redvar]	1✔
106	gufunc_txt += " "	1✔
107	gufunc_txt += legal_redvar + " = 0\n"	1✔
108	gufunc_txt += " "	1✔
109	gufunc_txt += self._sentinel_name + " = 0\n"	1✔
110
111	# Generate local_sum[local_id0] = redvar, for each reduction variable
112	for redvar in self._redvars:	1✔
113	legal_redvar = self._redvars_dict[redvar]	1✔
114	gufunc_txt += (	1✔
115	" "
116	+ f"local_sums_{legal_redvar}[local_id0] = {legal_redvar}\n"
117	)
118
119	gufunc_txt += (	1✔
120	" stride0 = local_size0 // 2\n"
121	+ " while stride0 > 0:\n"
122	+ " dpex.barrier(dpex.LOCAL_MEM_FENCE)\n"
123	+ " if local_id0 < stride0:\n"
124	)
125
126	for redvar in self._redvars:	1✔
127	redop = self._parfor_reddict[redvar].redop	1✔
128	redvar_legal = self._redvars_dict[redvar]	1✔
129	if redop == operator.iadd:	1!
130	gufunc_txt += (	1✔
131	" "
132	f"local_sums_{redvar_legal}[local_id0] "
133	f"+= local_sums_{redvar_legal}[local_id0 + stride0]\n"
134	)
NEW 135	elif redop == operator.imul:	×
NEW 136	gufunc_txt += (	×
137	" "
138	f"local_sums_{redvar_legal}[local_id0] "
139	f"*= local_sums_{redvar_legal}[local_id0 + stride0]\n"
140	)
141	else:
142	raise NotImplementedError
143
144	gufunc_txt += " stride0 >>= 1\n"	1✔
145	gufunc_txt += " if local_id0 == 0:\n"	1✔
146	for redvar in self._redvars:	1✔
147	for i, arg in enumerate(self._parfor_args):	1✔
148	if arg == redvar:	1✔
149	partial_sum_var = self._kernel_params[i]	1✔
150	redvar_legal = self._redvars_dict[redvar]	1✔
151	gufunc_txt += (	1✔
152	" "
153	f"{partial_sum_var}[group_id0] = "
154	f"local_sums_{redvar_legal}[0]\n"
155	)
156
157	gufunc_txt += " return None\n"	1✔
158
159	return gufunc_txt	1✔
160
161	def _generate_kernel_ir(self):	1✔
162	"""Exec the kernel_txt string into a Python function object and then
163	compile it using Numba's compiler front end.
164
165	Returns: The Numba functionIR object for the compiled kernel_txt string.
166
167	"""
168	globls = {"dpnp": dpnp, "dpex": dpex}	1✔
169	locls = {}	1✔
170	exec(self._kernel_txt, globls, locls)	1✔
171	kernel_fn = locls[self._kernel_name]	1✔
172
173	return compiler.run_frontend(kernel_fn)	1✔
174
175	@property	1✔
176	def kernel_ir(self):	1✔
177	"""Returns the Numba IR generated for a RangeKernelTemplate.
178
179	Returns: The Numba functionIR object for the compiled kernel_txt string.
180	"""
181	return self._kernel_ir	1✔
182
183	@property	1✔
184	def kernel_string(self):	1✔
185	"""Returns the function string generated for a RangeKernelTemplate.
186
187	Returns:
188	str: A string representing a stub kernel function for the parfor.
189	"""
NEW 190	return self._kernel_txt	×
191
192	def dump_kernel_string(self):	1✔
193	"""Helper to print the kernel function string."""
NEW 194	print(self._kernel_txt)	×
NEW 195	sys.stdout.flush()	×
196
197	def dump_kernel_ir(self):	1✔
198	"""Helper to dump the Numba IR for the RangeKernelTemplate."""
NEW 199	self._kernel_ir.dump()	×
200
201
202	class RemainderReduceIntermediateKernelTemplate:	1✔
203	def __init__(	1✔
204	self,
205	kernel_name,
206	kernel_params,
207	sentinel_name,
208	redvars,
209	parfor_reddict,
210	redvars_dict,
211	typemap,
212	legal_loop_indices,
213	global_size_var_name,
214	global_size_mod_var_name,
215	partial_sum_size_var_name,
216	partial_sum_var_name,
217	final_sum_var_name,
218	reductionKernelVar,
219	) -> None:
220	self._kernel_name = kernel_name	1✔
221	self._kernel_params = kernel_params	1✔
222	self._sentinel_name = sentinel_name	1✔
223	self._redvars = redvars	1✔
224	self._parfor_reddict = parfor_reddict	1✔
225	self._redvars_dict = redvars_dict	1✔
226	self._typemap = typemap	1✔
227	self._legal_loop_indices = legal_loop_indices	1✔
228	self._global_size_var_name = global_size_var_name	1✔
229	self._global_size_mod_var_name = global_size_mod_var_name	1✔
230	self._partial_sum_size_var_name = partial_sum_size_var_name	1✔
231	self._partial_sum_var_name = partial_sum_var_name	1✔
232	self._final_sum_var_name = final_sum_var_name	1✔
233	self._reductionKernelVar = reductionKernelVar	1✔
234
235	self._kernel_txt = self._generate_remainder_kernel_as_string()	1✔
236	self._kernel_ir = self._generate_kernel_ir()	1✔
237
238	def _generate_remainder_kernel_as_string(self):	1✔
239	gufunc_txt = ""	1✔
240	gufunc_txt += "def " + self._kernel_name	1✔
241	gufunc_txt += "(" + (", ".join(self._kernel_params))	1✔
242
243	for i in range(len(self._redvars)):	1✔
244	gufunc_txt += (	1✔
245	", "
246	+ f"{self._global_size_var_name[i]}, "
247	+ f"{self._global_size_mod_var_name[i]}, "
248	)
249	gufunc_txt += (	1✔
250	f"{self._partial_sum_size_var_name[i]}, "
251	+ f"{self._final_sum_var_name[i]}"
252	)
253
254	gufunc_txt += "):\n"	1✔
255
256	gufunc_txt += " "	1✔
257
258	gufunc_txt += (	1✔
259	"for j" + f" in range({self._partial_sum_size_var_name[0]}):\n"
260	)
261
262	for i, redvar in enumerate(self._redvars):	1✔
263	gufunc_txt += f" {self._final_sum_var_name[i]}[0] += {self._partial_sum_var_name[i]}[j]\n"	1✔
264
265	gufunc_txt += (	1✔
266	f" for j in range ({self._global_size_mod_var_name[0]}) :\n"
267	)
268
269	for redvar in self._redvars:	1✔
270	legal_redvar = self._redvars_dict[redvar]	1✔
271	gufunc_txt += " "	1✔
272	gufunc_txt += legal_redvar + " = 0\n"	1✔
273
274	gufunc_txt += (	1✔
275	" "
276	+ self._legal_loop_indices[0]
277	+ " = "
278	+ f"{self._global_size_var_name[0]} + j\n"
279	)
280
281	for redvar in self._redvars:	1✔
282	rtyp = str(self._typemap[redvar])	1✔
283	redvar = self._redvars_dict[redvar]	1✔
284	gufunc_txt += (	1✔
285	" "
286	+ f"local_sums_{redvar} = "
287	+ f"dpex.local.array(1, dpnp.{rtyp})\n"
288	)
289
290	gufunc_txt += " " + self._sentinel_name + " = 0\n"	1✔
291
292	for i, redvar in enumerate(self._redvars):	1✔
293	legal_redvar = self._redvars_dict[redvar]	1✔
294	gufunc_txt += (	1✔
295	" " + f"local_sums_{legal_redvar}[0] = {legal_redvar}\n"
296	)
297
298	for i, redvar in enumerate(self._redvars):	1✔
299	legal_redvar = self._redvars_dict[redvar]	1✔
300	redop = self._parfor_reddict[redvar].redop	1✔
301	if redop == operator.iadd:	1!
302	gufunc_txt += f" {self._final_sum_var_name[i]}[0] += local_sums_{legal_redvar}[0]\n"	1✔
NEW 303	elif redop == operator.imul:	×
NEW 304	gufunc_txt += f" {self._final_sum_var_name[i]}[0] *= local_sums_{legal_redvar}[0]\n"	×
305	else:
306	raise NotImplementedError
307
308	return gufunc_txt	1✔
309
310	def _generate_kernel_ir(self):	1✔
311	"""Exec the kernel_txt string into a Python function object and then
312	compile it using Numba's compiler front end.
313
314	Returns: The Numba functionIR object for the compiled kernel_txt string.
315
316	"""
317	globls = {"dpnp": dpnp, "dpex": dpex}	1✔
318	locls = {}	1✔
319	exec(self._kernel_txt, globls, locls)	1✔
320	kernel_fn = locls[self._kernel_name]	1✔
321
322	return compiler.run_frontend(kernel_fn)	1✔
323
324	@property	1✔
325	def kernel_ir(self):	1✔
326	"""Returns the Numba IR generated for a RangeKernelTemplate.
327
328	Returns: The Numba functionIR object for the compiled kernel_txt string.
329	"""
330	return self._kernel_ir	1✔
331
332	@property	1✔
333	def kernel_string(self):	1✔
334	"""Returns the function string generated for a RangeKernelTemplate.
335
336	Returns:
337	str: A string representing a stub kernel function for the parfor.
338	"""
NEW 339	return self._kernel_txt	×
340
341	def dump_kernel_string(self):	1✔
342	"""Helper to print the kernel function string."""
NEW 343	print(self._kernel_txt)	×
NEW 344	sys.stdout.flush()	×
345
346	def dump_kernel_ir(self):	1✔
347	"""Helper to dump the Numba IR for the RangeKernelTemplate."""
NEW 348	self._kernel_ir.dump()	×

IntelPython / numba-dpex / 4981627980

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