xKga) ddlmZddlZddlZddlZddlZddlmZm Z ddl m Z m Z m Z ddlmZddlmZmZ daed Zd ZGd d ZGd deZGddeZGddej6ZGddeZy))contextmanagerN) FakeCUDAArrayFakeWithinKernelCUDAArray)Dim3FakeCUDAModuleswapped_cuda_module)normalize_kernel_dimensions)wrap_argArgHintc#FKtJd|a dday#dawxYww)z* Push the current kernel context. Nz)concurrent simulated kernel not supported_kernel_context)mods _/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/simulator/kernel.py_push_kernel_contextrs1  "O$OO "O $s!!!ctS)zT Get the current kernel context. This is usually done by a device function. rr_get_kernel_contextr$s  rceZdZdZdZy) FakeOverloadzE Used only to provide the max_cooperative_grid_blocks method cy)Nrr)selfblockdims rmax_cooperative_grid_blocksz(FakeOverload.max_cooperative_grid_blocks/srN)__name__ __module__ __qualname____doc__rrrrrr+s rrceZdZdZy)FakeOverloadDictctSN)r)rkeys r __getitem__zFakeOverloadDict.__getitem__6s ~rN)rrr r'rrrr#r#5srr#c^eZdZdZdgdfdZdZdZdZdZd dZ e d Z e d Z y ) FakeCUDAKernelz( Wraps a @cuda.jit-ed function. Fc||_||_||_||_t ||_d|_d|_d|_d|_ tj||y)Nr) fn_device _fastmath_debuglist extensionsgrid_dim block_dimstreamdynshared_size functoolsupdate_wrapper)rr+devicefastmathr0debugs r__init__zFakeCUDAKernel.__init__AsY ! z*    r*rc  jr7tjt5j|cdddSt j j \}}t||j}t|5g fd}|Dcgc] }|| }}tj|5tj|D]7}tj||j} | j|g|9 ddd D] } |   dddy#1swYxYwcc}w#1swY2xYw#1swYyxYw)Nctjfdjd|f\}}t|tj r*|j dkDrt|j}nJt|tr|j}n(t|tjr t|}n|}t|tr t|S|S)Nc*|j|ddS)Nr)r3retr) prepare_args)ty_val extensionr>s rz;FakeCUDAKernel.__call__..fake_arg..bs.Di.D.D !/#rr) r5reducer0 isinstancenpndarrayndimr to_devicer voidrr)arg_retr>rs rfake_argz)FakeCUDAKernel.__call__..fake_arg_s"))#OO3K 3c2::.388a<"3-11$7CW----CRWW-',CCc=14S99 r)r,r r+rr r1r2rr4rrEndindex BlockManagerr.run) rargsr1r2fake_cuda_modulerMrJ fake_args grid_pointbmwbr>s ` @r__call__zFakeCUDAKernel.__call__Os& <<$TWW.A.CDtww~ED:$--:>..J)*(I*.*=*=? !"2 3D .377$3#$I7$TWW.>?"$**h"7J%dggxDKKPBBFF:2 2#8@ G4 3EDJ8??;4 3sBD- E D9'E ?A D> E -D69E >E E  Ecft|dd\|_|_t|dk(r |d|_|S)Nr )r r1r2lenr4)r configurations rr'zFakeCUDAKernel.__getitem__s? 'r): ; & t~ }  ""/"2D  rcyr%rrs rbindzFakeCUDAKernel.binds rc|Sr%r)rrQs r specializezFakeCUDAKernel.specializes rc:|dkrtd|z||d||fS)Nrz0Can't create ForAll with negative task count: %sr) ValueError)rntaskstpbr3 sharedmems rforallzFakeCUDAKernel.foralls5 A:O%&' 'FAvy011rctSr%)r#r^s r overloadszFakeCUDAKernel.overloadss !!rc|jSr%)r+r^s rpy_funczFakeCUDAKernel.py_funcs wwrN)rrr) rrr r!r:rWr'r_rargpropertyrirkrrrr)r)<sW-2b +/b 2 ""rr)cJeZdZdZfdZfdZdZdZdZdZ dZ xZ S) BlockThreadzG Manages the execution of a function for a single CUDA thread. c|rfd}|}n}tt| |tj|_d|_||_t||_ t||_ d|_ d|_ d|_ ||_t|jj}|jj |j |jj"|j"|jj$zzzz|_y)NcBtjd|i|y)Nraise)divide)rEseterr)rQkwargsfs r debug_wrapperz+BlockThread.__init__..debug_wrappers )4"6"r)targetFT)superrnr: threadingEventsyncthreads_eventsyncthreads_blocked_managerrblockIdx threadIdx exceptiondaemonabortr9 _block_dimxyz thread_id) rrumanagerr~rr9rvrwblockDim __class__s ` rr:zBlockThread.__init__s  ##FF k4))8!*!2#(  h y)   112))XZZ4>>;K;K;C::;?>>;K;K>>)r~rr^s r__str__zBlockThread.__str__s(, t~~FFr) rrr r!r:rPrrrrr __classcell__)rs@rrnrns.N0 0 F  GrrnceZdZdZdZdZy)rOa 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||_||_||_||_t j |tj |_y)N)dtype) _grid_dimr_fr.rEzerosbool_r)rrur1r2r9s rr:zBlockManager.__init__s5!# 88IRXX>rct}t}t}tjjD]S}fd}t |||j }|j |j||j|U|r|D]}|jr|j|!|js.|D]*} d| _ d| _| jj,|jdj|jd||k(r2|D]#}d|_|jj%t}t|Dcgc]}|js|c}}|r|D]9}|js|jdj|jdycc}w)Nc$jyr%)r)rQrsrrwz BlockManager.run..targetsrTFrr)setrErNrrnr.startaddr|rrr{with_tracebackis_alive) rrTrQthreads livethreadsblockedthreads block_pointrwtt_others ` ` rrPzBlockManager.run s~%e ::t7K FD*k4;;OA GGI KKN OOA  8 (("&&q)[[$+(, 6;311557$+ ++a.77 AGG!n,'A,1A)''++-("%;H;a!**,;HIK',A{{kk!n33AKKNCCIs G6GN)rrr r!r:rPrrrrOrOs"?(DrrO) contextlibrr5rrynumpyrEcudadrv.devicearrayrr kernelapirrr errorsr rQr r rrrrdictr#objectr)ThreadrnrOrrrrs% I@@0$   t_V_HPG)""PGfAD6ADr