+
    :i)                        ^ RI Ht ^ RIt^ RIt^ RIt^ RIt^RIHtH	t	 ^RI
HtHtHt ^RIHt ^RIHtHt  Rs]R 4       tR t ! R	 R
4      t ! R R]4      t ! R R]4      t ! R R]P6                  4      t ! R R]4      tR# )    )contextmanagerN)FakeCUDAArrayFakeWithinKernelCUDAArray)Dim3FakeCUDAModuleswapped_cuda_module)normalize_kernel_dimensions)wrap_argArgHintc              #  L   "   \         e   Q R4       hV s  Rx  Rs R#   Rs i ; i5i)z"
Push the current kernel context.
Nz)concurrent simulated kernel not supported_kernel_context)mods   &Y/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/simulator/kernel.py_push_kernel_contextr      s0      "O$OO"O$s   $ $!$c                     \         # )zL
Get the current kernel context. This is usually done by a device function.
r        r   _get_kernel_contextr   $   s
     r   c                   *   a  ] tR t^+t o RtR tRtV tR# )FakeOverloadz=
Used only to provide the max_cooperative_grid_blocks method
c                    ^#    r   )selfblockdims   &&r   max_cooperative_grid_blocks(FakeOverload.max_cooperative_grid_blocks/   s     r   r   N)__name__
__module____qualname____firstlineno____doc__r   __static_attributes____classdictcell____classdict__s   @r   r   r   +   s      r   r   c                   &   a  ] tR t^5t o R tRtV tR# )FakeOverloadDictc                    \        4       # N)r   )r   keys   &&r   __getitem__FakeOverloadDict.__getitem__6   s     ~r   r   N)r   r    r!   r"   r-   r$   r%   r&   s   @r   r)   r)   5   s      r   r)   c                   p   a  ] tR t^<t o RtRR ltR tR tR tR t	RR lt
]R	 4       t]R
 4       tRtV tR# )FakeCUDAKernelz 
Wraps a @cuda.jit-ed function.
Nc                    Vf   . pWn         W n        W0n        WPn        \	        V4      V n        R V n        R V n        ^ V n        ^ V n	        \        P                  ! W4       R # r+   )fn_device	_fastmath_debuglist
extensionsgrid_dim	block_dimstreamdynshared_size	functoolsupdate_wrapper)r   r2   devicefastmathr7   debugs   &&&&&&r   __init__FakeCUDAKernel.__init__A   s\     J!z*   *r   c                $  a a S P                   '       dB   \        S P                  \        4       4      ;_uu_ 4        S P                  ! V!  uuR R R 4       # \	        S P
                  S P                  4      w  r#\        W#S P                  4      p\        V4      ;_uu_ 4        . oVV 3R lpV Uu. uF
  qe! V4      NK  	  pp\        S P                  V4      ;_uu_ 4        \        P                  ! V!   F8  p\        S P                  W#S P                  4      p	V	P                  ! V.VO5!   K:  	  R R R 4       S F
  p
V
! 4        K  	  R R R 4       R #   + '       g   i     EL; iu upi   + '       g   i     LC; i  + '       g   i     R # ; i)Nc                   < \         P                  ! V3R  lSP                  RV 34      w  r\        V \        P
                  4      '       d-   V P                  ^ 8  d   \        V 4      P                  S4      pMV\        V \        4      '       d   V P                  S4      pM.\        V \        P                  4      '       d   \        V 4      pMT p\        V\        4      '       d   \        V4      # V# )c                 ,   < VP                   ! V R ^ RS/ # )r:   retr)prepare_args)ty_val	extensionrF   s   &&r   <lambda>;FakeCUDAKernel.__call__.<locals>.fake_arg.<locals>.<lambda>f   s%    i.D.D/# /# "/#r   N)r<   reducer7   
isinstancenpndarrayndimr
   	to_devicer   voidr   r   )arg_retrF   r   s   &  r   fake_arg)FakeCUDAKernel.__call__.<locals>.fake_argc   s    "))# OO3K c2::..388a<"3-11$7CW-----CRWW--',CCc=114S99
r   )r3   r   r2   r   r	   r8   r9   r   r;   r   rN   ndindexBlockManagerr5   run)r   argsr8   r9   fake_cuda_modulerV   rS   	fake_args
grid_pointbmwbrF   s   f*         @r   __call__FakeCUDAKernel.__call__S   s-   <<<$TWW.A.CDDww~ ED :$--:>..J *(*.*=*=?!"233 D. 377$3#$I7$TWW.>??"$**h"7J%dggxDKKPBFF:2	2 #8 @  G 43 EDDJ 8??; 433sB   E!E>/E&? E>AE+	0E>E#	&E>+E;6E>>F	c                |    \        VR ,          !  w  V n        V n        \        V4      ^8X  d   V^,          V n        V # ):N   N)r	   r8   r9   lenr;   )r   configurations   &&r   r-   FakeCUDAKernel.__getitem__   s=    'r):; 	&t~ }""/"2Dr   c                    R # r+   r   r   s   &r   bindFakeCUDAKernel.bind   s    r   c                    V # r+   r   )r   r[   s   &*r   
specializeFakeCUDAKernel.specialize   s    r   c                J    V^ 8  d   \        RV,          4      hW^W43,          # )r   z0Can't create ForAll with negative task count: %s)
ValueError)r   ntaskstpbr:   	sharedmems   &&&&&r   forallFakeCUDAKernel.forall   s1    A:O%& ' 'Av011r   c                    \        4       # r+   )r)   ri   s   &r   	overloadsFakeCUDAKernel.overloads   s    !!r   c                    V P                   # r+   )r2   ri   s   &r   py_funcFakeCUDAKernel.py_func   s    wwr   )	r5   r3   r4   r9   r;   r7   r2   r8   r:   )FNF)r   r   r   )r   r    r!   r"   r#   rA   ra   r-   rj   rm   rt   propertyrw   rz   r$   r%   r&   s   @r   r0   r0   <   sQ     +$/b2 " "  r   r0   c                   `   a a ] tR t^t oRtV 3R ltV 3R ltR tR tR t	R t
R tR	tVtV ;t# )
BlockThreadz?
Manages the execution of a function for a single CUDA thread.
c                4  <a V'       d
   V3R  lpTpMSp\         \        V `  VR7       \        P                  ! 4       V n        RV n        W n        \        V!  V n	        \        V!  V n
        RV n        RV n        RV n        WPn        \        V P                  P                  !  pV P                  P                   VP                   V P                  P"                  VP"                  V P                  P$                  ,          ,           ,          ,           V n        R# )c                  F   < \         P                  ! R R7       S! V / VB  R# )raise)divideN)rN   seterr)r[   kwargsfs   *,r   debug_wrapper+BlockThread.__init__.<locals>.debug_wrapper   s    		)4"6"r   )targetFNT)superr~   rA   	threadingEventsyncthreads_eventsyncthreads_blocked_managerr   blockIdx	threadIdx	exceptiondaemonabortr@   
_block_dimxyz	thread_id)
r   r   managerr   r   r@   r   r   blockDim	__class__s
   &f&&&&   r   rA   BlockThread.__init__   s    # #FFk4))8!*!2#( hy)

112))XZZ4>>;K;K;C::;?>>;K;K<L<L .M Nr   c                  <  \         \        V `  4        R#   \         d   pR \	        T P
                  4      ,          pR\	        T P                  4      ,          p\        T4      R8X  d   T: RT: 2pMT: RT: RT: 2p\        P                  ! 4       ^,          p\        T4      ! T4      T3T n         Rp?R# Rp?ii ; i)ztid=%szctaid=%s  z: N)r   r~   rZ   	Exceptionr6   r   r   strsysexc_infotyper   )r   etidctaidmsgtbr   s   &     r   rZ   BlockThread.run   s    	0+t(* 
	0T$..11Cdmm!44E1v|!$e,%(%3"B #1gclB/DNN
	0s    C BB;;C c                    V P                   '       d   \        R 4      hRV n        V P                  P	                  4        V P                  P                  4        V P                   '       d   \        R4      hR# )z"abort flag set on syncthreads callTz#abort flag set on syncthreads clearN)r   RuntimeErrorr   r   waitclearri   s   &r   syncthreadsBlockThread.syncthreads   s]    :::CDD#' ##%$$&:::DEE r   c                L   V P                   P                  V P                   P                  V P                   P                  3pWP                  P
                  V&   V P                  4        \        P                  ! V P                  P
                  4      pV P                  4        V# r+   )	r   r   r   r   r   block_stater   rN   count_nonzero)r   valueidxcounts   &&  r   syncthreads_countBlockThread.syncthreads_count   su    nn 0 0$..2B2BB).!!#&  !:!:;r   c                `   V P                   P                  V P                   P                  V P                   P                  3pWP                  P
                  V&   V P                  4        \        P                  ! V P                  P
                  4      pV P                  4        V'       d   ^# ^ # r   )	r   r   r   r   r   r   r   rN   allr   r   r   tests   &&  r   syncthreads_andBlockThread.syncthreads_and   {    nn 0 0$..2B2BB).!!#&vvdmm//0qar   c                `   V P                   P                  V P                   P                  V P                   P                  3pWP                  P
                  V&   V P                  4        \        P                  ! V P                  P
                  4      pV P                  4        V'       d   ^# ^ # r   )	r   r   r   r   r   r   r   rN   anyr   s   &&  r   syncthreads_orBlockThread.syncthreads_or   r   r   c                @    R V P                   : RV P                  : R2# )z
Thread <<<z, z>>>)r   r   ri   s   &r   __str__BlockThread.__str__   s    (,t~~FFr   )
r   r   r   r   r@   r   r   r   r   r   )r   r    r!   r"   r#   rA   rZ   r   r   r   r   r   r$   r%   __classcell__)r   r'   s   @@r   r~   r~      s6     N00
F  G Gr   r~   c                   0   a  ] tR t^t o RtR tR tRtV tR# )rY   a  
Manages the execution of a thread block.

When run() is called, all threads are started. Each thread executes until it
hits syncthreads(), at which point it sets its own syncthreads_blocked to
True so that the BlockManager knows it is blocked. It then waits on its
syncthreads_event.

The BlockManager polls threads to determine if they are blocked in
syncthreads(). If it finds a blocked thread, it adds it to the set of
blocked threads. When all threads are blocked, it unblocks all the threads.
The thread are unblocked by setting their syncthreads_blocked back to False
and setting their syncthreads_event.

The polling continues until no threads are alive, when execution is
complete.
c                    W n         W0n        Wn        W@n        \        P
                  ! V\        P                  R 7      V n        R# ))dtypeN)	_grid_dimr   _fr5   rN   zerosbool_r   )r   r   r8   r9   r@   s   &&&&&r   rA   BlockManager.__init__	  s-    !#88IRXX>r   c                  a a \        4       p\        4       p\        4       p\        P                  ! S P                  !   FT  pVV 3R  lp\	        VS WS P
                  4      pVP                  4        VP                  V4       VP                  V4       KV  	  V'       Ed   V F  pVP                  '       d   VP                  V4       K(  VP                  '       g   K<  V F+  p	RV	n
        RV	n        V	P                  P                  4        K-  	  VP                  ^ ,          P                  VP                  ^,          4      h	  WE8X  d5   V F$  pRVn        VP                  P                  4        K&  	  \        4       p\        V Uu. uF  qP                  4       '       g   K  VNK  	  up4      pEK  V FH  pVP                  '       g   K  VP                  ^ ,          P                  VP                  ^,          4      h	  R# u upi )c                  (   < SP                   ! S !   R # r+   )r   )r[   r   s   r   r    BlockManager.run.<locals>.target  s    r   TFN)setrN   rX   r   r~   r5   startaddr   r   r   r   with_tracebackis_alive)
r   r^   r[   threadslivethreadsblockedthreadsblock_pointr   tt_others
   f&j       r   rZ   BlockManager.run  s~   %e::t7KFD*4;;OAGGIKKNOOA 8 k ((("&&q)[[[ $+(,6;311557 $+
 ++a.77AGG ! ,'A,1A)''++- ( "%;H;a**,;HIK A{{{kk!n33AKKNCC  Is   H#H)r   r5   r   r   r   N)	r   r    r!   r"   r#   rA   rZ   r$   r%   r&   s   @r   rY   rY      s     "?(D (Dr   rY   )
contextlibr   r<   r   r   numpyrN   cudadrv.devicearrayr   r   	kernelapir   r   r   errorsr	   r[   r
   r   r   r   r   r   dictr)   objectr0   Threadr~   rY   r   r   r   <module>r      s    %  
   I @ @ 0 $
  
 
 t cV cPPG)"" PGfAD6 ADr   