
    shy                     "   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  Nf = 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    r/Users/tiagomarins/Projetos/claudeai/copy_bank/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_attr4verify_cuda_ndarray_interface.<locals>.requires_attr,   sD    sD!! &&'#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#'288$&#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     eS!!HEgs##jGJ	w<499$566 &&00E16A5\
W~
	((tzz1EF	99q=")"?"?JJdjj.A.A#C"..099$//J")"<"<X"F %%2215{,,W5H5H5J59CHDO 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2tdll8K%Lzz~~*.++*: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$K5tyy)9#::KYY!^FC%c**#d)s5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{{4f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ZZ'	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21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)zCopy `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     88q=$%%f-(.
38##C((c"%i:$$TO xx&__^<c# 6) #..55/35H &i:""44??*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1uls   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   . +dll+++FC%cjj&>??!#;%;;#%%f-;hhT__BGGDG%d0G??a""7$//*02 ;yyA~**4::TZZ,35
  **4::TZZ-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
[        XU R                  UU
S9v   M^     g7f)zSplit 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-99>455<<?djj111788DIIeDII.89:,,::&&uAKEeotyy1C[NE}}))%*:CNKHdjj)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3DMMB$**dll$**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-contiguousr   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     TZZ t||$::%..0'')) 6 
 $b	DJJ///NE"Is
 ax 6   ..GBK;;]]
 	
r   c                 H    U R                   R                  U R                  -  $ r   )r(   r<   r)   rf   s    r   nbytesDeviceNDArrayBase.nbytes  s    
 zz""TYY..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#    lD*55+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       DKK%%&&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	[        XxXU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-jj''---$$V,99?yy$##-46 6 ((1C0&&7,/LL.46 1: ,CII,0,/LLOSB "EE !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                  " XU
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 jj'',--$$V,4jsGD SYY^^E26B 	  399+=+=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 & & . .
00 , , ; ;
! !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    "gC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    99Q<
sxxA))A,&J !!  jjd)++   taxR(A 99Q</C1I!IIaL0SYYq\A5EFC1I1%H )
 -0s1vt0L,MN3q64()r   )numbar!  jit)r9   r#  r!  s   ` @r   _assign_kernelr.    sI     qy		 
		XX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    zz!}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*Xa[5$-"H"H{H4jzz!TZZ!ZZ$--A A ++--x?3?{{))**V\\6>>!ZZ$--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-4j++++%+8{{))**V\\6>>!ZZ$--$& &
 &&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-kk%%d+s11344jw<1mm(('!*5G<<::##/'djj18: : !hhq

;G**w040D0D28: qz!CKK!%gN N mm((#**5GSYY!ZZ'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  pUR"                  U
R"                  :  a(  [%        SUR"                  < SU
R"                  < S35      e[&        R(                  " U
R"                  [&        R*                  S9nUR                  XR"                  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X4-  5      e   [2        R4                  " [6        R8                  U
R                  S	5      n[;        U
R"                  5      R=                  UUS9" X5        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 kk%%c*--$$cjj1c:--..EGIIEkkG4j** UF88chh   GGCHHBHH5	*-))	((SXX%&'kk9%"3syy#))#<=IAv1Av!q&  "=ABAz"J K K > %%hllCIIqA
sxx ''
6'B3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    yy{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&    CKK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     ;;chh
J++v{!d< uZ !!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     99DXcii0(3;;2GH191~MD	 I
 r   zArray 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%djj.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     $$Ez	0	1	1zz'',e33c277##%c9F ((+f4U$C c"$S8F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zzTZZTZZ1 2 	2||v||#**djj2 3 	3 yyV^^v~~5,,6 7 	7 6yr   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 	;/5I3
 ;O/00 O/d
d!$ d!N ( (Vv!% v!r)V )X%rzz &

 ,E
"$ 3 3$N7  s   "D 	DD