7211502888

Committed 14 Dec 2023 04:09PM UTC coverage: 79.45% (+0.1%) from 79.335%

Build # 7211502888

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #1243 from IntelPython/feature/apply_builder_pattern_to_kernel_launcher

Apply builder pattern to kernel launcher

Run Details

1195 of 1762 branches covered (0.0%)

Branch coverage included in aggregate %.

126 of 127 new or added lines in 5 files covered. (99.21%)

1 existing line in 1 file now uncovered.

5339 of 6462 relevant lines covered (82.62%)

0.83 hits per line

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

96.3

/numba_dpex/core/datamodel/models.py

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

from numba.core import datamodel, types
from numba.core.datamodel.models import PrimitiveModel, StructModel
from numba.core.extending import register_model

from numba_dpex.core.exceptions import UnreachableError
from numba_dpex.utils import address_space

from ..types import (
    Array,
    DpctlSyclEvent,
    DpctlSyclQueue,
    DpnpNdArray,
    NdRangeType,
    RangeType,
    USMNdArray,
)


def _get_flattened_member_count(ty):
    """Return the number of fields in an instance of a given StructModel."""
    flattened_member_count = 0
    members = ty._members
    for member in members:
        if isinstance(member, types.UniTuple):
            flattened_member_count += member.count
        elif isinstance(
            member,
            (
                types.scalars.Integer,
                types.misc.PyObject,
                types.misc.RawPointer,
                types.misc.CPointer,
                types.misc.MemInfoPointer,
            ),
        ):
            flattened_member_count += 1
        else:
            raise UnreachableError

    return flattened_member_count


class GenericPointerModel(PrimitiveModel):
    def __init__(self, dmm, fe_type):
        adrsp = (
            fe_type.addrspace
            if fe_type.addrspace is not None
            else address_space.GLOBAL
        )
        be_type = dmm.lookup(fe_type.dtype).get_data_type().as_pointer(adrsp)
        super(GenericPointerModel, self).__init__(dmm, fe_type, be_type)


class USMArrayModel(StructModel):
    """A data model to represent a Dpex's array types in LLVM IR.

    Dpex's ArrayModel is based on Numba's ArrayModel for NumPy arrays. The
    dpex model adds an extra address space attribute to all pointer members
    in the array.
    """

    def __init__(self, dmm, fe_type):
        ndim = fe_type.ndim
        members = [
            (
                "meminfo",
                types.CPointer(fe_type.dtype, addrspace=fe_type.addrspace),
            ),
            (
                "parent",
                types.CPointer(types.pyobject, addrspace=fe_type.addrspace),
            ),
            ("nitems", types.intp),
            ("itemsize", types.intp),
            (
                "data",
                types.CPointer(fe_type.dtype, addrspace=fe_type.addrspace),
            ),
            (
                "sycl_queue",
                types.CPointer(types.void, addrspace=fe_type.addrspace),
            ),
            ("shape", types.UniTuple(types.intp, ndim)),
            ("strides", types.UniTuple(types.intp, ndim)),
        ]
        super(USMArrayModel, self).__init__(dmm, fe_type, members)

    @property
    def flattened_field_count(self):
        """Return the number of fields in an instance of a USMArrayModel."""
        return _get_flattened_member_count(self)


class DpnpNdArrayModel(StructModel):
    """Data model for the DpnpNdArray type.

    DpnpNdArrayModel is used by the numba_dpex.types.DpnpNdArray type and
    abstracts the usmarystruct_t C type defined in
    numba_dpex.core.runtime._usmarraystruct.h.

    The DpnpNdArrayModel differs from numba's ArrayModel by including an extra
    member sycl_queue that maps to _usmarraystruct.sycl_queue pointer. The
    _usmarraystruct.sycl_queue pointer stores the C++ sycl::queue pointer that
    was used to allocate the data for the dpnp.ndarray represented by an
    instance of _usmarraystruct.
    """

    def __init__(self, dmm, fe_type):
        ndim = fe_type.ndim
        members = [
            ("meminfo", types.MemInfoPointer(fe_type.dtype)),
            ("parent", types.pyobject),
            ("nitems", types.intp),
            ("itemsize", types.intp),
            ("data", types.CPointer(fe_type.dtype)),
            ("sycl_queue", types.voidptr),
            ("shape", types.UniTuple(types.intp, ndim)),
            ("strides", types.UniTuple(types.intp, ndim)),
        ]
        super(DpnpNdArrayModel, self).__init__(dmm, fe_type, members)

    @property
    def flattened_field_count(self):
        """Return the number of fields in an instance of a DpnpNdArrayModel."""
        return _get_flattened_member_count(self)


class SyclQueueModel(StructModel):
    """Represents the native data model for a dpctl.SyclQueue PyObject.

    Numba-dpex uses a C struct as defined in
    numba_dpex/core/runtime._queuestruct.h to store the required attributes for
    a ``dpctl.SyclQueue`` Python object.

        - ``queue_ref``: An opaque C pointer to an actual SYCL queue C++ object.
        - ``parent``: A PyObject* that stores a reference back to the original
                      ``dpctl.SyclQueue`` PyObject if the native struct is
                      created by unboxing the PyObject.
    """

    def __init__(self, dmm, fe_type):
        members = [
            (
                "meminfo",
                types.MemInfoPointer(types.pyobject),
            ),
            (
                "parent",
                types.pyobject,
            ),
            (
                "queue_ref",
                types.CPointer(types.int8),
            ),
        ]
        super(SyclQueueModel, self).__init__(dmm, fe_type, members)


class SyclEventModel(StructModel):
    """Represents the native data model for a dpctl.SyclEvent PyObject.

    Numba-dpex uses a C struct as defined in
    numba_dpex/core/runtime._eventstruct.h to store the required attributes for
    a ``dpctl.SyclEvent`` Python object.

        - ``event_ref``: An opaque C pointer to an actual SYCL event C++ object.
        - ``parent``: A PyObject* that stores a reference back to the original
                      ``dpctl.SyclEvent`` PyObject if the native struct is
                      created by unboxing the PyObject.
    """

    def __init__(self, dmm, fe_type):
        members = [
            (
                "meminfo",
                types.MemInfoPointer(types.pyobject),
            ),
            (
                "parent",
                types.pyobject,
            ),
            (
                "event_ref",
                types.CPointer(types.int8),
            ),
        ]
        super(SyclEventModel, self).__init__(dmm, fe_type, members)


class RangeModel(StructModel):
    """The native data model for a
    numba_dpex.core.kernel_interface.indexers.Range PyObject.
    """

    def __init__(self, dmm, fe_type):
        members = [
            ("ndim", types.int64),
            ("dim0", types.int64),
            ("dim1", types.int64),
            ("dim2", types.int64),
        ]
        super(RangeModel, self).__init__(dmm, fe_type, members)

    @property
    def flattened_field_count(self):
        """Return the number of fields in an instance of a RangeModel."""
        return _get_flattened_member_count(self)


class NdRangeModel(StructModel):
    """The native data model for a
    numba_dpex.core.kernel_interface.indexers.NdRange PyObject.
    """

    def __init__(self, dmm, fe_type):
        members = [
            ("ndim", types.int64),
            ("gdim0", types.int64),
            ("gdim1", types.int64),
            ("gdim2", types.int64),
            ("ldim0", types.int64),
            ("ldim1", types.int64),
            ("ldim2", types.int64),
        ]
        super(NdRangeModel, self).__init__(dmm, fe_type, members)

    @property
    def flattened_field_count(self):
        """Return the number of fields in an instance of a NdRangeModel."""
        return _get_flattened_member_count(self)


def _init_data_model_manager() -> datamodel.DataModelManager:
    """Initializes a DpexKernelTarget-specific data model manager.

    SPIRV kernel functions for certain types of devices require an explicit
    address space qualifier for pointers. For OpenCL HD Graphics
    devices, defining a kernel function (spir_kernel calling convention) with
    pointer arguments that have no address space qualifier causes a run time
    crash. For this reason, numba-dpex defines two separate data
    models: USMArrayModel and DpnpNdArrayModel. When a dpnp.ndarray object is
    passed as an argument to a ``numba_dpex.kernel`` decorated function it uses
    the USMArrayModel and when passed to a ``numba_dpex.dpjit`` decorated
    function it uses the DpnpNdArrayModel. The difference is due to the fact
    that inside a ``dpjit`` decorated function a dpnp.ndarray object can be
    passed to any other regular function.

    Returns:
        DataModelManager: A numba-dpex DpexKernelTarget-specific data model
        manager
    """
    dmm = datamodel.default_manager.copy()
    dmm.register(types.CPointer, GenericPointerModel)
    dmm.register(Array, USMArrayModel)

    # Register the USMNdArray type to USMArrayModel in numba_dpex's data model
    # manager. The dpex_data_model_manager is used by the DpexKernelTarget
    dmm.register(USMNdArray, USMArrayModel)

    # Register the DpnpNdArray type to USMArrayModel in numba_dpex's data model
    # manager. The dpex_data_model_manager is used by the DpexKernelTarget
    dmm.register(DpnpNdArray, USMArrayModel)

    # Register the DpctlSyclQueue type to SyclQueueModel in numba_dpex's data
    # model manager. The dpex_data_model_manager is used by the DpexKernelTarget
    dmm.register(DpctlSyclQueue, SyclQueueModel)

    return dmm


dpex_data_model_manager = _init_data_model_manager()


# Register the USMNdArray type to USMArrayModel in numba's default data model
# manager
register_model(USMNdArray)(USMArrayModel)

# Register the DpnpNdArray type to DpnpNdArrayModel in numba's default data
# model manager
register_model(DpnpNdArray)(DpnpNdArrayModel)

# Register the DpctlSyclQueue type
register_model(DpctlSyclQueue)(SyclQueueModel)

# Register the DpctlSyclEvent type
register_model(DpctlSyclEvent)(SyclEventModel)

# Register the RangeType type
register_model(RangeType)(RangeModel)

# Register the NdRangeType type
register_model(NdRangeType)(NdRangeModel)

1	# SPDX-FileCopyrightText: 2022 - 2023 Intel Corporation
2	#
3	# SPDX-License-Identifier: Apache-2.0
4
5	from numba.core import datamodel, types	1✔
6	from numba.core.datamodel.models import PrimitiveModel, StructModel	1✔
7	from numba.core.extending import register_model	1✔
8
9	from numba_dpex.core.exceptions import UnreachableError	1✔
10	from numba_dpex.utils import address_space	1✔
11
12	from ..types import (	1✔
13	Array,
14	DpctlSyclEvent,
15	DpctlSyclQueue,
16	DpnpNdArray,
17	NdRangeType,
18	RangeType,
19	USMNdArray,
20	)
21
22
23	def _get_flattened_member_count(ty):	1✔
24	"""Return the number of fields in an instance of a given StructModel."""
25	flattened_member_count = 0	1✔
26	members = ty._members	1✔
27	for member in members:	1✔
28	if isinstance(member, types.UniTuple):	1✔
29	flattened_member_count += member.count	1✔
30	elif isinstance(	1!
31	member,
32	(
33	types.scalars.Integer,
34	types.misc.PyObject,
35	types.misc.RawPointer,
36	types.misc.CPointer,
37	types.misc.MemInfoPointer,
38	),
39	):
40	flattened_member_count += 1	1✔
41	else:
42	raise UnreachableError	×
43
44	return flattened_member_count	1✔
45
46
47	class GenericPointerModel(PrimitiveModel):	1✔
48	def __init__(self, dmm, fe_type):	1✔
49	adrsp = (	1✔
50	fe_type.addrspace
51	if fe_type.addrspace is not None
52	else address_space.GLOBAL
53	)
54	be_type = dmm.lookup(fe_type.dtype).get_data_type().as_pointer(adrsp)	1✔
55	super(GenericPointerModel, self).__init__(dmm, fe_type, be_type)	1✔
56
57
58	class USMArrayModel(StructModel):	1✔
59	"""A data model to represent a Dpex's array types in LLVM IR.
60
61	Dpex's ArrayModel is based on Numba's ArrayModel for NumPy arrays. The
62	dpex model adds an extra address space attribute to all pointer members
63	in the array.
64	"""
65
66	def __init__(self, dmm, fe_type):	1✔
67	ndim = fe_type.ndim	1✔
68	members = [	1✔
69	(
70	"meminfo",
71	types.CPointer(fe_type.dtype, addrspace=fe_type.addrspace),
72	),
73	(
74	"parent",
75	types.CPointer(types.pyobject, addrspace=fe_type.addrspace),
76	),
77	("nitems", types.intp),
78	("itemsize", types.intp),
79	(
80	"data",
81	types.CPointer(fe_type.dtype, addrspace=fe_type.addrspace),
82	),
83	(
84	"sycl_queue",
85	types.CPointer(types.void, addrspace=fe_type.addrspace),
86	),
87	("shape", types.UniTuple(types.intp, ndim)),
88	("strides", types.UniTuple(types.intp, ndim)),
89	]
90	super(USMArrayModel, self).__init__(dmm, fe_type, members)	1✔
91
92	@property	1✔
93	def flattened_field_count(self):	1✔
94	"""Return the number of fields in an instance of a USMArrayModel."""
95	return _get_flattened_member_count(self)	1✔
96
97
98	class DpnpNdArrayModel(StructModel):	1✔
99	"""Data model for the DpnpNdArray type.
100
101	DpnpNdArrayModel is used by the numba_dpex.types.DpnpNdArray type and
102	abstracts the usmarystruct_t C type defined in
103	numba_dpex.core.runtime._usmarraystruct.h.
104
105	The DpnpNdArrayModel differs from numba's ArrayModel by including an extra
106	member sycl_queue that maps to _usmarraystruct.sycl_queue pointer. The
107	_usmarraystruct.sycl_queue pointer stores the C++ sycl::queue pointer that
108	was used to allocate the data for the dpnp.ndarray represented by an
109	instance of _usmarraystruct.
110	"""
111
112	def __init__(self, dmm, fe_type):	1✔
113	ndim = fe_type.ndim	1✔
114	members = [	1✔
115	("meminfo", types.MemInfoPointer(fe_type.dtype)),
116	("parent", types.pyobject),
117	("nitems", types.intp),
118	("itemsize", types.intp),
119	("data", types.CPointer(fe_type.dtype)),
120	("sycl_queue", types.voidptr),
121	("shape", types.UniTuple(types.intp, ndim)),
122	("strides", types.UniTuple(types.intp, ndim)),
123	]
124	super(DpnpNdArrayModel, self).__init__(dmm, fe_type, members)	1✔
125
126	@property	1✔
127	def flattened_field_count(self):	1✔
128	"""Return the number of fields in an instance of a DpnpNdArrayModel."""
UNCOV 129	return _get_flattened_member_count(self)	×
130
131
132	class SyclQueueModel(StructModel):	1✔
133	"""Represents the native data model for a dpctl.SyclQueue PyObject.
134
135	Numba-dpex uses a C struct as defined in
136	numba_dpex/core/runtime._queuestruct.h to store the required attributes for
137	a ``dpctl.SyclQueue`` Python object.
138
139	- ``queue_ref``: An opaque C pointer to an actual SYCL queue C++ object.
140	- ``parent``: A PyObject* that stores a reference back to the original
141	``dpctl.SyclQueue`` PyObject if the native struct is
142	created by unboxing the PyObject.
143	"""
144
145	def __init__(self, dmm, fe_type):	1✔
146	members = [	1✔
147	(
148	"meminfo",
149	types.MemInfoPointer(types.pyobject),
150	),
151	(
152	"parent",
153	types.pyobject,
154	),
155	(
156	"queue_ref",
157	types.CPointer(types.int8),
158	),
159	]
160	super(SyclQueueModel, self).__init__(dmm, fe_type, members)	1✔
161
162
163	class SyclEventModel(StructModel):	1✔
164	"""Represents the native data model for a dpctl.SyclEvent PyObject.
165
166	Numba-dpex uses a C struct as defined in
167	numba_dpex/core/runtime._eventstruct.h to store the required attributes for
168	a ``dpctl.SyclEvent`` Python object.
169
170	- ``event_ref``: An opaque C pointer to an actual SYCL event C++ object.
171	- ``parent``: A PyObject* that stores a reference back to the original
172	``dpctl.SyclEvent`` PyObject if the native struct is
173	created by unboxing the PyObject.
174	"""
175
176	def __init__(self, dmm, fe_type):	1✔
177	members = [	1✔
178	(
179	"meminfo",
180	types.MemInfoPointer(types.pyobject),
181	),
182	(
183	"parent",
184	types.pyobject,
185	),
186	(
187	"event_ref",
188	types.CPointer(types.int8),
189	),
190	]
191	super(SyclEventModel, self).__init__(dmm, fe_type, members)	1✔
192
193
194	class RangeModel(StructModel):	1✔
195	"""The native data model for a
196	numba_dpex.core.kernel_interface.indexers.Range PyObject.
197	"""
198
199	def __init__(self, dmm, fe_type):	1✔
200	members = [	1✔
201	("ndim", types.int64),
202	("dim0", types.int64),
203	("dim1", types.int64),
204	("dim2", types.int64),
205	]
206	super(RangeModel, self).__init__(dmm, fe_type, members)	1✔
207
208	@property	1✔
209	def flattened_field_count(self):	1✔
210	"""Return the number of fields in an instance of a RangeModel."""
211	return _get_flattened_member_count(self)	1✔
212
213
214	class NdRangeModel(StructModel):	1✔
215	"""The native data model for a
216	numba_dpex.core.kernel_interface.indexers.NdRange PyObject.
217	"""
218
219	def __init__(self, dmm, fe_type):	1✔
220	members = [	1✔
221	("ndim", types.int64),
222	("gdim0", types.int64),
223	("gdim1", types.int64),
224	("gdim2", types.int64),
225	("ldim0", types.int64),
226	("ldim1", types.int64),
227	("ldim2", types.int64),
228	]
229	super(NdRangeModel, self).__init__(dmm, fe_type, members)	1✔
230
231	@property	1✔
232	def flattened_field_count(self):	1✔
233	"""Return the number of fields in an instance of a NdRangeModel."""
234	return _get_flattened_member_count(self)	1✔
235
236
237	def _init_data_model_manager() -> datamodel.DataModelManager:	1✔
238	"""Initializes a DpexKernelTarget-specific data model manager.
239
240	SPIRV kernel functions for certain types of devices require an explicit
241	address space qualifier for pointers. For OpenCL HD Graphics
242	devices, defining a kernel function (spir_kernel calling convention) with
243	pointer arguments that have no address space qualifier causes a run time
244	crash. For this reason, numba-dpex defines two separate data
245	models: USMArrayModel and DpnpNdArrayModel. When a dpnp.ndarray object is
246	passed as an argument to a ``numba_dpex.kernel`` decorated function it uses
247	the USMArrayModel and when passed to a ``numba_dpex.dpjit`` decorated
248	function it uses the DpnpNdArrayModel. The difference is due to the fact
249	that inside a ``dpjit`` decorated function a dpnp.ndarray object can be
250	passed to any other regular function.
251
252	Returns:
253	DataModelManager: A numba-dpex DpexKernelTarget-specific data model
254	manager
255	"""
256	dmm = datamodel.default_manager.copy()	1✔
257	dmm.register(types.CPointer, GenericPointerModel)	1✔
258	dmm.register(Array, USMArrayModel)	1✔
259
260	# Register the USMNdArray type to USMArrayModel in numba_dpex's data model
261	# manager. The dpex_data_model_manager is used by the DpexKernelTarget
262	dmm.register(USMNdArray, USMArrayModel)	1✔
263
264	# Register the DpnpNdArray type to USMArrayModel in numba_dpex's data model
265	# manager. The dpex_data_model_manager is used by the DpexKernelTarget
266	dmm.register(DpnpNdArray, USMArrayModel)	1✔
267
268	# Register the DpctlSyclQueue type to SyclQueueModel in numba_dpex's data
269	# model manager. The dpex_data_model_manager is used by the DpexKernelTarget
270	dmm.register(DpctlSyclQueue, SyclQueueModel)	1✔
271
272	return dmm	1✔
273
274
275	dpex_data_model_manager = _init_data_model_manager()	1✔
276
277
278	# Register the USMNdArray type to USMArrayModel in numba's default data model
279	# manager
280	register_model(USMNdArray)(USMArrayModel)	1✔
281
282	# Register the DpnpNdArray type to DpnpNdArrayModel in numba's default data
283	# model manager
284	register_model(DpnpNdArray)(DpnpNdArrayModel)	1✔
285
286	# Register the DpctlSyclQueue type
287	register_model(DpctlSyclQueue)(SyclQueueModel)	1✔
288
289	# Register the DpctlSyclEvent type
290	register_model(DpctlSyclEvent)(SyclEventModel)	1✔
291
292	# Register the RangeType type
293	register_model(RangeType)(RangeModel)	1✔
294
295	# Register the NdRangeType type
296	register_model(NdRangeType)(NdRangeModel)	1✔

IntelPython / numba-dpex / 7211502888

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