ó
    „ñ:i“y  ã                   ó"  • S r SSKrSSKrSSKrSSKrSSKJr  SSKrSSK	r	SSK	J
r
  SSKJrJr  SSKJr  SSKJrJr  SSKJr  SS	KJr  SS
KJr  SSKJr  SSKJr  SSKJr   \" \S5      " S5      r S r"S r#S r$ " S S\
RJ                  5      r& " S S\&5      r'\ S 5       r( " S S\&5      r) " S S\*5      r+ " S S\&\RX                  5      r- " S S\&\RX                  5      r.S(S  jr/S(S! jr0S" r1S# r2S$r3S% r4S)S& jr5S' r6g! \! a    S r  Nœf = f)*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                 ó   • U $ ©N© )Úfuncs    Úa/srv/projetos/modelo_ml_acdoc/venv/lib/python3.13/site-packages/numba/cuda/cudadrv/devicearray.pyr   r      s   € Øˆó    c                 ó   • [        U SS5      $ )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 5        U 4S jnU" S[        5        U" S[        5        U" S[        R                  5        U" S[        5        g)z,Verify the CUDA ndarray interface for an objc                 ó˜   >• [        TU 5      (       d  [        U 5      e[        [        TU 5      U5      (       d  [        U < SU< 35      eg )Nz must be of type )ÚhasattrÚAttributeErrorÚ
isinstancer   )ÚattrÚtypr   s     €r   Úrequires_attrÚ4verify_cuda_ndarray_interface.<locals>.requires_attr,   sD   ø€ Ü�s˜D×!Ñ!Ü  Ó&Ð&Üœ' # tÓ,¨c×2Ñ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                 ó:   • [        U 5      (       d  [        S5      eg)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                   ó
  • \ rS rSrSrSrSrSS jr\S 5       r	SS jr
\S 5       rSS	 jrS
 r\S 5       r\S 5       r\R"                  SS j5       r\R"                  SS j5       rSS jrS rS rSS jrS r\S 5       rSrg)ÚDeviceNDArrayBaseé>   z$A on GPU NDArray representation
    TNc                 ó^  • [        U[        5      (       a  U4n[        U[        5      (       a  U4n[        R                  " U5      n[	        U5      U l        [	        U5      U R
                  :w  a  [        S5      e[        R                  R                  SXUR                  5      U l        [        U5      U l        [        U5      U l        X0l        [        [        R                   " ["        R$                  U R                  S5      5      U l        U R&                  S:”  a“  Uct  [(        R*                  " U R                  U R                  U R                  R                  5      U l        [.        R0                  " 5       R3                  U R,                  5      nOŒ[(        R4                  " U5      U l        Op[(        R6                  (       a   [(        R8                  R;                  S5      nO[=        S5      n[(        R>                  " [.        R0                  " 5       USS9nSU l        XPl         X@l!        g)zÍ
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Úndimr0   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(   rM   rL   Únulls          r   Ú__init__ÚDeviceNDArrayBase.__init__D   sˆ  € ô  �eœS×!Ñ!Ø�HˆEÜ�gœs×#Ñ#Ø�jˆGÜ—’˜“ˆÜ˜“JˆŒ	Üˆw‹<˜4Ÿ9™9Ó$ÜÐ5Ó6Ð6Ü ×&Ñ&×0Ñ0°°EØ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                 óÖ  • [         R                  (       a&  U R                  b  [        U R                  5      nO3SnO0U R                  R                  b  U R                  R                  nOSn[        U R                  5      [        U 5      (       a  S O[        U R                  5      US4U R                  R                  U R                  S:w  a  [        U R                  5      SS.$ S SS.$ )Nr   Fé   )r&   r'   ÚdataÚtypestrrM   Úversion)rB   rH   Údevice_ctypes_pointerr-   Úvaluer+   r&   Úis_contiguousr'   r(   ÚstrrM   )rN   Úptrs     r   Ú__cuda_array_interface__Ú*DeviceNDArrayBase.__cuda_array_interface__w   sÄ   € ä×!×!Ø×)Ñ)Ñ5Ü˜$×4Ñ4Ó5‘à‘à×)Ñ)×/Ñ/Ñ;Ø×0Ñ0×6Ñ6‘à�ô ˜4Ÿ:™:Ó&Ü,¨T×2Ñ2‘t¼¸d¿l¹lÓ8KØ˜%�LØ—z‘z—~‘~Ø*.¯+©+¸Ó*:”c˜$Ÿ+™+Ó&Øñ
ð 	
ð
 AEØñ
ð 	
r   c                 ó>   • [         R                   " U 5      nXl        U$ )zoBind a CUDA stream to this object so that all subsequent operation
on this array defaults to the given stream.
)ÚcopyrM   )rN   rM   Úclones      r   ÚbindÚDeviceNDArrayBase.bind�   s   € ô —	’	˜$“ˆØŒØˆr   c                 ó"   • U R                  5       $ r   ©Ú	transpose©rN   s    r   ÚTÚDeviceNDArrayBase.T•   s   € à�~‰~ÓÐr   c                 ó:  • U(       a-  [        U5      [        [        U R                  5      5      :X  a  U $ U R                  S:w  a  Sn[        U5      eUb:  [	        U5      [	        [        U R                  5      5      :w  a  [        SU< 35      eSSKJn  U" U 5      $ )Né   z2transposing a non-2D DeviceNDArray isn't supportedzinvalid axes list r   rd   )r+   Úranger9   ÚNotImplementedErrorÚsetr0   Únumba.cuda.kernels.transposere   )rN   ÚaxesÚmsgre   s       r   re   Ú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                 ó,   • U(       d  U R                   $ U$ r   ©rM   )rN   rM   s     r   Ú_default_streamÚ!DeviceNDArrayBase._default_stream¥   s   € Þ"(ˆt�{‰{Ð4¨fÐ4r   c                 ó  • SU R                   ;   nU R                  S   (       a
  U(       d  SnO U R                  S   (       a
  U(       d  SnOSn[        R                  " U R                  5      n[
        R                  " X0R                  U5      $ )úV
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   r:   r9   )rN   Ú	broadcastÚlayoutr(   s       r   Ú_numba_type_ÚDeviceNDArrayBase._numba_type_¨   sf   € ð( ˜Ÿ™Ñ%ˆ	Ø�:‰:�n×%®iØ‰FØ�Z‰Z˜×'¶	Ø‰FàˆFä×(Ò(¨¯©Ó4ˆÜ�{Š{˜5§)¡)¨VÓ4Ð4r   c                 óÆ   • U R                   c?  [        R                  (       a  [        R                  R	                  S5      $ [        S5      $ U R                   R                  $ )z:Returns the ctypes pointer to the GPU data buffer
        r   )rL   rB   rH   rI   rJ   r   rW   rf   s    r   rW   Ú'DeviceNDArrayBase.device_ctypes_pointerÇ   sF   € ð �=‰=Ñ Ü×%×%Ü—‘×2Ñ2°1Ó5Ð5ä “{Ð"à—=‘=×6Ñ6Ð6r   c                 ó  • UR                   S:X  a  g[        U 5        U R                  U5      n[        U 5      [        U5      pC[        R
                  " U5      (       a7  [        U5        [        X45        [        R                  " XU R                  US9  g[        R                  " UUR                  S   (       a  SOSS[        S:  a  UR                  S	   (       + OSS
9n[        X45        [        R                  " XU R                  US9  g)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   Nrs   rx   ry   r{   T©rj   r   Ú	WRITEABLE)ÚorderÚsubokr_   )r)   Úsentry_contiguousrt   Ú
array_corerB   Úis_device_memoryÚcheck_array_compatibilityÚdevice_to_devicerD   r,   Úarrayr}   r   Úhost_to_device)rN   ÚaryrM   Ú	self_coreÚary_cores        r   Úcopy_to_deviceÚ DeviceNDArrayBase.copy_to_deviceÓ   sÖ   € ð �8‰8�q‹=àä˜$ÔØ×%Ñ% fÓ-ˆä(¨Ó.´
¸3³�8Ü×#Ò# C×(Ñ(Ü˜cÔ"Ü% iÔ:Ü×$Ò$ T°·±ÈÓOô —x’xØØ&Ÿ_™_¨^×<‘cÀ#Øä  6Ó)ð #Ÿ.™.¨Ñ5Õ5Ø/3ñ5ˆHô & iÔ:Ü×"Ò" 4°4·?±?Ø*0ó2r   c                 ó¨  • [        S U R                   5       5      (       a&  Sn[        UR                  U R                  5      5      eU R                  S:¼  d   S5       eU R                  U5      nUc.  [        R                  " U R                  [        R                  S9nO[        X5        UnU R                  S:w  a   [        R                  " X@U R                  US9  Ucq  U R                  S:X  a,  [        R                  " U R                  U R                  US9nU$ [        R                  " U R                  U R                  U R                  US9nU$ )	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   ó*   #   • U  H	  oS :  v •  M     g7f©r   Nr   )Ú.0Úss     r   Ú	<genexpr>Ú1DeviceNDArrayBase.copy_to_host.<locals>.<genexpr>
  s   é € Ð+šl˜�1Žušlùs   ‚z2D->H copy not implemented for negative strides: {}r   zNegative memory size©r&   r(   rs   )r&   r(   Úbuffer)r&   r(   r'   rž   )Úanyr'   rl   ÚformatrD   rt   r,   ÚemptyÚbyter�   rB   Údevice_to_hostr)   Úndarrayr&   r(   )rN   r‘   rM   rp   Úhostarys        r   Úcopy_to_hostÚDeviceNDArrayBase.copy_to_hostó   s  € ô. Ñ+˜dŸlšlÓ+×+Ñ+ØFˆCÜ% c§j¡j°·±Ó&>Ó?Ð?Ø�‰ !Ó#Ð;Ð%;Ó;Ð#Ø×%Ñ% fÓ-ˆØ‰;Ü—h’h T§_¡_¼B¿G¹GÑD‰Gä% dÔ0ØˆGà�?‰?˜aÓÜ×"Ò" 7°$·/±/Ø*0ò2ð ‰;Ø�y‰y˜A‹~ÜŸ*š*¨4¯:©:¸T¿Z¹ZØ,3ñ5�ð
 ˆô Ÿ*š*¨4¯:©:¸T¿Z¹ZØ-1¯\©\À'ñK�àˆr   c           	   #   óL  #   • U R                  U5      nU R                  S:w  a  [        S5      eU R                  S   U R                  R
                  :w  a  [        S5      e[        [        R                  " [        U R                  5      U-  5      5      nU R                  nU R                  R
                  n[        U5       H\  nXa-  n[        Xq-   U R                  5      nX‡-
  4n	U R                  R                  Xu-  X…-  5      n
[        X”U R                  UU
S9v •  M^     g7f)z†Split the array into equal partition of the `section` size.
If the array cannot be equally divided, the last section will be
smaller.
r5   zonly support 1d arrayr   zonly support unit stride©r(   rM   rL   N)rt   r9   r0   r'   r(   r<   r-   ÚmathÚceilÚfloatr)   rk   ÚminrL   ÚviewÚDeviceNDArray)rN   ÚsectionrM   Únsectr'   r<   ÚiÚbeginÚendr&   rL   s              r   ÚsplitÚ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–ˆAØ‘KˆEÜ�e‘o t§y¡yÓ1ˆCØ‘[�NˆEØ—}‘}×)Ñ)¨%Ñ*:¸C¹NÓKˆHÜ °d·j±jÈØ)1ñ3ô 3ò ùs   ‚D"D$c                 ó   • U R                   $ )zEReturns a device memory object that is used as the argument.
        )rL   rf   s    r   Úas_cuda_argÚDeviceNDArrayBase.as_cuda_arg7  s   € ð �}‰}Ðr   c                 óÂ   • [         R                  " 5       R                  U R                  5      n[	        U R
                  U R                  U R                  S9n[        XS9$ )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   rE   Úget_ipc_handlerL   Údictr&   r'   r(   ÚIpcArrayHandle)rN   ÚipchÚdescs      r   r½   Ú DeviceNDArrayBase.get_ipc_handle<  sF   € ô ×"Ò"Ó$×3Ñ3°D·M±MÓBˆÜ˜$Ÿ*™*¨d¯l©lÀ$Ç*Á*ÑMˆÜ¨Ñ?Ð?r   c                 óÀ   • U R                   R                  US9u  p4[        UR                  UR                  U R
                  U R                  U5      U R                  S9$ )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(   rM   rL   )r=   Úsqueezer¯   r&   r'   r(   rt   rL   )rN   rÄ   rM   Ú	new_dummyÚ_s        r   rÆ   ÚDeviceNDArrayBase.squeezeG  sV   € ð& —{‘{×*Ñ*°Ð*Ð5‰ˆ	ÜØ—/‘/Ø×%Ñ%Ø—*‘*Ø×'Ñ'¨Ó/Ø—]‘]ñ
ð 	
r   c                 óø  • [         R                  " U5      n[        U R                  5      n[        U R                  5      nU R                  R
                  UR
                  :w  av  U R                  5       (       d  [        S5      e[        US   U R                  R
                  -  UR
                  5      u  US'   nUS:w  a  [        S5      eUR
                  US'   [        UUUU R                  U R                  S9$ )zUReturns 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_contiguousr0   Údivmodr¯   rM   rL   )rN   r(   r&   r'   Úrems        r   r®   Ú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   • U R                   R                  U R                  -  $ r   )r(   r<   r)   rf   s    r   ÚnbytesÚDeviceNDArrayBase.nbytesˆ  s   € ð
 �z‰z×"Ñ" T§Y¡YÑ.Ð.r   )	r=   rD   r(   rL   r9   r&   r)   rM   r'   r˜   ©r   r   ©Nr   )Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__Ú__cuda_memory__r   rP   Úpropertyr\   ra   rg   re   rt   r�   rW   r   Úrequire_contextr”   r¦   rµ   r¸   r½   rÆ   r®   rÒ   Ú__static_attributes__r   r   r   r2   r2   >   sÛ   † ñà€OØÐô1ðf ñ
ó ð
ô*ð ñ ó ð ô
#ò5ð ñ5ó ð5ð< ñ	7ó ð	7ð ×Ñó2ó ð2ð> ×Ñó,ó ð,ô\3ò*ò
	@ô
ò8#
ðJ ñ/ó ó/r   r2   c                   óü   ^ • \ rS rSrSrSU 4S jjr\S 5       r\S 5       r\	R                  S 5       r\	R                  SS j5       rSS jr\	R                  S	 5       r\	R                  SS
 j5       rSS jrSrU =r$ )ÚDeviceRecordi�  z
An on-GPU record type
c                 ó8   >• SnSn[         [        U ]  XEXU5        g ©Nr   )Úsuperrà   rP   )rN   r(   rM   rL   r&   r'   Ú	__class__s         €r   rP   ÚDeviceRecord.__init__”  s#   ø€ ØˆØˆÜŒl˜DÑ*¨5¸5Ø+3õ	5r   c                 ó@   • [        U R                  R                  5      $ ©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=   r}   rf   s    r   r}   ÚDeviceRecord.flagsš  ó   € ô �D—K‘K×%Ñ%Ó&Ð&r   c                 óB   • [         R                  " U R                  5      $ )rw   )r   r~   r(   rf   s    r   r�   ÚDeviceRecord._numba_type_¤  s   € ô ×'Ò'¨¯
©
Ó3Ð3r   c                 ó$   • U R                  U5      $ r   ©Ú_do_getitem©rN   Úitems     r   Ú__getitem__ÚDeviceRecord.__getitem__¬  ó   € à×Ñ Ó%Ð%r   c                 ó$   • U R                  X5      $ ©z0Do `__getitem__(item)` with CUDA stream
        rî   ©rN   rñ   rM   s      r   ÚgetitemÚDeviceRecord.getitem°  ó   € ð ×Ñ Ó-Ð-r   c                 ó¸  • U R                  U5      nU R                  R                  U   u  p4U R                  R	                  U5      nUR
                  S:X  aQ  UR                  b
  [        X2US9$ [        R                  " SUS9n[        R                  " XeUR                  US9  US   $ [        UR
                  S UR                  S   S5      u  pxn	[        XxX•US9$ )	Nr   r©   r5   ©r(   ©ÚdstÚsrcr)   rM   r   ry   ©r&   r'   r(   rL   rM   )rt   r(   ÚfieldsrL   r®   r&   Únamesrà   r,   r¡   rB   r£   r<   r   Úsubdtyper¯   )
rN   rñ   rM   r#   ÚoffsetÚnewdatar¥   r&   r'   r(   s
             r   rï   ÚDeviceRecord._do_getitem¶  sÜ   € Ø×%Ñ% fÓ-ˆØ—j‘j×'Ñ'¨Ñ-‰ˆØ—-‘-×$Ñ$ VÓ,ˆà�9‰9˜‹?Ø�y‰yÑ$Ü#¨#Ø-4ñ6ð 6ô Ÿ(š( 1¨CÑ0�Ü×&Ò&¨7Ø,/¯L©LØ.4ò6ð ˜1‘:Ðô ,¨C¯I©IØ,0Ø,/¯L©L¸©O¸SóBñ "ˆE˜Eô ! uØ',Ø(.ñ0ð 0r   c                 ó$   • U R                  X5      $ r   ©Ú_do_setitem©rN   ÚkeyrX   s      r   Ú__setitem__ÚDeviceRecord.__setitem__Î  ó   € à×Ñ Ó+Ð+r   c                 ó"   • U R                  XUS9$ ©z6Do `__setitem__(key, value)` with CUDA stream
        rs   r  ©rN   r  rX   rM   s       r   ÚsetitemÚDeviceRecord.setitemÒ  ó   € ð ×Ñ °6ÐÐ:Ð:r   c                 óê  • U R                  U5      nU(       + nU(       a%  [        R                  " 5       nUR                  5       nU R                  R
                  U   u  pgU R                  R                  U5      n[        U 5      " XcUS9n	[        U	R                  R                  U5      US9u  p«[        R                  " XšU
R                  R                  U5        U(       a  UR                  5         g g )Nr©   rs   )rt   r   rE   Úget_default_streamr(   r  rL   r®   ÚtypeÚauto_devicerB   rŽ   r<   Úsynchronize)rN   r  rX   rM   ÚsynchronousÚctxr#   r  r  ÚlhsÚrhsrÈ   s               r   r	  ÚDeviceRecord._do_setitemØ  sÅ   € à×%Ñ% fÓ-ˆð
 !”jˆÞÜ×%Ò%Ó'ˆCØ×+Ñ+Ó-ˆFð —j‘j×'Ñ'¨Ñ,‰ˆØ—-‘-×$Ñ$ VÓ,ˆä�4Œj˜s¸GÑDˆô ˜SŸY™YŸ^™^¨EÓ2¸6ÑB‰ˆô 	× Ò  ¨3¯9©9×+=Ñ+=¸vÔFæØ×ÑÕ ð r   r   r˜   rÔ   )rÖ   r×   rØ   rÙ   rÚ   rP   rÜ   r}   r�   r   rÝ   rò   rø   rï   r  r  r	  rÞ   Ú__classcell__)rä   s   @r   rà   rà   �  s¬   ø† ñ÷5ð ñ'ó ð'ð ñ4ó ð4ð ×Ññ&ó ð&ð ×Ñó.ó ð.ô
0ð0 ×Ññ,ó ð,ð ×Ñó;ó ð;÷
!ò !r   rà   c                 óv   ^ ^• SSK Jm  T S:X  a  TR                  S 5       nU$ TR                  UU 4S j5       nU$ )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                 ó   • US   U S'   g râ   r   )r  r  s     r   ÚkernelÚ_assign_kernel.<locals>.kernel  s   € à˜"‘gˆC�ŠGr   c                 óø  >• TR                  S5      nSn[        U R                  5       H  nX0R                  U   -  nM     X#:¼  a  g TR                  R                  ST4[        R                  S9n[        TS-
  SS5       HS  nX R                  U   -  USU4'   X R                  U   -  UR                  U   S:„  -  USU4'   X R                  U   -  nMU     U[        US   T5         U [        US   T5      '   g )Nr5   rj   r�   rË   r   )	Úgridrk   r9   r&   Úlocalr�   r   Úint64r
   )r  r  ÚlocationÚ
n_elementsr²   Úidxr!  r9   s         €€r   r#  r$    sý   ø€ à—9‘9˜Q“<ˆàˆ
Ü�s—x‘x–ˆAØŸ)™) A™,Ñ&ŠJñ !àÓ!ð ð �j‰j×ÑØ�d�)Ü—+‘+ð ð ˆô �t˜a‘x  RÖ(ˆAØ §9¡9¨Q¡<Ñ/ˆC��1�‰IØ!§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)r9   r#  r!  s   ` @r   Ú_assign_kernelr.  ÷  sI   ù€ õ àˆqƒyà	�‰ñ	ó 
ð	àˆà	‡X�XõNó ðNð. €Mr   c                   ó   • \ rS rSrSrS r\S 5       rS rSS jr	S r
S	 rSS
 jr\R                  S 5       r\R                  SS j5       rSS jr\R                  S 5       r\R                  SS j5       rSS jrSrg)r¯   i#  z
An on-GPU array type
c                 ó.   • U R                   R                  $ )z1
Return true if the array is Fortran-contiguous.
)r=   Úis_f_contigrf   s    r   Úis_f_contiguousÚDeviceNDArray.is_f_contiguous'  ó   € ð �{‰{×&Ñ&Ð&r   c                 ó@   • [        U R                  R                  5      $ rç   rè   rf   s    r   r}   ÚDeviceNDArray.flags-  rê   r   c                 ó.   • U R                   R                  $ )z+
Return true if the array is C-contiguous.
)r=   Úis_c_contigrf   s    r   rÍ   ÚDeviceNDArray.is_c_contiguous7  r4  r   Nc                 óŠ   • U(       a  U R                  5       R                  U5      $ U R                  5       R                  5       $ )z5
:return: an `numpy.ndarray`, so copies to the host.
)r¦   Ú	__array__)rN   r(   s     r   r;  ÚDeviceNDArray.__array__=  s9   € ö Ø×$Ñ$Ó&×0Ñ0°Ó7Ð7à×$Ñ$Ó&×0Ñ0Ó2Ð2r   c                 ó    • U R                   S   $ rÕ   )r&   rf   s    r   Ú__len__ÚDeviceNDArray.__len__F  s   € Ø�z‰z˜!‰}Ðr   c                 óä  • [        U5      S:X  a#  [        US   [        [        45      (       a  US   n[	        U 5      nXR
                  :X  a1  U" U R
                  U R                  U R                  U R                  S9$ U R                  R                  " U0 UD6u  pEXPR                  R                  /:X  a1  U" UR
                  UR                  U R                  U R                  S9$ [        S5      e)z–
Reshape the array without changing its contents, similarly to
:meth:`numpy.ndarray.reshape`. Example::

    d_arr = d_arr.reshape(20, 50, order='F')
r5   r   )r&   r'   r(   rL   úoperation requires copying)r8   r!   r+   rÌ   r  r&   r'   r(   rL   r=   ÚreshapeÚextentrl   )rN   ÚnewshapeÚkwsÚclsÚnewarrÚextentss         r   rB  ÚDeviceNDArray.reshapeI  sÎ   € ô ˆx‹=˜AÓ¤*¨X°a©[¼5Ä$¸-×"HÑ"HØ ‘{ˆHä�4‹jˆØ—z‘zÓ!á˜TŸZ™Z°·±Ø!ŸZ™Z°$·-±-ñAð Að Ÿ+™+×-Ò-¨xÐ?¸3Ñ?‰ˆà—{‘{×)Ñ)Ð*Ó*Ù˜VŸ\™\°6·>±>Ø!ŸZ™Z°$·-±-ñAð Aô &Ð&BÓCÐCr   c                 ó  • U R                  U5      n[        U 5      nU R                  R                  US9u  pEXPR                  R                  /:X  a2  U" UR
                  UR                  U R                  U R                  US9$ [        S5      e)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   rA  )
rt   r  r=   ÚravelrC  r&   r'   r(   rL   rl   )rN   rˆ   rM   rF  rG  rH  s         r   rK  ÚDeviceNDArray.ravela  s€   € ð ×%Ñ% fÓ-ˆÜ�4‹jˆØŸ+™+×+Ñ+°%Ð+Ð8‰ˆà—{‘{×)Ñ)Ð*Ó*Ù˜VŸ\™\°6·>±>Ø!ŸZ™Z°$·-±-Ø$ñ&ð &ô
 &Ð&BÓCÐCr   c                 ó$   • U R                  U5      $ r   rî   rð   s     r   rò   ÚDeviceNDArray.__getitem__s  rô   r   c                 ó$   • U R                  X5      $ rö   rî   r÷   s      r   rø   ÚDeviceNDArray.getitemw  rú   r   c                 óð  • U R                  U5      nU R                  R                  U5      n[        UR	                  5       5      n[        U 5      n[        U5      S:X  aÎ  U R                  R                  " US   6 nUR                  (       dz  U R                  R                  b  [        U R                  UUS9$ [        R                  " SU R                  S9n[        R                   " XvU R                  R"                  US9  US   $ U" UR$                  UR&                  U R                  XbS9$ U R                  R                  " UR(                  6 nU" UR$                  UR&                  U R                  XbS9$ )Nr5   r   r©   rü   rý   r   )rt   r=   rò   rÌ   Úiter_contiguous_extentr  r8   rL   r®   Úis_arrayr(   r  rà   r,   r¡   rB   r£   r<   r&   r'   rC  )rN   rñ   rM   ÚarrrH  rF  r  r¥   s           r   rï   ÚDeviceNDArray._do_getitem}  s9  € Ø×%Ñ% fÓ-ˆà�k‰k×%Ñ% dÓ+ˆÜ�s×1Ñ1Ó3Ó4ˆÜ�4‹jˆÜˆw‹<˜1ÓØ—m‘m×(Ò(¨'°!©*Ð5ˆGà—<—<à—:‘:×#Ñ#Ñ/Ü'¨d¯j©jÀØ18ñ:ð :ô !Ÿhšh q°·
±
Ñ;�GÜ×*Ò*¨wØ04·±×0DÑ0DØ28ò:ð ˜q‘zÐ!á §¡°C·K±KØ!%§¡°gñNð Nð —m‘m×(Ò(¨#¯*©*Ð5ˆGÙ˜SŸY™Y°·±Ø!ŸZ™Z°'ñJð Jr   c                 ó$   • U R                  X5      $ r   r  r
  s      r   r  ÚDeviceNDArray.__setitem__š  r  r   c                 ó"   • U R                  XUS9$ r  r  r  s       r   r  ÚDeviceNDArray.setitemž  r  r   c                 ó¢  • U R                  U5      nU(       + nU(       a%  [        R                  " 5       nUR                  5       nU R                  R                  U5      nU R                  R                  " UR                  6 n[        U[        R                  5      (       a  SnSn	OUR                  nUR                  n	[        U 5      " UU	U R                  UUS9n
[!        X#SS9u  p¼UR"                  U
R"                  :”  a(  [%        SUR"                  < SU
R"                  < S35      e[&        R(                  " U
R"                  [&        R*                  S9nUR                  XÚR"                  UR"                  -
  S & UR,                  " U6 n[/        [1        U
R                  UR                  5      5       H'  u  nu  nnUS	:w  d  M  UU:w  d  M  [%        S
UXï4-  5      e   [2        R4                  " [6        R8                  U
R                  S	5      n[;        U
R"                  5      R=                  UUS9" X«5        U(       a  UR?                  5         g g )Nr   r   T)rM   Úuser_explicitzCan't assign z-D array to z-D selfrü   r5   zCCan't copy sequence with size %d to array axis %d with dimension %drs   ) rt   r   rE   r  r=   rò   rL   r®   rC  r!   r   ÚElementr&   r'   r  r(   r  r9   r0   r,   Úonesr(  rB  Ú	enumerateÚzipr>   r?   r@   rA   r.  Úforallr  )rN   r  rX   rM   r  r  rT  r  r&   r'   r  r  rÈ   Ú	rhs_shaper²   ÚlÚrr*  s                     r   r	  ÚDeviceNDArray._do_setitem¤  sã  € à×%Ñ% fÓ-ˆð
 !”jˆÞÜ×%Ò%Ó'ˆCØ×+Ñ+Ó-ˆFð �k‰k×%Ñ% cÓ*ˆØ—-‘-×$Ò$ c§j¡jÐ1ˆä�cœ:×-Ñ-×.Ñ.àˆEØ‰Gà—I‘IˆEØ—k‘kˆGä�4ŒjØØØ—*‘*ØØñˆô ˜UÀÑF‰ˆØ�8‰8�c—h‘hÓÝØ—”Ø—”ðó ð ô —G’G˜CŸH™H¬B¯H©HÑ5ˆ	à*-¯)©)ˆ	—(‘(˜SŸX™XÑ%Ð&Ð'Ø�kŠk˜9Ð%ˆÜ"¤3 s§y¡y°#·)±)Ó#<Ö=‰IˆA‰v��1Ø�A�v˜!˜q�&Ü ð "=ØABÀA¸zñ"Jó Kð Kñ >ô ×%Ò%¤h§l¡l°C·I±I¸qÓAˆ
Ü�s—x‘xÓ ×'Ñ'¨
¸6Ð'ÑBÀ3ÔLÞØ×ÑÕ ð r   r   r   )ry   r   rÔ   )rÖ   r×   rØ   rÙ   rÚ   r2  rÜ   r}   rÍ   r;  r>  rB  rK  r   rÝ   rò   rø   rï   r  r  r	  rÞ   r   r   r   r¯   r¯   #  s®   † ñò'ð ñ'ó ð'ò'ô3òòDô0Dð$ ×Ññ&ó ð&ð ×Ñó.ó ð.ô
Jð: ×Ññ,ó ð,ð ×Ñó;ó ð;÷
5!r   r¯   c                   ó6   • \ rS rSrSrS rS rS rS rS r	Sr
g	)
r¿   iÜ  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                 ó   • X l         Xl        g r   ©Ú_array_descÚ_ipc_handle)rN   r»   r¼   s      r   rP   ÚIpcArrayHandle.__init__ï  s   € Ø%ÔØ%Õr   c                 óŒ   • U R                   R                  [        R                  " 5       5      n[	        SSU0U R
                  D6$ )z€
Returns a new *DeviceNDArray* that shares the allocation from the
original process.  Must not be used on the original process.
rL   r   )ri  Úopenr   rE   r¯   rh  )rN   Údptrs     r   rl  ÚIpcArrayHandle.openó  s<   € ð
 ×Ñ×$Ñ$¤W×%8Ò%8Ó%:Ó;ˆÜÑ? dÐ?¨d×.>Ñ.>Ñ?Ð?r   c                 ó8   • U R                   R                  5         g)z%
Closes the IPC handle to the array.
N)ri  Úcloserf   s    r   rp  ÚIpcArrayHandle.closeû  s   € ð 	×Ñ×ÑÕ r   c                 ó"   • U R                  5       $ r   )rl  rf   s    r   Ú	__enter__ÚIpcArrayHandle.__enter__  s   € Ø�y‰y‹{Ðr   c                 ó$   • U R                  5         g r   )rp  )rN   r  rX   Ú	tracebacks       r   Ú__exit__ÚIpcArrayHandle.__exit__  s   € Ø�
‰
�r   rg  N)rÖ   r×   rØ   rÙ   rÚ   rP   rl  rp  rs  rw  rÞ   r   r   r   r¿   r¿   Ü  s!   † ñò$&ò@ò!òõr   r¿   c                   ó"   • \ rS rSrSrSS jrSrg)ÚMappedNDArrayi  z,
A host array that uses CUDA mapped memory.
c                 ó   • Xl         X l        g r   ©rL   rM   ©rN   rL   rM   s      r   Údevice_setupÚMappedNDArray.device_setup  ó   € Ø ŒØ�r   r|  NrÔ   ©rÖ   r×   rØ   rÙ   rÚ   r~  rÞ   r   r   r   rz  rz    ó   † ñ÷r   rz  c                   ó"   • \ rS rSrSrSS jrSrg)ÚManagedNDArrayi  z-
A host array that uses CUDA managed memory.
c                 ó   • Xl         X l        g r   r|  r}  s      r   r~  ÚManagedNDArray.device_setup  r€  r   r|  NrÔ   r�  r   r   r   r„  r„    r‚  r   r„  c                 óX   • [        U R                  U R                  U R                  UUS9$ )z/Create a DeviceNDArray object that is like ary.©rM   rL   )r¯   r&   r'   r(   )r‘   rM   rL   s      r   Úfrom_array_liker‰    s&   € ä˜Ÿ™ C§K¡K°·±À6Ø"*ñ,ð ,r   c                 ó*   • [        U R                  XS9$ )z.Create a DeviceRecord object that is like rec.rˆ  )rà   r(   )ÚrecrM   rL   s      r   Úfrom_record_likerŒ  "  s   € ä˜Ÿ	™	¨&ÑDÐDr   c                 óÔ   • U R                   (       a  U R                  (       d  U $ / nU R                    H%  nUR                  US:X  a  SO
[        S5      5        M'     U [	        U5         $ )a/  
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‹   '  sS   € ð �;�;˜cŸhŸhØˆ
Ø€JØ—+”+ˆØ×Ñ˜v¨›{™!´°d³Ö<ñ àŒu�ZÓ Ñ!Ð!r   c                 óÜ   • U R                   R                  n[        [        U R                  5      [        U R
                  5      5       H   u  p#US:”  d  M  US:w  d  M  X:w  a    gX-  nM"     g)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.
r5   r   FT)r(   r<   r_  Úreversedr&   r'   )r‘   r)   r&   r‘  s       r   rY   rY   9  sZ   € ð �9‰9×Ñ€DÜœX c§i¡iÓ0´(¸3¿;¹;Ó2GÖH‰ˆØ�1�9˜ 1�Ø‹~ÙØ‰MŠDñ	 Ið
 r   z™Array contains non-contiguous buffer and cannot be transferred as a single memory region. Please ensure contiguous buffer with numpy .ascontiguousarray()c                 óŠ   • [        U 5      nUR                  S   (       d$  UR                  S   (       d  [        [        5      eg g )Nrx   rz   )r‹   r}   r0   Úerrmsg_contiguous_buffer)r‘   Úcores     r   rŠ   rŠ   O  s7   € Ü�c‹?€DØ�:‰:�n×%¨d¯j©j¸×.HÜÔ1Ó2Ð2ð /IÐ%r   c                 ób  • [         R                  " U 5      (       a  U S4$ [        U S5      (       a!  [        R                  R                  U 5      S4$ [        U [        R                  5      (       a
  [        XS9nO6[        R                  " U [        S:  a  SOSSS9n [        U 5        [        XS9nU(       au  [        R                  (       aQ  U(       dJ  [        U [         5      (       d5  [        U [        R"                  5      (       a  Sn[%        ['        U5      5        UR)                  XS9  US4$ )	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.
Fr\   rs   r†   NT)r_   r‰   zGHost array used in CUDA kernel will incur copy overhead to/from device.)rB   rŒ   r   r,  r!  Úas_cuda_arrayr!   r,   ÚvoidrŒ  r�   r   rŠ   r‰  r	   ÚCUDA_WARN_ON_IMPLICIT_COPYr¯   r¤   r   r   r”   )r   rM   r_   r[  Údevobjrp   s         r   r  r  U  sõ   € ô ×Ò ×$Ñ$Ø�EˆzÐÜ	�Ð0×	1Ñ	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Ñ7Ü# C¬¯©×4Ñ4ð;�CäÔ0°Ó5Ô6Ø×!Ñ! #Ð!Ñ5Ø�tˆ|Ðr   c                 óì  • U R                  5       UR                  5       p2U R                  UR                  :w  a'  [        SU R                  < SUR                  < 35      eUR                  UR                  :w  a'  [	        SU R                  < SUR                  < 35      eU R
                  (       aB  UR                  UR                  :w  a'  [	        SU R                  < SUR                  < 35      eg g )Nzincompatible dtype: z vs. zincompatible shape: zincompatible strides: )rÆ   r(   Ú	TypeErrorr&   r0   r)   r'   )Úary1Úary2Úary1sqÚary2sqs       r   r�   r�   |  s±   € Ø—\‘\“^ T§\¡\£^ˆ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   rB   Ú
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*   ÚDeviceArrayr2   rà   r.  r¯   Úobjectr¿   r¤   rz  r„  r‰  rŒ  r‹   rY   r•  rŠ   r  r�   r   r   r   Ú<module>r®     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   Á"D Ä	DÄD