\ K]c@@s~ddlmZmZddlZddlmZmZddlZddlZddl Z ddl Z ddl Z ddl mZddl mZmZmZmZddlmZmZddl mZmZmZmZddlmZd d lmZd d lm Z d d l!m"Z"m#Z#m$Z$d d l%m&Z&d dl'm(Z(d dl)m*Z*edZ+ee,e,e,ge-dZ.de/fdYZ0e,e,dZ1e2e,dZ3dZ4de/fdYZ5de/fdYZ6de/fdYZ7de/fdYZ8de/fd YZ9d!e/fd"YZ:d#e8fd$YZ;d%Z<d&e8fd'YZ=dS((i(tabsolute_importtprint_functionN(treducetwraps(tctypes_support(tconfigtcompilerttypestsigutils(tAbstractTemplatetConcreteTemplate(tfuncdescttypingtutilst serialize(tglobal_compiler_locki(t AutoTuner(t get_context(tnvvmt devicearraytdriver(tnormalize_kernel_dimensions(tget_current_device(twrap_argc C@sddlm}|j}|j}tj}|jd|jd|rk|jd|jdn|r|jdntjd|d |d |d |d |d |di} | j} | j | S(Ni(tCUDATargetDesct no_compiletno_cpython_wrappert boundcheckt debuginfot forceinlinet typingctxt targetctxtfunctargst return_typetflagstlocals( t descriptorRRRRtFlagstsett compile_extratlibrarytfinalize( tpyfuncR"R!tdebugtinlineRRRR#tcresR)((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt compile_cudas*         c C@st|tj|d|d|}|jj} |jj|j| |jj d|\} } t d| j d| j d|jj d|jj d|jd|d|d |jd |d |d | } | S( NR,R-t llvm_moduletnamet pretty_nametargtypesttype_annotationtlinkt call_helpertfastmatht extensionst max_registers(R/Rtvoidtfndesctllvm_func_namettarget_contexttprepare_cuda_kernelR)t signatureR!t CUDAKernelt _final_moduleR1tqualnameR4R6( R+R!R5R,R-R7R8R9R.tfnametlibtkerneltcukern((s2lib/python2.7/site-packages/numba/cuda/compiler.pytcompile_kernel;s"!        tDeviceFunctionTemplatecB@sMeZdZdZdZedZdZdZidZ RS(s#Unmaterialized device function cC@s(||_||_||_i|_dS(N(tpy_funcR,R-t _compileinfos(tselfR+R,R-((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt__init__Us   cC@sOtj|j}tj|j|}|j||j|jf}tj|fS(N(Rt#_get_function_globals_for_reductionRIt_reduce_functiont __class__R,R-t_rebuild_reduction(RKtglblst func_reducedR!((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt __reduce__[scC@s%tj|}t|d|d|S(NR,R-(Rt_rebuild_functiontcompile_device_template(tclsRRR,R-R ((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt_rebuildascC@s||jkrt|jd|d|jd|j}|j }||j|<|jg}|r{|jj||j |q|jj ||j |n |j|}|S(sCompile the function for the given argument types. Each signature is compiled once by caching the compiled function inside this object. Returns the `CompileResult`. R,R-N( RJR/RItNoneR,R-R)R=tinsert_user_functionR;tadd_user_function(RKR!R.tfirst_definitiontlibs((s2lib/python2.7/site-packages/numba/cuda/compiler.pytcompilefs      cC@s#|j|}|jj}t|S(sReturns the LLVM-IR text compiled for *args*. Parameters ---------- args: tuple[Type] Argument types. Returns ------- llvmir : str (RJR)RAtstr(RKR!R.tmod((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt inspect_llvms  c C@s[|j|}t}|j}|j}tj|}tj|ddd||}|S(sLReturns the PTX compiled for *args* for the currently active GPU Parameters ---------- args: tuple[Type] Argument types. nvvm_options : dict; optional See `CompilationUnit.compile` in `numba/cuda/cudadrv/nvvm.py`. Returns ------- ptx : bytes toptitarch(R`Rtdevicetcompute_capabilityRtget_arch_optiont llvm_to_ptx( RKR!t nvvm_optionstllvmirtcuctxRctccRbtptx((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt inspect_ptxs   ( t__name__t __module__t__doc__RLRSt classmethodRWR]R`Rl(((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRHRs    c@saddlm}t|d|d|dtffdY}|j}|j|S(scCreate a DeviceFunctionTemplate object and register the object to the CUDA typing context. i(RR,R-tdevice_function_templatec@seZZfdZRS(c@s| s tj|jS(N(tAssertionErrorR]R?(RKR!tkws(tdft(s2lib/python2.7/site-packages/numba/cuda/compiler.pytgenerics (RmRntkeyRu((Rt(s2lib/python2.7/site-packages/numba/cuda/compiler.pyRqs(R%RRHR RRY(R+R,R-RRqR((Rts2lib/python2.7/site-packages/numba/cuda/compiler.pyRUs  cC@st|||dtdtS(NR-R,(tDeviceFunctiontTruetFalse(R+R"R!R-R,((s2lib/python2.7/site-packages/numba/cuda/compiler.pytcompile_devicesc@sddlm}|j}|j}tj||t|dtffdY}tj d|d|d|}|j ||j |S(Ni(RRqc@seZZgZRS((RmRnRvtcases((textfntsig(s2lib/python2.7/site-packages/numba/cuda/compiler.pyRqsR1trestypeR3( R%RRRR R?tExternFunctionR R tExternalFunctionDescriptorRY(R1R~R3RRRRqR;((R|R}s2lib/python2.7/site-packages/numba/cuda/compiler.pytdeclare_device_functions   RwcB@s2eZdZdZedZdZRS(c@s|_|_|_t_t_tjjjdjdj_dt ffdY}j j |j j j jgdS(NR,R-Rqc@seZZjgZRS((RmRnRvR?R{((R.RK(s2lib/python2.7/site-packages/numba/cuda/compiler.pyRqs(RIR"R!RxR-RyR,R/R.R ttyping_contextRYR=R;R)(RKR+R"R!R-R,Rq((R.RKs2lib/python2.7/site-packages/numba/cuda/compiler.pyRLs        cC@s[tj|j}tj|j|}|j||j|j|j|jf}tj |fS(N( RRMRIRNROR"R!R-R,RP(RKtglobsRRR!((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRSs cC@s|tj|||||S(N(RRT(RVRRR"R!R-R,((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRWscC@sd}|j|j|jjS(Ns*(tformatRIR.R?(RKtfmt((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt__repr__s(RmRnRLRSRpRWR(((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRws  RcB@seZdZRS(cC@s||_||_dS(N(R1R}(RKR1R}((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRLs (RmRnRL(((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRstForAllcB@s,eZdZdZdZdZRS(cC@s1||_||_||_||_||_dS(N(REtntaskstthread_per_blocktstreamt sharedmem(RKRERttpbRR((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRLs     cG@st|jtr'|jj|}n |j}|j|}|d}|j||}|j||d|jd|j|S(NiRR( t isinstanceREtAutoJitCUDAKernelt specializet_compute_thread_per_blockRt configureRR(RKR!RERttpbm1tblkct((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt__call__ s  c C@s|j}|dkr|St}td|jjddd|jdd}y|j|\}}Wn#tk r|j|}nX|SdS(NiR tb2d_funccS@sdS(Ni((R((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt ttmemsizetblocksizelimiti( RRtdictt_functgetRtget_max_potential_block_sizetAttributeErrort_fallback_autotune_best(RKRERtctxtkwargst_((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRs       cC@s>y|jj}Wn$tk r9tjdd}nX|S(Ns,Could not autotune, using default tpb of 128i(tautotunetbestt ValueErrortwarningstwarn(RKRER((s2lib/python2.7/site-packages/numba/cuda/compiler.pyR-s    (RmRnRLRRR(((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRs  tCUDAKernelBasecB@s\eZdZdZdZdddZdZddddZdZdZ RS( s.Define interface for configurable kernels cC@s(d|_d|_d|_d|_dS(Nii(ii(iii(tgriddimtblockdimRR(RK((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRL;s   cC@s/|j}|j|}|jj|j|S(s+ Shallow copy the instance (ROt__new__t__dict__tupdate(RKRVtnew((s2lib/python2.7/site-packages/numba/cuda/compiler.pytcopyAs icC@sUt||\}}|j}t||_t||_||_||_|S(N(RRttupleRRRR(RKRRRRtclone((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRMs   cC@s.t|dkr!tdn|j|S(Niiis.must specify at least the griddim and blockdim(iii(tlenRR(RKR!((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt __getitem__Wsc C@st||d|d|d|S(sReturns a configured kernel for 1D kernel of given number of tasks ``ntasks``. This assumes that: - the kernel 1-to-1 maps global thread id ``cuda.grid(1)`` to tasks. - the kernel must check if the thread id is valid.RRR(R(RKRRRR((s2lib/python2.7/site-packages/numba/cuda/compiler.pytforall\scC@s|j|j|jfS(s Helper for serializing the grid, block and shared memory configuration. CUDA stream config is not serialized. (RRR(RK((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt_serialize_configfscC@s|\|_|_|_dS(sc Helper for deserializing the grid, block and shared memory configuration. N(RRR(RKR((s2lib/python2.7/site-packages/numba/cuda/compiler.pyt_deserialize_configms( RmRnRoRLRRRRRR(((s2lib/python2.7/site-packages/numba/cuda/compiler.pyR7s    t CachedPTXcB@s eZdZdZdZRS(s<A PTX cache that uses compute capability as a cache key cC@s.||_||_i|_|j|_dS(N(R1RhtcacheRt_extra_options(RKR1Rhtoptions((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRLxs   cC@st}|j}|j}|jj|}|d krtj|}tj|j ddd||j }||j| IsJThe .{} attribute is is deprecated and will be removed in a future releaseRcB@seZdZdZedZdZdZdZd d dZ d d dZ d dZ e d Zd ZRS( s CUDA Kernel object. When called, the kernel object will specialize itself for the given arguments (if no suitable specialized version already exists) & compute capability, and launch on the device associated with the current context. Kernel objects are not to be constructed by the user, but instead are created using the :func:`numba.cuda.jit` decorator. cC@sytt|j||_||_i|_||_t|jjdg|jdtg|D]}|jj|^q }|j|}|S(s| Compile and bind to the current context a version of this kernel specialized for the given *args*. (RRtresolve_argument_typeR](RKR!taR3RE((s2lib/python2.7/site-packages/numba/cuda/compiler.pyR*s(cC@stj|\}}|dks'ttj}|jj||f}|dkrd|jkrvd|jdPsN(RRdRXRTR`Rtitems(RKR?RdRj((s2lib/python2.7/site-packages/numba/cuda/compiler.pyR`Gs   cC@sV|ptj}|dk r5|j||fjStd|jjDSdS(s Return the generated assembly code for all signatures encountered thus far, or the LLVM IR for a specific signature and compute_capability if given. cs@s'|]\}}||jfVqdS(N(R(RZR}R[((s2lib/python2.7/site-packages/numba/cuda/compiler.pys ]sN(RRdRXRTRRR\(RKR?RdRj((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRSs   cC@sO|dkrtj}nx0tj|jD]\}}|jd|q+WdS(s Produce a dump of the Python source of this function annotated with the corresponding Numba IR and type information. The dump is written to *file*, or *sys.stdout* if *file* is *None*. RN(RXRRR t iteritemsRTR(RKRRR[((s2lib/python2.7/site-packages/numba/cuda/compiler.pyR`s  cC@s2tj|}||||}|j||S(s& Rebuild an instance. (RRTR(RVRRRRURR R((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRWls cC@s^tj|j}tj|j|}|j}|j||j|j|f}tj|fS(sd Reduce the instance for serialization. Compiled definitions are discarded. ( RRMRIRNRRORRURP(RKRQRRRR!((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRSvs   N(RmRnRoRLRSR8RRR]RXR`RRRpRWRS(((s2lib/python2.7/site-packages/numba/cuda/compiler.pyRs       (>t __future__RRRt functoolsRRRPRt threadingRtnumpyR6tnumbaRRRRRRtnumba.typing.templatesR R R R R Rtnumba.compiler_lockRtcudadrv.autotuneRtcudadrv.devicesRtcudadrvRRRterrorsRtapiRR!RR/RyRXRGtobjectRHRURxRzRRwRRRRRR@RHR(((s2lib/python2.7/site-packages/numba/cuda/compiler.pytsH      "" W &6>A