+
    :i"                         ^ RI Ht ^ RIHt ^ RIHt ^ RIHtH	t	H
t
  ! R R]4      t ! R R]
4      t ! R	 R
]	4      t ! R R]4      tRt ! R R]P"                  4      tRt ! R R]P(                  4      tR# )    )cuda)array)deviceufunc)UFuncMechanismGeneralizedUFuncGUFuncCallStepsc                   @   a  ] tR t^t o RtR tR tRR ltR tRt	V t
R# )	CUDAUFuncDispatcherz<
Invoke the CUDA ufunc specialization for the given inputs.
c                4    Wn         VP                  V n        R # N)	functions__name__)selftypes_to_retty_kernelspyfuncs   &&&T/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/vectorizers.py__init__CUDAUFuncDispatcher.__init__   s    /    c                B    \         P                  V P                  W4      # )af  
*args: numpy arrays or DeviceArrayBase (created by cuda.to_device).
       Cannot mix the two types in one call.

**kws:
    stream -- cuda stream; when defined, asynchronous mode is used.
    out    -- output array. Can be a numpy array or DeviceArrayBase
              depending on the input arguments.  Type must match
              the input arguments.
)CUDAUFuncMechanismcallr   )r   argskwss   &*,r   __call__CUDAUFuncDispatcher.__call__   s     "&&t~~tAAr   c                   \        \        V P                  P                  4       4      ^ ,          4      ^8X  g   Q R4       hVP                  ^8X  g   Q R4       hVP
                  ^ ,          p. pV^ 8X  d   \        R4      hV^8X  d
   V^ ,          # T;'       g    \        P                  ! 4       pVP                  4       ;_uu_ 4        \        P                  P                  P                  V4      '       d   TpM\        P                  ! W4      pV P                  WTV4      p\        RVP                   R7      pVP#                  WrR7       RRR4       V^ ,          #   + '       g   i     X^ ,          # ; i)r   zmust be a binary ufunczmust use 1d arrayzReduction on an empty array.)dtypestreamN)   )lenlistr   keysndimshape	TypeErrorr   r    auto_synchronizecudadrvdevicearrayis_cuda_ndarray	to_device_CUDAUFuncDispatcher__reducenp_arrayr   copy_to_host)r   argr    ngpu_memsmemoutbufs   &&&     r   reduceCUDAUFuncDispatcher.reduce   s-   4++-.q12a7 	A :A 	A7xx1}111}IIaL6:;;!Vq6M ((4;;=$$&&||''77<<nnS1--v6C4syy1CS0 ' 1v '& 1vs   BEE3	c                   VP                   ^ ,          pV^,          ^ 8w  dj   VP                  V^,
          4      w  rVVP                  V4       VP                  V4       V P                  WRV4      pVP                  V4       V ! WvWsR7      # VP                  V^,          4      w  rVP                  V4       VP                  V	4       V ! WWR7       V^,          ^8  d   V P                  WV4      # V# )r   )r4   r    )r&   splitappendr-   )
r   r3   r2   r    r1   fatcutthincutr4   leftrights
   &&&&      r   __reduceCUDAUFuncDispatcher.__reduce;   s    IIaLq5A:!iiA.OFOOF#OOG$--&9COOC #==))AF+KDOOD!OOE"$6Avz}}TV<<r   )r   r   Nr   )r   
__module____qualname____firstlineno____doc__r   r   r6   r-   __static_attributes____classdictcell____classdict__s   @r   r
   r
      s$     (B: r   r
   c                   \   a a ] tR t^St oR.tV 3R ltR tR tR tR t	R t
R tR	tVtV ;t# )
_CUDAGUFuncCallSteps_streamc                V   < \         SV `  WW44       VP                  R ^ 4      V n        R# )r    N)superr   getrL   )r   ninnoutr   kwargs	__class__s   &&&&&r   r   _CUDAGUFuncCallSteps.__init__X   s$    D1zz(A.r   c                .    \         P                  ! V4      # r   r   is_cuda_arrayr   objs   &&r   is_device_array$_CUDAGUFuncCallSteps.is_device_array\       !!#&&r   c                    \         P                  P                  P                  V4      '       d   V# \         P                  ! V4      # r   r   r)   r*   r+   as_cuda_arrayrX   s   &&r   as_device_array$_CUDAGUFuncCallSteps.as_device_array_   5     <<##33C88J!!#&&r   c                D    \         P                  ! WP                  R 7      # r   )r   r,   rL   )r   hostarys   &&r   r,   _CUDAGUFuncCallSteps.to_devicei   s    ~~gll;;r   c                >    VP                  W P                  R 7      pV# rd   )r/   rL   )r   devaryre   r4   s   &&& r   to_host_CUDAGUFuncCallSteps.to_hostl   s    !!',,!?
r   c                F    \         P                  ! WV P                  R 7      # )r&   r   r    )r   device_arrayrL   )r   r&   r   s   &&&r   allocate_device_array*_CUDAGUFuncCallSteps.allocate_device_arrayp   s      u$,,OOr   c                F    VP                  W P                  R 7      ! V!   R# r   N)forallrL   )r   kernelnelemr   s   &&&&r   launch_kernel"_CUDAGUFuncCallSteps.launch_kernels   s    eLL148r   )rL   )r   rB   rC   rD   	__slots__r   rZ   r`   r,   ri   rn   ru   rF   rG   __classcell__rS   rI   s   @@r   rK   rK   S   s8     I/''<P9 9r   rK   c                   N   a a ] tR t^wt oV 3R lt]R 4       tR tR tRt	Vt
V ;t# )CUDAGeneralizedUFuncc                H   < VP                   V n         \        SV `	  W4       R # r   )r   rN   r   )r   	kernelmapenginer   rS   s   &&&&r   r   CUDAGeneralizedUFunc.__init__x   s    +r   c                    \         # r   )rK   r   s   &r   _call_steps CUDAGeneralizedUFunc._call_steps|   s    ##r   c                    \         P                  P                  P                  VRVP                  VP
                  R7      # r   r&   stridesr   gpu_datarA   )r   r)   r*   DeviceNDArrayr   r   )r   aryr&   s   &&&r   _broadcast_scalar_input,CUDAGeneralizedUFunc._broadcast_scalar_input   s9    ||''55E>B<?II?B|| 6 M 	Mr   c                   \        V4      \        VP                  4      ,
          pRV,          VP                  ,           p\        P                  P
                  P                  VVVP                  VP                  R7      # r   )	r"   r&   r   r   r)   r*   r   r   r   )r   r   newshapenewax
newstridess   &&&  r   _broadcast_add_axis(CUDAGeneralizedUFunc._broadcast_add_axis   sa    HCII.E\CKK/
||''55H>H<?II?B|| 6 M 	Mr   )r   )r   rB   rC   rD   r   propertyr   r   r   rF   rG   rx   ry   s   @@r   r{   r{   w   s0     , $ $MM Mr   r{   c                   R   a  ] tR t^t o Rt^ tR tR tR tR t	R t
R tR tR	tV tR
# )r   z
Provide CUDA specialization
c                2    VP                  W#R 7      ! V!   R# rq   )rr   )r   funccountr    r   s   &&&&&r   launchCUDAUFuncMechanism.launch   s    E)40r   c                .    \         P                  ! V4      # r   rV   rX   s   &&r   rZ   "CUDAUFuncMechanism.is_device_array   r\   r   c                    \         P                  P                  P                  V4      '       d   V# \         P                  ! V4      # r   r^   rX   s   &&r   r`   "CUDAUFuncMechanism.as_device_array   rb   r   c                0    \         P                  ! WR 7      # rd   )r   r,   )r   re   r    s   &&&r   r,   CUDAUFuncMechanism.to_device   s    ~~g55r   c                &    VP                  VR 7      # rd   )r/   )r   rh   r    s   &&&r   ri   CUDAUFuncMechanism.to_host   s    ""&"11r   c                2    \         P                  ! WVR 7      # rl   )r   rm   )r   r&   r   r    s   &&&&r   rn   (CUDAUFuncMechanism.allocate_device_array   s      u&IIr   c                   \        \        V4      4       Uu. uF5  pW1P                  8  g!   VP                  V,          W#,          8w  g   K3  VNK7  	  pp\        V4      \        VP                  4      ,
          p^ .V,          \	        VP
                  4      ,           pV F  p^ Wc&   K	  	  \        P                  P                  P                  VVVP                  VP                  R7      # u upi )r   r   )ranger"   r%   r&   r#   r   r   r)   r*   r   r   r   )r   r   r&   ax
ax_differs
missingdimr   s   &&&    r   broadcast_device#CUDAUFuncMechanism.broadcast_device   s    #(U#4 5#4Rxx2%)3 b#4
 5 Z#cii.0
#
"T#++%66BGK  ||''55E>E<?II?B|| 6 M 	M5s   0C+C+ N)r   rB   rC   rD   rE   DEFAULT_STREAMr   rZ   r`   r,   ri   rn   r   rF   rG   rH   s   @r   r   r      s<      N1''62JM Mr   r   z
def __vectorized_{name}({args}, __out__):
    __tid__ = __cuda__.grid(1)
    if __tid__ < __out__.shape[0]:
        __out__[__tid__] = __core__({argitems})
c                   H   a  ] tR t^t o R tR tR tR t]R 4       t	Rt
V tR# )CUDAVectorizec                    \         P                  ! VR R R7      ! V P                  4      pW"P                  VP                  ,          P
                  P                  3# )T)deviceinline)r   jitr   	overloadsr   	signaturereturn_type)r   sigcudevfns   && r   _compile_coreCUDAVectorize._compile_core   sA    ((3tD9$++F))#((3==IIIIr   c                    V P                   P                  P                  4       pVP                  R \        RV/4       V# )__cuda____core__)r   __globals__copyupdater   )r   corefnglbls   && r   _get_globalsCUDAVectorize._get_globals   s9    {{&&++-Z) 	*r   c                .    \         P                  ! V4      # r   r   r   r   fnobjr   s   &&&r   _compile_kernelCUDAVectorize._compile_kernel   s    xxr   c                B    \        V P                  V P                  4      # r   )r
   r}   r   r   s   &r   build_ufuncCUDAVectorize.build_ufunc   s    "4>>4;;??r   c                    \         # r   )vectorizer_stager_sourcer   s   &r   _kernel_templateCUDAVectorize._kernel_template   s    ''r   r   N)r   rB   rC   rD   r   r   r   r   r   r   rF   rG   rH   s   @r   r   r      s/     J@ ( (r   r   zy
def __gufunc_{name}({args}):
    __tid__ = __cuda__.grid(1)
    if __tid__ < {checkedarg}:
        __core__({argitems})
c                   B   a  ] tR t^t o R tR t]R 4       tR tRt	V t
R# )CUDAGUFuncVectorizec                    \         P                  ! V P                  V P                  4      p\	        V P
                  VV P                  R 7      # ))r}   r~   r   )r   GUFuncEngineinputsig	outputsigr{   r}   r   )r   r~   s   & r   r   CUDAGUFuncVectorize.build_ufunc   s9    ))$--H#dnn+1+/;;8 	8r   c                :    \         P                  ! V4      ! V4      # r   r   r   s   &&&r   r   #CUDAGUFuncVectorize._compile_kernel   s    xx}U##r   c                    \         # r   )_gufunc_stager_sourcer   s   &r   r   $CUDAGUFuncVectorize._kernel_template   s    $$r   c                    \         P                  ! VR R7      ! V P                  4      pV P                  P                  P                  4       pVP                  R\         RV/4       V# )T)r   r   r   )r   r   r   py_funcr   r   r   )r   r   r   glblss   &&  r   r    CUDAGUFuncVectorize._get_globals   sR    #d+DKK8((--/j$ &* 	+r   r   N)r   rB   rC   rD   r   r   r   r   r   rF   rG   rH   s   @r   r   r      s-     8$ % % r   r   N)numbar   numpyr   r.   
numba.cudar   numba.cuda.deviceufuncr   r   r   objectr
   rK   r{   r   r   DeviceVectorizer   r   DeviceGUFuncVectorizer   r   r   r   <module>r      s     # "5 5H& HV!9? !9HM+ M2-M -M` (K// (2 +;; r   