o X۷i 7@s(ddlmZddlmZddlmZddlmZddlm Z ddlm Z ddlm ZddlmZgd Zd d d Zd dZGdddejZGdddeZGdddeZGddde ZGddde ZGddde ZGddde ZGddde ZedZed ZeZeZ eZ!eZ"d!S)") annotations)runtime)_compile) _cuda_types) BuiltinFunc)Constant)Data)wraps_class_method) this_gridthis_thread_blocksyncwait wait_prior memcpy_asyncz[#include #include namespace cg = cooperative_groups; z,#include )cgcg_memcpy_asynccCsHt|jd|}|dur"|jjt|t|jd|ddSdS)Ninclude_FT)getattr generatedcodesappend_header_to_codesetattr)envheaderflagrB/home/ubuntu/vllm_env/lib/python3.10/site-packages/cupyx/jit/cg.py_check_includes rc@s,eZdZdZdZddZddZddZdS) _ThreadGroupz( Base class for all cooperative groups. NcCstN)NotImplementedErrorselfrrr__init__(sz_ThreadGroup.__init__cCs|jSr  child_typer"rrr__str__+sz_ThreadGroup.__str__cCt|dt|jdtjS)Nrz.sync()r_Datacodervoidr#rinstancerrr_sync.s z_ThreadGroup._sync)__name__ __module__ __qualname____doc__r&r$r'r/rrrrr#s  rcseZdZdZddZeddZefddZedd Zed d Z ed d Z eddZ eddZ eddZ eddZeddZZS) _GridGroupaA handle to the current grid group. Must be created via :func:`this_grid`. .. seealso:: `CUDA Grid Group API`_, :class:`numba.cuda.cg.GridGroup` .. _CUDA Grid Group API: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#grid-group-cg cC d|_dS)Nzcg::grid_groupr%r"rrrr$< z_GridGroup.__init__cCr()zU is_valid() Returns whether the grid_group can synchronize. rz .is_valid())rr*r+rbool_r-rrris_valid? z_GridGroup.is_validcsd|j_t||S)z sync() Synchronize the threads named in the group. .. seealso:: :meth:`numba.cuda.cg.GridGroup.sync` T)r enable_cgsuperr/r- __class__rrr Is z_GridGroup.synccCr(z` thread_rank() Rank of the calling thread within ``[0, num_threads)``. rz.thread_rank()rr*r+ruint64r-rrr thread_rankWr9z_GridGroup.thread_rankcC2tdkr tdt|dt|jdtjS)z] block_rank() Rank of the calling block within ``[0, num_blocks)``. 4+z'block_rank() is supported on CUDA 11.6+rz .block_rank()_runtime_getLocalRuntimeVersion RuntimeErrorrr*r+rr@r-rrr block_ranka  z_GridGroup.block_rankcCrBzN num_threads() Total number of threads in the group. rCz(num_threads() is supported on CUDA 11.6+rz.num_threads()rDr-rrr num_threadsmrIz_GridGroup.num_threadscCrB)zL num_blocks() Total number of blocks in the group. rCz'num_blocks() is supported on CUDA 11.6+rz .num_blocks()rDr-rrr num_blocksyrIz_GridGroup.num_blockscCrB)z[ dim_blocks() Dimensions of the launched grid in units of blocks. rCz'dim_blocks() is supported on CUDA 11.6+rz .dim_blocks()rErFrGrr*r+rdim3r-rrr dim_blocksrIz_GridGroup.dim_blockscCrB)zc block_index() 3-Dimensional index of the block within the launched grid. rCz(block_index() is supported on CUDA 11.6+rz.block_index()rMr-rrr block_indexrIz_GridGroup.block_indexcCr(zG size() Total number of threads in the group. rz.size()r?r-rrrsize z_GridGroup.sizecCr()zZ group_dim() Dimensions of the launched grid in units of blocks. r .group_dim()rr*r+rrNr-rrr group_dimrSz_GridGroup.group_dim)r0r1r2r3r$_wraps_class_methodr8r rArHrKrLrOrPrRrV __classcell__rrr<rr43s.          r4cseZdZdZddZefddZeddZedd Zed d Z ed d Z eddZ eddZ eddZ ZS)_ThreadBlockGroupaA handle to the current thread block group. Must be created via :func:`this_thread_block`. .. seealso:: `CUDA Thread Block Group API`_ .. _CUDA Thread Block Group API: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#thread-block-group-cg cCr5)Nzcg::thread_blockr%r"rrrr$r6z_ThreadBlockGroup.__init__cst||S)zM sync() Synchronize the threads named in the group. )r;r/r-r<rrr sz_ThreadBlockGroup.synccCr(r>rr*r+ruint32r-rrrrAr9z_ThreadBlockGroup.thread_rankcCr()zc group_index() 3-Dimensional index of the block within the launched grid. rz.group_index()rUr-rrr group_indexr9z_ThreadBlockGroup.group_indexcCr()zf thread_index() 3-Dimensional index of the thread within the launched block. rz.thread_index()rUr-rrr thread_indexr9z_ThreadBlockGroup.thread_indexcCrB)z^ dim_threads() Dimensions of the launched block in units of threads. rCz(dim_threads() is supported on CUDA 11.6+rz.dim_threads()rMr-rrr dim_threadsrIz_ThreadBlockGroup.dim_threadscCrBrJ)rErFrGrr*r+rr[r-rrrrKrIz_ThreadBlockGroup.num_threadscCr(rQrZr-rrrrRrSz_ThreadBlockGroup.sizecCr()z\ group_dim() Dimensions of the launched block in units of threads. rrTrUr-rrrrV rSz_ThreadBlockGroup.group_dim)r0r1r2r3r$rWr rAr\r]r^rKrRrVrXrrr<rrYs&        rYcs,eZdZddZfddZddZZS) _ThisCgGroupcCsf|dkr d}d}n |dkrd}d}nt||_d|d|d |d |_|dkr1|jd 7_dSdS) Ngridz grid groupr4 thread_blockzthread block grouprYz Returns the current z (:class:`~cupyx.jit.cg.z/`). .. seealso:: :class:`cupyx.jit.cg.`z!, :func:`numba.cuda.cg.this_grid`)r! group_typer3)r#rcnametypenamerrrr$s$ z_ThisCgGroup.__init__cstdSr r;__call__r"r<rrrg*sz_ThisCgGroup.__call__cCsDtjrtd|jdkrt}n|jdkrt}td|jd|S)Nz)cooperative group is not supported on HIPr`raz cg::this_z())rEis_hiprGrcr4rYr*)r#rcg_typerrr call_const-s  z_ThisCgGroup.call_const)r0r1r2r$rgrjrXrrr<rr_s r_c$eZdZfddZddZZS)_SyncctdS)aCalls ``cg::sync()``. Args: group: a valid cooperative group .. seealso:: `cg::sync`_ .. _cg::sync: https://docs.nvidia.com/cuda/archive/11.6.0/cuda-c-programming-guide/index.html#collectives-cg-sync Nrfr#groupr<rrrg9 z_Sync.__call__cCs4t|jts tdt|dtd|jdtjS)Nz'group must be a valid cooperative grouprz cg::sync()) isinstancectyper ValueErrorrr*r+rr,r#rrorrrcallFs  z _Sync.callr0r1r2rgrvrXrrr<rrl7  rlcs0eZdZddfdd ZddddZZS) _MemcpySyncN) aligned_sizecrm)aCalls ``cg::memcpy_sync()``. Args: group: a valid cooperative group dst: the destination array that can be viewed as a 1D C-contiguous array dst_idx: the start index of the destination array element src: the source array that can be viewed as a 1D C-contiguous array src_idx: the start index of the source array element size (int): the number of bytes to be copied from ``src[src_index]`` to ``dst[dst_idx]`` aligned_size (int): Use ``cuda::aligned_size_t`` to guarantee the compiler that ``src``/``dst`` are at least N-bytes aligned. The behavior is undefined if the guarantee is not held. .. seealso:: `cg::memcpy_sync`_ .. _cg::memcpy_sync: https://docs.nvidia.com/cuda/archive/11.6.0/cuda-c-programming-guide/index.html#collectives-cg-memcpy-async Nrf)r#rodstdst_idxsrcsrc_idxrRrzr<rrrgOsz_MemcpySync.__call__c Cst|dt|dt||}t||}||fD]} t| jtjtjfs*tdqt |||}t |||}t |tj d|}t||}|j } |rat|tsWtdd|jd| d} td |j d |j d |j d | d tjS) Nrrzdst/src must be of array type. same_kindz,aligned_size must be a compile-time constantzcuda::aligned_size_t<>(rqzcg::memcpy_async(z, &(z), &(z), )rr*initrrrsrCArrayPtr TypeErrorr _indexing_astype_scalarr[r+ _Constantrtobjr,) r#rror{r|r}r~rRrzarr size_coderrrrvhs@         z_MemcpySync.callrwrrr<rryMs rycrk)_Waitcrm)aCalls ``cg::wait()``. Args: group: a valid cooperative group .. seealso: `cg::wait`_ .. _cg::wait: https://docs.nvidia.com/cuda/archive/11.6.0/cuda-c-programming-guide/index.html#collectives-cg-wait Nrfrnr<rrrgrpz_Wait.__call__cCs t|dtd|jdtjS)Nrz cg::wait(rqr)rurrrrvs z _Wait.callrwrrr<rrrxrcrk) _WaitPriorcrm)aXCalls ``cg::wait_prior()``. Args: group: a valid cooperative group step (int): wait for the first ``N`` steps to finish .. seealso: `cg::wait_prior`_ .. _cg::wait_prior: https://docs.nvidia.com/cuda/archive/11.6.0/cuda-c-programming-guide/index.html#collectives-cg-wait Nrfrnr<rrrgs z_WaitPrior.__call__cCs:t|dt|tstdtd|jd|jdtjS)Nrz$step must be a compile-time constantzcg::wait_priors8       c :