`i6ddlmZddlmZddlmZddlmZddlm Z ddlm Z ddlm ZgdZd d d Zd ZGd dejZGddeZGddeZGddeZGddeZGddeZGddeZGddeZedZedZeZeZeZeZ dS) )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_asyncct|jd|}|durE|jjt|t |jd|ddSdS)Ninclude_FT)getattr generatedcodesappend_header_to_codesetattr)envheaderflags `/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/cupyx/jit/cg.py_check_includersm 3="5V"5"5 6 6D u}} ""?6#:;;; 2&22D99999}c(eZdZdZdZdZdZdZdS) _ThreadGroupz( Base class for all cooperative groups. NctN)NotImplementedErrorselfs r__init__z_ThreadGroup.__init__&s!!rc|jSr# child_typer%s r__str__z_ThreadGroup.__str__)s /##rcft|dt|jdtjS)Nrz.sync()r_Datacodervoidr&rinstances r_syncz_ThreadGroup._sync,s0sD!!! ... 0@AAAr)__name__ __module__ __qualname____doc__r*r'r+r3rrr!r!!sQ22J"""$$$BBBBBrr!ceZdZdZdZedZefdZedZedZ edZ edZ ed Z ed Z ed Zed ZxZS) _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 cd|_dS)Nzcg::grid_groupr)r%s rr'z_GridGroup.__init__:s *rcft|dt|jdtjS)zU is_valid() Returns whether the grid_group can synchronize. rz .is_valid())rr.r/rbool_r1s ris_validz_GridGroup.is_valid=s2 sD!!! 222K4EFFFrc`d|j_t||S)z sync() Synchronize the threads named in the group. .. seealso:: :meth:`numba.cuda.cg.GridGroup.sync` T)r enable_cgsuperr3r&rr2 __class__s rr z_GridGroup.syncGs'#' ww}}S(+++rcft|dt|jdtjSz` thread_rank() Rank of the calling thread within ``[0, num_threads)``. rz.thread_rank()rr.r/ruint64r1s r thread_rankz_GridGroup.thread_rankU2 sD!!! 555{7IJJJrctjdkrtdt|dt |jdt jS)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/rrGr1s r block_rankz_GridGroup.block_rank_U  + - - 5 5HII IsD!!! 444k6HIIIrctjdkrtdt|dt |jdt jSzN num_threads() Total number of threads in the group. rKz(num_threads() is supported on CUDA 11.6+rz.num_threads()rLr1s r num_threadsz_GridGroup.num_threadskU  + - - 5 5IJJ JsD!!! 555{7IJJJrctjdkrtdt|dt |jdt jS)zL num_blocks() Total number of blocks in the group. rKz'num_blocks() is supported on CUDA 11.6+rz .num_blocks()rLr1s r num_blocksz_GridGroup.num_blockswrQrctjdkrtdt|dt |jdt jS)z[ dim_blocks() Dimensions of the launched grid in units of blocks. rKz'dim_blocks() is supported on CUDA 11.6+rz .dim_blocks()rMrNrOrr.r/rdim3r1s r dim_blocksz_GridGroup.dim_blockssU  + - - 5 5HII IsD!!! 444k6FGGGrctjdkrtdt|dt |jdt jS)zc block_index() 3-Dimensional index of the block within the launched grid. rKz(block_index() is supported on CUDA 11.6+rz.block_index()rYr1s r block_indexz_GridGroup.block_indexU  + - - 5 5IJJ JsD!!! 555{7GHHHrcft|dt|jdtjSzG size() Total number of threads in the group. rz.size()rFr1s rsizez_GridGroup.size2 sD!!! ... 0BCCCrcft|dt|jdtjS)zZ group_dim() Dimensions of the launched grid in units of blocks. r .group_dim()rr.r/rrZr1s r group_dimz_GridGroup.group_dim2 sD!!! 333[5EFFFr)r4r5r6r7r'_wraps_class_methodr>r rHrPrTrWr[r]rarf __classcell__rCs@rr:r:1sy+++GGG , , , , ,KKK J J J K K K J J J H H H I I IDDDGGGGGGGrr:ceZdZdZdZefdZedZedZedZ edZ edZ ed Z ed Z xZS) _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 cd|_dS)Nzcg::thread_blockr)r%s rr'z_ThreadBlockGroup.__init__s ,rcHt||S)zM sync() Synchronize the threads named in the group. )rAr3rBs rr z_ThreadBlockGroup.syncsww}}S(+++rcft|dt|jdtjSrErr.r/ruint32r1s rrHz_ThreadBlockGroup.thread_rankrIrcft|dt|jdtjS)zc group_index() 3-Dimensional index of the block within the launched grid. rz.group_index()rer1s r group_indexz_ThreadBlockGroup.group_indexs2 sD!!! 555{7GHHHrcft|dt|jdtjS)zf thread_index() 3-Dimensional index of the thread within the launched block. rz.thread_index()rer1s r thread_indexz_ThreadBlockGroup.thread_indexs2 sD!!! 666 8HIIIrctjdkrtdt|dt |jdt jS)z^ dim_threads() Dimensions of the launched block in units of threads. rKz(dim_threads() is supported on CUDA 11.6+rz.dim_threads()rYr1s r dim_threadsz_ThreadBlockGroup.dim_threadsr^rctjdkrtdt|dt |jdt jSrS)rMrNrOrr.r/rrqr1s rrTz_ThreadBlockGroup.num_threadsrUrcft|dt|jdtjSr`rpr1s rraz_ThreadBlockGroup.sizerbrcft|dt|jdtjS)z\ group_dim() Dimensions of the launched block in units of threads. rrdrer1s rrfz_ThreadBlockGroup.group_dim rgr)r4r5r6r7r'rhr rHrsrurwrTrarfrirjs@rrlrls;---,,,,,KKKIIIJJJ I I I K K KDDDGGGGGGGrrlc*eZdZdZfdZdZxZS) _ThisCgGroupc|dkrd}d}n|dkrd}d}nt||_d|d|d |d |_|dkr|xjd z c_dSdS) Ngridz grid groupr: thread_blockzthread block grouprlz Returns the current z (:class:`~cupyx.jit.cg.z/`). .. seealso:: :class:`cupyx.jit.cg.`z!, :func:`numba.cuda.cg.this_grid`)r$ group_typer7)r&rnametypenames rr'z_ThisCgGroup.__init__s   D#HH > ) )'D*HH% %$9!99;C99,4999     LL? ?LLLL rcHtdSr#rA__call__)r&rCs rrz_ThisCgGroup.__call__(s rctjrtd|jdkrt }n|jdkrt }t d|jd|S)Nz)cooperative group is not supported on HIPr~rz cg::this_z())rMis_hiprOrr:rlr.)r&rcg_types r call_constz_ThisCgGroup.call_const+sj ? LJKK K ?f $ $ llGG _ . .'))G4444g>>>r)r4r5r6r'rrrirjs@rr|r|sY@@@"???????rr|c$eZdZfdZdZxZS)_SynccHtdS)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 Nrr&grouprCs rrz_Sync.__call__7! rct|jtstdt |dt d|jdtjS)Nz'group must be a valid cooperative grouprz cg::sync()) isinstancectyper! ValueErrorrr.r/rr0r&rrs rcallz _Sync.callDsW%+|44 HFGG GsD!!!.... 0@AAArr4r5r6rrrirjs@rrr5sN     BBBBBBBrrc0eZdZddfd ZdddZxZS) _MemcpySyncN) aligned_sizecHtdS)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 Nr) r&rdstdst_idxsrcsrc_idxrarrCs rrz_MemcpySync.__call__Ms!. rc t|dt|dtj||}tj||}||fD]<} t| jt jt jfstd=tj |||}tj |||}tj |t j d|}tj||}|j } |r2t|tstdd|jd| d} td |j d |j d |j d | d t jS) Nrrzdst/src must be of array type. same_kindz,aligned_size must be a compile-time constantzcuda::aligned_size_t<>(rzcg::memcpy_async(z, &(z), &(z), )rr.initrrrCArrayPtr TypeErrorr _indexing_astype_scalarrqr/ _Constantrobjr0) r&rrrrrrrararr size_codes rrz_MemcpySync.callfssD!!!s-...jc""jc"": B BCI 2KODFF B @AAA B gs33 gs33& +$k388z$$$yN  +lI66 D BDDD*1A**&***I44444(44'04445@5EGG Grrrjs@rrrKsn#4GGGGGGGGGrrc$eZdZfdZdZxZS)_WaitcHtdS)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 Nrrs rrz_Wait.__call__rrcht|dtd|jdtjS)Nrz cg::wait(rr-rs rrz _Wait.calls3sD!!!.... 0@AAArrrjs@rrrsN     BBBBBBBrrc$eZdZfdZdZxZS) _WaitPriorcHtdS)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 Nrrs rrz_WaitPrior.__call__s! rct|dt|tstdt d|jd|jdtjS)Nrz$step must be a compile-time constantzcg::wait_priorrs))))))!!!!!!AAAAAA;;;;;;333333OOOOOO 9 9 9 2F ::: B B B B B;' B B B ~G~G~G~G~G~G~G~GB`G`G`G`G`G `G`G`GF?????<???@BBBBBLBBB,7G7G7G7G7G,7G7G7GtBBBBBLBBB(''''''''0 L  L00 uww uww Z\\ {}} r