Ë
    üÇ:j“y  ã                   ó  — d Z ddlZddlZddlZddlZddlmZ ddlZddl	Z	ddl	m
Z
 ddlmZmZ ddlmZ ddlmZmZ ddlmZ dd	lmZ dd
lmZ ddlmZ ddlmZ ddlmZ 	   eed«      d«      Z d„ Z"d„ Z#d„ Z$ G d„ de
jJ                  «      Z& G d„ de&«      Z'e d„ «       Z( G d„ de&«      Z) G d„ de*«      Z+ G d„ de&ejX                  «      Z- G d„ de&ejX                  «      Z.d(d „Z/d(d!„Z0d"„ Z1d#„ Z2d$Z3d%„ Z4d)d&„Z5d'„ Z6y# e!$ r d„ Z Y Œ˜w xY w)*zÑ
A CUDA ND Array is recognized by checking the __cuda_memory__ attribute
on the object.  If it exists and evaluate to True, it must define shape,
strides, dtype and size attributes similar to a NumPy ndarray.
é    N)Úc_void_p)Ú_devicearray)ÚdevicesÚ
dummyarray)Údriver)ÚtypesÚconfig)Úto_fixed_tuple)Únumpy_version)Únumpy_support)Úprepare_shape_strides_dtype)ÚNumbaPerformanceWarning)ÚwarnÚ	lru_cachec                 ó   — | S ©N© )Úfuncs    ús/home/mcse/projects/srt_converter/srt-converter-venv/lib/python3.12/site-packages/numba/cuda/cudadrv/devicearray.pyr   r      s   € Øˆó    c                 ó   — t        | dd«      S )z$Check if an object is a CUDA ndarrayÚ__cuda_ndarray__F)Úgetattr©Úobjs    r   Úis_cuda_ndarrayr   #   s   € ä�3Ð*¨EÓ2Ð2r   c                 ó¢   ‡ — t        ‰ «       ˆ fd„} |dt        «        |dt        «        |dt        j                  «        |dt        «       y)z,Verify the CUDA ndarray interface for an objc                 ó€   •— t        ‰| «      st        | «      ‚t        t        ‰| «      |«      st        | ›d|›�«      ‚y )Nz must be of type )ÚhasattrÚAttributeErrorÚ
isinstancer   )ÚattrÚtypr   s     €r   Úrequires_attrz4verify_cuda_ndarray_interface.<locals>.requires_attr,   s>   ø€ Ü�s˜DÔ!Ü  Ó&Ð&Üœ' # tÓ,¨cÔ2Ü ºDÁ#Ð!FÓGÐGð 3r   ÚshapeÚstridesÚdtypeÚsizeN)Úrequire_cuda_ndarrayÚtupleÚnpr'   Úint)r   r$   s   ` r   Úverify_cuda_ndarray_interfacer-   (   s?   ø€ ä˜ÔôHñ �'œ5Ô!Ù�)œUÔ#Ù�'œ2Ÿ8™8Ô$Ù�&œ#Õr   c                 ó0   — t        | «      st        d«      ‚y)z9Raises ValueError is is_cuda_ndarray(obj) evaluates Falsezrequire an cuda ndarray objectN)r   Ú
ValueErrorr   s    r   r)   r)   8   s   € ä˜3ÔÜÐ9Ó:Ð:ð  r   c                   óô   — e Zd ZdZdZdZdd„Zed„ «       Zdd„Z	ed„ «       Z
dd„Zd	„ Zed
„ «       Zed„ «       Zej                   dd„«       Zej                   dd„«       Zdd„Zd„ Zd„ Zdd„Zd„ Zed„ «       Zy)ÚDeviceNDArrayBasez$A on GPU NDArray representation
    TNc                 ó:  — t        |t        «      r|f}t        |t        «      r|f}t        j                  |«      }t	        |«      | _        t	        |«      | j
                  k7  rt        d«      ‚t        j                  j                  d|||j                  «      | _        t        |«      | _        t        |«      | _        || _        t        t        j                   t"        j$                  | j                  d«      «      | _        | j&                  dkD  r�|€rt)        j*                  | j                  | j                  | j                  j                  «      | _        t/        j0                  «       j3                  | j,                  «      }n‡t)        j4                  |«      | _        nlt(        j6                  r t(        j8                  j;                  d«      }nt=        d«      }t)        j>                  t/        j0                  «       |d¬«      }d| _        || _         || _!        y)a5  
        Args
        ----

        shape
            array shape.
        strides
            array strides.
        dtype
            data type as np.dtype coercible object.
        stream
            cuda stream.
        gpu_data
            user provided device memory for the ndarray data buffer
        zstrides not match ndimr   é   N)ÚcontextÚpointerr(   )"r!   r,   r+   r'   ÚlenÚndimr/   r   ÚArrayÚ	from_descÚitemsizeÚ_dummyr*   r%   r&   Ú	functoolsÚreduceÚoperatorÚmulr(   Ú_driverÚmemory_size_from_infoÚ
alloc_sizer   Úget_contextÚmemallocÚdevice_memory_sizeÚUSE_NV_BINDINGÚbindingÚCUdeviceptrr   ÚMemoryPointerÚgpu_dataÚstream)Úselfr%   r&   r'   rK   rJ   Únulls          r   Ú__init__zDeviceNDArrayBase.__init__D   sŠ  € ô  �eœSÔ!Ø�HˆEÜ�gœsÔ#Ø�jˆGÜ—‘˜“ˆÜ˜“JˆŒ	Üˆw‹<˜4Ÿ9™9Ò$ÜÐ5Ó6Ð6Ü ×&Ñ&×0Ñ0°°E¸7Ø16·±óAˆŒä˜5“\ˆŒ
Ü˜W“~ˆŒØˆŒ
Üœ	×(Ñ(¬¯©°t·z±zÀ1ÓEÓFˆŒ	à�9‰9�qŠ=ØÐÜ")×"?Ñ"?Ø—J‘J §¡¨d¯j©j×.AÑ.Aó#C�”ä"×.Ñ.Ó0×9Ñ9¸$¿/¹/ÓJ‘ä")×"<Ñ"<¸XÓ"F�•ô ×%Ò%Ü—‘×2Ñ2°1Ó5‘ä “{�Ü×,Ñ,´W×5HÑ5HÓ5JØ59ÀôCˆHàˆDŒOà ˆŒØˆ�r   c                 ó¼  — t         j                  r%| j                  �t        | j                  «      }n2d}n/| j                  j                  �| j                  j                  }nd}t        | j                  «      t        | «      rd nt        | j                  «      |df| j                  j                  | j                  dk7  rt        | j                  «      ddœS d ddœS )Nr   Fé   )r%   r&   ÚdataÚtypestrrK   Úversion)r@   rF   Údevice_ctypes_pointerr,   Úvaluer*   r%   Úis_contiguousr&   r'   ÚstrrK   )rL   Úptrs     r   Ú__cuda_array_interface__z*DeviceNDArrayBase.__cuda_array_interface__w   sÁ   € ä×!Ò!Ø×)Ñ)Ð5Ü˜$×4Ñ4Ó5‘à‘à×)Ñ)×/Ñ/Ð;Ø×0Ñ0×6Ñ6‘à�ô ˜4Ÿ:™:Ó&Ü,¨TÔ2‘t¼¸d¿l¹lÓ8KØ˜%�LØ—z‘z—~‘~Ø*.¯+©+¸Ò*:”c˜$Ÿ+™+Ó&Øñ
ð 	
ð
 AEØñ
ð 	
r   c                 ó>   — t        j                   | «      }||_        |S )zBind a CUDA stream to this object so that all subsequent operation
        on this array defaults to the given stream.
        )ÚcopyrK   )rL   rK   Úclones      r   ÚbindzDeviceNDArrayBase.bind�   s   € ô —	‘	˜$“ˆØˆŒØˆr   c                 ó"   — | j                  «       S r   ©Ú	transpose©rL   s    r   ÚTzDeviceNDArrayBase.T•   s   € à�~‰~ÓÐr   c                 ó&  — |r,t        |«      t        t        | j                  «      «      k(  r| S | j                  dk7  rd}t        |«      ‚|�8t	        |«      t	        t        | j                  «      «      k7  rt        d|›�«      ‚ddlm}  || «      S )Né   z2transposing a non-2D DeviceNDArray isn't supportedzinvalid axes list r   r_   )r*   Úranger7   ÚNotImplementedErrorÚsetr/   Únumba.cuda.kernels.transposer`   )rL   ÚaxesÚmsgr`   s       r   r`   zDeviceNDArrayBase.transpose™   sx   € Ù”E˜$“K¤5¬¨t¯y©yÓ)9Ó#:Ò:ØˆKØ�Y‰Y˜!Š^ØFˆCÜ% cÓ*Ð*ØÐ¤# d£)¬s´5¸¿¹Ó3CÓ/DÒ"DÝ±tÐ=Ó>Ð>å>Ù˜T“?Ð"r   c                 ó"   — |s| j                   S |S r   ©rK   )rL   rK   s     r   Ú_default_streamz!DeviceNDArrayBase._default_stream¥   s   € Ù"(ˆt�{‰{Ð4¨fÐ4r   c                 óò   — d| j                   v }| j                  d   r|sd}n| j                  d   r|sd}nd}t        j                  | j                  «      }t        j                  || j                  |«      S )ún
        Magic attribute expected by Numba to get the numba type that
        represents this object.
        r   ÚC_CONTIGUOUSÚCÚF_CONTIGUOUSÚFÚA)r&   Úflagsr   Ú
from_dtyper'   r   r8   r7   )rL   Ú	broadcastÚlayoutr'   s       r   Ú_numba_type_zDeviceNDArrayBase._numba_type_¨   sh   € ð( ˜Ÿ™Ð%ˆ	Ø�:‰:�nÒ%©iØ‰FØ�Z‰Z˜Ò'±	Ø‰FàˆFä×(Ñ(¨¯©Ó4ˆÜ�{‰{˜5 $§)¡)¨VÓ4Ð4r   c                 óº   — | j                   €:t        j                  rt        j                  j	                  d«      S t        d«      S | j                   j                  S )z:Returns the ctypes pointer to the GPU data buffer
        r   )rJ   r@   rF   rG   rH   r   rT   ra   s    r   rT   z'DeviceNDArrayBase.device_ctypes_pointerÇ   sF   € ð �=‰=Ð Ü×%Ò%Ü—‘×2Ñ2°1Ó5Ð5ä “{Ð"à—=‘=×6Ñ6Ð6r   c                 ó  — |j                   dk(  ryt        | «       | j                  |«      }t        | «      t        |«      }}t	        j
                  |«      r;t        |«       t        ||«       t	        j                  | || j                  |¬«       yt        j                  ||j                  d   rdnddt        dk  r|j                  d	    nd¬
«      }t        ||«       t	        j                  | || j                  |¬«       y)zŸCopy `ary` to `self`.

        If `ary` is a CUDA memory, perform a device-to-device transfer.
        Otherwise, perform a a host-to-device transfer.
        r   Nrl   rp   rq   rs   T©rd   r   Ú	WRITEABLE)ÚorderÚsubokr[   )r(   Úsentry_contiguousrm   Ú
array_corer@   Úis_device_memoryÚcheck_array_compatibilityÚdevice_to_devicerB   r+   Úarrayru   r   Úhost_to_device)rL   ÚaryrK   Ú	self_coreÚary_cores        r   Úcopy_to_devicez DeviceNDArrayBase.copy_to_deviceÓ   sÝ   € ð �8‰8�qŠ=àä˜$ÔØ×%Ñ% fÓ-ˆä(¨Ó.´
¸3³�8ˆ	Ü×#Ñ# CÔ(Ü˜cÔ"Ü% i°Ô:Ü×$Ñ$ T¨3°·±ÈÖOô —x‘xØØ&Ÿ_™_¨^Ò<‘cÀ#Øä  6Ò)ð #Ÿ.™.¨Ñ5Ñ5Ø/3ô5ˆHô & i°Ô:Ü×"Ñ" 4¨°4·?±?Ø*0ö2r   c                 ó¨  — t        d„ | j                  D «       «      r&d}t        |j                  | j                  «      «      ‚| j                  dk\  sJ d«       ‚| j                  |«      }|€0t        j                  | j                  t        j                  ¬«      }nt        | |«       |}| j                  dk7  r#t        j                  || | j                  |¬«       |€t| j                  dk(  r.t        j                  | j                  | j                  |¬«      }|S t        j                  | j                  | j                  | j                  |¬«      }|S )	a^  Copy ``self`` to ``ary`` or create a new Numpy ndarray
        if ``ary`` is ``None``.

        If a CUDA ``stream`` is given, then the transfer will be made
        asynchronously as part as the given stream.  Otherwise, the transfer is
        synchronous: the function returns after the copy is finished.

        Always returns the host array.

        Example::

            import numpy as np
            from numba import cuda

            arr = np.arange(1000)
            d_arr = cuda.to_device(arr)

            my_kernel[100, 100](d_arr)

            result_array = d_arr.copy_to_host()
        c              3   ó&   K  — | ]	  }|d k  –— Œ y­w©r   Nr   )Ú.0Úss     r   ú	<genexpr>z1DeviceNDArrayBase.copy_to_host.<locals>.<genexpr>
  s   è ø€ Ò+˜ˆq�1�uÑ+ùs   ‚z2D->H copy not implemented for negative strides: {}r   zNegative memory size©r%   r'   rl   )r%   r'   Úbuffer)r%   r'   r&   r’   )Úanyr&   rf   ÚformatrB   rm   r+   ÚemptyÚbyterƒ   r@   Údevice_to_hostr(   Úndarrayr%   r'   )rL   r‡   rK   rj   Úhostarys        r   Úcopy_to_hostzDeviceNDArrayBase.copy_to_hostó   s  € ô. Ñ+˜dŸl™lÔ+Ô+ØFˆCÜ% c§j¡j°·±Ó&>Ó?Ð?Ø�‰ !Ò#Ð;Ð%;Ó;Ð#Ø×%Ñ% fÓ-ˆØˆ;Ü—h‘h T§_¡_¼B¿G¹GÔD‰Gä% d¨CÔ0ØˆGà�?‰?˜aÒÜ×"Ñ" 7¨D°$·/±/Ø*0õ2ð ˆ;Ø�y‰y˜AŠ~ÜŸ*™*¨4¯:©:¸T¿Z¹ZØ,3ô5�ð
 ˆô Ÿ*™*¨4¯:©:¸T¿Z¹ZØ-1¯\©\À'ôK�àˆr   c              #   óT  K  — | j                  |«      }| j                  dk7  rt        d«      ‚| j                  d   | j                  j
                  k7  rt        d«      ‚t        t        j                  t        | j                  «      |z  «      «      }| j                  }| j                  j
                  }t        |«      D ]d  }||z  }t        ||z   | j                  «      }||z
  f}	| j                  j                  ||z  ||z  «      }
t        |	|| j                  ||
¬«      –— Œf y­w)zžSplit the array into equal partition of the `section` size.
        If the array cannot be equally divided, the last section will be
        smaller.
        r3   zonly support 1d arrayr   zonly support unit stride©r'   rK   rJ   N)rm   r7   r/   r&   r'   r:   r,   ÚmathÚceilÚfloatr(   re   ÚminrJ   ÚviewÚDeviceNDArray)rL   ÚsectionrK   Únsectr&   r:   ÚiÚbeginÚendr%   rJ   s              r   ÚsplitzDeviceNDArrayBase.split"  s  è ø€ ð
 ×%Ñ% fÓ-ˆØ�9‰9˜Š>ÜÐ4Ó5Ð5Ø�<‰<˜‰?˜dŸj™j×1Ñ1Ò1ÜÐ7Ó8Ð8Ü”D—I‘Iœe D§I¡IÓ.°Ñ8Ó9Ó:ˆØ—,‘,ˆØ—:‘:×&Ñ&ˆÜ�u“ò 	3ˆAØ˜‘KˆEÜ�e˜g‘o t§y¡yÓ1ˆCØ˜5‘[�NˆEØ—}‘}×)Ñ)¨%°(Ñ*:¸CÀ(¹NÓKˆHÜ  w°d·j±jÈØ)1ô3ó 3ñ	3ùs   ‚D&D(c                 ó   — | j                   S )zEReturns a device memory object that is used as the argument.
        )rJ   ra   s    r   Úas_cuda_argzDeviceNDArrayBase.as_cuda_arg7  s   € ð �}‰}Ðr   c                 óÎ   — t        j                  «       j                  | j                  «      }t	        | j
                  | j                  | j                  ¬«      }t        ||¬«      S )zÌ
        Returns a *IpcArrayHandle* object that is safe to serialize and transfer
        to another process to share the local allocation.

        Note: this feature is only available on Linux.
        )r%   r&   r'   )Ú
ipc_handleÚ
array_desc)	r   rC   Úget_ipc_handlerJ   Údictr%   r&   r'   ÚIpcArrayHandle)rL   ÚipchÚdescs      r   r®   z DeviceNDArrayBase.get_ipc_handle<  sH   € ô ×"Ñ"Ó$×3Ñ3°D·M±MÓBˆÜ˜$Ÿ*™*¨d¯l©lÀ$Ç*Á*ÔMˆÜ¨¸$Ô?Ð?r   c                 óÎ   — | j                   j                  |¬«      \  }}t        |j                  |j                  | j
                  | j                  |«      | j                  ¬«      S )a(  
        Remove axes of size one from the array shape.

        Parameters
        ----------
        axis : None or int or tuple of ints, optional
            Subset of dimensions to remove. A `ValueError` is raised if an axis
            with size greater than one is selected. If `None`, all axes with
            size one are removed.
        stream : cuda stream or 0, optional
            Default stream for the returned view of the array.

        Returns
        -------
        DeviceNDArray
            Squeezed view into the array.

        )Úaxis©r%   r&   r'   rK   rJ   )r;   Úsqueezer¢   r%   r&   r'   rm   rJ   )rL   r´   rK   Ú	new_dummyÚ_s        r   r¶   zDeviceNDArrayBase.squeezeG  sX   € ð& —{‘{×*Ñ*°Ð*Ó5‰ˆ	�1ÜØ—/‘/Ø×%Ñ%Ø—*‘*Ø×'Ñ'¨Ó/Ø—]‘]ô
ð 	
r   c                 óî  — t        j                  |«      }t        | j                  «      }t        | j                  «      }| j                  j
                  |j
                  k7  rp| j                  «       st        d«      ‚t        |d   | j                  j
                  z  |j
                  «      \  |d<   }|dk7  rt        d«      ‚|j
                  |d<   t        |||| j                  | j                  ¬«      S )zeReturns a new object by reinterpretting the dtype without making a
        copy of the data.
        zHTo change to a dtype of a different size, the array must be C-contiguouséÿÿÿÿr   zuWhen changing to a larger dtype, its size must be a divisor of the total size in bytes of the last axis of the array.rµ   )r+   r'   Úlistr%   r&   r:   Úis_c_contiguousr/   Údivmodr¢   rK   rJ   )rL   r'   r%   r&   Úrems        r   r¡   zDeviceNDArrayBase.viewc  sà   € ô —‘˜“ˆÜ�T—Z‘ZÓ ˆÜ�t—|‘|Ó$ˆà�:‰:×Ñ %§.¡.Ò0Ø×'Ñ'Ô)Ü ð6óð ô
 $Ø�b‘	˜DŸJ™J×/Ñ/Ñ/Ø—‘ó‰NˆE�"‰I�sð
 �aŠxÜ ð6óð ð  Ÿ.™.ˆG�B‰KäØØØØ—;‘;Ø—]‘]ô
ð 	
r   c                 óH   — | j                   j                  | j                  z  S r   )r'   r:   r(   ra   s    r   ÚnbyteszDeviceNDArrayBase.nbytesˆ  s   € ð
 �z‰z×"Ñ" T§Y¡YÑ.Ð.r   r�   ©r   r   ©Nr   )Ú__name__Ú
__module__Ú__qualname__Ú__doc__Ú__cuda_memory__r   rN   ÚpropertyrY   r]   rb   r`   rm   ry   rT   r   Úrequire_contextrŠ   rš   r¨   rª   r®   r¶   r¡   rÀ   r   r   r   r1   r1   >   sÛ   „ ñà€OØÐó1ðf ñ
ó ð
ó*ð ñ ó ð ó
#ò5ð ñ5ó ð5ð< ñ	7ó ð	7ð ×Ñò2ó ð2ð> ×Ñò,ó ð,ó\3ò*ò
	@ó
ò8#
ðJ ñ/ó ñ/r   r1   c                   óè   ‡ — e Zd ZdZdˆ fd„	Zed„ «       Zed„ «       Zej                  d„ «       Z
ej                  dd„«       Zdd„Zej                  d„ «       Zej                  dd	„«       Zdd
„Zˆ xZS )ÚDeviceRecordz
    An on-GPU record type
    c                 ó<   •— d}d}t         t        | �  |||||«       y ©Nr   )ÚsuperrË   rN   )rL   r'   rK   rJ   r%   r&   Ú	__class__s         €r   rN   zDeviceRecord.__init__”  s'   ø€ ØˆØˆÜŒl˜DÑ*¨5°'¸5À&Ø+3õ	5r   c                 ó@   — t        | j                  j                  «      S ©zÿ
        For `numpy.ndarray` compatibility. Ideally this would return a
        `np.core.multiarray.flagsobj`, but that needs to be constructed
        with an existing `numpy.ndarray` (as the C- and F- contiguous flags
        aren't writeable).
        ©r¯   r;   ru   ra   s    r   ru   zDeviceRecord.flagsš  ó   € ô �D—K‘K×%Ñ%Ó&Ð&r   c                 ó@   — t        j                  | j                  «      S )ro   )r   rv   r'   ra   s    r   ry   zDeviceRecord._numba_type_¤  s   € ô ×'Ñ'¨¯
©
Ó3Ð3r   c                 ó$   — | j                  |«      S r   ©Ú_do_getitem©rL   Úitems     r   Ú__getitem__zDeviceRecord.__getitem__¬  ó   € à×Ñ Ó%Ð%r   c                 ó&   — | j                  ||«      S ©z0Do `__getitem__(item)` with CUDA stream
        rÖ   ©rL   rÙ   rK   s      r   ÚgetitemzDeviceRecord.getitem°  ó   € ð ×Ñ  fÓ-Ð-r   c                 óÔ  — | j                  |«      }| j                  j                  |   \  }}| j                  j	                  |«      }|j
                  dk(  rY|j                  �t        |||¬«      S t        j                  d|¬«      }t        j                  |||j                  |¬«       |d   S t        |j
                  d |j                  d   d«      \  }}}	t        |||	||¬«      S )	Nr   rœ   r3   ©r'   ©ÚdstÚsrcr(   rK   r   rq   ©r%   r&   r'   rJ   rK   )rm   r'   ÚfieldsrJ   r¡   r%   ÚnamesrË   r+   r•   r@   r—   r:   r   Úsubdtyper¢   )
rL   rÙ   rK   r#   ÚoffsetÚnewdatar™   r%   r&   r'   s
             r   r×   zDeviceRecord._do_getitem¶  sè   € Ø×%Ñ% fÓ-ˆØ—j‘j×'Ñ'¨Ñ-‰ˆˆVØ—-‘-×$Ñ$ VÓ,ˆà�9‰9˜Š?Ø�y‰yÐ$Ü#¨#°fØ-4ô6ð 6ô Ÿ(™( 1¨CÔ0�Ü×&Ñ&¨7¸Ø,/¯L©LØ.4õ6ð ˜1‘:Ðô ,¨C¯I©IØ,0Ø,/¯L©L¸©O¸SóBñ "ˆE�7˜Eô ! u°gØ',°wØ(.ô0ð 0r   c                 ó&   — | j                  ||«      S r   ©Ú_do_setitem©rL   ÚkeyrU   s      r   Ú__setitem__zDeviceRecord.__setitem__Î  ó   € à×Ñ  UÓ+Ð+r   c                 ó*   — | j                  |||¬«      S ©z6Do `__setitem__(key, value)` with CUDA stream
        rl   rí   ©rL   rð   rU   rK   s       r   ÚsetitemzDeviceRecord.setitemÒ  ó   € ð ×Ñ  U°6ÐÓ:Ð:r   c                 óÞ  — | j                  |«      }| }|r$t        j                  «       }|j                  «       }| j                  j
                  |   \  }}| j                  j                  |«      } t        | «      |||¬«      }	t        |	j                  j                  |«      |¬«      \  }
}t        j                  |	|
|
j                  j                  |«       |r|j                  «        y y )Nrœ   rl   )rm   r   rC   Úget_default_streamr'   rç   rJ   r¡   ÚtypeÚauto_devicer@   r„   r:   Úsynchronize)rL   rð   rU   rK   ÚsynchronousÚctxr#   rê   rë   ÚlhsÚrhsr¸   s               r   rî   zDeviceRecord._do_setitemØ  sÏ   € à×%Ñ% fÓ-ˆð
 !�jˆÙÜ×%Ñ%Ó'ˆCØ×+Ñ+Ó-ˆFð —j‘j×'Ñ'¨Ñ,‰ˆˆVØ—-‘-×$Ñ$ VÓ,ˆàŒd�4‹j˜s¨6¸GÔDˆô ˜SŸY™YŸ^™^¨EÓ2¸6ÔB‰ˆˆQô 	× Ñ   c¨3¯9©9×+=Ñ+=¸vÔFáØ×ÑÕ ð r   r�   rÁ   )rÃ   rÄ   rÅ   rÆ   rN   rÈ   ru   ry   r   rÉ   rÚ   rß   r×   rñ   rö   rî   Ú__classcell__)rÏ   s   @r   rË   rË   �  s§   ø„ ñõ5ð ñ'ó ð'ð ñ4ó ð4ð ×Ññ&ó ð&ð ×Ñò.ó ð.ó
0ð0 ×Ññ,ó ð,ð ×Ñò;ó ð;÷
!r   rË   c                 ór   ‡ ‡— ddl mŠ ‰ dk(  r‰j                  d„ «       }|S ‰j                  ˆˆ fd„«       }|S )zÙ
    A separate method so we don't need to compile code every assignment (!).

    :param ndim: We need to have static array sizes for cuda.local.array, so
        bake in the number of dimensions into the kernel
    r   )Úcudac                 ó   — |d   | d<   y rÍ   r   )rÿ   r   s     r   Úkernelz_assign_kernel.<locals>.kernel  s   € à˜"‘gˆC�ŠGr   c                 óþ  •— ‰j                  d«      }d}t        | j                  «      D ]  }|| j                  |   z  }Œ ||k\  ry ‰j                  j                  d‰ft        j                  ¬«      }t        ‰dz
  dd«      D ]U  }|| j                  |   z  |d|f<   || j                  |   z  |j                  |   dkD  z  |d|f<   || j                  |   z  }ŒW |t        |d   ‰«         | t        |d   ‰«      <   y )Nr3   rd   r‘   rº   r   )	Úgridre   r7   r%   Úlocalr…   r   Úint64r
   )rÿ   r   ÚlocationÚ
n_elementsr¥   Úidxr  r7   s         €€r   r  z_assign_kernel.<locals>.kernel  s  ø€ à—9‘9˜Q“<ˆàˆ
Ü�s—x‘x“ò 	'ˆAØ˜#Ÿ)™) A™,Ñ&‰Jð	'à�zÒ!ð ð �j‰j×ÑØ�d�)Ü—+‘+ð ó ˆô �t˜a‘x  RÓ(ò 	&ˆAØ  3§9¡9¨Q¡<Ñ/ˆC��1�‰IØ! C§I¡I¨a¡LÑ0°S·Y±Y¸q±\ÀAÑ5EÑFˆC��1�‰IØ˜Ÿ™ 1™Ñ%‰Hð	&ð
 -0´¸sÀ1¹vÀtÓ0LÑ,MˆŒN˜3˜q™6 4Ó(Ò)r   )Únumbar  Újit)r7   r  r  s   ` @r   Ú_assign_kernelr  ÷  sI   ù€ õ àˆq‚yà	�‰ñ	ó 
ð	àˆà	‡X�XôNó ðNð. €Mr   c                   óì   — e Zd ZdZd„ Zed„ «       Zd„ Zdd„Zd„ Z	d„ Z
dd	„Zej                  d
„ «       Zej                  dd„«       Zdd„Zej                  d„ «       Zej                  dd„«       Zdd„Zy)r¢   z
    An on-GPU array type
    c                 ó.   — | j                   j                  S )zA
        Return true if the array is Fortran-contiguous.
        )r;   Úis_f_contigra   s    r   Úis_f_contiguouszDeviceNDArray.is_f_contiguous'  ó   € ð �{‰{×&Ñ&Ð&r   c                 ó@   — t        | j                  j                  «      S rÑ   rÒ   ra   s    r   ru   zDeviceNDArray.flags-  rÓ   r   c                 ó.   — | j                   j                  S )z;
        Return true if the array is C-contiguous.
        )r;   Úis_c_contigra   s    r   r¼   zDeviceNDArray.is_c_contiguous7  r  r   Nc                 ó€   — |r| j                  «       j                  |«      S | j                  «       j                  «       S )zE
        :return: an `numpy.ndarray`, so copies to the host.
        )rš   Ú	__array__)rL   r'   s     r   r  zDeviceNDArray.__array__=  s9   € ñ Ø×$Ñ$Ó&×0Ñ0°Ó7Ð7à×$Ñ$Ó&×0Ñ0Ó2Ð2r   c                 ó    — | j                   d   S rÂ   )r%   ra   s    r   Ú__len__zDeviceNDArray.__len__F  s   € Ø�z‰z˜!‰}Ðr   c                 óæ  — t        |«      dk(  rt        |d   t        t        f«      r|d   }t	        | «      }|| j
                  k(  r4 || j
                  | j                  | j                  | j                  ¬«      S  | j                  j                  |i |¤Ž\  }}|| j                  j                  gk(  r4 ||j
                  |j                  | j                  | j                  ¬«      S t        d«      ‚)z¶
        Reshape the array without changing its contents, similarly to
        :meth:`numpy.ndarray.reshape`. Example::

            d_arr = d_arr.reshape(20, 50, order='F')
        r3   r   )r%   r&   r'   rJ   úoperation requires copying)r6   r!   r*   r»   rú   r%   r&   r'   rJ   r;   ÚreshapeÚextentrf   )rL   ÚnewshapeÚkwsÚclsÚnewarrÚextentss         r   r  zDeviceNDArray.reshapeI  sÓ   € ô ˆx‹=˜AÒ¤*¨X°a©[¼5Ä$¸-Ô"HØ ‘{ˆHä�4‹jˆØ�t—z‘zÒ!á˜TŸZ™Z°·±Ø!ŸZ™Z°$·-±-ôAð Að .˜$Ÿ+™+×-Ñ-¨xÐ?¸3Ñ?‰ˆ�à�t—{‘{×)Ñ)Ð*Ò*Ù˜VŸ\™\°6·>±>Ø!ŸZ™Z°$·-±-ôAð Aô &Ð&BÓCÐCr   c                 ó,  — | j                  |«      }t        | «      }| j                  j                  |¬«      \  }}|| j                  j                  gk(  r5 ||j
                  |j                  | j                  | j                  |¬«      S t        d«      ‚)z¹
        Flattens a contiguous array without changing its contents, similar to
        :meth:`numpy.ndarray.ravel`. If the array is not contiguous, raises an
        exception.
        )r~   ræ   r  )
rm   rú   r;   Úravelr  r%   r&   r'   rJ   rf   )rL   r~   rK   r"  r#  r$  s         r   r&  zDeviceNDArray.ravela  s„   € ð ×%Ñ% fÓ-ˆÜ�4‹jˆØŸ+™+×+Ñ+°%Ð+Ó8‰ˆ�à�t—{‘{×)Ñ)Ð*Ò*Ù˜VŸ\™\°6·>±>Ø!ŸZ™Z°$·-±-Ø$ô&ð &ô
 &Ð&BÓCÐCr   c                 ó$   — | j                  |«      S r   rÖ   rØ   s     r   rÚ   zDeviceNDArray.__getitem__s  rÛ   r   c                 ó&   — | j                  ||«      S rÝ   rÖ   rÞ   s      r   rß   zDeviceNDArray.getitemw  rà   r   c                 ó  — | j                  |«      }| j                  j                  |«      }t        |j	                  «       «      }t        | «      }t        |«      dk(  rÔ | j                  j                  |d   Ž }|j                  s�| j                  j                  �t        | j                  ||¬«      S t        j                  d| j                  ¬«      }t        j                   ||| j                  j"                  |¬«       |d   S  ||j$                  |j&                  | j                  ||¬«      S  | j                  j                  |j(                  Ž } ||j$                  |j&                  | j                  ||¬«      S )Nr3   r   rœ   râ   rã   ræ   )rm   r;   rÚ   r»   Úiter_contiguous_extentrú   r6   rJ   r¡   Úis_arrayr'   rè   rË   r+   r•   r@   r—   r:   r%   r&   r  )rL   rÙ   rK   Úarrr$  r"  rë   r™   s           r   r×   zDeviceNDArray._do_getitem}  sC  € Ø×%Ñ% fÓ-ˆà�k‰k×%Ñ% dÓ+ˆÜ�s×1Ñ1Ó3Ó4ˆÜ�4‹jˆÜˆw‹<˜1ÒØ(�d—m‘m×(Ñ(¨'°!©*Ð5ˆGà—<’<à—:‘:×#Ñ#Ð/Ü'¨d¯j©jÀØ18ô:ð :ô !Ÿh™h q°·
±
Ô;�GÜ×*Ñ*¨w¸GØ04·±×0DÑ0DØ28õ:ð ˜q‘zÐ!á §¡°C·K±KØ!%§¡°gÀfôNð Nð )�d—m‘m×(Ñ(¨#¯*©*Ð5ˆGÙ˜SŸY™Y°·±Ø!ŸZ™Z°'À&ôJð Jr   c                 ó&   — | j                  ||«      S r   rí   rï   s      r   rñ   zDeviceNDArray.__setitem__š  rò   r   c                 ó*   — | j                  |||¬«      S rô   rí   rõ   s       r   rö   zDeviceNDArray.setitemž  r÷   r   c                 óˆ  — | j                  |«      }| }|r$t        j                  «       }|j                  «       }| j                  j                  |«      } | j                  j                  |j                  Ž }t        |t        j                  «      rd}d}	n|j                  }|j                  }	 t        | «      ||	| j                  ||¬«      }
t!        ||d¬«      \  }}|j"                  |
j"                  kD  r&t%        d|j"                  ›d|
j"                  ›d�«      ‚t'        j(                  |
j"                  t&        j*                  ¬«      }|j                  ||
j"                  |j"                  z
  d   |j,                  |Ž }t/        t1        |
j                  |j                  «      «      D ]$  \  }\  }}|d	k7  sŒ||k7  sŒt%        d
|||fz  «      ‚ t3        j4                  t6        j8                  |
j                  d	«      } t;        |
j"                  «      j=                  ||¬«      |
|«       |r|j?                  «        y y )Nr   ræ   T)rK   Úuser_explicitzCan't assign z-D array to z-D selfrâ   r3   zCCan't copy sequence with size %d to array axis %d with dimension %drl   ) rm   r   rC   rù   r;   rÚ   rJ   r¡   r  r!   r   ÚElementr%   r&   rú   r'   rû   r7   r/   r+   Úonesr	  r  Ú	enumerateÚzipr<   r=   r>   r?   r  Úforallrü   )rL   rð   rU   rK   rý   rþ   r,  rë   r%   r&   rÿ   r   r¸   Ú	rhs_shaper¥   ÚlÚrr  s                     r   rî   zDeviceNDArray._do_setitem¤  sù  € à×%Ñ% fÓ-ˆð
 !�jˆÙÜ×%Ñ%Ó'ˆCØ×+Ñ+Ó-ˆFð �k‰k×%Ñ% cÓ*ˆØ$�$—-‘-×$Ñ$ c§j¡jÐ1ˆä�cœ:×-Ñ-Ô.àˆEØ‰Gà—I‘IˆEØ—k‘kˆGàŒd�4‹jØØØ—*‘*ØØôˆô ˜U¨6ÀÔF‰ˆˆQØ�8‰8�c—h‘hÒÝØ—“Ø—“ðó ð ô —G‘G˜CŸH™H¬B¯H©HÔ5ˆ	à*-¯)©)ˆ	�#—(‘(˜SŸX™XÑ%Ð&Ð'Øˆc�k‰k˜9Ð%ˆÜ"¤3 s§y¡y°#·)±)Ó#<Ó=ò 	K‰IˆA‰v��1Ø�A‹v˜!˜q›&Ü ð "=ØABÀAÀq¸zñ"Jó Kð Kð	Kô ×%Ñ%¤h§l¡l°C·I±I¸qÓAˆ
ØBŒ�s—x‘xÓ ×'Ñ'¨
¸6Ð'ÓBÀ3ÈÔLÙØ×ÑÕ ð r   r   )rq   r   rÁ   )rÃ   rÄ   rÅ   rÆ   r  rÈ   ru   r¼   r  r  r  r&  r   rÉ   rÚ   rß   r×   rñ   rö   rî   r   r   r   r¢   r¢   #  s®   „ ñò'ð ñ'ó ð'ò'ó3òòDó0Dð$ ×Ññ&ó ð&ð ×Ñò.ó ð.ó
Jð: ×Ññ,ó ð,ð ×Ñò;ó ð;ô
5!r   r¢   c                   ó.   — e Zd ZdZd„ Zd„ Zd„ Zd„ Zd„ Zy)r°   a"  
    An IPC array handle that can be serialized and transfer to another process
    in the same machine for share a GPU allocation.

    On the destination process, use the *.open()* method to creates a new
    *DeviceNDArray* object that shares the allocation from the original process.
    To release the resources, call the *.close()* method.  After that, the
    destination can no longer use the shared array object.  (Note: the
    underlying weakref to the resource is now dead.)

    This object implements the context-manager interface that calls the
    *.open()* and *.close()* method automatically::

        with the_ipc_array_handle as ipc_array:
            # use ipc_array here as a normal gpu array object
            some_code(ipc_array)
        # ipc_array is dead at this point
    c                 ó    — || _         || _        y r   )Ú_array_descÚ_ipc_handle)rL   r¬   r­   s      r   rN   zIpcArrayHandle.__init__ï  s   € Ø%ˆÔØ%ˆÕr   c                 óŠ   — | j                   j                  t        j                  «       «      }t	        dd|i| j
                  ¤ŽS )z˜
        Returns a new *DeviceNDArray* that shares the allocation from the
        original process.  Must not be used on the original process.
        rJ   r   )r<  Úopenr   rC   r¢   r;  )rL   Údptrs     r   r>  zIpcArrayHandle.openó  s<   € ð
 ×Ñ×$Ñ$¤W×%8Ñ%8Ó%:Ó;ˆÜÑ? dÐ?¨d×.>Ñ.>Ñ?Ð?r   c                 ó8   — | j                   j                  «        y)z5
        Closes the IPC handle to the array.
        N)r<  Úclosera   s    r   rA  zIpcArrayHandle.closeû  s   € ð 	×Ñ×ÑÕ r   c                 ó"   — | j                  «       S r   )r>  ra   s    r   Ú	__enter__zIpcArrayHandle.__enter__  s   € Ø�y‰y‹{Ðr   c                 ó$   — | j                  «        y r   )rA  )rL   rú   rU   Ú	tracebacks       r   Ú__exit__zIpcArrayHandle.__exit__  s   € Ø�
‰
�r   N)	rÃ   rÄ   rÅ   rÆ   rN   r>  rA  rC  rF  r   r   r   r°   r°   Ü  s!   „ ñò$&ò@ò!òór   r°   c                   ó   — e Zd ZdZdd„Zy)ÚMappedNDArrayz4
    A host array that uses CUDA mapped memory.
    c                 ó    — || _         || _        y r   ©rJ   rK   ©rL   rJ   rK   s      r   Údevice_setupzMappedNDArray.device_setup  ó   € Ø ˆŒØˆ�r   NrÁ   ©rÃ   rÄ   rÅ   rÆ   rL  r   r   r   rH  rH    ó   „ ñôr   rH  c                   ó   — e Zd ZdZdd„Zy)ÚManagedNDArrayz5
    A host array that uses CUDA managed memory.
    c                 ó    — || _         || _        y r   rJ  rK  s      r   rL  zManagedNDArray.device_setup  rM  r   NrÁ   rN  r   r   r   rQ  rQ    rO  r   rQ  c                 ó^   — t        | j                  | j                  | j                  ||¬«      S )z/Create a DeviceNDArray object that is like ary.©rK   rJ   )r¢   r%   r&   r'   )r‡   rK   rJ   s      r   Úfrom_array_likerU    s&   € ä˜Ÿ™ C§K¡K°·±À6Ø"*ô,ð ,r   c                 ó2   — t        | j                  ||¬«      S )z.Create a DeviceRecord object that is like rec.rT  )rË   r'   )ÚrecrK   rJ   s      r   Úfrom_record_likerX  "  s   € ä˜Ÿ	™	¨&¸8ÔDÐDr   c                 óº   — | j                   r| j                  s| S g }| j                   D ]#  }|j                  |dk(  rdn
t        d«      «       Œ% | t	        |«         S )aG  
    Extract the repeated core of a broadcast array.

    Broadcast arrays are by definition non-contiguous due to repeated
    dimensions, i.e., dimensions with stride 0. In order to ascertain memory
    contiguity and copy the underlying data from such arrays, we must create
    a view without the repeated dimensions.

    r   N)r&   r(   ÚappendÚslicer*   )r‡   Ú
core_indexÚstrides      r   r�   r�   '  sX   € ð �;Š;˜cŸhšhØˆ
Ø€JØ—+‘+ò =ˆØ×Ñ˜v¨š{™!´°d³Õ<ð=àŒu�ZÓ Ñ!Ð!r   c                 óÔ   — | j                   j                  }t        t        | j                  «      t        | j
                  «      «      D ]  \  }}|dkD  sŒ|dk7  sŒ||k7  r y||z  }Œ y)zÓ
    Returns True iff `ary` is C-style contiguous while ignoring
    broadcasted and 1-sized dimensions.
    As opposed to array_core(), it does not call require_context(),
    which can be quite expensive.
    r3   r   FT)r'   r:   r4  Úreversedr%   r&   )r‡   r(   r%   r]  s       r   rV   rV   9  sd   € ð �9‰9×Ñ€DÜœX c§i¡iÓ0´(¸3¿;¹;Ó2GÓHò ‰ˆˆvØ�1‹9˜ 1›Ø�vŠ~ÙØ�E‰M‰Dð	ð
 r   z™Array contains non-contiguous buffer and cannot be transferred as a single memory region. Please ensure contiguous buffer with numpy .ascontiguousarray()c                 óv   — t        | «      }|j                  d   s|j                  d   st        t        «      ‚y y )Nrp   rr   )r�   ru   r/   Úerrmsg_contiguous_buffer)r‡   Úcores     r   r€   r€   O  s7   € Ü�c‹?€DØ�:‰:�nÒ%¨d¯j©j¸Ò.HÜÔ1Ó2Ð2ð /IÐ%r   c                 ó*  — t        j                  | «      r| dfS t        | d«      r!t        j                  j                  | «      dfS t        | t        j                  «      rt        | |¬«      }n;t        j                  | t        dk  rdndd¬«      } t        | «       t        | |¬«      }|ret        j                  rB|s@t        | t         «      s0t        | t        j"                  «      rd}t%        t'        |«      «       |j)                  | |¬«       |dfS )	zº
    Create a DeviceRecord or DeviceArray like obj and optionally copy data from
    host to device. If obj already represents device memory, it is returned and
    no copy is made.
    FrY   rl   r|   NT)r[   r   zGHost array used in CUDA kernel will incur copy overhead to/from device.)r@   r‚   r   r  r  Úas_cuda_arrayr!   r+   ÚvoidrX  r…   r   r€   rU  r	   ÚCUDA_WARN_ON_IMPLICIT_COPYr¢   r˜   r   r   rŠ   )r   rK   r[   r0  Údevobjrj   s         r   rû   rû   U  sì   € ô ×Ñ Ô$Ø�EˆzÐÜ	�Ð0Ô	1Ü�z‰z×'Ñ'¨Ó,¨eÐ3Ð3ä�cœ2Ÿ7™7Ô#Ü% c°&Ô9‰Fô —(‘(ØÜ+¨fÒ4‘U¸$ØôˆCô ˜cÔ"Ü$ S°Ô8ˆFÙÜ×0Ò0á%Ü# C¬Ô7Ü# C¬¯©Ô4ð;�CäÔ0°Ó5Ô6Ø×!Ñ! #¨fÐ!Ô5Ø�tˆ|Ðr   c                 óÒ  — | j                  «       |j                  «       }}| j                  |j                  k7  r%t        d| j                  ›d|j                  ›�«      ‚|j                  |j                  k7  r%t	        d| j                  ›d|j                  ›�«      ‚| j
                  r?|j                  |j                  k7  r%t	        d| j                  ›d|j                  ›�«      ‚y y )Nzincompatible dtype: z vs. zincompatible shape: zincompatible strides: )r¶   r'   Ú	TypeErrorr%   r/   r(   r&   )Úary1Úary2Úary1sqÚary2sqs       r   rƒ   rƒ   |  s³   € Ø—\‘\“^ T§\¡\£^ˆF€FØ‡z�z�T—Z‘ZÒÝØŸ› T§Z¢Zð1ó 2ð 	2à‡|�|�v—|‘|Ò#ÝØŸ*›* d§j¢jð2ó 3ð 	3ð ‡y‚y�V—^‘^ v§~¡~Ò5ÝØŸ,›,¨¯ªð6ó 7ð 	7ð 6€yr   r�   )r   TF)7rÆ   r�   r<   r>   r[   Úctypesr   Únumpyr+   r  r   Únumba.cuda.cudadrvr   r   r   r@   Ú
numba.corer   r	   Únumba.np.unsafe.ndarrayr
   Únumba.np.numpy_supportr   Únumba.npr   Únumba.cuda.api_utilr   Únumba.core.errorsr   Úwarningsr   r   r   r    r   r-   r)   ÚDeviceArrayr1   rË   r  r¢   Úobjectr°   r˜   rH  rQ  rU  rX  r�   rV   ra  r€   rû   rƒ   r   r   r   ú<module>rz     s'  ðñó Û Û Û Ý ã ã Ý ß 2Ý 0ß $Ý 2Ý 0Ý "Ý ;Ý 5Ý ðØ/‘˜	 ;Ó/°Ó5€Iò3ò
ò ;ôO/˜×0Ñ0ô O/ôd
d!Ð$ô d!ðN ñ(ó ð(ôVv!Ð%ô v!ôr)�Vô )ôXÐ% r§z¡zô ôÐ&¨¯
©
ô ó,óEò
"ò$ð 3Ð ò3ó$óN7øð ò ôðús   Á"C? Ã?D
Ä	D
