1@j*(ddlmZddlZddlZddlZddlZddlmZm Z ddl m Z m Z m Z ddlmZddlmZmZ daed Zd ZGd d ZGd deZGddeZGddejZGddeZdS))contextmanagerN) FakeCUDAArrayFakeWithinKernelCUDAArray)Dim3FakeCUDAModuleswapped_cuda_module)normalize_kernel_dimensions)wrap_argArgHintc#LKt Jd|a dVdadS#dawxYw)z* Push the current kernel context. Nz)concurrent simulated kernel not supported_kernel_context)mods LC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/simulator/kernel.py_push_kernel_contextrsJ  " "$O " " "O $s#ctS)zT Get the current kernel context. This is usually done by a device function. rr_get_kernel_contextr$s  rceZdZdZdZdS) FakeOverloadzE Used only to provide the max_cooperative_grid_blocks method cdS)Nrr)selfblockdims rmax_cooperative_grid_blocksz(FakeOverload.max_cooperative_grid_blocks/s qrN)__name__ __module__ __qualname____doc__rrrrrr+s-rrceZdZdZdS)FakeOverloadDictctSN)r)rkeys r __getitem__zFakeOverloadDict.__getitem__6s~~rN)rrr r'rrrr#r#5s#rr#cheZdZdZ d dZdZdZdZdZdd Z e d Z e d Z dS)FakeCUDAKernelz( Wraps a @cuda.jit-ed function. FNc|g}||_||_||_||_t ||_d|_d|_d|_d|_ tj ||dS)Nr) fn_device _fastmath_debuglist extensionsgrid_dim block_dimstreamdynshared_size functoolsupdate_wrapper)rr+devicefastmathr0debugs r__init__zFakeCUDAKernel.__init__Asr  J ! z**   r*****rc t jrCtjt5j|cdddS#1swxYwYt jj\}}t||j}t|5g fd fd|D}tj|5tj |D],}tj||j }|j|g|R- dddn #1swxYwY D] }| ddddS#1swxYwYdS)Nctjfdjd|f\}}t|tjr.|jdkr#t|}nWt|tr|}n,t|tj rt|}n|}t|trt|S|S)Nc |j|ddS)Nr)r3retr) prepare_args)ty_val extensionr>s rz;FakeCUDAKernel.__call__..fake_arg..fs%.Di.D !/#/#/#rr) r5reducer0 isinstancenpndarrayndimr to_devicer voidrr)arg_retr>rs rfake_argz)FakeCUDAKernel.__call__..fake_argcs")####O3K 3c2:..38a<<"3--11$77CCW------CCRW--',,CCCc=11:4S999 rc&g|] }|Srr).0rJrMs r z+FakeCUDAKernel.__call__..zs!7773#777r)r,r r+rr r1r2rr4rrEndindex BlockManagerr.run) rargsr1r2fake_cuda_module fake_args grid_pointbmwbrMr>s ` @@r__call__zFakeCUDAKernel.__call__SsX < &$TW.A.C.CDD & &tw~ & & & & & & & & & & & & & & & &:$-:>.JJ)*(I*.*=?? !"2 3 3$ $ D      .8777$777I$TW.>?? 3 3"$*h"733J%dgxDKPPBBF:2 222223 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3    G$ $ $ $ $ $ $ $ $ $ $ $ $ $ $ $ $ $ sG AA A,D-<.debug_wrappers1 ))))4"6"""""r)targetFT)superrqr: threadingEventsyncthreads_eventsyncthreads_blocked_managerrblockIdx threadIdx exceptiondaemonabortr9 _block_dimxyz thread_id) rrxmanagerrrr9ryrzblockDim __class__s ` rr:zBlockThread.__init__s   # # # # ##FFF k4  )))888!*!2!2#(  h y)   12)XZ4>;K;C:;?>;K>>)rrras r__str__zBlockThread.__str__s(, t~~~FFr) rrr r!r:rSrrrrr __classcell__)rs@rrqrqsNNNNN0 0 0 0 0 0 F F F      GGGGGGGrrqceZdZdZdZdZdS)rRa 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|_dS)N)dtype) _grid_dimr_fr.rEzerosbool_r)rrxr1r2r9s rr:zBlockManager.__init__ s=!# 8IRX>>>rc\t}t}t}tjjD]^}fd}t |||j}|||||_|r|D]}}|jr|||j rW|D])} d| _ d| _| j *|j d |j d~||kr3|D]"}d|_|j #t}td|D}||D]4}|j r+|j d |j d5dS)NcjdSr%)r)rTrsrrzz BlockManager.run..targetsrTFrrc:g|]}||Sr)is_alive)rOts rrPz$BlockManager.run..3s%HHHa!**,,HHHHr) setrErQrrqr.startaddrrrr~with_traceback) rrWrTthreads livethreadsblockedthreads block_pointrzrt_others ` ` rrSzBlockManager.runs%%ee :t7  K      FD*k4;OOA GGIII KKNNN OOA      J H H( H"&&q))))[ H$+88(, 6;31557777+a.77 AGGG Hn,,'..A,1A)'++----!$HH;HHHIIK' J, D DA{ Dk!n33AKNCCC D D DrN)rrr r!r:rSrrrrRrRsA"???(D(D(D(D(DrrR) contextlibrr5rr|numpyrEcudadrv.devicearrayrr kernelapirrr errorsr rTr r rrrrdictr#objectr)ThreadrqrRrrrrs%%%%%% IIIIIIII@@@@@@@@@@000000$$$$$$$$    tcccccVcccPPGPGPGPGPG)"PGPGPGfADADADADAD6ADADADADADr