\ K]c@sdZddlmZddlZddlZddlZddlZddlm Z de fdYZ de fdYZ d e fd YZ d e fd YZejZejZejZejZd e fdYZde fdYZedZdS(sf Implements the cuda module as called from within an executing kernel (@cuda.jit-decorated function). i(tcontextmanagerN(t numpy_supporttDim3cBs2eZdZdZdZdZdZRS(s; Used to implement thread/block indices/dimensions cCs||_||_||_dS(N(txtytz(tselfRRR((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyt__init__s  cCsd|j|j|jfS(Ns (%s, %s, %s)(RRR(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyt__str__scCsd|j|j|jfS(NsDim3(%s, %s, %s)(RRR(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyt__repr__sccs|jV|jV|jVdS(N(RRR(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyt__iter__s(t__name__t __module__t__doc__RRR R (((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyRs    t FakeCUDALocalcBseZdZdZRS(s CUDA Local arrays cCstj|}tj||S(N(Rtas_dtypetnptempty(Rtshapetdtype((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytarray)s(R R R R(((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyR%st FakeCUDAConstcBseZdZdZRS(s CUDA Const arrays cCs|S(N((Rtary((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyt array_like2s(R R R R(((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyR.stFakeCUDASharedcBs eZdZdZdZRS(s CUDA Shared arrays. Limitations: assumes that only one call to cuda.shared.array is on a line, and that that line is only executed once per thread. i.e.:: a = cuda.shared.array(...); b = cuda.shared.array(...) will erroneously alias a and b, and:: for i in range(10): sharedarrs[i] = cuda.shared.array(...) will alias all arrays created at that point (though it is not certain that this would be supported by Numba anyway). cCs1i|_||_tj|dtj|_dS(NR(t _allocationst_dynshared_sizeRtzerostbytet _dynshared(Rtdynshared_size((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyRHs  cCstj|}|dkrJ|j|j}tj|jjd|d|Stj t j }|ddd!}|j j |}|dkrtj||}||j |(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytlocalscCs|jS(N(R?(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytsharedscCs|jS(N(R@(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytconstscCs|jS(N(RA(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytatomicscCs tjjS(N(t threadingtcurrent_threadt threadIdx(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyRJscCs tjjS(N(RHRItblockIdx(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyRKscCsdS(Ni ((R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytwarpsizescCstjjdS(Ni (RHRIt thread_id(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytlaneidscCstjjdS(N(RHRIt syncthreads(R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyROscCsdS(N((R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyt threadfencescCsdS(N((R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytthreadfence_blockscCsdS(N((R((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytthreadfence_systemscCstjj|S(N(RHRItsyncthreads_count(RR/((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyRSscCstjj|S(N(RHRItsyncthreads_and(RR/((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyRTscCstjj|S(N(RHRItsyncthreads_or(RR/((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyRUscCst|jdS(Nt1(tbinR(RR/((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytpopcscCs |||S(N((Rtatbtc((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytfmascCs#tdj|ddddS(Ns{:032b}ii(tinttformat(RR/((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytbrevscCs,dj|}t|t|jdS(Ns{:032b}t0(R^tlentlstrip(RR/ts((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytclzscCs,dj|}t|t|jdS(Ns{:032b}R`(R^Ratrstrip(RR/Rc((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytffsscCs|r |S|S(N((RRYRZR[((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytselpscCs|j}|j}|j}|j|j|j}|dkrB|S|j|j|j}|dkro||fS|j|j|j}|dkr|||fStd|dS(Niiis*Global ID has 1-3 dimensions. %d requested(R=RKRJRRRt RuntimeError(RtntbdimtbidttidRRR((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytgrids        cCs|j}|j}|j|j}|dkr2|S|j|j}|dkrX||fS|j|j}|dkr|||fStd|dS(Niiis,Global grid has 1-3 dimensions. %d requested(R=R<RRRRh(RRiRjtgdimRRR((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytgridsizes       (R R R RtpropertyRDRERFRGRJRKRLRNRORPRQRRRSRTRURXR\R_RdRfRgRmRo(((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyR;s2                c#sddlm|j}tfd|jD}tfd|jD}|j|z dVWd|j|XdS(Ni(tcudac3s-|]#\}}|kr||fVqdS(N((t.0tktv(Rq(s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pys sc3s!|]\}}|fVqdS(N((RrRsRt(tfake_cuda_module(s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pys s(tnumbaRqt __globals__tdicttitemstupdate(tfnRutfn_globstorigtrepl((RqRus=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pytswapped_cuda_module s ""  (R t contextlibRR%RHR#tnumpyRRvRtobjectRRRRtLockR-R1R4R6R,R;R(((s=lib/python2.7/site-packages/numba/cuda/simulator/kernelapi.pyts"     -    $