|
| 1 | +# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation |
| 2 | +# |
| 3 | +# SPDX-License-Identifier: Apache-2.0 |
| 4 | + |
| 5 | +import operator |
| 6 | + |
| 7 | +from llvmlite.ir import IRBuilder |
| 8 | +from numba.core import cgutils, errors, imputils, types |
| 9 | +from numba.core.imputils import impl_ret_borrowed |
| 10 | +from numba.extending import intrinsic, overload_attribute |
| 11 | +from numba.np.arrayobj import _getitem_array_generic as np_getitem_array_generic |
| 12 | +from numba.np.arrayobj import make_array |
| 13 | + |
| 14 | +from numba_dpex.core.types import DpnpNdArray, USMNdArray |
| 15 | +from numba_dpex.core.types.dpctl_types import DpctlSyclQueue |
| 16 | +from numba_dpex.kernel_api_impl.spirv.arrayobj import ( |
| 17 | + _getitem_array_generic as kernel_getitem_array_generic, |
| 18 | +) |
| 19 | +from numba_dpex.kernel_api_impl.spirv.target import SPIRVTargetContext |
| 20 | + |
| 21 | +from .dpctlimpl import lower_builtin |
| 22 | + |
| 23 | +# can't import name because of the circular import |
| 24 | +DPEX_TARGET_NAME = "dpex" |
| 25 | + |
| 26 | +# ========================================================================= |
| 27 | +# Helps to parse dpnp constructor arguments |
| 28 | +# ========================================================================= |
| 29 | + |
| 30 | + |
| 31 | +# TODO: target specific |
| 32 | +@lower_builtin(operator.getitem, USMNdArray, types.Integer) |
| 33 | +@lower_builtin(operator.getitem, USMNdArray, types.SliceType) |
| 34 | +def getitem_arraynd_intp(context, builder, sig, args): |
| 35 | + """ |
| 36 | + Overrding the numba.np.arrayobj.getitem_arraynd_intp to support dpnp.ndarray |
| 37 | +
|
| 38 | + The data model for numba.types.Array and numba_dpex.types.DpnpNdArray |
| 39 | + are different. DpnpNdArray has an extra attribute to store a sycl::queue |
| 40 | + pointer. For that reason, np_getitem_arraynd_intp needs to be overriden so |
| 41 | + that when returning a view of a dpnp.ndarray the sycl::queue pointer |
| 42 | + member in the LLVM IR struct gets properly updated. |
| 43 | + """ |
| 44 | + getitem_call_in_kernel = isinstance(context, SPIRVTargetContext) |
| 45 | + _getitem_array_generic = np_getitem_array_generic |
| 46 | + |
| 47 | + if getitem_call_in_kernel: |
| 48 | + _getitem_array_generic = kernel_getitem_array_generic |
| 49 | + |
| 50 | + aryty, idxty = sig.args |
| 51 | + ary, idx = args |
| 52 | + |
| 53 | + assert aryty.ndim >= 1 |
| 54 | + ary = make_array(aryty)(context, builder, ary) |
| 55 | + |
| 56 | + res = _getitem_array_generic( |
| 57 | + context, builder, sig.return_type, aryty, ary, (idxty,), (idx,) |
| 58 | + ) |
| 59 | + ret = impl_ret_borrowed(context, builder, sig.return_type, res) |
| 60 | + |
| 61 | + if isinstance(sig.return_type, USMNdArray) and not getitem_call_in_kernel: |
| 62 | + array_val = args[0] |
| 63 | + array_ty = sig.args[0] |
| 64 | + sycl_queue_attr_pos = context.data_model_manager.lookup( |
| 65 | + array_ty |
| 66 | + ).get_field_position("sycl_queue") |
| 67 | + sycl_queue_attr = builder.extract_value(array_val, sycl_queue_attr_pos) |
| 68 | + ret = builder.insert_value(ret, sycl_queue_attr, sycl_queue_attr_pos) |
| 69 | + |
| 70 | + return ret |
| 71 | + |
| 72 | + |
| 73 | +@intrinsic(target=DPEX_TARGET_NAME) |
| 74 | +def ol_usm_nd_array_sycl_queue( |
| 75 | + ty_context, |
| 76 | + ty_dpnp_nd_array: DpnpNdArray, |
| 77 | +): |
| 78 | + if not isinstance(ty_dpnp_nd_array, DpnpNdArray): |
| 79 | + raise errors.TypingError("Argument must be DpnpNdArray") |
| 80 | + |
| 81 | + ty_queue: DpctlSyclQueue = ty_dpnp_nd_array.queue |
| 82 | + |
| 83 | + sig = ty_queue(ty_dpnp_nd_array) |
| 84 | + |
| 85 | + def codegen(context, builder: IRBuilder, sig, args: list): |
| 86 | + array_proxy = cgutils.create_struct_proxy(ty_dpnp_nd_array)( |
| 87 | + context, |
| 88 | + builder, |
| 89 | + value=args[0], |
| 90 | + ) |
| 91 | + |
| 92 | + queue_ref = array_proxy.sycl_queue |
| 93 | + |
| 94 | + queue_struct_proxy = cgutils.create_struct_proxy(ty_queue)( |
| 95 | + context, builder |
| 96 | + ) |
| 97 | + |
| 98 | + queue_struct_proxy.queue_ref = queue_ref |
| 99 | + queue_struct_proxy.meminfo = array_proxy.meminfo |
| 100 | + |
| 101 | + # Warning: current implementation prevents whole object from being |
| 102 | + # destroyed as long as sycl_queue attribute is being used. It should be |
| 103 | + # okay since anywere we use it as an argument callee creates a copy |
| 104 | + # so it does not steel reference. |
| 105 | + # |
| 106 | + # We can avoid it by: |
| 107 | + # queue_ref_copy = sycl.dpctl_queue_copy(builder, queue_ref) #noqa E800 |
| 108 | + # queue_struct_proxy.queue_ref = queue_ref_copy #noqa E800 |
| 109 | + # queue_struct->meminfo = |
| 110 | + # nrt->manage_memory(queue_ref_copy, DPCTLEvent_Delete); |
| 111 | + # but it will allocate new meminfo object which can negatively affect |
| 112 | + # performance. |
| 113 | + # Speaking philosophically attribute is a part of the object and as long |
| 114 | + # as nobody can still the reference it is a part of the owner object |
| 115 | + # and lifetime is tied to it. |
| 116 | + # TODO: we want to have queue: queuestruct_t instead of |
| 117 | + # queue_ref: QueueRef as an attribute for DPNPNdArray. |
| 118 | + |
| 119 | + queue_value = queue_struct_proxy._getvalue() |
| 120 | + |
| 121 | + # We need to incref meminfo so that queue model is preventing parent |
| 122 | + # ndarray from being destroyed, that can destroy queue that we are |
| 123 | + # using. |
| 124 | + return imputils.impl_ret_borrowed( |
| 125 | + context, builder, ty_queue, queue_value |
| 126 | + ) |
| 127 | + |
| 128 | + return sig, codegen |
| 129 | + |
| 130 | + |
| 131 | +@overload_attribute(USMNdArray, "sycl_queue", target=DPEX_TARGET_NAME) |
| 132 | +def dpnp_nd_array_sycl_queue(arr): |
| 133 | + """Returns :class:`dpctl.SyclQueue` object associated with USM data. |
| 134 | +
|
| 135 | + This is an overloaded attribute implementation for dpnp.sycl_queue. |
| 136 | +
|
| 137 | + Args: |
| 138 | + arr (numba_dpex.core.types.DpnpNdArray): Input array from which to |
| 139 | + take sycl_queue. |
| 140 | +
|
| 141 | + Returns: |
| 142 | + function: Local function `ol_dpnp_nd_array_sycl_queue()`. |
| 143 | + """ |
| 144 | + |
| 145 | + def get(arr): |
| 146 | + return ol_usm_nd_array_sycl_queue(arr) |
| 147 | + |
| 148 | + return get |
0 commit comments