o ia)@sddlmZddlZddlZddlZddlZddlmZm Z ddl m Z m Z m Z ddlmZddlmZmZ daed d Zd d ZGd ddZGdddeZGdddeZGdddejZGdddeZdS))contextmanagerN) FakeCUDAArrayFakeWithinKernelCUDAArray)Dim3FakeCUDAModuleswapped_cuda_module)normalize_kernel_dimensions)wrap_argArgHintccs.tdus Jd|azdVWdadSdaw)z* Push the current kernel context. Nz)concurrent simulated kernel not supported_kernel_context)modrj/var/www/html/eduruby.in/lip-sync/lip-sync-env/lib/python3.10/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+s rc@seZdZddZdS)FakeOverloadDictcCtSN)r)rkeyrrr __getitem__6szFakeOverloadDict.__getitem__N)rrrrrrrrr5s rc@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 s jrtjt j|WdS1swYtjj\}}t||j}t |Sgfddfdd|D}tj|"t j |D]}t j||j }|j|g|RqRWdn1sqwYD]}|qxWddS1swYdS)Ncstfddjd|f\}}t|tjr#|jdkr#t|}nt|t r.|}n t|tj r9t |}n|}t|t rDt |S|S)Ncs|j|ddS)Nr)r(retr)Z prepare_args)Zty_val extension)r.rrbs z;FakeCUDAKernel.__call__..fake_arg..r) r*reducer% isinstancenpZndarrayndimr Z to_devicer voidrr)arg_ret)r.rrrfake_arg_s       z)FakeCUDAKernel.__call__..fake_argcsg|]}|qSrr).0r6)r9rr vsz+FakeCUDAKernel.__call__..)r"rr!rr r&r'rr)rr3ndindex BlockManagerr#run) rargsr&r'Zfake_cuda_moduleZ fake_args grid_pointZbmwbr)r9r.rr__call__Os0  "zFakeCUDAKernel.__call__cCs2t|dd\|_|_t|dkr|d|_|S)Nr )r r&r'lenr))r configurationrrrrs   zFakeCUDAKernel.__getitem__cCsdSrrrrrrbindzFakeCUDAKernel.bindcGs|Srr)rr?rrr specializerIzFakeCUDAKernel.specializercCs$|dkr td|||d||fS)Nrz0Can't create ForAll with negative task count: %sr) ValueError)rZntasksZtpbr(Z sharedmemrrrforalls zFakeCUDAKernel.forallcCrr)rrGrrr overloadszFakeCUDAKernel.overloadscCs|jSr)r!rGrrrpy_funcrNzFakeCUDAKernel.py_funcN)rrr) rrrrr-rBrrHrJrLpropertyrMrOrrrrr <s1   r csTeZdZdZfddZfddZddZdd Zd d Zd d Z ddZ Z S) BlockThreadzG Manages the execution of a function for a single CUDA thread. c s|r fdd}|}n}tt|j|dt|_d|_||_t||_ t||_ d|_ d|_ d|_ ||_t|jj}|j j|j|j j|j|j j|_dS)Ncstjdd|i|dS)Nraise)divide)r3Zseterr)r?kwargsfrr debug_wrappers z+BlockThread.__init__..debug_wrapper)targetFT)superrQr- threadingEventsyncthreads_eventsyncthreads_blocked_managerrblockIdx threadIdx exceptiondaemonabortr, _block_dimxyz thread_id) rrVmanagerr_r`r,rWrXZblockDim __class__rUrr-s(     zBlockThread.__init__c sz tt|WdStyN}z8dt|j}dt|j}t|dkr-d||f}nd|||f}t d}t |||f|_ WYd}~dSd}~ww)Nztid=%szctaid=%sz%s %sz %s %s: %sr ) rYrQr> Exceptionr$r`r_strsysexc_infotypera)retidZctaidmsgtbrjrrr>s   zBlockThread.runcCs:|jrtdd|_|j|j|jrtddS)Nz"abort flag set on syncthreads callTz#abort flag set on syncthreads clear)rc RuntimeErrorr]r\waitclearrGrrr syncthreadss  zBlockThread.syncthreadscCsD|jj|jj|jjf}||jj|<|t|jj}||Sr) r`rerfrgr^ block_stateryr3Z count_nonzero)rvalueidxcountrrrsyncthreads_counts  zBlockThread.syncthreads_countcCL|jj|jj|jjf}||jj|<|t|jj}||r$dSdSNrr) r`rerfrgr^rzryr3allrr{r|testrrrsyncthreads_and   zBlockThread.syncthreads_andcCrr) r`rerfrgr^rzryr3anyrrrrsyncthreads_orrzBlockThread.syncthreads_orcCsd|j|jfS)NzThread <<<%s, %s>>>)r_r`rGrrr__str__szBlockThread.__str__) rrrrr-r>ryr~rrr __classcell__rrrjrrQs   rQc@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_dimrd_fr#r3ZzerosZbool_rz)rrVr&r'r,rrrr-s zBlockManager.__init__c s$t}t}t}tjjD] }fdd}t|||j}|||||q|r}|D])}|jr?||q4|j r]|D] } d| _ d| _| j qD|j d |j dq4||krr|D] }d|_|j qdt}tdd|D}|s2|D]}|j r|j d |j dqdS) NcsjdSr)rrr?rrrrXsz BlockManager.run..targetTFrrcSsg|]}|r|qSr)is_alive)r:trrrr;/sz$BlockManager.run..) setr3r<rdrQr#startaddr]rarcr\with_traceback) rr@r?threadsZ livethreadsZblockedthreadsZ block_pointrXrZt_otherrrrr> s@      zBlockManager.runN)rrrrr-r>rrrrr=s r=) contextlibrr*rorZnumpyr3Zcudadrv.devicearrayrrZ kernelapirrrerrorsr r?r r rrrrdictrobjectr ThreadrQr=rrrrs&     dS