8290948108

Committed 15 Mar 2024 03:35AM UTC coverage: 82.266% (-0.3%) from 82.555%

Build # 8290948108

Build Type

Pull #1378

github

Committed by

web-flow

Commit Message

Merge 7a8b6cad7 into af24e66d6

Pull Request Pull Request #1378: Extend indexing function tests to kernel simulator

Run Details

1557 of 2165 branches covered (71.92%)

Branch coverage included in aggregate %.

6436 of 7551 relevant lines covered (85.23%)

0.85 hits per line

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

52.21

/numba_dpex/kernel_api/index_space_ids.py

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

"""Implements a mock Python classes to represent ``sycl::item`` and
``sycl::nd_item``for prototyping numba_dpex kernel functions before they are JIT
compiled.
"""

from .ranges import Range


class Group:
    """Analogue to the ``sycl::group`` type."""

    def __init__(
        self,
        global_range: Range,
        local_range: Range,
        group_range: Range,
        index: list,
    ):
        self._global_range = global_range
        self._local_range = local_range
        self._group_range = group_range
        self._index = index
        self._leader = False

    def get_group_id(self, dim):
        """Returns the index of the work-group within the global nd-range for
        specified dimension.

        Since the work-items in a work-group have a defined position within the
        global nd-range, the returned group id can be used along with the local
        id to uniquely identify the work-item in the global nd-range.
        """
        if dim > len(self._index) - 1:
            raise ValueError(
                "Dimension value is out of bounds for the group index"
            )
        return self._index[dim]

    def get_group_linear_id(self):
        """Returns a linearized version of the work-group index."""
        if self.dimensions == 1:
            return self.get_group_id(0)
        if self.dimensions == 2:
            return self.get_group_id(0) * self.get_group_range(
                1
            ) + self.get_group_id(1)
        return (
            (
                self.get_group_id(0)
                * self.get_group_range(1)
                * self.get_group_range(2)
            )
            + (self.get_group_id(1) * self.get_group_range(2))
            + (self.get_group_id(2))
        )

    def get_group_range(self, dim):
        """Returns a the extent of the range representing the number of groups
        in the nd-range for a specified dimension.
        """
        return self._group_range[dim]

    def get_group_linear_range(self):
        """Return the total number of work-groups in the nd_range."""
        num_wg = 1
        for i in range(self.dimensions):
            num_wg *= self.get_group_range(i)

        return num_wg

    def get_local_range(self, dim):
        """Returns the extent of the SYCL range representing all dimensions
        of the local range for a specified dimension. This local range may
        have been provided by the programmer, or chosen by the SYCL runtime.
        """
        return self._local_range[dim]

    def get_local_linear_range(self):
        """Return the total number of work-items in the work-group."""
        num_wi = 1
        for i in range(self.dimensions):
            num_wi *= self.get_local_range(i)

        return num_wi

    @property
    def leader(self):
        """Return true for exactly one work-item in the work-group, if the
        calling work-item is the leader of the work-group, and false for all
        other work-items in the work-group.

        The leader of the work-group is determined during construction of the
        work-group, and is invariant for the lifetime of the work-group. The
        leader of the work-group is guaranteed to be the work-item with a
        local id of 0.


        Returns:
            bool: If the work item is the designated leader of the
        """
        return self._leader

    @property
    def dimensions(self) -> int:
        """Returns the rank of a Group object.
        Returns:
            int: Number of dimensions in the Group object
        """
        return self._global_range.ndim

    @leader.setter
    def leader(self, work_item_id):
        """Sets the leader attribute for the group."""
        self._leader = work_item_id


class Item:
    """Analogue to the ``sycl::item`` class.

    Identifies an instance of the function object executing at each point in an
    :class:`.Range`.
    """

    def __init__(self, extent: Range, index: list):
        self._extent = extent
        self._index = index

    def get_linear_id(self):
        """Get the linear id associated with this item for all dimensions.
        Original implementation could be found at ``sycl::item_base`` class.

        Returns:
            int: The linear id.
        """
        if self.dimensions == 1:
            return self.get_id(0)
        if self.dimensions == 2:
            return self.get_id(0) * self.get_range(1) + self.get_id(1)
        return (
            (self.get_id(0) * self.get_range(1) * self.get_range(2))
            + (self.get_id(1) * self.get_range(2))
            + (self.get_id(2))
        )

    def get_id(self, idx):
        """Get the id for a specific dimension.

        Returns:
            int: The id
        """
        return self._index[idx]

    def get_linear_range(self):
        """Return the total number of work-items in the work-group."""
        num_wi = 1
        for i in range(self.dimensions):
            num_wi *= self.get_range(i)

        return num_wi

    def get_range(self, idx):
        """Get the range size for a specific dimension.

        Returns:
            int: The size
        """
        return self._extent[idx]

    @property
    def dimensions(self) -> int:
        """Returns the number of dimensions of a Item object.

        Returns:
            int: Number of dimensions in the Item object
        """
        return self._extent.ndim


class NdItem:
    """Analogue to the ``sycl::nd_item`` class.

    Identifies an instance of the function object executing at each point in an
    :class:`.NdRange`.
    """

    # TODO: define group type
    def __init__(self, global_item: Item, local_item: Item, group: Group):
        # TODO: assert offset and dimensions
        self._global_item = global_item
        self._local_item = local_item
        self._group = group
        if self.get_local_linear_id() == 0:
            self._group.leader = True

    def get_global_id(self, idx):
        """Get the global id for a specific dimension.

        Returns:
            int: The global id
        """
        return self._global_item.get_id(idx)

    def get_global_linear_id(self):
        """Get the linearized global id for the item for all dimensions.

        Returns:
            int: The global linear id.
        """
        # Instead of calling self._global_item.get_linear_id(), the linearization
        # logic is duplicated here so that the method can be JIT compiled by
        # numba-dpex and works in both Python and Numba nopython modes.
        if self.dimensions == 1:
            return self.get_global_id(0)
        if self.dimensions == 2:
            return self.get_global_id(0) * self.get_global_range(
                1
            ) + self.get_global_id(1)
        return (
            (
                self.get_global_id(0)
                * self.get_global_range(1)
                * self.get_global_range(2)
            )
            + (self.get_global_id(1) * self.get_global_range(2))
            + (self.get_global_id(2))
        )

    def get_local_id(self, idx):
        """Get the local id for a specific dimension.

        Returns:
            int: The local id
        """
        return self._local_item.get_id(idx)

    def get_local_linear_id(self):
        """Get the local linear id associated with this item for all
        dimensions.

        Returns:
            int: The local linear id.
        """
        # Instead of calling self._local_item.get_linear_id(), the linearization
        # logic is duplicated here so that the method can be JIT compiled by
        # numba-dpex and works in both Python and Numba nopython modes.
        if self.dimensions == 1:
            return self.get_local_id(0)
        if self.dimensions == 2:
            return self.get_local_id(0) * self.get_local_range(
                1
            ) + self.get_local_id(1)
        return (
            (
                self.get_local_id(0)
                * self.get_local_range(1)
                * self.get_local_range(2)
            )
            + (self.get_local_id(1) * self.get_local_range(2))
            + (self.get_local_id(2))
        )

    def get_global_range(self, idx):
        """Get the global range size for a specific dimension.

        Returns:
            int: The size
        """
        return self._global_item.get_range(idx)

    def get_local_range(self, idx):
        """Get the local range size for a specific dimension.

        Returns:
            int: The size
        """
        return self._local_item.get_range(idx)

    def get_local_linear_range(self):
        """Return the total number of work-items in the work-group."""
        num_wi = 1
        for i in range(self.dimensions):
            num_wi *= self.get_local_range(i)

        return num_wi

    def get_global_linear_range(self):
        """Return the total number of work-items in the work-group."""
        num_wi = 1
        for i in range(self.dimensions):
            num_wi *= self.get_global_range(i)

        return num_wi

    def get_group(self):
        """Returns the group.

        Returns:
            A group object."""
        return self._group

    @property
    def dimensions(self) -> int:
        """Returns the rank of a NdItem object.

        Returns:
            int: Number of dimensions in the NdItem object
        """
        return self._global_item.dimensions

1	# SPDX-FileCopyrightText: 2024 Intel Corporation
2	#
3	# SPDX-License-Identifier: Apache-2.0
4
5	"""Implements a mock Python classes to represent ``sycl::item`` and	1✔
6	``sycl::nd_item``for prototyping numba_dpex kernel functions before they are JIT
7	compiled.
8	"""
9
10	from .ranges import Range	1✔
11
12
13	class Group:	1✔
14	"""Analogue to the ``sycl::group`` type."""
15
16	def __init__(	1✔
17	self,
18	global_range: Range,
19	local_range: Range,
20	group_range: Range,
21	index: list,
22	):
23	self._global_range = global_range	1✔
24	self._local_range = local_range	1✔
25	self._group_range = group_range	1✔
26	self._index = index	1✔
27	self._leader = False	1✔
28
29	def get_group_id(self, dim):	1✔
30	"""Returns the index of the work-group within the global nd-range for
31	specified dimension.
32
33	Since the work-items in a work-group have a defined position within the
34	global nd-range, the returned group id can be used along with the local
35	id to uniquely identify the work-item in the global nd-range.
36	"""
37	if dim > len(self._index) - 1:	×
38	raise ValueError(	×
39	"Dimension value is out of bounds for the group index"
40	)
41	return self._index[dim]	×
42
43	def get_group_linear_id(self):	1✔
44	"""Returns a linearized version of the work-group index."""
45	if self.dimensions == 1:	×
46	return self.get_group_id(0)	×
47	if self.dimensions == 2:	×
48	return self.get_group_id(0) * self.get_group_range(	×
49	1
50	) + self.get_group_id(1)
51	return (	×
52	(
53	self.get_group_id(0)
54	* self.get_group_range(1)
55	* self.get_group_range(2)
56	)
57	+ (self.get_group_id(1) * self.get_group_range(2))
58	+ (self.get_group_id(2))
59	)
60
61	def get_group_range(self, dim):	1✔
62	"""Returns a the extent of the range representing the number of groups
63	in the nd-range for a specified dimension.
64	"""
65	return self._group_range[dim]	×
66
67	def get_group_linear_range(self):	1✔
68	"""Return the total number of work-groups in the nd_range."""
69	num_wg = 1	×
70	for i in range(self.dimensions):	×
71	num_wg *= self.get_group_range(i)	×
72
73	return num_wg	×
74
75	def get_local_range(self, dim):	1✔
76	"""Returns the extent of the SYCL range representing all dimensions
77	of the local range for a specified dimension. This local range may
78	have been provided by the programmer, or chosen by the SYCL runtime.
79	"""
80	return self._local_range[dim]	×
81
82	def get_local_linear_range(self):	1✔
83	"""Return the total number of work-items in the work-group."""
84	num_wi = 1	×
85	for i in range(self.dimensions):	×
86	num_wi *= self.get_local_range(i)	×
87
88	return num_wi	×
89
90	@property	1✔
91	def leader(self):	1✔
92	"""Return true for exactly one work-item in the work-group, if the
93	calling work-item is the leader of the work-group, and false for all
94	other work-items in the work-group.
95
96	The leader of the work-group is determined during construction of the
97	work-group, and is invariant for the lifetime of the work-group. The
98	leader of the work-group is guaranteed to be the work-item with a
99	local id of 0.
100
101
102	Returns:
103	bool: If the work item is the designated leader of the
104	"""
105	return self._leader	×
106
107	@property	1✔
108	def dimensions(self) -> int:	1✔
109	"""Returns the rank of a Group object.
110	Returns:
111	int: Number of dimensions in the Group object
112	"""
113	return self._global_range.ndim	×
114
115	@leader.setter	1✔
116	def leader(self, work_item_id):	1✔
117	"""Sets the leader attribute for the group."""
118	self._leader = work_item_id	1✔
119
120
121	class Item:	1✔
122	"""Analogue to the ``sycl::item`` class.
123
124	Identifies an instance of the function object executing at each point in an
125	:class:`.Range`.
126	"""
127
128	def __init__(self, extent: Range, index: list):	1✔
129	self._extent = extent	1✔
130	self._index = index	1✔
131
132	def get_linear_id(self):	1✔
133	"""Get the linear id associated with this item for all dimensions.
134	Original implementation could be found at ``sycl::item_base`` class.
135
136	Returns:
137	int: The linear id.
138	"""
139	if self.dimensions == 1:	1!
140	return self.get_id(0)	1✔
141	if self.dimensions == 2:	×
142	return self.get_id(0) * self.get_range(1) + self.get_id(1)	×
143	return (	×
144	(self.get_id(0) * self.get_range(1) * self.get_range(2))
145	+ (self.get_id(1) * self.get_range(2))
146	+ (self.get_id(2))
147	)
148
149	def get_id(self, idx):	1✔
150	"""Get the id for a specific dimension.
151
152	Returns:
153	int: The id
154	"""
155	return self._index[idx]	1✔
156
157	def get_linear_range(self):	1✔
158	"""Return the total number of work-items in the work-group."""
159	num_wi = 1	×
160	for i in range(self.dimensions):	×
161	num_wi *= self.get_range(i)	×
162
163	return num_wi	×
164
165	def get_range(self, idx):	1✔
166	"""Get the range size for a specific dimension.
167
168	Returns:
169	int: The size
170	"""
171	return self._extent[idx]	1✔
172
173	@property	1✔
174	def dimensions(self) -> int:	1✔
175	"""Returns the number of dimensions of a Item object.
176
177	Returns:
178	int: Number of dimensions in the Item object
179	"""
180	return self._extent.ndim	1✔
181
182
183	class NdItem:	1✔
184	"""Analogue to the ``sycl::nd_item`` class.
185
186	Identifies an instance of the function object executing at each point in an
187	:class:`.NdRange`.
188	"""
189
190	# TODO: define group type
191	def __init__(self, global_item: Item, local_item: Item, group: Group):	1✔
192	# TODO: assert offset and dimensions
193	self._global_item = global_item	1✔
194	self._local_item = local_item	1✔
195	self._group = group	1✔
196	if self.get_local_linear_id() == 0:	1✔
197	self._group.leader = True	1✔
198
199	def get_global_id(self, idx):	1✔
200	"""Get the global id for a specific dimension.
201
202	Returns:
203	int: The global id
204	"""
205	return self._global_item.get_id(idx)	1✔
206
207	def get_global_linear_id(self):	1✔
208	"""Get the linearized global id for the item for all dimensions.
209
210	Returns:
211	int: The global linear id.
212	"""
213	# Instead of calling self._global_item.get_linear_id(), the linearization
214	# logic is duplicated here so that the method can be JIT compiled by
215	# numba-dpex and works in both Python and Numba nopython modes.
216	if self.dimensions == 1:	×
217	return self.get_global_id(0)	×
218	if self.dimensions == 2:	×
219	return self.get_global_id(0) * self.get_global_range(	×
220	1
221	) + self.get_global_id(1)
222	return (	×
223	(
224	self.get_global_id(0)
225	* self.get_global_range(1)
226	* self.get_global_range(2)
227	)
228	+ (self.get_global_id(1) * self.get_global_range(2))
229	+ (self.get_global_id(2))
230	)
231
232	def get_local_id(self, idx):	1✔
233	"""Get the local id for a specific dimension.
234
235	Returns:
236	int: The local id
237	"""
238	return self._local_item.get_id(idx)	1✔
239
240	def get_local_linear_id(self):	1✔
241	"""Get the local linear id associated with this item for all
242	dimensions.
243
244	Returns:
245	int: The local linear id.
246	"""
247	# Instead of calling self._local_item.get_linear_id(), the linearization
248	# logic is duplicated here so that the method can be JIT compiled by
249	# numba-dpex and works in both Python and Numba nopython modes.
250	if self.dimensions == 1:	1✔
251	return self.get_local_id(0)	1✔
252	if self.dimensions == 2:	1✔
253	return self.get_local_id(0) * self.get_local_range(	1✔
254	1
255	) + self.get_local_id(1)
256	return (	1✔
257	(
258	self.get_local_id(0)
259	* self.get_local_range(1)
260	* self.get_local_range(2)
261	)
262	+ (self.get_local_id(1) * self.get_local_range(2))
263	+ (self.get_local_id(2))
264	)
265
266	def get_global_range(self, idx):	1✔
267	"""Get the global range size for a specific dimension.
268
269	Returns:
270	int: The size
271	"""
272	return self._global_item.get_range(idx)	×
273
274	def get_local_range(self, idx):	1✔
275	"""Get the local range size for a specific dimension.
276
277	Returns:
278	int: The size
279	"""
280	return self._local_item.get_range(idx)	1✔
281
282	def get_local_linear_range(self):	1✔
283	"""Return the total number of work-items in the work-group."""
284	num_wi = 1	×
285	for i in range(self.dimensions):	×
286	num_wi *= self.get_local_range(i)	×
287
288	return num_wi	×
289
290	def get_global_linear_range(self):	1✔
291	"""Return the total number of work-items in the work-group."""
292	num_wi = 1	×
293	for i in range(self.dimensions):	×
294	num_wi *= self.get_global_range(i)	×
295
296	return num_wi	×
297
298	def get_group(self):	1✔
299	"""Returns the group.
300
301	Returns:
302	A group object."""
303	return self._group	×
304
305	@property	1✔
306	def dimensions(self) -> int:	1✔
307	"""Returns the rank of a NdItem object.
308
309	Returns:
310	int: Number of dimensions in the NdItem object
311	"""
312	return self._global_item.dimensions	1✔

IntelPython / numba-dpex / 8290948108

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