o װ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)envheaderflagr@/home/ubuntu/.local/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 Gs z_GridGroup.synccCr'z` thread_rank() Rank of the calling thread within ``[0, num_threads)``. rz.thread_rank()rr)r*ruint64r,rrr thread_rankUr8z_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_rankcCrAzN num_threads() Total number of threads in the group. rBz(num_threads() is supported on CUDA 11.6+rz.num_threads()rCr,rrr num_threadskrHz_GridGroup.num_threadscCrA)zL num_blocks() Total number of blocks in the group. rBz'num_blocks() is supported on CUDA 11.6+rz .num_blocks()rCr,rrr num_blockswrHz_GridGroup.num_blockscCrA)z[ dim_blocks() Dimensions of the launched grid in units of blocks. rBz'dim_blocks() is supported on CUDA 11.6+rz .dim_blocks()rDrErFrr)r*rdim3r,rrr dim_blocksrHz_GridGroup.dim_blockscCrA)zc block_index() 3-Dimensional index of the block within the launched grid. rBz(block_index() is supported on CUDA 11.6+rz.block_index()rLr,rrr block_indexrHz_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*rrMr,rrr group_dimrRz_GridGroup.group_dim)r/r0r1r2r#_wraps_class_methodr7r r@rGrJrKrNrOrQrU __classcell__rrr;rr31s.          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. )r:r.r,r;rrr sz_ThreadBlockGroup.synccCr'r=rr)r*ruint32r,rrrr@r8z_ThreadBlockGroup.thread_rankcCr')zc group_index() 3-Dimensional index of the block within the launched grid. rz.group_index()rTr,rrr group_indexr8z_ThreadBlockGroup.group_indexcCr')zf thread_index() 3-Dimensional index of the thread within the launched block. rz.thread_index()rTr,rrr thread_indexr8z_ThreadBlockGroup.thread_indexcCrA)z^ dim_threads() Dimensions of the launched block in units of threads. rBz(dim_threads() is supported on CUDA 11.6+rz.dim_threads()rLr,rrr dim_threadsrHz_ThreadBlockGroup.dim_threadscCrArI)rDrErFrr)r*rrZr,rrrrJrHz_ThreadBlockGroup.num_threadscCr'rPrYr,rrrrQrRz_ThreadBlockGroup.sizecCr')z\ group_dim() Dimensions of the launched block in units of threads. rrSrTr,rrrrU rRz_ThreadBlockGroup.group_dim)r/r0r1r2r#rVr r@r[r\r]rJrQrUrWrrr;rrXs&        rXcs,eZdZddZfddZddZZS) _ThisCgGroupcCsf|dkr d}d}n |dkrd}d}nt||_d|d|d |d |_|dkr1|jd 7_dSdS) Ngridz grid groupr3 thread_blockzthread block grouprXz 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"rbnametypenamerrrr#s$ z_ThisCgGroup.__init__cstdSrr:__call__r!r;rrrf(sz_ThisCgGroup.__call__cCsDtjrtd|jdkrt}n|jdkrt}td|jd|S)Nz)cooperative group is not supported on HIPr_r`z cg::this_z())rDis_hiprFrbr3rXr))r"rcg_typerrr call_const+s  z_ThisCgGroup.call_const)r/r0r1r#rfrirWrrr;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 Nrer"groupr;rrrf7 z_Sync.__call__cCs4t|jts tdt|dtd|jdtjS)Nz'group must be a valid cooperative grouprz cg::sync()) isinstancectyper ValueErrorrr)r*rr+r"rrnrrrcallDs  z _Sync.callr/r0r1rfrurWrrr;rrk5  rkcs0eZdZddfdd ZddddZZS) _MemcpySyncN) aligned_sizecrl)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 Nre)r"rndstdst_idxsrcsrc_idxrQryr;rrrfMsz_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<>(rpzcg::memcpy_async(z, &(z), &(z), )rr)initrqrrrCArrayPtr TypeErrorr _indexing_astype_scalarrZr* _Constantrsobjr+) r"rrnrzr{r|r}rQryarr size_coderrrrufs@         z_MemcpySync.callrvrrr;rrxKs rxcrj)_Waitcrl)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 Nrermr;rrrfroz_Wait.__call__cCs t|dtd|jdtjS)Nrz cg::wait(rpr(rtrrrrus z _Wait.callrvrrr;rrrwrcrj) _WaitPriorcrl)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 Nrermr;rrrfs z_WaitPrior.__call__cCs:t|dt|tstdtd|jd|jdtjS)Nrz$step must be a compile-time constantzcg::wait_priors6      c :