o 8ήc6)@sddlmZddlmZddlZddlZddlZddlm Z m Z ddl m Z m Z mZddlmZdd lmZmZ daed d Zd d ZGdddZGdddeZGdddeZGdddejZGdddeZdS))contextmanager)reduceN) FakeCUDAArrayFakeWithinKernelCUDAArray)Dim3FakeCUDAModuleswapped_cuda_module)normalize_kernel_dimensions)wrap_argArgHintccs.tdus Jd|azdVWdadSdaw)z* Push the current kernel context. Nz'conrrent simulated kernel not supported_kernel_context)modrB/tmp/pip-target-vg8gfxp4/lib/python/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)selfblockdimrrrmax_cooperative_grid_blocks/sz(FakeOverload.max_cooperative_grid_blocksN)__name__ __module__ __qualname____doc__rrrrrr+s rc@seZdZddZdS)FakeOverloadDictcCtSN)r)rkeyrrr __getitem__6szFakeOverloadDict.__getitem__N)rrrr!rrrrr5s rc@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_device _fastmath_debuglist extensionsgrid_dim block_dimstreamdynshared_size)rr#devicefastmathr(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|tr-|}n t|tj r8t |}n|}t|t rCt |S|S)Ncs|j|ddS)Nr)r+retr) prepare_args)ty_val extension)r1rras z;FakeCUDAKernel.__call__..fake_arg..r) rr( isinstancenpndarrayndimr to_devicer voidrr)arg_ret)r1rrrfake_arg^s       z)FakeCUDAKernel.__call__..fake_argcsg|]}|qSrr).0r<)r?rr usz+FakeCUDAKernel.__call__..)r$r r#rr r)r*rr,rr7ndindex BlockManagerr&run) rargsr)r*fake_cuda_module fake_args grid_pointbmwbr)r?r1rr__call__Ns0  "zFakeCUDAKernel.__call__cCs2t|dd\|_|_t|dkr|d|_|S)Nr )r r)r*lenr,)r configurationrrrr!s   zFakeCUDAKernel.__getitem__cCsdSrrrrrrbindzFakeCUDAKernel.bindcGs|Srr)rrErrr specializerRzFakeCUDAKernel.specializercCs$|dkr td|||d||fS)Nrz0Can't create ForAll with negative task count: %sr) ValueError)rntaskstpbr+ sharedmemrrrforalls zFakeCUDAKernel.forallcCrr)rrPrrr overloadszFakeCUDAKernel.overloadscCs|jSr)r#rPrrrpy_funcrZzFakeCUDAKernel.py_funcN)rrr) rrrrr0rKr!rQrSrXpropertyrYr[rrrrr"<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|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)r7seterr)rEkwargsfrr debug_wrappers z+BlockThread.__init__..debug_wrapper)targetFT)superr]r0 threadingEventsyncthreads_eventsyncthreads_blocked_managerrblockIdx threadIdx exceptiondaemonabortr/ _block_dimxyz thread_id) rrcmanagerrlrmr/rdreblockDim __class__rbrr0s(     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 ) rfr]rD Exceptionr'rmrlstrsysexc_infotypern)retidctaidmsgtbrxrrrDs   zBlockThread.runcCs:|jrtdd|_|j|j|jrtddS)Nz"abort flag set on syncthreads callTz#abort flag set on syncthreads clear)rp RuntimeErrorrjriwaitclearrPrrr syncthreadss  zBlockThread.syncthreadscCsD|jj|jj|jjf}||jj|<|t|jj}||Sr) rmrrrsrtrk block_staterr7 count_nonzero)rvalueidxcountrrrsyncthreads_counts  zBlockThread.syncthreads_countcCL|jj|jj|jjf}||jj|<|t|jj}||r$dSdSNrr) rmrrrsrtrkrrr7allrrrtestrrrsyncthreads_and   zBlockThread.syncthreads_andcCrr) rmrrrsrtrkrrr7anyrrrrsyncthreads_orrzBlockThread.syncthreads_orcCsd|j|jfS)NzThread <<<%s, %s>>>)rlrmrPrrr__str__szBlockThread.__str__) rrrrr0rDrrrrr __classcell__rrrxrr]s   r]c@s eZdZdZddZddZdS)rCa 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)dtype) _grid_dimrq_fr&r7zerosbool_r)rrcr)r*r/rrrr0s 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)rrrErrrresz BlockManager.run..targetTFrrcSsg|]}|r|qSr)is_alive)r@trrrrA.sz$BlockManager.run..) setr7rBrqr]r&startaddrjrnrpriwith_traceback) rrHrEthreads livethreadsblockedthreads block_pointrert_otherrrrrD s@      zBlockManager.runN)rrrrr0rDrrrrrCs rC) contextlibr functoolsrr}rgnumpyr7cudadrv.devicearrayrr kernelapirrr errorsr rEr r rrrrdictrobjectr"Threadr]rCrrrrs&     cS