8299321732

Committed 15 Mar 2024 04:25PM UTC coverage: 83.018% (+0.2%) from 82.822%

Build # 8299321732

Build Type

Pull #1331

github

Committed by

web-flow

Commit Message

Merge b7de69cf6 into 9fd6cdc1c

Pull Request Pull Request #1331: A sycl::local_accessor-like API for numba-dpex kernel

Run Details

1625 of 2225 branches covered (73.03%)

Branch coverage included in aggregate %.

184 of 205 new or added lines in 12 files covered. (89.76%)

1 existing line in 1 file now uncovered.

6612 of 7697 relevant lines covered (85.9%)

0.86 hits per line

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

87.18

/numba_dpex/kernel_api/launcher.py

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

"""Implementation of mock kernel launcher functions
"""

from inspect import signature
from itertools import product

from .index_space_ids import Group, Item, NdItem
from .local_accessor import LocalAccessor, _LocalAccessorMock
from .ranges import NdRange, Range


def _range_kernel_launcher(kernel_fn, index_range, *kernel_args):
    """Executes a function that mocks a range kernel.

    Converts the range into a set of index tuple that represent an element in
    the iteration domain over which the kernel will be executed. Then the
    kernel is called sequentially over that set of indices after each index
    value is used to construct an Item object.

    Args:
        kernel_fn : A callable function object
        index_range (numba_dpex.Range): An instance of a Range object

    Raises:
        ValueError: If the number of passed in kernel arguments is not the
        number of function parameters subtracted by one. The first kernel
        argument is expected to be an Item object.
    """

    range_sets = [range(ir) for ir in index_range]
    index_tuples = list(product(*range_sets))

    for karg in kernel_args:
        if isinstance(karg, LocalAccessor):
            raise TypeError(
                "LocalAccessor arguments are only supported for NdRange kernels"
            )

    for idx in index_tuples:
        it = Item(extent=index_range, index=idx)

        if len(signature(kernel_fn).parameters) - len(kernel_args) != 1:
            raise ValueError(
                "Required number of kernel function arguments do not "
                "match provided number of kernel args"
            )

        kernel_fn(it, *kernel_args)


def _ndrange_kernel_launcher(kernel_fn, index_range, *kernel_args):
    """Executes a function that mocks a range kernel.

    Args:
        kernel_fn : A callable function object
        index_range (numba_dpex.NdRange): An instance of a NdRange object

    Raises:
        ValueError: If the number of passed in kernel arguments is not the
        number of function parameters subtracted by one. The first kernel
        argument is expected to be an Item object.
    """
    group_range = tuple(
        gr // lr
        for gr, lr in zip(index_range.global_range, index_range.local_range)
    )
    local_range_sets = [range(ir) for ir in index_range.local_range]
    group_range_sets = [range(gr) for gr in group_range]
    local_index_tuples = list(product(*local_range_sets))
    group_index_tuples = list(product(*group_range_sets))

    modified_kernel_args = []
    for karg in kernel_args:
        if isinstance(karg, LocalAccessor):
            karg = _LocalAccessorMock(karg)
        modified_kernel_args.append(karg)

    # Loop over the groups (parallel loop)
    for gidx in group_index_tuples:
        # loop over work items in the group (parallel loop)
        for lidx in local_index_tuples:
            global_id = []
            # to calculate global indices
            for dim, gidx_val in enumerate(gidx):
                global_id.append(
                    gidx_val * index_range.local_range[dim] + lidx[dim]
                )
            if len(signature(kernel_fn).parameters) - len(kernel_args) != 1:
                raise ValueError(
                    "Required number of kernel function arguments do not "
                    "match provided number of kernel args"
                )

            kernel_fn(
                NdItem(
                    global_item=Item(
                        extent=index_range.global_range, index=global_id
                    ),
                    local_item=Item(extent=index_range.local_range, index=lidx),
                    group=Group(
                        index_range.global_range,
                        index_range.local_range,
                        group_range,
                        gidx,
                    ),
                ),
                *modified_kernel_args
            )


def call_kernel(kernel_fn, index_range, *kernel_args):
    """Mocks the launching of a kernel function over either a Range or NdRange.

    Args:
        kernel_fn : A callable function object
        index_range (numba_dpex.Range): An instance of a Range object

    Raises:
        ValueError: If the first positional argument is not callable
        ValueError: If the second positional argument is not a Range or an
        Ndrange object
    """
    if not callable(kernel_fn):
        raise ValueError(
            "Expected the first positional argument to be a function object"
        )
    if isinstance(index_range, Range):
        _range_kernel_launcher(kernel_fn, index_range, *kernel_args)
    elif isinstance(index_range, NdRange):
        _ndrange_kernel_launcher(kernel_fn, index_range, *kernel_args)
    else:
        raise ValueError(
            "Expected second positional argument to be Range or NdRange object"
        )

1	# SPDX-FileCopyrightText: 2024 Intel Corporation
2	#
3	# SPDX-License-Identifier: Apache-2.0
4
5	"""Implementation of mock kernel launcher functions	1✔
6	"""
7
8	from inspect import signature	1✔
9	from itertools import product	1✔
10
11	from .index_space_ids import Group, Item, NdItem	1✔
12	from .local_accessor import LocalAccessor, _LocalAccessorMock	1✔
13	from .ranges import NdRange, Range	1✔
14
15
16	def _range_kernel_launcher(kernel_fn, index_range, *kernel_args):	1✔
17	"""Executes a function that mocks a range kernel.
18
19	Converts the range into a set of index tuple that represent an element in
20	the iteration domain over which the kernel will be executed. Then the
21	kernel is called sequentially over that set of indices after each index
22	value is used to construct an Item object.
23
24	Args:
25	kernel_fn : A callable function object
26	index_range (numba_dpex.Range): An instance of a Range object
27
28	Raises:
29	ValueError: If the number of passed in kernel arguments is not the
30	number of function parameters subtracted by one. The first kernel
31	argument is expected to be an Item object.
32	"""
33
34	range_sets = [range(ir) for ir in index_range]	1✔
35	index_tuples = list(product(*range_sets))	1✔
36
37	for karg in kernel_args:	1✔
38	if isinstance(karg, LocalAccessor):	1!
NEW 39	raise TypeError(	×
40	"LocalAccessor arguments are only supported for NdRange kernels"
41	)
42
43	for idx in index_tuples:	1✔
44	it = Item(extent=index_range, index=idx)	1✔
45
46	if len(signature(kernel_fn).parameters) - len(kernel_args) != 1:	1!
47	raise ValueError(	×
48	"Required number of kernel function arguments do not "
49	"match provided number of kernel args"
50	)
51
52	kernel_fn(it, *kernel_args)	1✔
53
54
55	def _ndrange_kernel_launcher(kernel_fn, index_range, *kernel_args):	1✔
56	"""Executes a function that mocks a range kernel.
57
58	Args:
59	kernel_fn : A callable function object
60	index_range (numba_dpex.NdRange): An instance of a NdRange object
61
62	Raises:
63	ValueError: If the number of passed in kernel arguments is not the
64	number of function parameters subtracted by one. The first kernel
65	argument is expected to be an Item object.
66	"""
67	group_range = tuple(	1✔
68	gr // lr
69	for gr, lr in zip(index_range.global_range, index_range.local_range)
70	)
71	local_range_sets = [range(ir) for ir in index_range.local_range]	1✔
72	group_range_sets = [range(gr) for gr in group_range]	1✔
73	local_index_tuples = list(product(*local_range_sets))	1✔
74	group_index_tuples = list(product(*group_range_sets))	1✔
75
76	modified_kernel_args = []	1✔
77	for karg in kernel_args:	1✔
78	if isinstance(karg, LocalAccessor):	1✔
79	karg = _LocalAccessorMock(karg)	1✔
80	modified_kernel_args.append(karg)	1✔
81
82	# Loop over the groups (parallel loop)
83	for gidx in group_index_tuples:	1✔
84	# loop over work items in the group (parallel loop)
85	for lidx in local_index_tuples:	1✔
86	global_id = []	1✔
87	# to calculate global indices
88	for dim, gidx_val in enumerate(gidx):	1✔
89	global_id.append(	1✔
90	gidx_val * index_range.local_range[dim] + lidx[dim]
91	)
92	if len(signature(kernel_fn).parameters) - len(kernel_args) != 1:	1!
UNCOV 93	raise ValueError(	×
94	"Required number of kernel function arguments do not "
95	"match provided number of kernel args"
96	)
97
98	kernel_fn(	1✔
99	NdItem(
100	global_item=Item(
101	extent=index_range.global_range, index=global_id
102	),
103	local_item=Item(extent=index_range.local_range, index=lidx),
104	group=Group(
105	index_range.global_range,
106	index_range.local_range,
107	group_range,
108	gidx,
109	),
110	),
111	*modified_kernel_args
112	)
113
114
115	def call_kernel(kernel_fn, index_range, *kernel_args):	1✔
116	"""Mocks the launching of a kernel function over either a Range or NdRange.
117
118	Args:
119	kernel_fn : A callable function object
120	index_range (numba_dpex.Range): An instance of a Range object
121
122	Raises:
123	ValueError: If the first positional argument is not callable
124	ValueError: If the second positional argument is not a Range or an
125	Ndrange object
126	"""
127	if not callable(kernel_fn):	1!
128	raise ValueError(	×
129	"Expected the first positional argument to be a function object"
130	)
131	if isinstance(index_range, Range):	1✔
132	_range_kernel_launcher(kernel_fn, index_range, *kernel_args)	1✔
133	elif isinstance(index_range, NdRange):	1!
134	_ndrange_kernel_launcher(kernel_fn, index_range, *kernel_args)	1✔
135	else:
136	raise ValueError(	×
137	"Expected second positional argument to be Range or NdRange object"
138	)

IntelPython / numba-dpex / 8299321732

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