o ְi8@sddlZddlmZddlmZddlmZddlmZddlmZd0dd Z d0d d Z d0d d Z d0ddZ ddZ ddZedddZd1ddZddZdZdZdZejdded eed!Zd"Zejd#d$d%ed&d'eed!Zejd#d(d)ed&d*eed!Zd2d,d-Zd.d/ZdS)3N)_core)fusion)_util)_routines_indexing)_routines_statisticsFcC|j||||dS)aReturns the indices of the maximum along an axis. Args: a (cupy.ndarray): Array to take argmax. axis (int): Along which axis to find the maximum. ``a`` is flattened by default. dtype: Data type specifier. out (cupy.ndarray): Output array. keepdims (bool): If ``True``, the axis ``axis`` is preserved as an axis of length one. Returns: cupy.ndarray: The indices of the maximum of ``a`` along an axis. .. note:: ``dtype`` and ``keepdim`` arguments are specific to CuPy. They are not in NumPy. .. note:: ``axis`` argument accepts a tuple of ints, but this is specific to CuPy. NumPy does not support it. .. seealso:: :func:`numpy.argmax` axisdtypeoutkeepdims)argmaxar r r r rH/home/ubuntu/.local/lib/python3.10/site-packages/cupy/_sorting/search.pyr r cC0|jjdvrt|||||dSt|||||S)aReturn the indices of the maximum values in the specified axis ignoring NaNs. For all-NaN slice ``-1`` is returned. Subclass cannot be passed yet, subok=True still unsupported Args: a (cupy.ndarray): Array to take nanargmax. axis (int): Along which axis to find the maximum. ``a`` is flattened by default. Returns: cupy.ndarray: The indices of the maximum of ``a`` along an axis ignoring NaN values. .. note:: For performance reasons, ``cupy.nanargmax`` returns ``out of range values`` for all-NaN slice whereas ``numpy.nanargmax`` raises ``ValueError`` .. seealso:: :func:`numpy.nanargmax` biur)r kindr _statistics _nanargmaxrrrr nanargmax( rcCr)aReturns the indices of the minimum along an axis. Args: a (cupy.ndarray): Array to take argmin. axis (int): Along which axis to find the minimum. ``a`` is flattened by default. dtype: Data type specifier. out (cupy.ndarray): Output array. keepdims (bool): If ``True``, the axis ``axis`` is preserved as an axis of length one. Returns: cupy.ndarray: The indices of the minimum of ``a`` along an axis. .. note:: ``dtype`` and ``keepdim`` arguments are specific to CuPy. They are not in NumPy. .. note:: ``axis`` argument accepts a tuple of ints, but this is specific to CuPy. NumPy does not support it. .. seealso:: :func:`numpy.argmin` r)argminrrrrrArrcCr)aReturn the indices of the minimum values in the specified axis ignoring NaNs. For all-NaN slice ``-1`` is returned. Subclass cannot be passed yet, subok=True still unsupported Args: a (cupy.ndarray): Array to take nanargmin. axis (int): Along which axis to find the minimum. ``a`` is flattened by default. Returns: cupy.ndarray: The indices of the minimum of ``a`` along an axis ignoring NaN values. .. note:: For performance reasons, ``cupy.nanargmin`` returns ``out of range values`` for all-NaN slice whereas ``numpy.nanargmin`` raises ``ValueError`` .. seealso:: :func:`numpy.nanargmin` rr)r rrr _nanargminrrrr nanargmin_rrcCstj|dd|S)aReturn the indices of the elements that are non-zero. Returns a tuple of arrays, one for each dimension of a, containing the indices of the non-zero elements in that dimension. Args: a (cupy.ndarray): array Returns: tuple of arrays: Indices of elements that are non-zero. .. warning:: This function may synchronize the device. .. seealso:: :func:`numpy.nonzero` rarg_name)r check_arraynonzerorrrrr xsr cCstj|dd|dS)aReturn indices that are non-zero in the flattened version of a. This is equivalent to a.ravel().nonzero()[0]. Args: a (cupy.ndarray): input array Returns: cupy.ndarray: Output array, containing the indices of the elements of a.ravel() that are non-zero. .. warning:: This function may synchronize the device. .. seealso:: :func:`numpy.flatnonzero` rrr)rrravelr r!rrr flatnonzerosr# cupy_where)z???->?z?bb->bz?BB->Bz?hh->hz?HH->Hz?ii->iz?II->Iz?ll->lz?LL->Lz?qq->qz?QQ->Qz?ee->ez?ff->fz?hd->dz?Hd->dz?dd->dz?FF->Fz?DD->Dzout0 = in0 ? in1 : in2cCs`|du|dufd}|dkrtd|dkrt|Str'tt|||St|d||S)aReturn elements, either from x or y, depending on condition. If only condition is given, return ``condition.nonzero()``. Args: condition (cupy.ndarray): When True, take x, otherwise take y. x (cupy.ndarray): Values from which to choose on ``True``. y (cupy.ndarray): Values from which to choose on ``False``. Returns: cupy.ndarray: Each element of output contains elements of ``x`` when ``condition`` is ``True``, otherwise elements of ``y``. If only ``condition`` is given, return the tuple ``condition.nonzero()``, the indices where ``condition`` is True. .. warning:: This function may synchronize the device if both ``x`` and ``y`` are omitted. .. seealso:: :func:`numpy.where` NTz)Must provide both 'x' and 'y' or neither.?)count ValueErrorr r _is_fusing _call_ufunc _where_ufuncastype) conditionxymissingrrrwheresr2cCstj|ddt|S)aReturn the indices of the elements that are non-zero. Returns a (N, ndim) dimantional array containing the indices of the non-zero elements. Where `N` is number of non-zero elements and `ndim` is dimension of the given array. Args: a (cupy.ndarray): array Returns: cupy.ndarray: Indices of elements that are non-zero. .. seealso:: :func:`numpy.argwhere` rr)rr _indexing_ndarray_argwherer!rrrargwheres r5zO template __device__ bool _isnan(T val) { return val != val; } a #ifdef __HIP_DEVICE_COMPILE__ #define no_thread_divergence(do_work, to_return) \ if (!is_done) { \ do_work; \ is_done = true; \ } #else #define no_thread_divergence(do_work, to_return) \ do_work; \ if (to_return) { return; } #endif a #ifdef __HIP_DEVICE_COMPILE__ bool is_done = false; #endif // Array is assumed to be monotonically // increasing unless a check is requested with the // `assume_increasing = False` parameter. // `digitize` allows increasing and decreasing arrays. bool inc = true; if (!assume_increasing && n_bins >= 2) { // In the case all the bins are nan the array is considered // to be decreasing in numpy inc = (bins[0] <= bins[n_bins-1]) || (!_isnan(bins[0]) && _isnan(bins[n_bins-1])); } if (_isnan(x)) { long long pos = (inc ? n_bins : 0); if (!side_is_right) { if (inc) { while (pos > 0 && _isnan(bins[pos-1])) { --pos; } } else { while (pos < n_bins && _isnan(bins[pos])) { ++pos; } } } no_thread_divergence( y = pos , true ) } bool greater = false; if (side_is_right) { greater = inc && x >= bins[n_bins-1]; } else { greater = (inc ? x > bins[n_bins-1] : x <= bins[n_bins-1]); } if (greater) { no_thread_divergence( y = n_bins , true ) } long long left = 0; // In the case the bins is all NaNs, digitize // needs to place all the valid values to the right if (!inc) { while (_isnan(bins[left]) && left < n_bins) { ++left; } if (left == n_bins) { no_thread_divergence( y = n_bins , true ) } if (side_is_right && !_isnan(bins[n_bins-1]) && !_isnan(x) && bins[n_bins-1] > x) { no_thread_divergence( y = n_bins , true ) } } long long right = n_bins-1; while (left < right) { long long m = left + (right - left) / 2; bool look_right = true; if (side_is_right) { look_right = (inc ? bins[m] <= x : bins[m] > x); } else { look_right = (inc ? bins[m] < x : bins[m] >= x); } if (look_right) { left = m + 1; } else { right = m; } } no_thread_divergence( y = right , false ) zIS x, raw T bins, int64 n_bins, bool side_is_right, bool assume_increasingzint64 ycupy_searchsorted_kernel)namepreambleaJ #ifdef __HIP_DEVICE_COMPILE__ #define no_thread_divergence(do_work, to_return) \ if (!is_done) { \ do_work; \ is_done = true; \ } #else #define no_thread_divergence(do_work, to_return) \ do_work; \ if (to_return) { \ out = (y == n_bins ? false : bins[y] == x); \ if (invert) out = !out; \ return; \ } #endif z*S x, raw T bins, int64 n_bins, bool invertzbool outze const bool assume_increasing = true; const bool side_is_right = false; long long y; zQ out = (y == n_bins ? false : bins[y] == x); if (invert) out = !out; cupy_exists_kernelzbool out, int64 yzT const bool assume_increasing = true; const bool side_is_right = false; #cupy_exists_and_searchsorted_kernelleftcCst||||dS)aFinds indices where elements should be inserted to maintain order. Find the indices into a sorted array ``a`` such that, if the corresponding elements in ``v`` were inserted before the indices, the order of ``a`` would be preserved. Args: a (cupy.ndarray): Input array. If ``sorter`` is ``None``, then it must be sorted in ascending order, otherwise ``sorter`` must be an array of indices that sort it. v (cupy.ndarray): Values to insert into ``a``. side : {'left', 'right'} If ``left``, return the index of the first suitable location found If ``right``, return the last such index. If there is no suitable index, return either 0 or length of ``a``. sorter : 1-D array_like Optional array of integer indices that sort array ``a`` into ascending order. They are typically the result of :func:`~cupy.argsort`. Returns: cupy.ndarray: Array of insertion points with the same shape as ``v``. .. note:: When a is not in ascending order, behavior is undefined. .. seealso:: :func:`numpy.searchsorted` T) _searchsorted)rvsidesorterrrr searchsortedsr@cCst|tjs tdt|tjstd|jdkrtd|jdkr&td|jdkr4tj|jtj dS|j j dk}|j j dk}|rK|sK| |j }n |rU|sU| |j }|d urr|j j d vrct d |j|jkrmtd ||}tj|jtj d}t|||j|d k|||S)z`assume_increasing` is used in the kernel to skip monotonically increasing or decreasing verification inside the cuda kernel. z'Only int or ndarray are supported for az'Only int or ndarray are supported for vr%z!object too deep for desired arrayz+object of too small depth for desired arrayr)r cN)iuzsorter must be of integer typezsorter.size must equal a.sizeright) isinstancecupyndarrayNotImplementedErrorndimr)sizezerosshapeint64r rr- TypeErrortake_searchsorted_kernel)rr=r>r?assume_increasing a_iscomplex v_iscomplexr0rrrr<s2           r<)NNNF)NN)r;N)rFr cupy._corerrrr3rrr rrrr r# create_ufuncr,r2r5 _preamble _hip_preamble_searchsorted_codeElementwiseKernelrP_exists_kernel_exists_and_searchsorted_kernelr@r<rrrrsh          %O