o W۷iAB@sddlmZddlZddlmZddlmZddlmZddlm Z ddl m Z ddl Z e dd d d d d dZddZddZdddZddZddZdddZejZZejddge ddRd d!d"Zejd#d$d%d&d'd(d)d*d+d,d-dg e ddd.d/Rd0d1d"ZejZejd2dge ddRd3d4d"Zd5Zd6Zejd7d$d%efd&d'efd(d)efd*d+efd,d-efdg e ddd.efd/efRd8d9d"Zd:Zed7d$d%efd&d'efd(d)efd*d+efd,d-efdg e ddd.efd/efRd8Z ejd;dd?Rd@dAd"Z!dBZ"dCZ#ejdDdEdFdGdHdIdJdKdLdMdNdOde#fd?e#fdQe#fdRe#fRdSe"dTdUdVdWZ$dXZ%ejdYdEdFdGdHdIdJdKdLdMdNdOde%fd?e%fdQe%fdRe%fRdZe"d[d\d]dWZ&ejd^dEdFdGdHdIdJdKdLdMdNdOd_g e d=d`dPdadbdQdRRdSdcd"Z'ejdddEdFdGdHdIdJdKdLdMdNdOdeg e d=dfdPdgdhdQdRRdZdid"Z(djZ)ejdkdldmdndodpdqdrdsdtdudvdwg e dxdydzdPd{d|d}d~Rde)ddZ*dddZ+dddZ,dddZ-ej.j/ddddZ0dddZ1dS)) annotationsN)_core)_routines_math)fusion) stride_tricks) bf16_loopz T x1, T x2zT yzx1 * x2za + bzy = a0 dot_productcCsL|jdks |jdkrtd|jjdvs|jjdvrdSt|||r$dSdS)Nz&Only 1d inputs are supported currentlybuidirectfft)ndimNotImplementedErrordtypekind_fftconv_faster)in1in2moderE/home/ubuntu/vllm_env/lib/python3.10/site-packages/cupy/_math/misc.py_choose_conv_methods rcCsdS)zJ .. seealso:: :func: `scipy.signal._signaltools._fftconv_faster` Tr)xhrrrrr&srfullcCs|jdkr td|jdkrtd|jdkrtd|j|jkr&||}}|}|}t|||}|dkr@t|||}|S|dkrLt|||}|Std) arReturns the discrete, linear convolution of two one-dimensional sequences. Args: a (cupy.ndarray): first 1-dimensional input. v (cupy.ndarray): second 1-dimensional input. mode (str, optional): `valid`, `same`, `full` Returns: cupy.ndarray: Discrete, linear convolution of a and v. .. seealso:: :func:`numpy.convolve` rza cannot be emptyzv cannot be emptyr z"v cannot be multidimensional arrayr r zUnsupported method)size ValueErrorrravelr _dot_convolve _fft_convolve)avrmethodoutrrrconvolve/s$        r%cCsdd}|jd|jdkr||}}d|jdd}|jjdks&|jjdkr0tjjtjj}}n tjjtjj}}t||}|jd|jd}}ddl } | j j ||d} ||| } ||| } || | | } |dkrxd||d}}n!|dkr|dd|}||}n|dkr|d|}}nt d | d ||f} |jd vrt | } | j|d d S)Nrr crsamevalidz5acceptable mode flags are `valid`, `same`, or `full`..iuFcopy)shaperrcupyr ifftrfftirfft result_typecupyx.scipy.fftscipy next_fast_lenraroundastype)a1a2roffsetr r0rn1n2cupyxout_sizefa1fa2r$startendrrrr Rs8       r c Cs d}|j|jkr||}}d|jd}t||}|j|j}}|j|dd}|j|dd}|dkrB||d}t||d}n%|dkr]|}|dd|}t||d||f}n |dkrg||d}|jd} t|||f| | f}t||ddd dd } | S) Nrr r'Fr,rr)r*r&)axis) rr/r3r8padstridesr as_strided _dot_kernel) r9r:rr;rr<r=r?pad_sizestrideoutputrrrrys*      rcCs.trtjtj||||dS|j|||dS)aClips the values of an array to a given interval. This is equivalent to ``maximum(minimum(a, a_max), a_min)``, while this function is more efficient. Args: a (cupy.ndarray): The source array. a_min (scalar, cupy.ndarray or None): The left side of the interval. When it is ``None``, it is ignored. a_max (scalar, cupy.ndarray or None): The right side of the interval. When it is ``None``, it is ignored. out (cupy.ndarray): Output array. Returns: cupy.ndarray: Clipped array. .. seealso:: :func:`numpy.clip` Notes ----- When `a_min` is greater than `a_max`, `clip` returns an array in which all values are equal to `a_max`. r$)r _is_fusing _call_ufunc_mathclip)r!a_mina_maxr$rrrrPs rP cupy_cbrtze->ezf->fzd->dzout0 = cbrt(in0)zJElementwise cube root function. .. seealso:: :data:`numpy.cbrt` )doc cupy_squarezb->bzB->Bzh->hzH->Hzi->izI->Izl->lzL->Lzq->qzQ->QzF->FzD->Dzout0 = in0 * in0zIElementwise square function. .. seealso:: :data:`numpy.square` cupy_fabszout0 = abs(in0)zuCalculates absolute values element-wise. Only real values are handled. .. seealso:: :data:`numpy.fabs` zout0 = in0 > 0z out0 = in0 / abs(in0) cupy_signzout0 = (in0 > 0) - (in0 < 0)zElementwise sign function. It returns -1, 0, or 1 depending on the sign of the input. .. seealso:: :data:`numpy.sign` z if (in0.real() == 0) { out0 = (in0.imag() > 0) - (in0.imag() < 0); } else { out0 = (in0.real() > 0) - (in0.real() < 0); } cupy_heavisideee->er'ff->fdd->dz if (isnan(in0)) { out0 = in0; } else if (in0 == 0) { out0 = in1; } else { out0 = (in0 > 0); } zTCompute the Heaviside step function. .. seealso:: :data:`numpy.heaviside` z; #ifndef NAN #define NAN __int_as_float(0x7fffffff) #endif zLout0 = (isnan(in0) | isnan(in1)) ? out0_type(NAN) : out0_type(max(in0, in1)) cupy_maximumz??->?zbb->bzBB->Bzhh->hzHH->Hzii->izII->Izll->lzLL->Lzqq->qzQQ->Q)codezFF->FzDD->Dzout0 = max(in0, in1)zTakes the maximum of two arrays elementwise. If NaN appears, it returns the NaN. .. seealso:: :data:`numpy.maximum` )OP_MAXr r max)preamblerT cutensor_op scatter_opzLout0 = (isnan(in0) | isnan(in1)) ? out0_type(NAN) : out0_type(min(in0, in1)) cupy_minimumzout0 = min(in0, in1)zTakes the minimum of two arrays elementwise. If NaN appears, it returns the NaN. .. seealso:: :data:`numpy.minimum` )OP_MINr r min cupy_fmax)rYout0 = fmax(in0, in1)rg)rZrg)r[rgzTakes the maximum of two arrays elementwise. If NaN appears, it returns the other operand. .. seealso:: :data:`numpy.fmax` cupy_fmin)rYout0 = fmin(in0, in1)ri)rZri)r[rizTakes the minimum of two arrays elementwise. If NaN appears, it returns the other operand. .. seealso:: :data:`numpy.fmin` a template __device__ T nan_to_num(T x, T nan, T posinf, T neginf) { if (isnan(x)) return nan; if (isinf(x)) return x > 0 ? posinf : neginf; return x; } template __device__ complex nan_to_num(complex x, T nan, T posinf, T neginf) { T re = nan_to_num(x.real(), nan, posinf, neginf); T im = nan_to_num(x.imag(), nan, posinf, neginf); return complex(re, im); } cupy_nan_to_num_z????->?zbbbb->bzBBBB->Bzhhhh->hzHHHH->Hziiii->izIIII->Izllll->lzLLLL->Lzqqqq->qzQQQQ->Q)zeeee->e%out0 = nan_to_num(in0, in1, in2, in3)r rk)zffff->frk)zdddd->drk)zFfff->Frk)zDddd->Drkz out0 = in0zQElementwise nan_to_num function. .. seealso:: :func:`numpy.nan_to_num` )r`rTcCsb|jdvr t|j}|jdvrd}n|dur+|dur+|r%t|jnt|j}t||S)NFDefdr)charr/rlowerfinforer_ asanyarray)rrnegrrr_check_nan_infs   rtTcCsft|tjs td|}n |rt|n|}|j}t||}t||d}t||d}t|||||dS)zReplace NaN with zero and infinity with large finite numbers (default behaviour) or with the numbers defined by the user using the `nan`, `posinf` and/or `neginf` keywords. .. seealso:: :func:`numpy.nan_to_num` rFTrL) isinstancer/ndarrayr empty_likerrt _nan_to_num)rr-nanposinfneginfr$rrrr nan_to_nums    r}dcCsTt|jjtjs |S|dkrt|jj}|j|}tt |j |kr(|j }|S)a#If input is complex with all imaginary parts close to zero, return real parts. "Close to zero" is defined as `tol` * (machine epsilon of the type for `a`). .. warning:: This function may synchronize the device. .. seealso:: :func:`numpy.real_if_close` r ) issubclassrtyper/complexfloatingnumpyrqepsallabsoluteimagreal)r!tolfrrr real_if_closes  r)for_each_devicecCsLd}|d7}d}|r d}nd}|d7}|tjjj7}d}tj|||d|d S) Nzraw V x, raw U idx, z2raw W fx, raw Y fy, U len, raw Y left, raw Y rightzZ yztypedef double real_t; ztypedef Z real_t; ztypedef Z value_t; a U x_idx = idx[i] - 1; if ( _isnan(x[i]) ) { y = x[i]; } else if (x_idx < 0) { y = left[0]; } else if (x[i] == fx[len - 1]) { // searchsorted cannot handle both of the boundary points, // so we must detect and correct ourselves... y = fy[len - 1]; } else if (x_idx >= len - 1) { y = right[0]; } else { const Z slope = (value_t)(fy[x_idx+1] - fy[x_idx]) / \ ((real_t)fx[x_idx+1] - (real_t)fx[x_idx]); Z out = slope * ((real_t)x[i] - (real_t)fx[x_idx]) \ + (value_t)fy[x_idx]; if (_isnan(out)) { out = slope * ((real_t)x[i] - (real_t)fx[x_idx+1]) \ + (value_t)fy[x_idx+1]; if (_isnan(out) && (fy[x_idx] == fy[x_idx+1])) { out = fy[x_idx]; } } y = out; } cupy_interp)r`)r/_sortingsearch _preambleElementwiseKernel) is_complex in_params out_paramsr`r]rrr_get_interp_kernels rc Cs|jdks |jdkrtd|j|jkrtd|jdkr!td|jjs)tdt||}t|tj s?t d |tj |dur|dkrKtd t |}d}d}| tj }| tj }||;}||;}t|}||}||}t|d d|||dd|f}t|d d||ddf}|jjsJ|jjsJ|jjd krd nd }tj|j|d} tj||dd} |dur|dnt||j}|dur|d nt||j}t|d k} | || |||j||| | S)a One-dimensional linear interpolation. Args: x (cupy.ndarray): a 1D array of points on which the interpolation is performed. xp (cupy.ndarray): a 1D array of points on which the function values (``fp``) are known. fp (cupy.ndarray): a 1D array containing the function values at the the points ``xp``. left (float or complex): value to return if ``x < xp[0]``. Default is ``fp[0]``. right (float or complex): value to return if ``x > xp[-1]``. Default is ``fp[-1]``. period (None or float): a period for the x-coordinates. Parameters ``left`` and ``right`` are ignored if ``period`` is specified. Default is ``None``. Returns: cupy.ndarray: The interpolated values, same shape as ``x``. .. note:: This function may synchronize if ``left`` or ``right`` is not already on the device. .. seealso:: :func:`numpy.interp` r zxp and fp must be 1D arraysz$fp and xp are not of the same lengthrzarray of sample points is emptyz-Non-C-contiguous x is currently not supportedzACannot cast array data from {} to {} according to the rule 'safe'Nzperiod must be a non-zero valuer&r(Dd)rright)side)rrrflags c_contiguousrr/ common_typecan_castfloat64 TypeErrorformatabsr8argsort concatenaterremptyr. searchsortedarrayr) rxpfpleftrperiodx_dtypeasort_xp out_dtyperKidxkernrrrinterpsL       (    r)r)N)TruNN)r~)NNN)2 __future__rr/r cupy._corerrOrcupy.libr cupy._utilrrReductionKernelrHrrr%r rrPsqrt sqrt_fixed create_ufunccbrtsquarerfabs_unsigned_sign _complex_signsign_legacy_complex_sign _legacy_sign heaviside_float_preamble_float_maximummaximum_float_minimumminimumfmaxfmin_nan_to_num_preambleryrtr}r_utilmemoizerrrrrrs        #'  "                    *