ó Û³Y]c@s#ddlmZddlmZddlZddlZddlZddlm Z ddl m Z ddl m Z mZddlmZmZmZd d lmZd d lmZmZdaed „ƒZd „Zdefd„ƒYZdejfd„ƒYZdefd„ƒYZ dS(iÿÿÿÿ(tprint_function(tcontextmanagerN(tsix(treraisei(t to_devicet auto_device(tDim3tFakeCUDAModuletswapped_cuda_modulei(tnormalize_kernel_dimensions(twrap_argtArgHintccs5tdkstdƒ‚|az dVWddaXdS(s* Push the current kernel context. s'conrrent simulated kernel not supportedN(t_kernel_contexttNonetAssertionError(tmod((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyt_push_kernel_contexts  cCstS(sT Get the current kernel context. This is usually done by a device function. (R (((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyt_get_kernel_context'stFakeCUDAKernelcBsbeZdZegd„Zd„Zd„Zd„Zd„Zdddd„Z e d„ƒZ RS( s( Wraps a @cuda.jit-ed function. cCs6||_||_||_t|ƒ|_|ddS(Nii(iiii(tfnt_devicet _fastmathtlistt extensions(tselfRtdevicetfastmathR((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyt__init__3s    c sˆjr5tˆjtƒƒˆj|ŒSWdQXntˆjˆjˆjƒ}t|ƒµg‰‡‡fd†}g|D]}||ƒ^q|}tˆj|ƒMxEt j ˆjŒD]1}t ˆjˆjˆjƒ}|j ||ŒqºWWdQXxˆD] }|ƒqüWWdQXdS(Ncsˆtjj‡fd†ˆjd|fƒ\}}t|tjƒrd|jdkrdt |ƒj ˆƒSt|t ƒr€|j ˆƒS|SdS(Ncs|jdddˆ|ŒS(Ntstreamitretr(t prepare_args(tty_valt extension(R(s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pytKs i( RtmovestreduceRR t isinstancetnptndarraytndimR RR (targt_(RR(s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pytfake_argHs  ! ( RRRRRtgrid_dimt block_dimtdynshared_sizeRR%tndindext BlockManagertrun( Rtargstfake_cuda_moduleR*R(t fake_argst grid_pointtbmtwb((RRs:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyt__call__<s    cCsBt|d Œ\|_|_t|ƒdkr>|d|_n|S(Niii(R R+R,tlenR-(Rt configuration((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyt __getitem__escCsdS(N((R((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pytbindnscGs|S(N((RR1((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyt specializersicCs||d||fS(Ni((RtntasksttpbRt sharedmem((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pytforalluscCs*d}|d7}|jr&|d7}n|S(sl Required in order to proceed through some tests, but serves no functional purpose. s.consts .locals div.full.ftz.f32(R(Rtres((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pytptxxs    ( t__name__t __module__t__doc__tFalseRR7R:R;R<R@tpropertyRB(((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyR.s )  t BlockThreadcBsMeZdZd„Zd„Zd„Zd„Zd„Zd„Zd„Z RS(sG Manages the execution of a function for a single CUDA thread. cCsºtt|ƒjd|ƒtjƒ|_t|_||_t |Œ|_ t |Œ|_ d|_ t|_t|_t |jjŒ}|j j|j|j j|j|j j|_dS(Nttarget(tsuperRHRt threadingtEventtsyncthreads_eventRFtsyncthreads_blockedt_managerRtblockIdxt threadIdxR t exceptiontTruetdaemontabortt _block_dimtxtytzt thread_id(RtftmanagerRPRQtblockDim((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyRŽs     cCsÃytt|ƒjƒWn¥tk r¾}dt|jƒ}dt|jƒ}t|ƒdkrtd||f}nd|||f}tj ƒd}t |ƒt |ƒ|ƒ|f|_ nXdS(Nstid=%ssctaid=%sts%s %ss %s %s: %si( RJRHR0t ExceptionRRQRPtstrtsystexc_infottypeRR(Rtettidtctaidtmsgttb((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyR0›scCsW|jrtdƒ‚nt|_|jjƒ|jjƒ|jrStdƒ‚ndS(Ns"abort flag set on syncthreads calls#abort flag set on syncthreads clear(RUt RuntimeErrorRSRNRMtwaittclear(R((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyt syncthreads¨s     cCsX||jj|jj|jj|jjf<|jƒtj|jjƒ}|jƒ|S(N( ROt block_stateRQRWRXRYRlR%t count_nonzero(Rtvaluetcount((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pytsyncthreads_count´s +  cCsb||jj|jj|jj|jjf<|jƒtj|jjƒ}|jƒ|r^dSdS(Nii( RORmRQRWRXRYRlR%tall(RRottest((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pytsyncthreads_and»s +  cCsb||jj|jj|jj|jjf<|jƒtj|jjƒ}|jƒ|r^dSdS(Nii( RORmRQRWRXRYRlR%tany(RRoRs((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pytsyncthreads_orÂs +  cCsd|j|jfS(NsThread <<<%s, %s>>>(RPRQ(R((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyt__str__És( RCRDRERR0RlRqRtRvRw(((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyRHŠs   R/cBs eZdZd„Zd„ZRS(sç 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|dtjƒ|_dS(Ntdtype(t _grid_dimRVt_fR%tzerostboolRm(RR[R+R,((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyRßs   c s›tƒ}tƒ}tƒ}xetjˆjŒD]Q}‡‡fd†}t|ˆ||ƒ}|jƒ|j|ƒ|j|ƒq.Wxç|rlxp|D]h}|jr²|j|ƒq“|jr“x-|D]%} t | _ t | _| j jƒqÂWt |jŒq“q“W||kr>x$|D]}t |_|j jƒqWtƒ}ntg|D]}|jƒrH|^qHƒ}q†Wx'|D]}|jrtt |jŒqtqtWdS(NcsˆjˆŒdS(N(Rz((R1R(s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyRIës(tsetR%R.RVRHtstarttaddRNRRRSRURFRMRtis_alive( RR4R1tthreadst livethreadstblockedthreadst block_pointRItttt_other((R1Rs:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyR0ås8                /  (RCRDRERR0(((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyR/Ís (!t __future__Rt contextlibRRaRKtnumpyR%tnumbaRt numba.sixRtcudadrv.devicearrayRRt kernelapiRRRterrorsR R1R R R R RRtobjectRtThreadRHR/(((s:lib/python2.7/site-packages/numba/cuda/simulator/kernel.pyts     \C