o 0 i6@sddlmZddlmZddlmZddlmZddlm Z ddlm Z ddlm ZgdZd d d Zd d ZGdddejZGdddeZGdddeZGdddeZGdddeZGdddeZGdddeZGdddeZedZedZeZeZeZeZ d S)!)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)envheaderflagrX/home/app/PaddleOCR-VL-test/.venv_paddleocr/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*rZbool_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)rZ enable_cgsuperr.r, __class__rrr Gs z_GridGroup.synccCr'z` thread_rank() Rank of the calling thread within ``[0, num_threads)``. rz.thread_rank()rr)r*ruint64r,rrr thread_rankUr7z_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_rank_  z_GridGroup.block_rankcCr?zN num_threads() Total number of threads in the group. r@z(num_threads() is supported on CUDA 11.6+rz.num_threads()rAr,rrr num_threadskrFz_GridGroup.num_threadscCr?)zL num_blocks() Total number of blocks in the group. r@z'num_blocks() is supported on CUDA 11.6+rz .num_blocks()rAr,rrr num_blockswrFz_GridGroup.num_blockscCr?)z[ dim_blocks() Dimensions of the launched grid in units of blocks. r@z'dim_blocks() is supported on CUDA 11.6+rz .dim_blocks()rBrCrDrr)r*rdim3r,rrr dim_blocksrFz_GridGroup.dim_blockscCr?)zc block_index() 3-Dimensional index of the block within the launched grid. r@z(block_index() is supported on CUDA 11.6+rz.block_index()rJr,rrr block_indexrFz_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*rrKr,rrr group_dimrPz_GridGroup.group_dim)r/r0r1r2r#_wraps_class_methodr6r r>rErHrIrLrMrOrS __classcell__rrr9rr31s.          r3cseZdZdZddZefddZeddZedd 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 cCr4)Nzcg::thread_blockr$r!rrrr#r5z_ThreadBlockGroup.__init__cst||S)zM sync() Synchronize the threads named in the group. )r8r.r,r9rrr sz_ThreadBlockGroup.synccCr'r;rr)r*ruint32r,rrrr>r7z_ThreadBlockGroup.thread_rankcCr')zc group_index() 3-Dimensional index of the block within the launched grid. rz.group_index()rRr,rrr group_indexr7z_ThreadBlockGroup.group_indexcCr')zf thread_index() 3-Dimensional index of the thread within the launched block. rz.thread_index()rRr,rrr thread_indexr7z_ThreadBlockGroup.thread_indexcCr?)z^ dim_threads() Dimensions of the launched block in units of threads. r@z(dim_threads() is supported on CUDA 11.6+rz.dim_threads()rJr,rrr dim_threadsrFz_ThreadBlockGroup.dim_threadscCr?rG)rBrCrDrr)r*rrXr,rrrrHrFz_ThreadBlockGroup.num_threadscCr'rNrWr,rrrrOrPz_ThreadBlockGroup.sizecCr')z\ group_dim() Dimensions of the launched block in units of threads. rrQrRr,rrrrS rPz_ThreadBlockGroup.group_dim)r/r0r1r2r#rTr r>rYrZr[rHrOrSrUrrr9rrVs&        rVcs,eZdZddZfddZddZZS) _ThisCgGroupcCsf|dkr d}d}n |dkrd}d}nt||_d|d|d |d |_|dkr1|jd 7_dSdS) Ngridz grid groupr3 thread_blockzthread block grouprVz Returns the current z (:class:`~cupyx.jit.cg.z/`). .. seealso:: :class:`cupyx.jit.cg.`z!, :func:`numba.cuda.cg.this_grid`)r group_typer2)r"r`nametypenamerrrr#s$ z_ThisCgGroup.__init__cstdSrr8__call__r!r9rrrd(sz_ThisCgGroup.__call__cCsDtjrtd|jdkrt}n|jdkrt}td|jd|S)Nz)cooperative group is not supported on HIPr]r^z cg::this_z())rBZis_hiprDr`r3rVr))r"rZcg_typerrr call_const+s  z_ThisCgGroup.call_const)r/r0r1r#rdrerUrrr9rr\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 Nrcr"groupr9rrrd7 z_Sync.__call__cCs4t|jts tdt|dtd|jdtjS)Nz'group must be a valid cooperative grouprz cg::sync()) isinstancectyper ValueErrorrr)r*rr+r"rrjrrrcallDs  z _Sync.callr/r0r1rdrqrUrrr9rrg5  rgcs0eZdZddfdd ZddddZZS) _MemcpySyncN) aligned_sizecrh)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 Nrc)r"rjdstdst_idxsrcsrc_idxrOrur9rrrdMsz_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.Z same_kindz,aligned_size must be a compile-time constantzcuda::aligned_size_t<>(rlzcg::memcpy_async(z, &(z), &(z), )rr)initrmrnrZCArrayZPtr TypeErrorrZ _indexingZ_astype_scalarrXr* _Constantroobjr+) r"rrjrvrwrxryrOruZarrZ size_coderrrrqfs@         z_MemcpySync.callrrrrr9rrtKs rtcrf)_Waitcrh)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 Nrcrir9rrrdrkz_Wait.__call__cCs t|dtd|jdtjS)Nrz cg::wait(rlr(rprrrrqs z _Wait.callrrrrr9rrrsrcrf) _WaitPriorcrh)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 Nrcrir9rrrds z_WaitPrior.__call__cCs:t|dt|tstdtd|jd|jdtjS)Nrz$step must be a compile-time constantzcg::wait_priors6      c :