+
    :im0                        R t ^ RIHt ^ RIt^ RIt^ RIt^ RIHt ^ RIt	^ RI
Ht ^RIHt  ! R R]4      t ! R R	4      t ! R
 R4      t ! R R]4      t ! R R]4      t ! R R]4      t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t]P(                  ! 4       t  ! R R]4      t! ! R R]4      t" ! R R]4      t#]R 4       t$R# )zf
Implements the cuda module as called from within an executing kernel
(@cuda.jit-decorated function).
)contextmanagerN)types)numpy_support)vector_typesc                   <   a  ] tR t^t o RtR tR tR tR tRt	V t
R# )Dim3z3
Used to implement thread/block indices/dimensions
c                *    Wn         W n        W0n        R # Nxyz)selfr   r   r   s   &&&&\/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/simulator/kernelapi.py__init__Dim3.__init__   s        c                \    R V P                   : RV P                  : RV P                  : R2# )(, )r
   r   s   &r   __str__Dim3.__str__   s    !%88r   c                \    R V P                   : RV P                  : RV P                  : R2# )zDim3(r   r   r
   r   s   &r   __repr__Dim3.__repr__   s    %)VVTVVTVV<<r   c              #  b   "   V P                   x  V P                  x  V P                  x  R # 5ir	   r
   r   s   &r   __iter__Dim3.__iter__!   s      ffffffs   -/r
   N)__name__
__module____qualname____firstlineno____doc__r   r   r   r   __static_attributes____classdictcell____classdict__s   @r   r   r      s#     
9= r   r   c                   *   a  ] tR t^'t o RtR tRtV tR# )	GridGroupz#
Used to implement the grid group.
c                L    \         P                  ! 4       P                  4        R # r	   	threadingcurrent_threadsyncthreadsr   s   &r   syncGridGroup.sync,   s     	  "..0r    N)r    r!   r"   r#   r$   r0   r%   r&   r'   s   @r   r*   r*   '   s     1 1r   r*   c                   *   a  ] tR t^3t o RtR tRtV tR# )
FakeCUDACgz
CUDA Cooperative Groups
c                    \        4       # r	   )r*   r   s   &r   	this_gridFakeCUDACg.this_grid7   s
    {r   r2   N)r    r!   r"   r#   r$   r6   r%   r&   r'   s   @r   r4   r4   3   s      r   r4   c                   *   a  ] tR t^;t o RtR tRtV tR# )FakeCUDALocalz
CUDA Local arrays
c                    \        V\        P                  4      '       d   \        P                  ! V4      p\
        P                  ! W4      # r	   )
isinstancer   Typer   as_dtypenpempty)r   shapedtypes   &&&r   arrayFakeCUDALocal.array?   s2    eUZZ((!**51Exx%%r   r2   N)r    r!   r"   r#   r$   rB   r%   r&   r'   s   @r   r9   r9   ;   s     & &r   r9   c                   *   a  ] tR t^Et o RtR tRtV tR# )FakeCUDAConstz
CUDA Const arrays
c                    V# r	   r2   )r   arys   &&r   
array_likeFakeCUDAConst.array_likeI   s    
r   r2   N)r    r!   r"   r#   r$   rH   r%   r&   r'   s   @r   rE   rE   E   s      r   rE   c                   0   a  ] tR t^Mt o RtR tR tRtV tR# )FakeCUDAShareda  
CUDA Shared arrays.

Limitations: assumes that only one call to cuda.shared.array is on a line,
and that that line is only executed once per thread. i.e.::

    a = cuda.shared.array(...); b = cuda.shared.array(...)

will erroneously alias a and b, and::

    for i in range(10):
        sharedarrs[i] = cuda.shared.array(...)

will alias all arrays created at that point (though it is not certain that
this would be supported by Numba anyway).
c                v    / V n         Wn        \        P                  ! V\        P                  R 7      V n        R# )rA   N)_allocations_dynshared_sizer>   zerosbyte
_dynshared)r   dynshared_sizes   &&r   r   FakeCUDAShared.__init___   s&    -((>Ar   c                
   \        V\        P                  4      '       d   \        P                  ! V4      pV^ 8X  dJ   V P
                  VP                  ,          p\        P                  ! V P                  P                  W#R7      # \        P                  ! \        P                  ! 4       4      pVR,          R,          pV P                  P!                  V4      pVf%   \        P"                  ! W4      pW`P                  V&   V# )    )rA   count:rV      N)r;   r   r<   r   r=   rO   itemsizer>   
frombufferrR   data	tracebackextract_stacksys	_getframerN   getr?   )r   r@   rA   rW   stackcallerress   &&&    r   rB   FakeCUDAShared.arrayd   s    eUZZ((!**51E A: ((ENN:E==!5!5UPP
 ''8r3##F+;((5(C(+f%
r   )rN   rR   rO   N)	r    r!   r"   r#   r$   r   rB   r%   r&   r'   s   @r   rK   rK   M   s     "B
 r   rK   c                   t   a  ] tR t^t o R tR tR tR tR tR t	R t
R tR	 tR
 tR tR tR tR tRtV tR# )FakeCUDAAtomicc                    \         ;_uu_ 4        W,          pW;;,          V,          uu&   R R R 4       V#   + '       g   i     X# ; ir	   )addlockr   rB   indexvalolds   &&&& r   addFakeCUDAAtomic.add   5    W,CLCL  
 W 
	   6A	c                    \         ;_uu_ 4        W,          pW;;,          V,          uu&   R R R 4       V#   + '       g   i     X# ; ir	   )sublockrj   s   &&&& r   subFakeCUDAAtomic.sub   rp   rq   c                    \         ;_uu_ 4        W,          pW;;,          V,          uu&   R R R 4       V#   + '       g   i     X# ; ir	   )andlockrj   s   &&&& r   and_FakeCUDAAtomic.and_   rp   rq   c                    \         ;_uu_ 4        W,          pW;;,          V,          uu&   R R R 4       V#   + '       g   i     X# ; ir	   )orlockrj   s   &&&& r   or_FakeCUDAAtomic.or_   s5    V,CLCL  
 V 
rq   c                    \         ;_uu_ 4        W,          pW;;,          V,          uu&   R R R 4       V#   + '       g   i     X# ; ir	   )xorlockrj   s   &&&& r   xorFakeCUDAAtomic.xor   rp   rq   c                    \         ;_uu_ 4        W,          pWC8  d   ^ W&   MW;;,          ^,          uu&   RRR4       V#   + '       g   i     X# ; irV   N)inclockrj   s   &&&& r   incFakeCUDAAtomic.inc   sA    W,Cz !  
 W 
s   (AA	c                    \         ;_uu_ 4        W,          pV^ 8X  g   WC8  d   W1V&   MW;;,          ^,          uu&   RRR4       V#   + '       g   i     X# ; ir   )declockrj   s   &&&& r   decFakeCUDAAtomic.dec   sG    W,Cqci"e!  
 W 
s   /AA	c                t    \         ;_uu_ 4        W,          pW1V&   R R R 4       V#   + '       g   i     X# ; ir	   )exchlockrj   s   &&&& r   exchFakeCUDAAtomic.exch   s0    X,C%L  
 X 
s   &7	c                    \         ;_uu_ 4        W,          p\        WC4      W&   R R R 4       V#   + '       g   i     X# ; ir	   )maxlockmaxrj   s   &&&& r   r   FakeCUDAAtomic.max   4    W,Cs=EL  
 W 
	   /A 	c                    \         ;_uu_ 4        W,          p\        WC4      W&   R R R 4       V#   + '       g   i     X# ; ir	   )minlockminrj   s   &&&& r   r   FakeCUDAAtomic.min   r   r   c                    \         ;_uu_ 4        W,          p\        P                  ! W,          V.4      W&   R R R 4       V#   + '       g   i     X# ; ir	   )r   r>   nanmaxrj   s   &&&& r   r   FakeCUDAAtomic.nanmax   @    W,C99elC%89EL  
 W 
   )AA	c                    \         ;_uu_ 4        W,          p\        P                  ! W,          V.4      W&   R R R 4       V#   + '       g   i     X# ; ir	   )r   r>   nanminrj   s   &&&& r   r   FakeCUDAAtomic.nanmin   r   r   c                    \         ;_uu_ 4        RVP                  ,          pW,          pWR8X  d   W1V&   VuuRRR4       #   + '       g   i     R# ; i)rV   N)rV   )compare_and_swaplockndim)r   rB   rm   rl   rk   loadeds   &&&&  r   compare_and_swapFakeCUDAAtomic.compare_and_swap   s;    !!5::%E\F}"e "!!!s   'AA	c                    \         ;_uu_ 4        W,          pWS8X  d   WAV&   VuuR R R 4       #   + '       g   i     R # ; ir	   )caslock)r   rB   rk   rm   rl   r   s   &&&&& r   casFakeCUDAAtomic.cas   s)    W\F}"e	 WWWs   .?	r2   N)r    r!   r"   r#   rn   rt   rx   r|   r   r   r   r   r   r   r   r   r   r   r%   r&   r'   s   @r   rg   rg      sP      r   rg   c                      a  ] tR t^t o R tR tR tR tR tR t	R t
R tR	 tR
 tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR t R t!Rt"V t#R # )!FakeCUDAFp16c                    W,           # r	   r2   r   abs   &&&r   haddFakeCUDAFp16.hadd   	    ur   c                    W,
          # r	   r2   r   s   &&&r   hsubFakeCUDAFp16.hsub   r   r   c                    W,          # r	   r2   r   s   &&&r   hmulFakeCUDAFp16.hmul   r   r   c                    W,          # r	   r2   r   s   &&&r   hdivFakeCUDAFp16.hdiv   r   r   c                     W,          V,           # r	   r2   r   r   r   cs   &&&&r   hfmaFakeCUDAFp16.hfma       uqyr   c                    V) # r	   r2   r   r   s   &&r   hnegFakeCUDAFp16.hneg   s	    r	r   c                    \        V4      # r	   )absr   s   &&r   habsFakeCUDAFp16.habs   s    1vr   c                N    \         P                  ! V\         P                  R 7      # rM   )r>   sinfloat16r   r   s   &&r   hsinFakeCUDAFp16.hsin       vvarzz**r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   cosr   r   s   &&r   hcosFakeCUDAFp16.hcos  r   r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   logr   r   s   &&r   hlogFakeCUDAFp16.hlog  r   r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   log2r   r   s   &&r   hlog2FakeCUDAFp16.hlog2      wwq

++r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   log10r   r   s   &&r   hlog10FakeCUDAFp16.hlog10      xx,,r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   expr   r   s   &&r   hexpFakeCUDAFp16.hexp  r   r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   exp2r   r   s   &&r   hexp2FakeCUDAFp16.hexp2  r   r   c                <    \         P                  ! ^
V,          4      # )
   r>   r   r   s   &&r   hexp10FakeCUDAFp16.hexp10  s    zz"'""r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   sqrtr   r   s   &&r   hsqrtFakeCUDAFp16.hsqrt  r   r   c                <    \         P                  ! VR,          4      # )g      ?g      r   r   s   &&r   hrsqrtFakeCUDAFp16.hrsqrt  s    zz!t)$$r   c                N    \         P                  ! V\         P                  R 7      # r   r>   ceilr   r   s   &&r   hceilFakeCUDAFp16.hceil  r   r   c                N    \         P                  ! V\         P                  R 7      # r   r   r   s   &&r   hfloorFakeCUDAFp16.hfloor   r   r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   
reciprocalr   r   s   &&r   hrcpFakeCUDAFp16.hrcp#  s    }}Qbjj11r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   truncr   r   s   &&r   htruncFakeCUDAFp16.htrunc&  r   r   c                N    \         P                  ! V\         P                  R 7      # r   )r>   rintr   r   s   &&r   hrintFakeCUDAFp16.hrint)  r   r   c                
    W8H  # r	   r2   r   s   &&&r   heqFakeCUDAFp16.heq,  	    vr   c                
    W8g  # r	   r2   r   s   &&&r   hneFakeCUDAFp16.hne/  r  r   c                
    W8  # r	   r2   r   s   &&&r   hgeFakeCUDAFp16.hge2  r  r   c                
    W8  # r	   r2   r   s   &&&r   hgtFakeCUDAFp16.hgt5  	    ur   c                
    W8*  # r	   r2   r   s   &&&r   hleFakeCUDAFp16.hle8  r  r   c                
    W8  # r	   r2   r   s   &&&r   hltFakeCUDAFp16.hlt;  r  r   c                    \        W4      # r	   )r   r   s   &&&r   hmaxFakeCUDAFp16.hmax>      1yr   c                    \        W4      # r	   )r   r   s   &&&r   hminFakeCUDAFp16.hminA  r$  r   r2   N)$r    r!   r"   r#   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r  r  r  r  r  r  r  r  r"  r&  r%   r&   r'   s   @r   r   r      s     +++,-+,#,%,,2-, r   r   c                   *  a  ] tR tRt o RtR t]R 4       t]R 4       t]R 4       t	]R 4       t
]R 4       t]R	 4       t]R
 4       t]R 4       t]R 4       t]R 4       tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR t Rt!V t"R# ) FakeCUDAModuleiE  a  
An instance of this class will be injected into the __globals__ for an
executing function in order to implement calls to cuda.*. This will fail to
work correctly if the user code does::

    from numba import cuda as something_else

In other words, the CUDA module must be called cuda.
c                   \        V!  V n        \        V!  V n        \        4       V n        \        4       V n        \        V4      V n        \        4       V n
        \        4       V n        \        4       V n        \        P                   ! 4        F0  w  rE\#        WV4       VP$                   F  p\#        WV4       K  	  K2  	  R # r	   )r   gridDimblockDimr4   _cgr9   _localrK   _sharedrE   _constrg   _atomicr   _fp16r   itemssetattraliases)r   grid_dim	block_dimrS   namesvtyaliass   &&&&   r   r   FakeCUDAModule.__init__P  s    Xi(<#o%n5#o%'!^
 ',,.JDD%T* & /r   c                    V P                   # r	   )r-  r   s   &r   cgFakeCUDAModule.cgc  s    xxr   c                    V P                   # r	   )r.  r   s   &r   localFakeCUDAModule.localg      {{r   c                    V P                   # r	   )r/  r   s   &r   sharedFakeCUDAModule.sharedk      ||r   c                    V P                   # r	   )r0  r   s   &r   constFakeCUDAModule.consto  rB  r   c                    V P                   # r	   )r1  r   s   &r   atomicFakeCUDAModule.atomics  rF  r   c                    V P                   # r	   )r2  r   s   &r   fp16FakeCUDAModule.fp16w  s    zzr   c                @    \         P                  ! 4       P                  # r	   )r-   r.   	threadIdxr   s   &r   rQ  FakeCUDAModule.threadIdx{  s    '')333r   c                @    \         P                  ! 4       P                  # r	   )r-   r.   blockIdxr   s   &r   rT  FakeCUDAModule.blockIdx  s    '')222r   c                    ^ #     r2   r   s   &r   warpsizeFakeCUDAModule.warpsize  s    r   c                N    \         P                  ! 4       P                  ^ ,          # rW  )r-   r.   	thread_idr   s   &r   laneidFakeCUDAModule.laneid  s    '')33b88r   c                L    \         P                  ! 4       P                  4        R # r	   r,   r   s   &r   r/   FakeCUDAModule.syncthreads  s      "..0r   c                    R # r	   r2   r   s   &r   threadfenceFakeCUDAModule.threadfence      r   c                    R # r	   r2   r   s   &r   threadfence_block FakeCUDAModule.threadfence_block  rd  r   c                    R # r	   r2   r   s   &r   threadfence_system!FakeCUDAModule.threadfence_system  rd  r   c                J    \         P                  ! 4       P                  V4      # r	   )r-   r.   syncthreads_countr   rl   s   &&r   rl   FakeCUDAModule.syncthreads_count  s    '');;C@@r   c                J    \         P                  ! 4       P                  V4      # r	   )r-   r.   syncthreads_andrm  s   &&r   rp  FakeCUDAModule.syncthreads_and  s    '')99#>>r   c                J    \         P                  ! 4       P                  V4      # r	   )r-   r.   syncthreads_orrm  s   &&r   rs  FakeCUDAModule.syncthreads_or  s    '')88==r   c                6    \        V4      P                  R 4      # )1)binrW   rm  s   &&r   popcFakeCUDAModule.popc  s    3x~~c""r   c                     W,          V,           # r	   r2   r   s   &&&&r   fmaFakeCUDAModule.fma  r   r   c                    VR,          # )   gUUUUUU?r2   r   s   &&r   cbrtFakeCUDAModule.cbrt  s    U|r   c                L    \        R P                  V4      RRR1,          ^4      # ){:032b}N)intformatrm  s   &&r   brevFakeCUDAModule.brev  s#    9##C(2.22r   c                x    R P                  V4      p\        V4      \        VP                  R4      4      ,
          # r  0)r  lenlstrip)r   rl   ss   && r   clzFakeCUDAModule.clz  s.    S!1vAHHSM***r   c                    R P                  V4      p\        V4      \        VP                  R4      4      ,
          ^,           ^!,          pV# r  )r  r  rstrip)r   rl   r  rs   &&  r   ffsFakeCUDAModule.ffs  s>     S!Vc!((3-((1,2r   c                    V'       d   V# T# r	   r2   r   s   &&&&r   selpFakeCUDAModule.selp  s    q1r   c                   V P                   pV P                  pV P                  pVP                  VP                  ,          VP                  ,           pV^8X  d   V# VP                  VP                  ,          VP                  ,           pV^8X  d   WV3# VP
                  VP
                  ,          VP
                  ,           pV^8X  d   WVV3# \        RV,          4      h)r~  z*Global ID has 1-3 dimensions. %d requested)r,  rT  rQ  r   r   r   RuntimeError)r   nbdimbidtidr   r   r   s   &&      r   gridFakeCUDAModule.grid  s    }}mmnnEEDFFNSUU"6HEEDFFNSUU"66MEEDFFNSUU"6!9G!KLLr   c                @   V P                   pV P                  pVP                  VP                  ,          pV^8X  d   V# VP                  VP                  ,          pV^8X  d   WE3# VP                  VP                  ,          pV^8X  d   WEV3# \        RV,          4      h)r~  z,Global grid has 1-3 dimensions. %d requested)r,  r+  r   r   r   r  )r   r  r  gdimr   r   r   s   &&     r   gridsizeFakeCUDAModule.gridsize  s    }}||FFTVVO6HFFTVVO66MFFTVVO6!9IAMNNr   )r1  r-  r0  r2  r.  r/  r,  r+  N)#r    r!   r"   r#   r$   r   propertyr=  r@  rD  rH  rK  rN  rQ  rT  rY  r]  r/   rb  rf  ri  rl  rp  rs  rx  r{  r  r  r  r  r  r  r  r%   r&   r'   s   @r   r)  r)  E  s0    +&             4 4 3 3   9 91A?>#3+M O Or   r)  c              #   <  aa"   ^ RI Ho V P                  p\        V3R lVP	                  4        4       4      p\        V3R lVP	                  4        4       4      pVP                  V4        Rx  VP                  V4       R#   TP                  T4       i ; i5i)rV   )cudac              3   >   <"   T F  w  rVSJ g   K  W3x  K  	  R # 5ir	   r2   ).0kvr  s   &  r   	<genexpr>&swapped_cuda_module.<locals>.<genexpr>  s     A#341qDy#3s   
c              3   0   <"   T F  w  rVS3x  K  	  R # 5ir	   r2   )r  r  r  fake_cuda_modules   &  r   r  r    s     ?,$!$%,s   N)numbar  __globals__dictr3  update)fnr  fn_globsorigreplr  s   &f   @r   swapped_cuda_moduler    sm     ~~HA8>>#3AAD?$**,??DOOD 	s   A*B/B 3BBB)%r$   
contextlibr   r_   r-   r]   
numba.corer   numpyr>   numba.npr   r   objectr   r*   r4   r9   rE   rK   Lockri   rs   rw   r{   r   r   r   r   r   r   r   r   rg   r   r)  r  r2   r   r   <module>r     s?  
 & 
     " &6 *	1 	1 &F &F ,V ,^ ..

..

..
		
..

..

..
 ~~' 
..

..

..
>>\V \~Y6 YxXOV XOv  r   