o 8ήcw@sddlZddlZddlZddlZddlZddlmZmZm Z m Z m Z m Z ddl mZmZddlmZddlmZddlmZddlmZmZddlmZdd lmZdd lmZmZdd l m!Z!dd l"m#Z#dd l$m%Z%ddl&m'Z'ddl(m)Z)m*Z*ddl+m Z,ddl-m.Z.ddl-m/Z/ddl0m1Z1Gdddej2Z3Gddde4Z5GdddZ6GdddeZ7GdddeZ8Gdddeej2Z9dS) N)config serializesigutilstypestypingutils)Cache CacheImpl)global_compiler_lock) Dispatcher)NumbaPerformanceWarning)Purposetypeofget_current_device)wrap_arg) compile_cuda CUDACompiler)driver) get_context) get_cudalib) cuda_target)missing_launch_config_msgnormalize_kernel_dimensions)rcuda) _dispatcher)warncseZdZdZe   d*fdd ZeddZed d Zd d Z ed dZ eddZ e fddZ ddZddZeddZeddZddZddZdd Zd+d!d"Zd,d$d%Zd-d&d'Zd(d)ZZS)._Kernelz CUDA Kernel specialized for a given set of argument types. When called, this object launches the kernel on the device. NFTc  sL| rtdtd|_d|_||_||_||_||_|p g|_ |j|j|| r+dndd} t |jt j |j|j|j||| d} | j }|jj}|j}|j}|| j| j|| ||| \}}|s`g}d|v|_|jrs|tdd d |D]}||qu|j|_| j|_| j|_||_| j|_||_ | j|_| j |_ g|_!g|_"g|_#dS) Nz,Cannot compile a device function as a kernelFr)debuglineinfofastmathopt)r r!inliner" nvvm_optionscudaCGGetIntrinsicHandle cudadevrtT)static)$ RuntimeErrorsuper__init__ objectmode entry_pointpy_funcargtypesr r! extensionsrrvoidtarget_context__code__ co_filenameco_firstlinenoprepare_cuda_kernellibraryfndesc get_asm_str cooperativeappendradd_linking_filename entry_name signaturetype_annotation_type_annotation _codelibrary call_helper environment_referenced_environmentslifted reload_init)selfr.r/linkr r!r$r"r0 max_registersr#devicer%crestgt_ctxcodefilenamelinenumlibkernelfilepath __class__r?rArBr r!rCr0) clsr:r=r? codelibraryr r!rCr0instancerTrVrW_rebuilds z_Kernel._rebuildc Cs(t|j|j|j|j|j|j|j|jdS)a Reduce the instance for serialization. Compiled definitions are serialized in PTX form. Type annotation are discarded. Thread, block and shared memory configuration are serialized. Stream information is discarded. )r:r=r?rfr r!rCr0) dictr:r>r?rBr r!rCr0rZrVrVrW_reduce_statess z_Kernel._reduce_statescCs|jdS)z7 Force binding to current CUDA context N)rB get_cufuncrZrVrVrWbindz _Kernel.bindcCstS)z, Get current active context rrZrVrVrWrKz_Kernel.devicecCs|jjjS)zN The number of registers used by each thread for this kernel. )rBrkattrsregsrZrVrVrWregs_per_threadsz_Kernel.regs_per_threadcCr])z6 Returns the LLVM IR for this kernel. )rB get_llvm_strrZrVrVrW inspect_llvms z_Kernel.inspect_llvmcCs|jj|dS)z7 Returns the PTX code for this kernel. cc)rBr9)rHrurVrVrW inspect_asmrmz_Kernel.inspect_asmcCr])zp Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )rBget_sassrZrVrVrW inspect_sasss z_Kernel.inspect_sasscCsb|jdur td|durtj}td|j|jf|dtd|dt|j|dtd|ddS) 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*. Nz Type annotation is not availablez%s %sfilezP--------------------------------------------------------------------------------zP================================================================================)rA ValueErrorsysstdoutprintr>rb)rHr{rVrVrW inspect_typess  z_Kernel.inspect_typesrcCsHt}|j}t|trtdd|}||||}|jj }||S)a Calculates the maximum number of blocks that can be launched for this kernel in a cooperative grid in the current context, for the given block and dynamic shared memory sizes. :param blockdim: Block dimensions, either as a scalar for a 1D block, or a tuple for 2D or 3D blocks. :param dynsmemsize: Dynamic shared memory size in bytes. :return: The maximum number of blocks in the grid. cSs||SrYrV)xyrVrVrWsz5_Kernel.max_cooperative_grid_blocks..) rrBrk isinstancer` functoolsreduce$get_active_blocks_per_multiprocessorrKMULTIPROCESSOR_COUNT)rHblockdim dynsmemsizectxcufunc active_per_smsm_countrVrVrWmax_cooperative_grid_blockss  z#_Kernel.max_cooperative_grid_blockscs|j|jr*jd}j|\}}|ttjksJt} |j d|dg} g} t |j |D] \} } | | | || | q4t jrLt jd}nd}|rS|jpT|}t jjg||||| Rd|ji|jrt t| ||| jdkr݇fddfddd D}fd dd D}| j}|j|\}}}|durd }n|\}}}tj|}d |||f}d |||f}|rd||dff|dd}|||f}||| D]}|qdS)N __errcode__r)streamr:cs<jdj|f\}}t}tt||||jS)Nz%s__%s__) moduleget_global_symbolr=ctypesc_intrdevice_to_host addressofvalue)r=memszval)rrVrW load_symbol#s z#_Kernel.launch..load_symbolcg|]}d|qS)tidrV.0irrVrW +z"_Kernel.launch..zyxcr)ctaidrVrrrVrWr,rz"In function %r, file %s, line %s, z%stid=%s ctaid=%sz%s: %s)rBrkr r=rrrsizeofrmemsetziprb _prepare_argsrUSE_NV_BINDINGbindingCUstreamhandle launch_kernelr:rrrrC get_exceptionospathabspath)rHragriddimrr sharedmemexcnameexcmemexcszexcvalretr kernelargstv zero_stream stream_handlerrrNexcclsexc_argsloclocinfosymrSlinenoprefixwbrV)rrrWlaunchsn        z_Kernel.launchcCsXt|jD] }|j||||d\}}qt|tjrt|||}tj }t d} t d} ||j } ||j j } t|} tjrEt| } t | }|| || || || ||t|jD] }|||j|qht|jD] }|||j|qzdSt|tjrttd||}||dS|tjkrtt|tj}||dS|tjkrt|}||dS|tj krt!|}||dS|tj"krt#t|}||dS|tj$kr|t!|j%|t!|j&dS|tj'kr |t|j%|t|j&dSt|tj(tj)fr8|t*|tj+dSt|tj,r\t|||}|j-} tjrUt t| } || dSt|tj.rt/|t/|ksnJt0||D]\}}|1|||||qsdSt|tj2rz|1|j |j3|||WdSt4yt4||wt4||)zF Convert arguments to ctypes and append to kernelargs )rrrzc_%sN)5reversedr0 prepare_argsrrArrayr to_devicer c_ssize_tc_void_psizedtypeitemsizerdevice_pointerrintr;rangendimshapestridesIntegergetattrfloat16c_uint16npviewuint16float64c_doublefloat32c_floatbooleanc_uint8 complex64realimag complex128 NPDatetime NPTimedeltac_int64int64Recorddevice_ctypes_pointer BaseTuplelenrr EnumMemberrNotImplementedError)rHtyrrrr extensiondevaryc_intpmeminfoparentnitemsrptrdataaxcvaldevrecrrrVrVrWrEs                        z_Kernel._prepare_args) NFFFFNNTFrY)rrr)__name__ __module__ __qualname____doc__r r+propertyr7r@r\r^rb classmethodrhrjrlrKrqrsrvrxrrrr __classcell__rVrVrTrWrs>R          Hrc@$eZdZddZddZddZdS)ForAllcCs6|dkr td|||_||_||_||_||_dS)Nrz0Can't create ForAll with negative task count: %s)r| dispatcherntasksthread_per_blockrr)rHr r tpbrrrVrVrWr+s zForAll.__init__cGs^|jdkrdS|jjr|j}n|jj|}||}|j|d|}||||j|jf|S)Nrr)r r  specialized specialize_compute_thread_per_blockrr)rHrarrrrVrVrW__call__s    zForAll.__call__cCsZ|j}|dkr |St}tt|j}t|jd|j dd}|j di|\}}|S)Nri)funcb2d_funcmemsizeblocksizelimitrV) r rnextiter overloadsvaluesrirBrkrget_max_potential_block_size)rHr rrrRkwargs_rVrVrWrsz ForAll._compute_thread_per_blockN)rrrr+rrrVrVrVrWr s r c@seZdZddZddZdS)_LaunchConfigurationc Csl||_||_||_||_||_tjr2d}|d|d|d}||kr4d|d}tt|dSdSdS)Nrrz Grid size zB will likely result in GPU under-utilization due to low occupancy.) r rrrrrCUDA_LOW_OCCUPANCY_WARNINGSrr ) rHr rrrr min_grid_size grid_sizemsgrVrVrWr+s   z_LaunchConfiguration.__init__cGs|j||j|j|j|jSrY)r callrrrrrHrarVrVrWrsz_LaunchConfiguration.__call__N)rrrr+rrVrVrVrWrs rc@r ) CUDACacheImplcCs|SrY)rj)rHrRrVrVrWrszCUDACacheImpl.reducecCstjdi|S)NrV)rrh)rHr2payloadrVrVrWrebuildzCUDACacheImpl.rebuildcCsdS)NTrV)rHrLrVrVrWcheck_cachableszCUDACacheImpl.check_cachableN)rrrrr)r+rVrVrVrWr's r'c@seZdZdZeZdS) CUDACachezS Implements a cache that saves and loads CUDA kernels and compile results. N)rrrrr' _impl_classrVrVrVrWr,sr,cseZdZdZdZeZeffdd Ze ddZ ddZ e j d d d9d d ZddZd:ddZe ddZddZddZddZddZddZe ddZd;d!d"Zd#d$Zd%d&Zd'd(Zd)d*Zd;d+d,Zd;d-d.Zd;d/d0Zd;d1d2Z d3d4Z!e"d5d6Z#d7d8Z$Z%S)<CUDADispatchera CUDA Dispatcher object. When configured and called, the dispatcher 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. Dispatcher objects are not to be constructed by the user, but instead are created using the :func:`numba.cuda.jit` decorator. Fcs"tj|||dd|_i|_dS)N) targetoptionspipeline_classF)r*r+ _specializedspecializations)rHr.r/r0rTrVrWr+s  zCUDADispatcher.__init__cCs t|SrY) cuda_typesr.rZrVrVrW _numba_type_*r_zCUDADispatcher._numba_type_cCst|j|_dSrY)r,r._cacherZrVrVrWenable_caching.r*zCUDADispatcher.enable_cachingr)maxsizercCst||\}}t|||||SrY)rr)rHrrrrrVrVrW configure1szCUDADispatcher.configurecCst|dvr td|j|S)N)r rz.must specify at least the griddim and blockdim)rr|r8r&rVrVrW __getitem__6s  zCUDADispatcher.__getitem__cCst|||||dS)a3Returns a 1D-configured dispatcher for a given number of tasks. This assumes that: - the kernel maps the Global Thread ID ``cuda.grid(1)`` to tasks on a 1-1 basis. - the kernel checks that the Global Thread ID is upper-bounded by ``ntasks``, and does nothing if it is not. :param ntasks: The number of tasks. :param tpb: The size of a block. An appropriate value is chosen if this parameter is not supplied. :param stream: The stream on which the configured dispatcher will be launched. :param sharedmem: The number of bytes of dynamic shared memory required by the kernel. :return: A configured dispatcher, ready to launch on a set of arguments.)rrr)r )rHr rrrrVrVrWforall;szCUDADispatcher.forallcCs |jdS)aS A list of objects that must have a `prepare_args` function. When a specialized kernel is called, each argument will be passed through to the `prepare_args` (from the last object in this list to the first). The arguments to `prepare_args` are: - `ty` the numba type of the argument - `val` the argument value itself - `stream` the CUDA stream used for the current call to the kernel - `retr` a list of zero-arg functions that you may want to append post-call cleanup work to. The `prepare_args` function must return a tuple `(ty, val)`, which will be passed in turn to the next right-most `extension`. After all the extensions have been called, the resulting `(ty, val)` will be passed into Numba's default argument marshalling logic. r0)r/getrZrVrVrWr0Qs zCUDADispatcher.extensionscOsttrY)r|r)rHrarrVrVrWrfszCUDADispatcher.__call__cCsD|jr tt|j}n tjj|g|R}||||||dS)zJ Compile if necessary and invoke this kernel with *args*. N) rrrrrrr _cuda_callr)rHrarrrrrRrVrVrWr%jszCUDADispatcher.callcs(|rJfdd|D}t|S)Ncsg|]}|qSrV) typeof_pyvalrarZrVrWrxsz4CUDADispatcher._compile_for_args..)compiler`)rHrakwsr/rVrZrW_compile_for_argsusz CUDADispatcher._compile_for_argscCsDzt|tjWSty!t|r ttj|ddtjYSw)NF)sync)rr argumentr|r is_cuda_array as_cuda_array)rHrrVrVrWr>{s  zCUDADispatcher.typeof_pyvalcstj}tfdd|D}jrtdj||f}|r"|Sj}tj |d}| || d|_ |j||f<|S)zd Create a new instance of this dispatcher specialized for the given *args*. csg|]}j|qSrV) typingctxresolve_argument_typer?rZrVrWrrz-CUDADispatcher.specialize..zDispatcher already specialized)r/T) rcompute_capabilityr`rr)r2r<r/r.r.rAdisable_compiler1)rHrarur/specializationr/rVrZrWrs$ zCUDADispatcher.specializecCrX)z> True if the Dispatcher has been specialized. )r1rZrVrVrWrrnzCUDADispatcher.specializedNcCsD|dur |j|jjS|jrtt|jjSdd|jDS)a Returns the number of registers used by each thread in this kernel for the device in the current context. :param signature: The signature of the compiled kernel to get register usage for. This may be omitted for a specialized kernel. :return: The number of registers used by the compiled variant of the kernel for the given signature and current device. NcSsi|]\}}||jqSrV)rqrsigoverloadrVrVrW s z6CUDADispatcher.get_regs_per_thread..)rrarqrrrritemsrHr?rVrVrWget_regs_per_threads z"CUDADispatcher.get_regs_per_threadcCsP|jr |t||jj}d|}tj|||jd}t |j}||||fS)z Get a typing.ConcreteTemplate for this dispatcher and the given *args* and *kws* types. This allows resolution of the return type. A (template, pysig, args, kws) tuple is returned. zCallTemplate({0}))key signatures) _can_compilecompile_devicer`r.rformatrmake_concrete_templatenopython_signaturesr pysignature)rHrarB func_namer= call_templatepysigrVrVrWget_call_templates   z CUDADispatcher.get_call_templatec Cs||jvrX|jF|jd}|jd}|jd}||jdr$dnd|d}t|jd|||||d }||j|<|j|j|j |j gWd|S1sQwY|S|j|}|S) zCompile the device function for the given argument types. Each signature is compiled once by caching the compiled function inside this object. Returns the `CompileResult`. r r$r"r#rr)r r#r"N)r r$r"r%) r_compiling_counterr/r<rr.r2insert_user_functionr-r8r7)rHrar r$r"r%rLrVrVrWrWs4         zCUDADispatcher.compile_devicecCs,dd|D}|j||dd||j|<dS)NcSsg|]}|jqSrV)_coder?rVrVrWrsz/CUDADispatcher.add_overload..Tr)_insertr)rHrRr/c_sigrVrVrW add_overloadszCUDADispatcher.add_overloadcCst|\}}|dus|tjksJ|jrtt|jS|j |}|dur*|S|j ||j }|dur@|j |d7<n&|j|d7<|jsPtdt|j|fi|j}||j ||||||S)z Compile and bind to the current context a version of this kernel specialized for the given signature. NrzCompilation disabled)rnormalize_signaturernonerrrrrr<r5 load_overload targetctx _cache_hits _cache_missesrVr)rr.r/rl save_overloadre)rHrNr/ return_typerRrVrVrWrAs$  zCUDADispatcher.compilecCsb|jd}|dur|r|j|jS|j|S|r'dd|jDSdd|jDS)z Return the LLVM IR for this kernel. :param signature: A tuple of argument types. :return: The LLVM IR for the given signature, or a dict of LLVM IR for all previously-encountered signatures. rKNcSsi|] \}}||jqSrV)r7rrrMrVrVrWrP4z/CUDADispatcher.inspect_llvm..cSi|] \}}||qSrV)rsrMrVrVrWrP7)r/r<rr7rrrsrQrHr?rKrVrVrWrs#s zCUDADispatcher.inspect_llvmcsvtj|jd}|dur!|r|j|jS|j|S|r/fdd|jDSfdd|jDS)a+ Return this kernel's PTX assembly code for for the device in the current context. :param signature: A tuple of argument types. :return: The PTX code for the given signature, or a dict of PTX codes for all previously-encountered signatures. rKNcsi|] \}}||jqSrV)r7r9rMrtrVrWrPLsz.CUDADispatcher.inspect_asm..csi|] \}}||qSrV)rvrMrtrVrWrPOrn) rrJr/r<rr7r9rvrQrqrVrtrWrv:s   zCUDADispatcher.inspect_asmcCs>|jdr td|dur|j|Sdd|jDS)a Return this kernel's SASS assembly code for for the device in the current context. :param signature: A tuple of argument types. :return: The SASS code for the given signature, or a dict of SASS codes for all previously-encountered signatures. SASS for the device in the current context is returned. Requires nvdisasm to be available on the PATH. rKz(Cannot inspect SASS of a device functionNcSrorV)rx)rrNdefnrVrVrWrPerpz/CUDADispatcher.inspect_sass..)r/r<r)rrxrQrRrVrVrWrxRs zCUDADispatcher.inspect_sasscCs2|durtj}|jD] \}}|j|dq dS)ryNrz)r}r~rrQr)rHr{rrrrVrVrWrhs zCUDADispatcher.inspect_typescCs|jD]}|qdSrY)rrrl)rHrrrVrVrWrlts zCUDADispatcher.bindcCs|||}|S)rcrV)rer.r/rgrVrVrWrhxs zCUDADispatcher._rebuildcCst|j|jdS)zd Reduce the instance for serialization. Compiled definitions are discarded. )r.r/)rir.r/rZrVrVrWrjszCUDADispatcher._reduce_statesr)rrrrY)&rrrr _fold_argsr targetdescrrr+rr4r6r lru_cacher8r:r;r0rr%rCr>rrrSr_rWrerArsrvrxrrlrrhrjrrVrVrTrWr.sD         $ $     r.):numpyrrr}rr numba.corerrrrrrnumba.core.cachingrr numba.core.compiler_lockr numba.core.dispatcherr numba.core.errorsr numba.core.typing.typeofr rnumba.cuda.apirnumba.cuda.argsrnumba.cuda.compilerrrnumba.cuda.cudadrvrnumba.cuda.cudadrv.devicesrnumba.cuda.cudadrv.libsrnumba.cuda.descriptorrnumba.cuda.errorsrr numba.cudar3numbarrwarningsr ReduceMixinrobjectr rr'r,r.rVrVrVrWs@              .