o iB@spUddlmZmZddlZddlZddlmZddlmZddl m Z ddl m Z ddl m Z ddl mZdd l mZdd l mZdd l mZdd lmZGd dde ZGddde ZGddde ZGddde ZGddde ZGddde ZGddde ZGddde ZGddde ZGdd d e ZGd!d"d"e Zeee ee!ee"eiZ#eee fe$d#<eZ%eZ&eZ'eZ(ed$Z)ed%Z*eZ+ed&d'ej,rd(nd)Z-ed*d+Z.ed,d-Z/ed.d/Z0ed0d/Z1ed1d2Z2ed3d2Z3ed4d/ej,rd(nd5Z4ed6d/Z5ed7d/Z6ed8d/Z7d9ej,rd(nd:Z8ed;e8Z9ede8Z@sz!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 allobjrintbooltype__name__warningswarnctypedtypekindr initmoderScalarr ) renvrrstartstopstepstep_is_positiver1rrrcall2sf               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(rr1rCArrayndimr coder6)rr7rkwdsargrrrr<os z LenFunc.callNr.r>r?r<rrrrrBm rBc@rA)MinFuncc<t|dkrtdt||rtdtfdd|S)Nrz(min() expects at least 2 arguments, got rCcttj||fdSr)r _call_ufunccupyminimumabr7rrzMinFunc.call..r'r(r rr7rrHrrUrr<   z MinFunc.callNrJrrrrrL}rKrLc@rA)MaxFunccrM)Nrz(max() expects at least 2 arguments, got rCcrNr)r rOrPmaximumrRrUrrrVrWzMaxFunc.call..rXrYrrUrr<rZz MaxFunc.callNrJrrrrr[rKr[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 rvoidrr7rrr call_consts zSyncThreads.call_constr.r>r?rrbr@rrrrr]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)rrfrrrr 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 HIPrremask is out of range same_kindz __syncwarp()z __syncwarp())ris_hipr/r0RuntimeWarningrr r* ValueErrorr _astype_scalarrint32r r4rGr`)rr7rfrGrrrr<s"    z SyncWarp.callr=rrrrrds rdcs(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)rr2size alignmentrrrrrgzSharedMemory.__call__cCsN|jdd}t|}t|t|||}||j|<||j|<t|t|S)N_smem)prefix) get_fresh_variable_namerto_ctyper r SharedMemdeclslocalsPtr)rr7r2rsrtnamer1varrrrrbs    zSharedMemory.call_constrrcrrrrrrs rrcs0eZdZddZdfdd 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.rkCASHFz5uint16 atomic operation is not supported before sm_70z(&, rl)rrr r4rr1rrEr|r(r _indexingr2r}rrpcharr+rget_compute_capability RuntimeErrorrG) rr7rrrvalue2r}rtargetr1rGrrrr<s0     $  z AtomicOp.callrr.r>r?rrr<r@rrrrrsrcs,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_coderor)rr5rrrrr s"  zGridFunc.__init__crrr)rrFrrrrCrzGridFunc.__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"Usz&GridFunc.call_const..zSTD::make_pair(rlzSTD::make_tuple() rr+r(r rrruint32rojoinTuple)rr7rFdims elts_coder1rr^rrbFs zGridFunc.call_const)r.r>r?rrrbr@rrrrrs #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 rrrrrr`s  zWarpShuffleOp.__init__ )widthcrrr)rrfr~val_idrrrrrlrzWarpShuffleOp.__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 integerrhrirrerj)updownrk)rrzwidth needs to be power of 2@r(rrl)rr r4r1r2r}rr(r* Exceptionrrmr/r0rnrorrrrqr rprr hexrG) rr7rfr~rrr}r1val_id_trGrrrr<osD         " zWarpShuffleOp.callrrrrrr^s 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(); } )rrm)rpreamblerrr _get_preambles  zLaneID._get_preamblecCs|j|tdtjS)NzLaneId()) generatedadd_coderr rrrarrrrbs zLaneID.call_const)r.r>r?rrrbr@rrrrrs  rbuiltin_functions_dictrrAdd)rqruint64float32float64r)float16Sub)rqrExch)rqrrrMin)rqrrMaxInc)rDecr)uint16AndOrXor)rqrint64rr)rrrrrxor)=typingrrr/rPcupy_backends.cuda.apir cupy.cudar cupyx.jitrrcupyx.jit._internal_typesrr r r r functoolsr rrBrLr[r]rdrrrrrrranger'minmaxr__annotations__range_ syncthreadssyncwarp shared_memoryrrlaneidrm 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  "?@:!