5005997591

Build Type

github

Committed by GitHub

Commit Message

Merge c9f9bc7cb into c0dd04f61

Pull Request Pull Request #999: Enable dpnp parfor reduction

Run Details

1791 of 3460 branches covered (51.76%)

Branch coverage included in aggregate %.

588 of 630 new or added lines in 11 files covered. (93.33%)

32 existing lines in 1 file now uncovered.

5990 of 9440 relevant lines covered (63.45%)

0.63 hits per line

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

86.27

/numba_dpex/core/utils/kernel_templates/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

from .kernel_template_iface import KernelTemplateInterface


class TreeReduceIntermediateKernelTemplate(KernelTemplateInterface):
    """The class to build reduction main kernel_txt template and
    compiled Numba functionIR."""

    def __init__(
        self,
        kernel_name,
        kernel_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
        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_kernel_stub_as_string()
        self._kernel_ir = self._generate_kernel_ir()

    def _generate_kernel_stub_as_string(self):
        """Generate reduction main kernel template"""

        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
            TreeReduceIntermediateKernelTemplate.

        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
            TreeReduceIntermediateKernelTemplate.

        Returns:
            str: A string representing a stub reduction 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 a
        TreeReduceIntermediateKernelTemplate."""
        self._kernel_ir.dump()


class RemainderReduceIntermediateKernelTemplate(KernelTemplateInterface):
    """The class to build reduction remainder kernel_txt template and
    compiled Numba functionIR.
    """

    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_kernel_stub_as_string()
        self._kernel_ir = self._generate_kernel_ir()

    def _generate_kernel_stub_as_string(self):
        """Generate reduction remainder kernel template"""

        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
            RemainderReduceIntermediateKernelTemplate.

        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
            RemainderReduceIntermediateKernelTemplate.

        Returns:
            str: A string representing a stub reduction 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
        RemainderReduceIntermediateKernelTemplate."""

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

IntelPython / numba-dpex / 5005997591

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