U ,d6)@sddlmZddlmZddlZddlZddlZddlm Z m Z ddl m Z m Z mZddlmZdd lmZmZdaed d Zd d ZGdddZGdddeZGdddeZGdddejZGdddeZdS))contextmanager)reduceN) FakeCUDAArrayFakeWithinKernelCUDAArray)Dim3FakeCUDAModuleswapped_cuda_module)normalize_kernel_dimensions)wrap_argArgHintccs*tdkstd|az dVW5daXdS)z* Push the current kernel context. Nz'conrrent simulated kernel not supported)_kernel_contextAssertionError)modr?/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/simulator/kernel.py_push_kernel_contexts  rcCstS)zT Get the current kernel context. This is usually done by a device function. )rrrrr_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. FcCs>||_||_||_||_t||_d|_d|_d|_d|_ dS)Nr) fn_deviceZ _fastmath_debuglist extensionsgrid_dim block_dimstreamdynshared_size)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|tjrD|jdkrDt|}n0t|trZ|}nt|tj rpt |}n|}t|t rt |S|S)Ncs|j|ddS)Nr)r(retr)Z prepare_args)Zty_val extension)r,rrasz;FakeCUDAKernel.__call__..fake_arg..r) rr% isinstancenpZndarrayndimr Z to_devicer voidrr)arg_ret)r,rrrfake_arg^s       z)FakeCUDAKernel.__call__..fake_argcsg|] }|qSrr).0r3)r6rr usz+FakeCUDAKernel.__call__..)r"r r!rr r&r'rr)rr0ndindex BlockManagerr#run) rargsr&r'Zfake_cuda_moduleZ fake_args grid_pointZbmwbr)r6r,rr__call__Ns& 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)rrDrrr overloadsszFakeCUDAKernel.overloadscCs|jSr)r!rDrrrpy_funcszFakeCUDAKernel.py_funcN)rrr) rrrrr+r?rrErFrHpropertyrIrJrrrrr <s 1   r csTeZdZdZfddZfddZddZdd 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)r0Zseterr)r<kwargsfrr debug_wrappers z+BlockThread.__init__..debug_wrapper)targetFT)superrLr+ threadingEventsyncthreads_eventsyncthreads_blocked_managerrblockIdx threadIdx exceptiondaemonabortr* _block_dimxyzZ thread_id) rrQmanagerrZr[r*rRrSZblockDim __class__rPrr+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 ) rTrLr; Exceptionr$r[rZstrsysexc_infotyper\)retidZctaidmsgtbrdrrr;s  zBlockThread.runcCs:|jrtdd|_|j|j|jr6tddS)Nz"abort flag set on syncthreads callTz#abort flag set on syncthreads clear)r^ RuntimeErrorrXrWwaitclearrDrrr syncthreadss  zBlockThread.syncthreadscCsD|jj|jj|jjf}||jj|<|t|jj}||Sr) r[r`rarbrY block_statersr0Z count_nonzero)rvalueidxcountrrrsyncthreads_counts  zBlockThread.syncthreads_countcCsL|jj|jj|jjf}||jj|<|t|jj}||rHdSdSNrr) r[r`rarbrYrtrsr0allrrurvtestrrrsyncthreads_ands  zBlockThread.syncthreads_andcCsL|jj|jj|jjf}||jj|<|t|jj}||rHdSdSry) r[r`rarbrYrtrsr0anyr{rrrsyncthreads_ors  zBlockThread.syncthreads_orcCsd|j|jfS)NzThread <<<%s, %s>>>)rZr[rDrrr__str__szBlockThread.__str__) rrrrr+r;rsrxr}rr __classcell__rrrdrrLs   rLc@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_dimr__fr#r0zerosZbool_rt)rrQr&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<rrrrSsz BlockManager.run..targetTFrrcSsg|]}|r|qSr)is_alive)r7trrrr8.sz$BlockManager.run..) setr0r9r_rLr#startaddrXr\r^rWwith_traceback) rr=r<threadsZ livethreadsZblockedthreadsZ block_pointrSrZt_otherrrrr; s8     zBlockManager.runN)rrrrr+r;rrrrr:sr:) contextlibr functoolsrrirUZnumpyr0Zcudadrv.devicearrayrrZ kernelapirrr errorsr r<r r rrrrdictrobjectr ThreadrLr:rrrrs"     cS