U hda)@sddlmZddlZddlZddlZddlZddlmZm Z ddl m Z m Z m Z ddlmZddlmZmZdaed d Zd d ZGd ddZGdddeZGdddeZGdddejZGdddeZdS))contextmanagerN) FakeCUDAArrayFakeWithinKernelCUDAArray)Dim3FakeCUDAModuleswapped_cuda_module)normalize_kernel_dimensions)wrap_argArgHintccs*tdkstd|az dVW5daXdS)z* Push the current kernel context. Nz)concurrent simulated kernel not supported)_kernel_contextAssertionError)modrT/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/simulator/kernel.py_push_kernel_contexts  rcCstS)zT Get the current kernel context. This is usually done by a device function. )r rrrr_get_kernel_context$src@seZdZdZddZdS) FakeOverloadzE Used only to provide the max_cooperative_grid_blocks method cCsdS)Nrr)selfZblockdimrrrmax_cooperative_grid_blocks/sz(FakeOverload.max_cooperative_grid_blocksN)__name__ __module__ __qualname____doc__rrrrrr+src@seZdZddZdS)FakeOverloadDictcCstSN)r)rkeyrrr __getitem__6szFakeOverloadDict.__getitem__N)rrrrrrrrr5src@sbeZdZdZdgdfddZddZddZd d Zd d ZdddZ e ddZ e ddZ dS)FakeCUDAKernelz( Wraps a @cuda.jit-ed function. FcCsJ||_||_||_||_t||_d|_d|_d|_d|_ t ||dS)Nr) fn_deviceZ _fastmath_debuglist extensionsgrid_dim block_dimstreamdynshared_size functoolsupdate_wrapper)rr ZdeviceZfastmathr$debugrrr__init__As zFakeCUDAKernel.__init__c sjr2tjtj|W5QRSQRXtjj\}}t||j}t |gfddfdd|D}tj|8t j |D]&}t j||j }|j|f|qW5QRXD] }|qW5QRXdS)Ncstfddjd|f\}}t|tjrF|jdkrFt|}n0t|t r\|}nt|tj rrt |}n|}t|t rt |S|S)Ncs|j|ddS)Nr)r'retr)Z prepare_args)Zty_val extension)r-rrbsz;FakeCUDAKernel.__call__..fake_arg..r) r)reducer$ isinstancenpZndarrayndimr Z to_devicer voidrr)arg_ret)r-rrrfake_arg_s       z)FakeCUDAKernel.__call__..fake_argcsg|] }|qSrr).0r5)r8rr vsz+FakeCUDAKernel.__call__..)r!rr rr r%r&rr(rr2ndindex BlockManagerr"run) rargsr%r&Zfake_cuda_moduleZ fake_args grid_pointZbmwbr)r8r-rr__call__Os& zFakeCUDAKernel.__call__cCs2t|dd\|_|_t|dkr.|d|_|S)Nr )r r%r&lenr()r configurationrrrrs    zFakeCUDAKernel.__getitem__cCsdSrrrrrrbindszFakeCUDAKernel.bindcGs|Srr)rr>rrr specializeszFakeCUDAKernel.specializercCs$|dkrtd|||d||fS)Nrz0Can't create ForAll with negative task count: %sr) ValueError)rZntasksZtpbr'Z sharedmemrrrforalls zFakeCUDAKernel.forallcCstSr)rrFrrr overloadsszFakeCUDAKernel.overloadscCs|jSr)r rFrrrpy_funcszFakeCUDAKernel.py_funcN)rrr) rrrrr,rArrGrHrJpropertyrKrLrrrrr<s1   rcsTeZdZdZfddZfddZddZdd Zd d Zd d Z ddZ Z S) BlockThreadzG Manages the execution of a function for a single CUDA thread. c s|rfdd}|}n}tt|j|dt|_d|_||_t||_ t||_ d|_ d|_ d|_ ||_t|jj}|j j|j|j j|j|j j|_dS)Ncstjdd||dS)Nraise)divide)r2Zseterr)r>kwargsfrr debug_wrappers z+BlockThread.__init__..debug_wrapper)targetFT)superrNr, threadingEventsyncthreads_eventsyncthreads_blocked_managerrblockIdx threadIdx exceptiondaemonabortr+ _block_dimxyz thread_id) rrSmanagerr\r]r+rTrUZblockDim __class__rRrr,s(     zBlockThread.__init__c sztt|Wntk r}zfdt|j}dt|j}t|dkrZd||f}nd|||f}t d}t |||f|_ W5d}~XYnXdS)Nztid=%szctaid=%sz%s %sz %s %s: %sr ) rVrNr= Exceptionr#r]r\strsysexc_infotyper^)retidZctaidmsgtbrgrrr=s  zBlockThread.runcCs:|jrtdd|_|j|j|jr6tddS)Nz"abort flag set on syncthreads callTz#abort flag set on syncthreads clear)r` RuntimeErrorrZrYwaitclearrFrrr syncthreadss  zBlockThread.syncthreadscCsD|jj|jj|jjf}||jj|<|t|jj}||Sr) r]rbrcrdr[ block_statervr2Z count_nonzero)rvalueidxcountrrrsyncthreads_counts  zBlockThread.syncthreads_countcCsL|jj|jj|jjf}||jj|<|t|jj}||rHdSdSNrr) r]rbrcrdr[rwrvr2allrrxrytestrrrsyncthreads_ands  zBlockThread.syncthreads_andcCsL|jj|jj|jjf}||jj|<|t|jj}||rHdSdSr|) r]rbrcrdr[rwrvr2anyr~rrrsyncthreads_ors  zBlockThread.syncthreads_orcCsd|j|jfS)NzThread <<<%s, %s>>>)r\r]rFrrr__str__szBlockThread.__str__) rrrrr,r=rvr{rrr __classcell__rrrgrrNs   rNc@s eZdZdZddZddZdS)r<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. cCs.||_||_||_||_tj|tjd|_dS)N)Zdtype)Z _grid_dimra_fr"r2ZzerosZbool_rw)rrSr%r&r+rrrr,s zBlockManager.__init__c s"t}t}t}tjjD]@}fdd}t|||j}|||||q|r|D]R}|jr~||qh|j rh|D]} d| _ d| _| j q|j d |j dqh||kr|D]}d|_|j qt}tdd|D}q`|D] }|j r|j d |j dqdS) NcsjdSr)rrr>rrrrUsz BlockManager.run..targetTFrrcSsg|]}|r|qSr)is_alive)r9trrrr:/sz$BlockManager.run..) setr2r;rarNr"startaddrZr^r`rYwith_traceback) rr?r>threadsZ livethreadsZblockedthreadsZ block_pointrUrZt_otherrrrr= s8     zBlockManager.runN)rrrrr,r=rrrrr<sr<) contextlibrr)rlrWnumpyr2Zcudadrv.devicearrayrrZ kernelapirrrerrorsr r>r r r rrrdictrobjectrThreadrNr<rrrrs"    dS