o Y۷iݜ@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'm(Z(ddl)m Z*ddl+m,Z,ddl+m-Z-ddl.m/Z/gdZ0Gdddej1Z2Gddde3Z4GdddZ5GdddeZ6GdddeZ7Gdddeej1Z8dS) N)config serializesigutilstypestypingutils)Cache CacheImpl)global_compiler_lock) Dispatcher)NumbaPerformanceWarning)Purposetypeof)get_current_device)wrap_arg) compile_cuda CUDACompiler)driver) get_context) cuda_target)missing_launch_config_msgnormalize_kernel_dimensions)rcuda) _dispatcher)warn)hsinhcoshloghlog10hlog2hexphexp10hexp2hsqrthrsqrthfloorhceilhrcphrinthtrunchdivcseZdZdZe   d2fdd ZeddZed d Zd d Z ed dZ eddZ e fddZ ddZddZeddZeddZeddZeddZedd Zd!d"Zd#d$Zd%d&Zd'd(Zd3d)d*Zd4d,d-Zd5d.d/Zd0d1ZZS)6_Kernelz CUDA Kernel specialized for a given set of argument types. When called, this object launches the kernel on the device. NFTc  s| rtdtd|_d|_||_||_||_||_|p g|_ || r'dndd} t j } t |jt j|j|j|||| | d }|j}|jj}|j}|j}||j|j||| ||| \}|sag}dv|_|jrnd_fd d tD}|rtjtjt}tj|d }| ||D]}!|q|j"|_#|j$|_$|j%|_&|_'|j(|_(||_|j|_|j)|_)g|_*g|_+g|_,dS) Nz,Cannot compile a device function as a kernelFr)fastmathoptdebuglineinfoinliner. nvvm_optionscccudaCGGetIntrinsicHandleTcs"g|] }d|vr|qS)__numba_wrapper_) get_asm_str).0fnlibK/home/ubuntu/vllm_env/lib/python3.10/site-packages/numba/cuda/dispatcher.py lsz$_Kernel.__init__..zcpp_function_wrappers.cu)- RuntimeErrorsuper__init__ objectmode entry_pointpy_funcargtypesr1r2 extensionsrcompute_capabilityrrvoidtarget_context__code__ co_filenameco_firstlinenoprepare_cuda_kernellibraryfndescr8 cooperativeneeds_cudadevrtcuda_fp16_math_funcsospathdirnameabspath__file__joinappendadd_linking_filename entry_name signaturetype_annotation_type_annotation _codelibrary call_helper environment_referenced_environmentslifted reload_init)selfrErFlinkr1r2r3r.rG max_registersr/devicer4r5crestgt_ctxcodefilenamelinenumkernelresbasedirfunctions_cu_pathfilepath __class__r;r>rB.sp       z_Kernel.__init__cC|jSN)rargr=r=r>rOz_Kernel.librarycCrwrx)r`ryr=r=r>r_rzz_Kernel.type_annotationcCrwrx)rdryr=r=r>_find_referenced_environmentssz%_Kernel._find_referenced_environmentscC |jSrx)rJcodegenryr=r=r>r} z_Kernel.codegencCs t|jjSrx)tupler^argsryr=r=r>argument_typess z_Kernel.argument_typesc sX||} t|| d| _|| _|| _|| _d| _|| _|| _ || _ || _ || _ | S)& Rebuild an instance. N) __new__rArBrDrQr]r^r`rar1r2rbrG) clsrQr\r^ codelibraryr1r2rbrGinstancerur=r>_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. )rQr\r^rr1r2rbrG) dictrQr]r^rar1r2rbrGryr=r=r>_reduce_statess z_Kernel._reduce_statescCs|jdS)z7 Force binding to current CUDA context N)ra get_cufuncryr=r=r>bindz _Kernel.bindcC|jjjS)zN The number of registers used by each thread for this kernel. )rarattrsregsryr=r=r>regs_per_threadz_Kernel.regs_per_threadcCr)zD The amount of constant memory used by this kernel. )rarrconstryr=r=r>const_mem_sizerz_Kernel.const_mem_sizecCr)zM The amount of shared memory used per block for this kernel. )rarrsharedryr=r=r>shared_mem_per_blockrz_Kernel.shared_mem_per_blockcCr)z: The maximum allowable threads per block. )rarr maxthreadsryr=r=r>max_threads_per_blockrz_Kernel.max_threads_per_blockcCr)zM The amount of local memory used per thread for this kernel. )rarrlocalryr=r=r>local_mem_per_threadrz_Kernel.local_mem_per_threadcCr|)z6 Returns the LLVM IR for this kernel. )ra get_llvm_strryr=r=r> inspect_llvms z_Kernel.inspect_llvmcCs|jj|dS)z7 Returns the PTX code for this kernel. r5)rar8)rgr5r=r=r> inspect_asmrz_Kernel.inspect_asmcCr|)zv Returns the CFG of the SASS for this kernel. Requires nvdisasm to be available on the PATH. )ra get_sass_cfgryr=r=r>inspect_sass_cfg z_Kernel.inspect_sass_cfgcCr|)zp Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )raget_sassryr=r=r> inspect_sassrz_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================================================================================)r` ValueErrorsysstdoutprintr]r)rgrr=r=r> 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||Srxr=)xyr=r=r>&sz5_Kernel.max_cooperative_grid_blocks..) rrar isinstancer functoolsreduce$get_active_blocks_per_multiprocessorrjMULTIPROCESSOR_COUNT)rgblockdim dynsmemsizectxcufunc active_per_smsm_countr=r=r>max_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)streamrQcs<jdj|f\}}t}tt||||jS)Nz%s__%s__) moduleget_global_symbolr\ctypesc_intrdevice_to_host addressofvalue)r\memszval)rr=r> load_symbolSs z#_Kernel.launch..load_symbolcg|]}d|qS)tidr=r9irr=r>r?[z"_Kernel.launch..zyxcr)ctaidr=rrr=r>r?\rz"In function %r, file %s, line %s, z%stid=%s ctaid=%sz%s: %s)rarr1r\rrrsizeofrmemsetzipr _prepare_argsrUSE_NV_BINDINGbindingCUstreamhandle launch_kernelrQrrrrb get_exceptionrTrUrW)rgrgriddimrr sharedmemexcnameexcmemexcszexcvalretr kernelargstv zero_stream stream_handlerrrmexcclsexc_argsloclocinfosymrtlinenoprefixwbr=)rrr>launch-sn        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)5reversedrG prepare_argsrrArrayr to_devicer c_ssize_tc_void_psizedtypeitemsizerdevice_pointerrintrZrangendimshapestridesIntegergetattrfloat16c_uint16npviewuint16float64c_doublefloat32c_floatbooleanc_uint8 complex64realimag complex128 NPDatetime NPTimedeltac_int64int64Recorddevice_ctypes_pointer BaseTuplelenrr EnumMemberrNotImplementedError)rgtyrrrr extensiondevaryc_intpmeminfoparentnitemsrptrdataaxcvaldevrecrrr=r=r>rus                        z_Kernel._prepare_args) NFFFFNNTFrx)rrr)__name__ __module__ __qualname____doc__r rBpropertyrOr_r{r}r classmethodrrrrrrrrrrrrrrrr __classcell__r=r=rur>r,(sL\             Hr,c@$eZdZddZddZddZdS)ForAllcCs6|dkr td|||_||_||_||_||_dS)Nrz0Can't create ForAll with negative task count: %s)r dispatcherntasksthread_per_blockrr)rgr0r1tpbrrr=r=r>rBs zForAll.__init__cGs^|jdkrdS|jjr|j}n|jj|}||}|j|d|}||||j|jf|S)Nrr)r1r0 specialized specialize_compute_thread_per_blockrr)rgrr4rrr=r=r>__call__s    zForAll.__call__cCsZ|j}|dkr |St}tt|j}t|jd|j dd}|j di|\}}|S)Nri)funcb2d_funcmemsizeblocksizelimitr=) r2rnextiter overloadsvaluesrrarrget_max_potential_block_size)rgr0r3rrpkwargs_r=r=r>r6sz ForAll._compute_thread_per_blockN)r'r(r)rBr7r6r=r=r=r>r/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.) r0rrrrrCUDA_LOW_OCCUPANCY_WARNINGSrr ) rgr0rrrr min_grid_size grid_sizemsgr=r=r>rBs   z_LaunchConfiguration.__init__cGs|j||j|j|j|jSrx)r0callrrrrrgrr=r=r>r7sz_LaunchConfiguration.__call__N)r'r(r)rBr7r=r=r=r>rCs rCc@r.) CUDACacheImplcCs|Srx)r)rgrpr=r=r>r szCUDACacheImpl.reducecCstjdi|S)Nr=)r,r)rgrJpayloadr=r=r>rebuild#zCUDACacheImpl.rebuildcCsdS)NTr=)rgrkr=r=r>check_cachable&szCUDACacheImpl.check_cachableN)r'r(r)rrNrPr=r=r=r>rLs rLcs$eZdZdZeZfddZZS) CUDACachezS Implements a cache that saves and loads CUDA kernels and compile results. csFddlm}|dt||WdS1swYdS)Nr)target_overrider)numba.core.target_extensionrRrA load_overload)rgsigrJrRrur=r>rT7s   $zCUDACache.load_overload)r'r(r)r*rL _impl_classrTr-r=r=rur>rQ1srQcsBeZdZdZdZeZeffdd Ze ddZ ddZ e j d d dAd d ZddZdBddZe ddZddZddZddZddZddZe ddZdCd!d"ZdCd#d$ZdCd%d&ZdCd'd(ZdCd)d*Zd+d,ZdCd-d.Zd/d0Zd1d2Z dCd3d4Z!dCd5d6Z"dCd7d8Z#dCd9d:Z$dCd;d<Z%e&d=d>Z'd?d@Z(Z)S)DCUDADispatchera 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)rArB _specializedspecializations)rgrErXrYrur=r>rBRs  zCUDADispatcher.__init__cCs t|Srx) cuda_typesrWryr=r=r> _numba_type_br~zCUDADispatcher._numba_type_cCst|j|_dSrx)rQrE_cacheryr=r=r>enable_cachingfrOzCUDADispatcher.enable_cachingrD)maxsizercCst||\}}t|||||Srx)rrC)rgrrrrr=r=r> configureiszCUDADispatcher.configurecCst|dvr td|j|S)N)rEr-z.must specify at least the griddim and blockdim)rrrarKr=r=r> __getitem__ns  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.)r3rr)r/)rgr1r3rrr=r=r>forallsszCUDADispatcher.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. rG)rXgetryr=r=r>rGs zCUDADispatcher.extensionscOsttrx)rr)rgrrAr=r=r>r7szCUDADispatcher.__call__cCsD|jr tt|j}n tjj|g|R}||||||dS)zJ Compile if necessary and invoke this kernel with *args*. N) r4r<r=r>r?rr _cuda_callr)rgrrrrrrpr=r=r>rJszCUDADispatcher.callcs(|rJfdd|D}t|S)Ncsg|]}|qSr= typeof_pyvalr9aryr=r>r?sz4CUDADispatcher._compile_for_args..)compiler)rgrkwsrFr=ryr>_compile_for_argssz CUDADispatcher._compile_for_argscCsDzt|tjWSty!t|r ttj|ddtjYSw)NF)sync)rr argumentrr is_cuda_array as_cuda_array)rgrr=r=r>rhs  zCUDADispatcher.typeof_pyvalcsjrtdtj}tfdd|D}j||f}|r"|Sj}tj |d}| || d|_ |j||f<|S)zd Create a new instance of this dispatcher specialized for the given *args*. zDispatcher already specializedc3s|]}|VqdSrxrgriryr=r> sz,CUDADispatcher.specialize..)rXT) r4r@rrHrr[rerXrWrErkdisable_compilerZ)rgrr5rFspecializationrXr=ryr>r5s  zCUDADispatcher.specializecCrw)z> True if the Dispatcher has been specialized. )rZryr=r=r>r4szCUDADispatcher.specializedNcCD|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. NcSi|]\}}||jqSr=)rr9rUoverloadr=r=r>  z6CUDADispatcher.get_regs_per_thread..)r>rrr4r<r=r?itemsrgr^r=r=r>get_regs_per_thread z"CUDADispatcher.get_regs_per_threadcCru)a Returns the size in bytes of constant memory used by this kernel for the device in the current context. :param signature: The signature of the compiled kernel to get constant memory usage for. This may be omitted for a specialized kernel. :return: The size in bytes of constant memory allocated by the compiled variant of the kernel for the given signature and current device. NcSrvr=)rrwr=r=r>ryrzz5CUDADispatcher.get_const_mem_size..)r>rrr4r<r=r?r{r|r=r=r>get_const_mem_sizes z!CUDADispatcher.get_const_mem_sizecCru)a Returns the size in bytes of statically allocated shared memory for this kernel. :param signature: The signature of the compiled kernel to get shared memory usage for. This may be omitted for a specialized kernel. :return: The amount of shared memory allocated by the compiled variant of the kernel for the given signature and current device. NcSrvr=)rrwr=r=r>ryrzz;CUDADispatcher.get_shared_mem_per_block..)r>rrr4r<r=r?r{r|r=r=r>get_shared_mem_per_blockr~z'CUDADispatcher.get_shared_mem_per_blockcCru)a( Returns the maximum allowable number of threads per block for this kernel. Exceeding this threshold will result in the kernel failing to launch. :param signature: The signature of the compiled kernel to get the max threads per block for. This may be omitted for a specialized kernel. :return: The maximum allowable threads per block for the compiled variant of the kernel for the given signature and current device. NcSrvr=)rrwr=r=r>ry,rzz.)r>rrr4r<r=r?r{r|r=r=r>get_max_threads_per_blocks z(CUDADispatcher.get_max_threads_per_blockcCru)a Returns the size in bytes of local memory per thread for this kernel. :param signature: The signature of the compiled kernel to get local memory usage for. This may be omitted for a specialized kernel. :return: The amount of local memory allocated by the compiled variant of the kernel for the given signature and current device. NcSrvr=)rrwr=r=r>ry?rzz;CUDADispatcher.get_local_mem_per_thread..)r>rrr4r<r=r?r{r|r=r=r>get_local_mem_per_thread/r~z'CUDADispatcher.get_local_mem_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_devicerrEr'formatrmake_concrete_templatenopython_signaturesr pysignature)rgrrl func_namer\ call_templatepysigr=r=r>get_call_templateBs   z CUDADispatcher.get_call_templatec Cs||jvrc|jQ|jd}|jd}|jd}|jd}|jdr)dnd|d}tj}t|j||||||||d } | |j|<| j | j | j | j gWd | S1s\wY| 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`. r1r2r3r.r/r-r)r/r.r0N) r>_compiling_counterrXrerrHrrErJinsert_user_functionrDrPrO) rgr return_typer1r2r3r.r4r5rkr=r=r>r]s:          zCUDADispatcher.compile_devicecCs,dd|D}|j||dd||j|<dS)NcSsg|]}|jqSr=)_coderir=r=r>r?sz/CUDADispatcher.add_overload..Tr)_insertr>)rgrprFc_sigr=r=r> 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_signaturernoner4r<r=r>r?rer^rT targetctx _cache_hits _cache_missesrr@r,rErXr save_overloadr)rgrUrFrrpr=r=r>rks$  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. rjNcSsi|] \}}||jqSr=)rOrrwr=r=r>ryz/CUDADispatcher.inspect_llvm..cSi|] \}}||qSr=)rrwr=r=r>ry)rXrer>rOrrr{rgr^rjr=r=r>rs 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. rjNcsi|] \}}||jqSr=)rOr8rwrr=r>rysz.CUDADispatcher.inspect_asm..csi|] \}}||qSr=)rrwrr=r>ryr) rrHrXrer>rOr8rr{rr=rr>rs   zCUDADispatcher.inspect_asmcC>|jdr td|dur|j|Sdd|jDS)a Return this kernel's CFG for the device in the current context. :param signature: A tuple of argument types. :return: The CFG for the given signature, or a dict of CFGs for all previously-encountered signatures. The CFG for the device in the current context is returned. Requires nvdisasm to be available on the PATH. rjz'Cannot get the CFG of a device functionNcSrr=)rr9rUdefnr=r=r>ryrz3CUDADispatcher.inspect_sass_cfg..)rXrer@r>rr{r|r=r=r>rs zCUDADispatcher.inspect_sass_cfgcCr)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. rjz(Cannot inspect SASS of a device functionNcSrr=)rrr=r=r>ryrz/CUDADispatcher.inspect_sass..)rXrer@r>rr{r|r=r=r>rs zCUDADispatcher.inspect_sasscCs2|durtj}|jD] \}}|j|dq dS)rNr)rrr>r{r)rgrrBrr=r=r>rs zCUDADispatcher.inspect_typescCs|||}|S)rr=)rrErXrr=r=r>rs zCUDADispatcher._rebuildcCst|j|jdS)zd Reduce the instance for serialization. Compiled definitions are discarded. )rErX)rrErXryr=r=r>rszCUDADispatcher._reduce_statesr&)rrrrx)*r'r(r)r* _fold_argsr targetdescrrrBr+r]r_r lru_cacherarcrdrGr7rJrmrhr5r4r}rrrrrrrrkrrrrrr,rrr-r=r=rur>rW@sL              ' $     rW)9numpyrrTrrr 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.descriptorrnumba.cuda.errorsrr numba.cudar\numbarrwarningsrrS ReduceMixinr,objectr/rCrLrQrWr=r=r=r>s@              /.