+
    :iy                     &   R t ^ RIt^ RIt^ RIt^ RIt^ RIHt ^ RIt^ RI	t	^ RI	H
t
 ^ RIHtHt ^ RIHt ^ RIHtHt ^ RIHt ^ RIHt ^ R	IHt ^ R
IHt ^ RIHt ^ RIHt  ]! ]R4      ! R4      t R t"R t#R t$ ! R R]
PJ                  4      t& ! R R]&4      t'] R 4       t( ! R R]&4      t) ! R R]*4      t+ ! R R]&]PX                  4      t- ! R R]&]PX                  4      t.R'R lt/R'R  lt0R! t1R" t2R#t3R$ t4R(R% lt5R& t6R#   ]! d    R t  Li ; i))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                     V # N )funcs   &\/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/cudadrv/devicearray.pyr   r      s        c                    \        V RR4      # )z$Check if an object is a CUDA ndarray__cuda_ndarray__F)getattrobjs   &r   is_cuda_ndarrayr   #   s    3*E22r   c                   a  \        S 4       V 3R lpV! R\        4       V! R\        4       V! R\        P                  4       V! R\        4       R# )z,Verify the CUDA ndarray interface for an objc                    < \        SV 4      '       g   \        V 4      h\        \        SV 4      V4      '       g   \        V : R V: 24      hR# )z must be of type N)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   f r   verify_cuda_ndarray_interfacer-   (   s?    H '5!)U#'288$&#r   c                >    \        V 4      '       g   \        R4      hR# )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                     a  ] tR t^>t o RtRtRtRR lt]R 4       t	RR lt
]R 4       tRR ltR	 t]R
 4       t]R 4       t]P"                  RR l4       t]P"                  RR l4       tRR ltR tR tRR ltR t]R 4       tRtV tR# )DeviceNDArrayBasez$A on GPU NDArray representation
    TNc                r   \        V\        4      '       d   V3p\        V\        4      '       d   V3p\        P                  ! V4      p\	        V4      V n        \	        V4      V P
                  8w  d   \        R4      h\        P                  P                  ^ WVP                  4      V n        \        V4      V n        \        V4      V n        W0n        \        \        P                   ! \"        P$                  V P                  ^4      4      V n        V P&                  ^ 8  d   Vfu   \(        P*                  ! V P                  V P                  V P                  P                  4      V n        \.        P0                  ! 4       P3                  V P,                  4      pM\(        P4                  ! V4      V n        Mt\(        P6                  '       d!   \(        P8                  P;                  ^ 4      pM\=        ^ 4      p\(        P>                  ! \.        P0                  ! 4       V^ R7      p^ V n        WPn         W@n!        R# )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 ndimN)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'   rJ   rI   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                   \         P                  '       d(   V P                  e   \        V P                  4      pM4^ pM1V P                  P                  e   V P                  P                  pM^ pR\        V P                  4      R\        V 4      '       d   R M\        V P                  4      RVR3RV P                  P                  RV P                  ^ 8w  d   \        V P                  4      R^/# R R^/# )Nr%   r&   dataFtypestrrJ   version)r?   rE   device_ctypes_pointerr,   valuer*   r%   is_contiguousr&   r'   strrJ   )rK   ptrs   & r   __cuda_array_interface__*DeviceNDArrayBase.__cuda_array_interface__w   s    !!!))5$445))//;0066 U4::&}T22tdll8KS%Ltzz~~$++*:c$++&q
 	

 AEq
 	
r   c                >    \         P                   ! V 4      pWn        V# )zoBind a CUDA stream to this object so that all subsequent operation
on this array defaults to the given stream.
)copyrJ   )rK   rJ   clones   && r   bindDeviceNDArrayBase.bind   s     		$r   c                "    V P                  4       # r   	transposerK   s   &r   TDeviceNDArrayBase.T   s    ~~r   c                D   V'       d/   \        V4      \        \        V P                  4      4      8X  d   V # V P                  ^8w  d   Rp\        V4      hVe<   \	        V4      \	        \        V P                  4      4      8w  d   \        RV: 24      h^ RIHp V! V 4      # )   z2transposing a non-2D DeviceNDArray isn't supportedzinvalid axes list r`   )r*   ranger6   NotImplementedErrorsetr/   numba.cuda.kernels.transposera   )rK   axesmsgra   s   &&  r   ra   DeviceNDArrayBase.transpose   sx    E$K5tyy)9#::KYY!^FC%c**#d)s53C/D"Dt=>>>T?"r   c                .    V'       g   V P                   # T# r   rJ   )rK   rJ   s   &&r   _default_stream!DeviceNDArrayBase._default_stream   s    "(t{{4f4r   c                6   ^ V P                   9   pV P                  R,          '       d   V'       g   RpM&V P                  R,          '       d   V'       g   RpMRp\        P                  ! V P                  4      p\
        P                  ! W0P                  V4      # )V
Magic attribute expected by Numba to get the numba type that
represents this object.
C_CONTIGUOUSCF_CONTIGUOUSFA)r&   flagsr   
from_dtyper'   r   r7   r6   )rK   	broadcastlayoutr'   s   &   r   _numba_type_DeviceNDArrayBase._numba_type_   sl    ( %	::n%%iFZZ''	FF((4{{5))V44r   c                    V P                   fA   \        P                  '       d    \        P                  P	                  ^ 4      # \        ^ 4      # V P                   P                  # )z:Returns the ctypes pointer to the GPU data buffer
        )rI   r?   rE   rF   rG   r   rS   rb   s   &r   rS   'DeviceNDArrayBase.device_ctypes_pointer   sI     == %%%22155{"==666r   c                H   VP                   ^ 8X  d   R# \        V 4       V P                  V4      p\        V 4      \        V4      rC\        P
                  ! V4      '       d<   \        V4       \        W44       \        P                  ! WV P                  VR7       R# \        P                  ! TVP                  R,          '       d   RMRR\        R	8  d   VP                  R,          '       * MRR7      p\        W44       \        P                  ! WV P                  VR7       R# )
zCopy `ary` to `self`.

If `ary` is a CUDA memory, perform a device-to-device transfer.
Otherwise, perform a a host-to-device transfer.
Nro   rt   ru   rw   T	WRITEABLE)ordersubokr[   rf       )r(   sentry_contiguousrp   
array_corer?   is_device_memorycheck_array_compatibilitydevice_to_devicerA   r+   arrayry   r
   host_to_device)rK   aryrJ   	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                $   \         ;QJ d&    R V P                   4       F  '       g   K   RM	  RM! R V P                   4       4      '       d'   Rp\        VP                  V P                  4      4      hV P                  ^ 8  g   Q R4       hV P                  V4      pVf2   \        P                  ! V P                  \        P                  R7      pM\        W4       TpV P                  ^ 8w  d$   \        P                  ! W@V P                  VR7       Vfy   V P                  ^ 8X  d0   \        P                  ! V P                  V P                  VR7      pV# \        P                  ! V P                  V P                  V P                  VR	7      pV# )
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   *   "   T F	  q^ 8  x  K  	  R# 5ir   Nr   ).0ss   & r   	<genexpr>1DeviceNDArrayBase.copy_to_host.<locals>.<genexpr>
  s     +l1uls   TFz2D->H copy not implemented for negative strides: {}zNegative memory sizer%   r'   ro   )r%   r'   buffer)r%   r'   r&   r   )anyr&   rh   formatrA   rp   r+   emptybyter   r?   device_to_hostr(   ndarrayr%   r'   )rK   r   rJ   rl   hostarys   &&&  r   copy_to_hostDeviceNDArrayBase.copy_to_host   s"   . 3+dll+333+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           	   #    "   V P                  V4      pV P                  ^8w  d   \        R4      hV P                  ^ ,          V P                  P
                  8w  d   \        R4      h\        \        P                  ! \        V P                  4      V,          4      4      pV P                  pV P                  P
                  p\        V4       Fs  pWa,          p\        Wq,           V P                  4      pW,
          3p	V P                  P                  Wu,          W,          4      p
\        WV P                  VV
R7      x  Ku  	  R# 5i)zSplit the array into equal partition of the `section` size.
If the array cannot be equally divided, the last section will be
smaller.
zonly support 1d arrayzonly support unit strider'   rJ   rI   N)rp   r6   r/   r&   r'   r9   r,   mathceilfloatr(   rg   minrI   viewDeviceNDArray)rK   sectionrJ   nsectr&   r9   ibeginendr%   rI   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   EEc                    V P                   # )zEReturns a device memory object that is used as the argument.
        )rI   rb   s   &r   as_cuda_argDeviceNDArrayBase.as_cuda_arg7  s     }}r   c                    \         P                  ! 4       P                  V P                  4      p\	        V P
                  V P                  V P                  R7      p\        WR7      # )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   rB   get_ipc_handlerI   dictr%   r&   r'   IpcArrayHandle)rK   ipchdescs   &  r   r    DeviceNDArrayBase.get_ipc_handle<  sF     ""$33DMMB$**dll$**M??r   c                    V P                   P                  VR7      w  r4\        VP                  VP                  V P
                  V P                  V4      V P                  R7      # )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'   rJ   rI   )r:   squeezer   r%   r&   r'   rp   rI   )rK   r   rJ   	new_dummy_s   &&&  r   r   DeviceNDArrayBase.squeezeG  sV    & {{***5	//%%**''/]]
 	
r   c                   \         P                  ! V4      p\        V P                  4      p\        V P                  4      pV P                  P
                  VP
                  8w  d   V P                  4       '       g   \        R4      h\        VR,          V P                  P
                  ,          VP
                  4      w  VR&   pV^ 8w  d   \        R4      hVP
                  VR&   \        VVVV P                  V P                  R7      # )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-contiguouszuWhen 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&   r9   is_c_contiguousr/   divmodr   rJ   rI   )rK   r'   r%   r&   rems   &&   r   r   DeviceNDArrayBase.viewc  s     TZZ t||$::%..0'')) 6 
 $b	DJJ///NE"Is
 ax 6   ..GBK;;]]
 	
r   c                P    V P                   P                  V P                  ,          # r   )r'   r9   r(   rb   s   &r   nbytesDeviceNDArrayBase.nbytes  s    
 zz""TYY..r   )	r:   rA   r'   rI   r6   r%   r(   rJ   r&   r   r   r   )Nr   )__name__
__module____qualname____firstlineno____doc____cuda_memory__r   rM   propertyrX   r]   rc   ra   rp   r}   rS   r   require_contextr   r   r   r   r   r   r   r   __static_attributes____classdictcell____classdict__s   @r   r1   r1   >   s     O1f 
 
*    
#5 5 5< 	7 	7 2 2> , ,\3*
	@
8#
J / /r   r1   c                     a a ] tR tRt oRtRV 3R llt]R 4       t]R 4       t]	P                  R 4       t]	P                  RR l4       tRR lt]	P                  R	 4       t]	P                  RR
 l4       tRR ltRtVtV ;t# )DeviceRecordi  z
An on-GPU record type
c                :   < RpRp\         \        V `  WEWV4       R # Nr   )superr   rM   )rK   r'   rJ   rI   r%   r&   	__class__s   &&&&  r   rM   DeviceRecord.__init__  s#    lD*55+3	5r   c                @    \        V P                  P                  4      # 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:   ry   rb   s   &r   ry   DeviceRecord.flags       DKK%%&&r   c                B    \         P                  ! V P                  4      # )rs   )r   rz   r'   rb   s   &r   r}   DeviceRecord._numba_type_  s     ''

33r   c                $    V P                  V4      # r   _do_getitemrK   items   &&r   __getitem__DeviceRecord.__getitem__      %%r   c                $    V P                  W4      # z0Do `__getitem__(item)` with CUDA stream
        r   rK   r   rJ   s   &&&r   getitemDeviceRecord.getitem       --r   c                   V P                  V4      pV P                  P                  V,          w  r4V P                  P	                  V4      pVP
                  R8X  d`   VP                  e   \        W2VR7      # \        P                  ! ^VR7      p\        P                  ! WeVP                  VR7       V^ ,          # \        VP
                  R VP                  ^ ,          R4      w  rxp	\        WxWVR7      # )Nr   r'   dstsrcr(   rJ   ru   r%   r&   r'   rI   rJ   r   )rp   r'   fieldsrI   r   r%   namesr   r+   r   r?   r   r9   r   subdtyper   )
rK   r   rJ   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                $    V P                  W4      # r   _do_setitemrK   keyrT   s   &&&r   __setitem__DeviceRecord.__setitem__      ++r   c                (    V P                  WVR7      # z6Do `__setitem__(key, value)` with CUDA stream
        ro   r  rK   r
  rT   rJ   s   &&&&r   setitemDeviceRecord.setitem       6::r   c                   V P                  V4      pV'       * pV'       d&   \        P                  ! 4       pVP                  4       pV P                  P
                  V,          w  rgV P                  P                  V4      p\        V 4      ! WcVR 7      p	\        V	P                  P                  V4      VR7      w  r\        P                  ! WV
P                  P                  V4       V'       d   VP                  4        R# R# )r   ro   N)rp   r   rB   get_default_streamr'   r   rI   r   typeauto_devicer?   r   r9   synchronize)rK   r
  rT   rJ   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   rM   r   ry   r}   r   r   r   r   r   r  r  r  r   r   __classcell__)r   r   s   @@r   r   r     s     5 ' ' 4 4 & & . .
00 , , ; ;
! !r   r   c                x   a a ^ RI Ho S ^ 8X  d   SP                  R 4       pV# SP                  VV 3R l4       pV# )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
)cudac                     VR,          V R&   R # r   r   )r  r  s   &&r   kernel_assign_kernel.<locals>.kernel  s    "gCGr   c                 t  < SP                  ^4      p^p\        V P                  4       F  pW0P                  V,          ,          pK  	  W#8  d   R# SP                  P                  ^S3\        P                  R7      p\        S^,
          RR4       Fs  pW P                  V,          ,          V^ V3&   W P                  V,          ,          VP                  V,          ^8  ,          V^V3&   W P                  V,          ,          pKu  	  V\        V^,          S4      ,          V \        V^ ,          S4      &   R# )   Nr   r   )	gridrg   r6   r%   localr   r   int64r	   )r  r  location
n_elementsr   idxr   r6   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)r6   r"  r   s   f @r   _assign_kernelr.    sI     qy		 
		XXN N. Mr   c                     a  ] tR tRt o RtR t]R 4       tR tRR lt	R t
R	 tRR
 lt]P                  R 4       t]P                  RR l4       tRR lt]P                  R 4       t]P                  RR l4       tRR ltRtV tR# )r   i#  z
An on-GPU array type
c                .    V P                   P                  # )z1
Return true if the array is Fortran-contiguous.
)r:   is_f_contigrb   s   &r   is_f_contiguousDeviceNDArray.is_f_contiguous'       {{&&&r   c                @    \        V P                  P                  4      # r   r   rb   s   &r   ry   DeviceNDArray.flags-  r   r   c                .    V P                   P                  # )z+
Return true if the array is C-contiguous.
)r:   is_c_contigrb   s   &r   r   DeviceNDArray.is_c_contiguous7  r4  r   Nc                    V'       d    V P                  4       P                  V4      # V P                  4       P                  4       # )z5
:return: an `numpy.ndarray`, so copies to the host.
)r   	__array__)rK   r'   s   &&r   r;  DeviceNDArray.__array__=  s9     $$&0077$$&0022r   c                (    V P                   ^ ,          # r   )r%   rb   s   &r   __len__DeviceNDArray.__len__F  s    zz!}r   c                   \        V4      ^8X  d-   \        V^ ,          \        \        34      '       d
   V^ ,          p\	        V 4      pWP
                  8X  d5   V! V P
                  V P                  V P                  V P                  R7      # V P                  P                  ! V/ VB w  rEWPP                  P                  .8X  d5   V! VP
                  VP                  V P                  V P                  R7      # \        R4      h)z
Reshape the array without changing its contents, similarly to
:meth:`numpy.ndarray.reshape`. Example::

    d_arr = d_arr.reshape(20, 50, order='F')
)r%   r&   r'   rI   operation requires copying)r5   r    r*   r   r  r%   r&   r'   rI   r:   reshapeextentrh   )rK   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                ,   V P                  V4      p\        V 4      pV P                  P                  VR7      w  rEWPP                  P                  .8X  d6   V! VP
                  VP                  V P                  V P                  VR7      # \        R4      h)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  )
rp   r  r:   ravelrC  r%   r&   r'   rI   rh   )rK   r   rJ   rF  rG  rH  s   &&&   r   rK  DeviceNDArray.ravela  s     %%f-4j++++%+8{{))**V\\6>>!ZZ$--$& &
 &&BCCr   c                $    V P                  V4      # r   r   r   s   &&r   r   DeviceNDArray.__getitem__s  r   r   c                $    V P                  W4      # r   r   r   s   &&&r   r   DeviceNDArray.getitemw  r   r   c                (   V P                  V4      pV P                  P                  V4      p\        VP	                  4       4      p\        V 4      p\        V4      ^8X  d   V P                  P                  ! V^ ,          !  pVP                  '       g   V P                  P                  e   \        V P                  VVR7      # \        P                  ! ^V P                  R7      p\        P                   ! WvV P                  P"                  VR7       V^ ,          # V! VP$                  VP&                  V P                  WbR7      # V P                  P                  ! VP(                  !  pV! VP$                  VP&                  V P                  WbR7      # )r%  r   r   r   r   )rp   r:   r   r   iter_contiguous_extentr  r5   rI   r   is_arrayr'   r  r   r+   r   r?   r   r9   r%   r&   rC  )rK   r   rJ   arrrH  rF  r  r   s   &&&     r   r   DeviceNDArray._do_getitem}  s;   %%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                $    V P                  W4      # r   r  r	  s   &&&r   r  DeviceNDArray.__setitem__  r  r   c                (    V P                  WVR7      # r  r  r  s   &&&&r   r  DeviceNDArray.setitem  r  r   c                   V P                  V4      pV'       * pV'       d&   \        P                  ! 4       pVP                  4       pV P                  P                  V4      pV P                  P                  ! VP                  !  p\        V\        P                  4      '       d   R
pR
p	MVP                  pVP                  p	\        V 4      ! VV	V P                  VVR 7      p
\!        W#RR7      w  rVP"                  V
P"                  8  d)   \%        RVP"                  : RV
P"                  : R24      h\&        P(                  ! V
P"                  \&        P*                  R7      pVP                  WP"                  VP"                  ,
          R% VP,                  ! V!  p\/        \1        V
P                  VP                  4      4       F-  w  pw  ppV^8w  g   K  VV8w  g   K  \%        RVW3,          4      h	  \2        P4                  ! \6        P8                  V
P                  ^4      p\;        V
P"                  4      P=                  VVR	7      ! W4       V'       d   VP?                  4        R# R# )r   T)rJ   user_explicitzCan't assign z-D array to z-D selfr   NzCCan't copy sequence with size %d to array axis %d with dimension %dro   r   ) rp   r   rB   r  r:   r   rI   r   rC  r    r   Elementr%   r&   r  r'   r  r6   r/   r+   onesr(  rB  	enumeratezipr;   r<   r=   r>   r.  forallr  )rK   r
  rT   rJ   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   )ru   r   r   )r   r   r   r   r   r2  r   ry   r   r;  r>  rB  rK  r   r   r   r   r   r  r  r  r   r   r   s   @r   r   r   #  s     ' ' ''3D0D$ & & . .
J: , , ; ;
5! 5!r   r   c                   B   a  ] tR tRt o RtR tR tR tR tR t	Rt
V tR	# )
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                    W n         Wn        R # r   _array_desc_ipc_handle)rK   r   r   s   &&&r   rM   IpcArrayHandle.__init__  s    %%r   c                    V P                   P                  \        P                  ! 4       4      p\	        RRV/V P
                  B # )z
Returns a new *DeviceNDArray* that shares the allocation from the
original process.  Must not be used on the original process.
rI   r   )ri  openr   rB   r   rh  )rK   dptrs   & r   rl  IpcArrayHandle.open  s<    
 $$W%8%8%:;?d?d.>.>??r   c                :    V P                   P                  4        R# )z%
Closes the IPC handle to the array.
N)ri  closerb   s   &r   rp  IpcArrayHandle.close  s     	 r   c                "    V P                  4       # r   )rl  rb   s   &r   	__enter__IpcArrayHandle.__enter__  s    yy{r   c                &    V P                  4        R # r   )rp  )rK   r  rT   	tracebacks   &&&&r   __exit__IpcArrayHandle.__exit__  s    

r   rg  N)r   r   r   r   r   rM   rl  rp  rs  rw  r   r   r   s   @r   r   r     s)     $&@! r   r   c                   .   a  ] tR tRt o RtRR ltRtV tR# )MappedNDArrayi  z,
A host array that uses CUDA mapped memory.
c                    Wn         W n        R # r   rI   rJ   rK   rI   rJ   s   &&&r   device_setupMappedNDArray.device_setup       r   r|  Nr   r   r   r   r   r   r~  r   r   r   s   @r   rz  rz          r   rz  c                   .   a  ] tR tRt o RtRR ltRtV tR# )ManagedNDArrayi  z-
A host array that uses CUDA managed memory.
c                    Wn         W n        R # r   r|  r}  s   &&&r   r~  ManagedNDArray.device_setup  r  r   r|  Nr   r  r   s   @r   r  r    r  r   r  c                ^    \        V P                  V P                  V P                  VVR7      # )z/Create a DeviceNDArray object that is like ary.rJ   rI   )r   r%   r&   r'   )r   rJ   rI   s   &&&r   from_array_liker    s&    CKK6"*, ,r   c                0    \        V P                  WR7      # )z.Create a DeviceRecord object that is like rec.r  )r   r'   )recrJ   rI   s   &&&r   from_record_liker  "  s    		&DDr   c                    V P                   '       d   V P                  '       g   V # . pV P                    F&  pTP                  V^ 8X  d   ^ M
\        R4      4       K(  	  V \	        V4      ,          # )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.

N)r&   r(   appendslicer*   )r   
core_indexstrides   &  r   r   r   '  sW     ;;;chhh
J++v{!d< uZ !!r   c                    V P                   P                  p\        \        V P                  4      \        V P
                  4      4       F(  w  r#V^8  g   K  V^ 8w  g   K  W8w  d    R# W,          pK*  	  R# )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.
FT)r'   r9   r_  reversedr%   r&   )r   r(   r%   r  s   &   r   rU   rU   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                     \        V 4      pVP                  R ,          '       g+   VP                  R,          '       g   \        \        4      hR# R# )rt   rv   N)r   ry   r/   errmsg_contiguous_buffer)r   cores   & r   r   r   O  s=    c?D::n%%djj.H.H122 /I%r   c                   \         P                  ! V 4      '       d   V R3# \        V R4      '       d"   \        P                  P                  V 4      R3# \        V \        P                  4      '       d   \        WR7      pM=\        P                  ! T \        R8  d   RMRRR7      p \        V 4       \        WR7      pV'       d}   \        P                  '       dU   V'       gM   \        V \         4      '       g7   \        V \        P"                  4      '       d   Rp\%        \'        V4      4       VP)                  WR7       VR3# )	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.
FrX   ro   NT)r[   r   zGHost array used in CUDA kernel will incur copy overhead to/from device.r   )r?   r   r   r,  r   as_cuda_arrayr    r+   voidr  r   r
   r   r  r   CUDA_WARN_ON_IMPLICIT_COPYr   r   r   r   r   )r   rJ   r[   r[  devobjrl   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0%#C77#C44;C056!!#!5t|r   c                    V P                  4       VP                  4       r2V P                  VP                  8w  d(   \        R V P                  : RVP                  : 24      hVP                  VP                  8w  d(   \	        RV P                  : RVP                  : 24      hV P
                  '       dE   VP                  VP                  8w  d(   \	        RV P                  : RVP                  : 24      hR# R# )zincompatible dtype: z vs. zincompatible shape: zincompatible strides: N)r   r'   	TypeErrorr%   r/   r(   r&   )ary1ary2ary1sqary2sqs   &&  r   r   r   |  s    \\^T\\^FzzTZZTZZ1 2 	2||v||#**djj2 3 	3 yyyV^^v~~5,,6 7 	7 6yr   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   rz  r  r  r  r   rU   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