o X۷iC@s|UddlmZddlmZddlmZddlZddlZddlm Z ddl m Z ddl m Z ddl mZdd lmZdd lmZdd lmZdd lmZdd l mZddlmZGdddeZGdddeZGdddeZGdddeZGdddeZGdddeZGdddeZGdddeZGdd d eZGd!d"d"eZ Gd#d$d$eZ!e"ee#ee$ee%eiZ&d%e'd&<eZ(eZ)eZ*eZ+ed'Z,ed(Z-e!Z.ed)d*e j/rd+nd,Z0ed-d.Z1ed/d0Z2ed1d2Z3ed3d2Z4ed4d5Z5ed6d5Z6ed7d2e j/r d+nd8Z7ed9d2Z8ed:d2Z9ed;d2Z:de;Ze dAe;Z?dS)B) annotations)Any)MappingN)runtime)device) _cuda_types)_cuda_typerules) BuiltinFunc)Data)Constant)Range)_compile)reducecs0eZdZddfdd ZddddZZS) RangeFuncNunrollctdS)aRange with loop unrolling support. Args: start (int): Same as that of built-in :obj:`range`. stop (int): Same as that of built-in :obj:`range`. step (int): Same as that of built-in :obj:`range`. unroll (int or bool or None): - If `True`, add ``#pragma unroll`` directive before the loop. - If `False`, add ``#pragma unroll(1)`` directive before the loop to disable unrolling. - If an `int`, add ``#pragma unroll(n)`` directive before the loop, where the integer ``n`` means the number of iterations to unroll. - If `None` (default), leave the control of loop unrolling to the compiler (no ``#pragma``). .. seealso:: `#pragma unroll`_ .. _#pragma unroll: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#pragma-unroll Nsuper__call__)selfrargs __class__N/home/ubuntu/vllm_env/lib/python3.10/site-packages/cupyx/jit/_builtin_funcs.pyrszRangeFunc.__call__c Gst|dkr tdt|dkrtd|dtd}}}n*t|dkr4|d|dtd}}}nt|dkr@|\}}}n tdt||durtdd||||fDs^td |j}t|tsut|tsutd t|j |d ur{d}|d usd|krd ksnt dt|tr|jdk}n |j j jdkrd }nd}t||}t||}t||}|j j jdvrtd|j j jdvrtd|j j jdvrtd|jdkrtt}n |jdkr|j }nJt||||||dS)Nrz)range expected at least 1 argument, got 0z'range expected at most 3 argument, got css|]}t|tVqdSN) isinstancer ).0xrrr Csz!RangeFunc.call..zCloop unrolling requires constant start, stop, step and unroll valuez-unroll value expected to be of type int, got FTlzUloop unrolling is ignored as the unroll value is non-positive or greater than INT_MAXuiuz%range supports only for integer type.numpycudar)len TypeErrorr allobjr intbooltype__name__warningswarnctypedtypekindr initmoderScalarr ) renvrrstartstopstepstep_is_positiver2rrrcall5sf               zRangeFunc.callr/ __module__ __qualname__rr= __classcell__rrrrrsrc@eZdZddZdS)LenFunccOsrt|dkrtdt||rtd|d}t|jtjs$td|jjs,tdtd|jdt d S) Nrz#len() expects only 1 argument, got #keyword arguments are not supportedrzlen() supports only array typezlen() of unsized arrayzstatic_cast(z .shape()[0])q) r(r)r r2rCArrayndimr coder7)rr8rkwdsargrrrr=rs z LenFunc.callNr/r?r@r=rrrrrCp rCc@rB)MinFuncc<t|dkrtdt||rtdtfdd|S)Nrz(min() expects at least 2 arguments, got rDcttj||fdSr)r _call_ufunccupyminimumabr8rrzMinFunc.call..r(r)rrr8rrIrrVrr=   z MinFunc.callNrKrrrrrMrLrMc@rB)MaxFunccrN)Nrz(max() expects at least 2 arguments, got rDcrOr)r rPrQmaximumrSrVrrrWrXzMaxFunc.call..rYrZrrVrr=r[z MaxFunc.callNrKrrrrr\rLr\cs$eZdZfddZddZZS) SyncThreadscr)zCalls ``__syncthreads()``. .. seealso:: `Synchronization functions`_ .. _Synchronization functions: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#synchronization-functions NrrrrrrzSyncThreads.__call__cCs tdtjS)Nz__syncthreads())r rvoidrr8rrr call_consts zSyncThreads.call_constr/r?r@rrcrArrrrr^s  r^cs0eZdZddfdd ZddddZZS)SyncWarp)maskcr)a:Calls ``__syncwarp()``. Args: mask (int): Active threads in a warp. Default is 0xffffffff. .. seealso:: `Synchronization functions`_ .. _Synchronization functions: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#synchronization-functions Nr)rrgrrrr zSyncWarp.__call__NcCstjr|durtd|dtd}|rDt|tr-d|jkr(dks-tdtdt |t j d|}t ||}d|jd}nd }t |t jS) Nmask  is ignored on HIPrrfmask is out of range same_kindz __syncwarp()z __syncwarp())ris_hipr0r1RuntimeWarningr r r+ ValueErrorr _astype_scalarrint32r r5rHra)rr8rgrHrrrr=s"    z SyncWarp.callr>rrrrres recs(eZdZdfdd ZdddZZS) SharedMemoryNcr)aAllocates shared memory and returns it as a 1-D array. Args: dtype (dtype): The dtype of the returned array. size (int or None): If ``int`` type, the size of static shared memory. If ``None``, declares the shared memory with extern specifier. alignment (int or None): Enforce the alignment via __align__(N). Nr)rr3size alignmentrrrrrhzSharedMemory.__call__cCsN|jdd}t|}t|t|||}||j|<||j|<t|t|S)N_smem)prefix) get_fresh_variable_namerto_ctyper r SharedMemdeclslocalsPtr)rr8r3rtrunamer2varrrrrcs    zSharedMemory.call_constrrdrrrrrss rscs0eZdZddZdfdd ZdddZZS) AtomicOpcCs.||_d||_||_d|jd}||_dS)Natomic Calls the ``a`` function to operate atomically on ``array[index]``. Please refer to `Atomic Functions`_ for detailed explanation. Args: array: A :class:`cupy.ndarray` to index over. index: A valid index such that the address to the corresponding array element ``array[index]`` can be computed. value: Represent the value to use for the specified operation. For the case of :obj:`atomic_cas`, this is the value for ``array[index]`` to compare with. alt_value: Only used in :obj:`atomic_cas` to represent the value to swap to. .. seealso:: `Numba's corresponding atomic functions`_ .. _Atomic Functions: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions .. _Numba's corresponding atomic functions: https://numba.readthedocs.io/en/stable/cuda-reference/kernel.html#synchronization-and-atomic-operations _op_name_dtypes__doc__ropdtypesdocrrr__init__s   zAtomicOp.__init__NctdSrr)rarrayindexvalue alt_valuerrrrzAtomicOp.__call__c Cs2|j}|j}t||}t|jtjtjfst dt |||}|j} | j j |jvr7t d|d| j dt || d|}t||}|dkr|dusOJ| j jdkrattdkratd t || d|}t||}|d |jd |jd |jd } n|dusJ|d |jd |jd } t| | S) Nz)The first argument must be of array type.`` does not support  input.rlCASHFz5uint16 atomic operation is not supported before sm_70z(&, rm)rrr r5r r2rrFr}r)r _indexingr3r~rrqcharr,rget_compute_capability RuntimeErrorrH) rr8rrrvalue2r~rtargetr2rHrrrr=s0     $  z AtomicOp.callrr/r?r@rrr=rArrrrrsrcs,eZdZddZfddZddZZS)GridFunccCst|dkrd|_d|_d|_d|_n|dkr"d|_d|_d |_d |_ntd d |jd |jd|jd}||_dS)Ngridz%Compute the thread index in the grid.z1jit.threadIdx.x + jit.blockIdx.x * jit.blockDim.xznumba.cuda.gridz+threadIdx.{n} + blockIdx.{n} * blockDim.{n}gridsizezCompute the grid size.zjit.blockDim.x * jit.gridDim.xznumba.cuda.gridsizezblockDim.{n} * gridDim.{n}zunsupported functionz zH Computation of the first integer is as follows:: a and for the other two integers the ``y`` and ``z`` attributes are used. Args: ndim (int): The dimension of the grid. Only 1, 2, or 3 is allowed. Returns: int or tuple: If ``ndim`` is 1, an integer is returned, otherwise a tuple. .. note:: This function follows the convention of Numba's :func:`z `. )_desc_eq_link_coderpr)rr6rrrrr#s"  zGridFunc.__init__crrr)rrGrrrrFrzGridFunc.__call__cst|ts td|dkrtjjddtjS|dkrd}n |dkr&d}ntd d fd d |D}t tjg|}|dkrLtd |d|Std|d|S)Nzndim must be an integerrr"nr)r"yr)r"rzzOnly ndim=1,2,3 are supportedrc3s|] }jj|dVqdS)rN)rformat)r!rr_rrr#Xsz&GridFunc.call_const..zSTD::make_pair(rmzSTD::make_tuple() r r,r)r rrruint32rpjoinTuple)rr8rGdims elts_coder2rr_rrcIs zGridFunc.call_const)r/r?r@rrrcrArrrrr!s #rcs8eZdZddZddfdd Zdddd ZZS) WarpShuffleOpcCs>||_d|r |dndd|_||_d|jd}||_dS)N__shfl__syncrz`` function. Please refer to `Warp Shuffle Functions`_ for detailed explanation. .. _Warp Shuffle Functions: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#warp-shuffle-functions rrrrrrcs  zWarpShuffleOp.__init__ )widthcrrr)rrgrval_idrrrrrorzWarpShuffleOp.__call__Nc Csh|j}t||}|j}|jj|jvrtd|d|jdz|j}Wn t y/tdwt j r>t d|dtnd|krKdksPtd td |jd vrYtj}ntj}t||d |}t||}|r{t|trz|jd vrztd n t j rtdntd}t|tjd |}t||}|dt|d|jd|j} | d|jd7} t| |S)Nrrrzmask must be an integerrirjrrfrk)updownrl)rrzwidth needs to be power of 2@r(rrm)rr r5r2r3r~rr)r+ Exceptionrrnr0r1rorprrrrrr rqr r hexrH) rr8rgrrrr~r2val_id_trHrrrr=rsD         " zWarpShuffleOp.callrrrrrras rcs,eZdZfddZddZddZZS)LaneIDcr)zReturns the lane ID of the calling thread, ranging in ``[0, jit.warpsize)``. .. note:: Unlike :obj:`numba.cuda.laneid`, this is a callable function instead of a property. Nrr_rrrrr`zLaneID.__call__cCs"d}tjs |d7}|S|d7}|S)Nz2__device__ __forceinline__ unsigned int LaneId() {z unsigned int ret; asm ("mov.u32 %0, %%laneid;" : "=r"(ret) ); return ret; } z3 return __lane_id(); } )rrn)rpreamblerrr _get_preambles  zLaneID._get_preamblecCs|j|tdtjS)NzLaneId()) generatedadd_coderr rrrbrrrrcs zLaneID.call_const)r/r?r@rrrcrArrrrrs  rzMapping[Any, BuiltinFunc]builtin_functions_dictrrAdd)rrruint64float32float64r)float16Sub)rrrExch)rrrrrMin)rrrrMaxInc)rDecr)uint16AndOrXor)rrrint64rr)rrrrrxor)@ __future__rtypingrcollections.abcrr0rQcupy_backends.cuda.apir cupy.cudar cupyx.jitrrcupyx.jit._internal_typesr r r r r functoolsrrrCrMr\r^rersrrrrranger(minmaxr__annotations__range_ syncthreadssyncwarp shared_memoryrrlaneidrn atomic_add atomic_sub atomic_exch atomic_min atomic_max atomic_inc atomic_dec atomic_cas atomic_and atomic_or atomic_xor _shfl_dtypes shfl_sync shfl_up_syncshfl_down_sync shfl_xor_syncrrrrs            Z  "?@:!