:j) 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 p/home/mcse/projects/srt_converter/srt-converter-venv/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#cZeZdZdZ d dZdZdZdZdZd dZ e d Z e d Z y) FakeCUDAKernelz( Wraps a @cuda.jit-ed function. Nc|g}||_||_||_||_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__Asf  J ! 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..fs.Di.D.D !/#rr) r5reducer0 isinstancenpndarrayndimr to_devicer voidrr)arg_retr>rs rfake_argz)FakeCUDAKernel.__call__..fake_argcs"))#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__SsG <<$TWW.A.CD &tww~ & &:$--:>..J)*(I*.*=*=? !"2 3$ D .3773#7I7$TWW.>? 3"$**h"73J%dggxDKKPBBFF:2 23 3   G$ $  & &J8 3 3;$ $ sBD- 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 wwr)FNF)rrr) rrr r!r:rWr'r_rargpropertyrirkrrrr)r)<sU BG+$/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.runs%e ::t7 K FD*k4;;OA GGI KKN OOA    H(("&&q)[[$+8(, 6;3115578 ++a.77 AGG Hn,'.A,1A)''++-."%;Ha!**,HIK', DA{{kk!n33AKKNCC DIs G6GN)rrr r!r:rPrrrrOrOs"?(DrrO) contextlibrr5rrynumpyrEcudadrv.devicearrayrr kernelapirrr errorsr rQr r rrrrdictr#objectr)ThreadrnrOrrrrs% I@@0$   tcVcPPG)""PGfAD6ADr