a Df@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)ZhsinZhcosZhlogZhlog10Zhlog2ZhexpZhexp10Zhexp2ZhsqrtZhrsqrtZhfloorZhceilZhrcpZhrintZhtruncZhdivc seZdZdZed2fdd 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| r tdtd|_d|_||_||_||_||_|p@g|_ || rNdndd} t j } t |jt j|j|j|||| | d }|j}|jj}|j}|j}||j|j||| ||| \}|sg}dv|_|jrd_fd d tD}|r tjtjt}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_optionsccZcudaCGGetIntrinsicHandleTcs"g|]}d|vr|qS)Z__numba_wrapper_) get_asm_str).0fnlibb/nfs/NAS7/SABIOD/METHODE/ermites/ermites_venv/lib/python3.9/site-packages/numba/cuda/dispatcher.py lsz$_Kernel.__init__..zcpp_function_wrappers.cu)- RuntimeErrorsuper__init__Z objectmode entry_pointpy_funcargtypesr!r" extensionsrcompute_capabilityrrvoidtarget_context__code__ co_filenameco_firstlinenoZprepare_cuda_kernellibraryfndescr& cooperativeZneeds_cudadevrtcuda_fp16_math_funcsospathdirnameabspath__file__joinappendZadd_linking_filename entry_name signaturetype_annotation_type_annotation _codelibrary call_helper environment_referenced_environmentsZliftedZ reload_init)selfr2r3linkr!r"r#rr4Z max_registersrdevicer$r%cresZtgt_ctxcodefilenamelinenumkernelresZbasedirZfunctions_cu_pathfilepath __class__r)r,r0.sp     z_Kernel.__init__cCs|jSN)rKrOr+r+r,r;sz_Kernel.librarycCs|jSr[)rJr\r+r+r,rIsz_Kernel.type_annotationcCs|jSr[)rNr\r+r+r,_find_referenced_environmentssz%_Kernel._find_referenced_environmentscCs |jSr[)r7codegenr\r+r+r,r^sz_Kernel.codegencCs t|jjSr[)tuplerHargsr\r+r+r,argument_typessz_Kernel.argument_typesc sX||} t|| d| _|| _|| _|| _d| _|| _|| _ || _ || _ || _ | S)& Rebuild an instance. N) __new__r/r0r1r=rGrHrJrKr!r"rLr4) clsr=rFrH codelibraryr!r"rLr4instancerYr+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. )r=rFrHrer!r"rLr4) dictr=rGrHrKr!r"rLr4r\r+r+r,_reduce_statess  z_Kernel._reduce_statescCs|jdS)z7 Force binding to current CUDA context N)rK get_cufuncr\r+r+r,bindsz _Kernel.bindcCs|jjjS)zN The number of registers used by each thread for this kernel. )rKrjattrsregsr\r+r+r,regs_per_threadsz_Kernel.regs_per_threadcCs|jjjS)zD The amount of constant memory used by this kernel. )rKrjrlconstr\r+r+r,const_mem_sizesz_Kernel.const_mem_sizecCs|jjjS)zM The amount of shared memory used per block for this kernel. )rKrjrlZsharedr\r+r+r,shared_mem_per_blocksz_Kernel.shared_mem_per_blockcCs|jjjS)z: The maximum allowable threads per block. )rKrjrlZ maxthreadsr\r+r+r,max_threads_per_blocksz_Kernel.max_threads_per_blockcCs|jjjS)zM The amount of local memory used per thread for this kernel. )rKrjrllocalr\r+r+r,local_mem_per_threadsz_Kernel.local_mem_per_threadcCs |jS)z6 Returns the LLVM IR for this kernel. )rK get_llvm_strr\r+r+r, inspect_llvmsz_Kernel.inspect_llvmcCs|jj|dS)z7 Returns the PTX code for this kernel. r%)rKr&)rOr%r+r+r, inspect_asmsz_Kernel.inspect_asmcCs |jS)zv Returns the CFG of the SASS for this kernel. Requires nvdisasm to be available on the PATH. )rKZ get_sass_cfgr\r+r+r,inspect_sass_cfgsz_Kernel.inspect_sass_cfgcCs |jS)zp Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )rKZget_sassr\r+r+r, inspect_sasssz_Kernel.inspect_sasscCsb|jdurtd|dur tj}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================================================================================)rJ ValueErrorsysstdoutprintrGra)rOr}r+r+r, inspect_typess  z_Kernel.inspect_typesrcCsHt}|j}t|tr*tdd|}||||}|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||Sr[r+)xyr+r+r,&z5_Kernel.max_cooperative_grid_blocks..) rrKrj isinstancer_ functoolsreduceZ$get_active_blocks_per_multiprocessorrQZMULTIPROCESSOR_COUNT)rOblockdimZ dynsmemsizectxcufuncZ active_per_smZsm_countr+r+r,max_cooperative_grid_blockss   z#_Kernel.max_cooperative_grid_blockscs|j|jrTjd}j|\}}|ttjks>Jt} |j d|dg} g} t |j |D]\} } | | | || | qht jrt jd}nd}|r|jp|}t jjg||||| Rd|ji|jrt t| ||| jdkrfddfddd D}fd dd D}| j}|j|\}}}|durVd }n$|\}}}tj|}d |||f}d |||f}|rd||dff|dd}n|f}||| D] }|qdS)NZ __errcode__r)streamr=cs<jdj|f\}}t}tt||||jS)Nz%s__%s__) moduleget_global_symbolrFctypesc_intrdevice_to_host addressofvalue)rFZmemszval)rr+r, load_symbolSs z#_Kernel.launch..load_symbolcsg|]}d|qS)tidr+r'irr+r,r-[rz"_Kernel.launch..Zzyxcsg|]}d|qS)ctaidr+rrr+r,r-\rz"In function %r, file %s, line %s, z%stid=%s ctaid=%sz%s: %s)rKrjr!rFrrrsizeofrmemsetzipra _prepare_argsrUSE_NV_BINDINGZbindingZCUstreamhandleZ launch_kernelr=rrrrLZ get_exceptionr?r@rB)rOr`griddimrr sharedmemexcnameZexcmemZexcszZexcvalretr kernelargstvZ zero_streamZ stream_handlerrrSZexcclsZexc_argslocZlocinfosymrXlinenoprefixwbr+)rrr,launch-sj         z_Kernel.launchcCsft|jD]}|j||||d\}}q t|tjrt|||}tj }t d} t d} ||j } ||j j } t|} tjrt| } t | }|| || || || ||t|jD]}|||j|qt|jD]}|||j|qnNt|tjrBttd||}||n |tjkrttt|tj}||n|tjkrt|}||n|tj krt!|}||n|tj"krt#t|}||n~|tj$kr|t!|j%|t!|j&nJ|tj'krL|t|j%|t|j&nt|tj(tj)frz|t*|tj+nt|tj,rt|||}|j-} tjrt t| } || nt|tj.rt/|t/|ksJt0||D]\}}|1|||||qnTt|tj2rXz|1|j |j3|||Wnt4yTt4||Yn0n t4||dS)zF Convert arguments to ctypes and append to kernelargs )rrrzc_%sN)5reversedr4Z prepare_argsrrArrayrZ to_devicer c_ssize_tc_void_psizeZdtypeitemsizerZdevice_pointerrintrErangendimshapestridesZIntegergetattrfloat16c_uint16npviewZuint16float64c_doublefloat32c_floatbooleanc_uint8Z complex64realimagZ complex128Z NPDatetimeZ NPTimedeltac_int64int64ZRecordZdevice_ctypes_pointerZ BaseTuplelenrrZ EnumMemberrNotImplementedError)rOtyrrrr extensionZdevaryZc_intpZmeminfoparentZnitemsrptrdataaxcvalZdevrecrrr+r+r,rus                     z_Kernel._prepare_args) NFFFFNNTF)N)r)rr)__name__ __module__ __qualname____doc__r r0propertyr;rIr]r^ra classmethodrgrirkrnrprqrrrtrvrxryrzrrrr __classcell__r+r+rYr,r(sJ\             Hrc@s$eZdZddZddZddZdS)ForAllcCs6|dkrtd|||_||_||_||_||_dS)Nrz0Can't create ForAll with negative task count: %s)r~ dispatcherntasksthread_per_blockrr)rOrrtpbrrr+r+r,r0szForAll.__init__cGs^|jdkrdS|jjr|j}n |jj|}||}|j|d|}||||j|jf|S)Nrr)rr specialized specialize_compute_thread_per_blockrr)rOr`rrrr+r+r,__call__s    zForAll.__call__cCs^|j}|dkr|St}tt|j}t|jd|j dd}|j fi|\}}|SdS)Nri)funcZb2d_funcZmemsizeZblocksizelimit) rrnextiter overloadsvaluesrhrKrjrZget_max_potential_block_size)rOrrrrVkwargs_r+r+r,rsz ForAll._compute_thread_per_blockN)rrrr0rrr+r+r+r,rs rc@seZdZddZddZdS)_LaunchConfigurationc Csd||_||_||_||_||_tjr`d}|d|d|d}||kr`d|d}tt|dS)Nrrz Grid size zB will likely result in GPU under-utilization due to low occupancy.) rrrrrrZCUDA_LOW_OCCUPANCY_WARNINGSrr ) rOrrrrrZ min_grid_sizeZ grid_sizemsgr+r+r,r0s  z_LaunchConfiguration.__init__cGs|j||j|j|j|jSr[)rcallrrrrrOr`r+r+r,rsz_LaunchConfiguration.__call__N)rrrr0rr+r+r+r,rsrc@s$eZdZddZddZddZdS) CUDACacheImplcCs|Sr[)ri)rOrVr+r+r,r szCUDACacheImpl.reducecCstjfi|Sr[)rrg)rOr7payloadr+r+r,rebuild#szCUDACacheImpl.rebuildcCsdS)NTr+)rOrRr+r+r,check_cachable&szCUDACacheImpl.check_cachableN)rrrrrrr+r+r+r,rsrcs$eZdZdZeZfddZZS) CUDACachezS Implements a cache that saves and loads CUDA kernels and compile results. csFddlm}|dt||WdS1s80YdS)Nr)target_overrider)Znumba.core.target_extensionrr/ load_overload)rOsigr7rrYr+r,r7s  zCUDACache.load_overload)rrrrrZ _impl_classrrr+r+rYr,r1srcsBeZdZdZdZeZeffdd Ze ddZ ddZ e j d d dAd d ZddZdBddZe ddZddZddZddZddZddZe ddZdCd!d"ZdDd#d$ZdEd%d&ZdFd'd(ZdGd)d*Zd+d,ZdHd-d.Zd/d0Zd1d2Z dId3d4Z!dJd5d6Z"dKd7d8Z#dLd9d:Z$dMd;d<Z%e&d=d>Z'd?d@Z(Z)S)NCUDADispatchera 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/r0 _specializedspecializations)rOr2rrrYr+r,r0Rs   zCUDADispatcher.__init__cCs t|Sr[) cuda_typesrr\r+r+r, _numba_type_bszCUDADispatcher._numba_type_cCst|j|_dSr[)rr2_cacher\r+r+r,enable_cachingfszCUDADispatcher.enable_cachingr)maxsizercCst||\}}t|||||Sr[)rr)rOrrrrr+r+r, configureiszCUDADispatcher.configurecCst|dvrtd|j|S)N)rrz.must specify at least the griddim and blockdim)rr~rrr+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.)rrr)r)rOrrrrr+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. r4)rgetr\r+r+r,r4szCUDADispatcher.extensionscOs ttdSr[)r~r)rOr`rr+r+r,rszCUDADispatcher.__call__cCsD|jrtt|j}ntjj|g|R}||||||dS)zJ Compile if necessary and invoke this kernel with *args*. N) rrrrrrr Z _cuda_callr)rOr`rrrrrVr+r+r,rszCUDADispatcher.callcs(|rJfdd|D}t|S)Ncsg|]}|qSr+) typeof_pyvalr'ar\r+r,r-rz4CUDADispatcher._compile_for_args..)compiler_)rOr`kwsr3r+r\r,_compile_for_argssz CUDADispatcher._compile_for_argscCsLzt|tjWStyFt|r@ttj|ddtjYSYn0dS)NF)sync)rr argumentr~rZ is_cuda_arrayZ as_cuda_array)rOrr+r+r,r s  zCUDADispatcher.typeof_pyvalcstj}tfdd|D}jr,tdj||f}|rD|Sj}tj |d}| || d|_ |j||f<|S)zd Create a new instance of this dispatcher specialized for the given *args*. csg|]}j|qSr+)Z typingctxZresolve_argument_typer r\r+r,r-rz-CUDADispatcher.specialize..zDispatcher already specialized)rT) rr5r_rr.rr rrr2r Zdisable_compiler)rOr`r%r3Zspecializationrr+r\r,rs$ zCUDADispatcher.specializecCs|jS)z> True if the Dispatcher has been specialized. )rr\r+r+r,rszCUDADispatcher.specializedNcCsH|dur|j|jjS|jr0tt|jjSdd|jDSdS)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|]\}}||jqSr+)rnr'roverloadr+r+r, sz6CUDADispatcher.get_regs_per_thread..)rr`rnrrrritemsrOrHr+r+r,get_regs_per_threads z"CUDADispatcher.get_regs_per_threadcCsH|dur|j|jjS|jr0tt|jjSdd|jDSdS)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. NcSsi|]\}}||jqSr+)rprr+r+r,rsz5CUDADispatcher.get_const_mem_size..)rr`rprrrrrrr+r+r,get_const_mem_sizes z!CUDADispatcher.get_const_mem_sizecCsH|dur|j|jjS|jr0tt|jjSdd|jDSdS)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. NcSsi|]\}}||jqSr+)rqrr+r+r,rsz;CUDADispatcher.get_shared_mem_per_block..)rr`rqrrrrrrr+r+r,get_shared_mem_per_blocks z'CUDADispatcher.get_shared_mem_per_blockcCsH|dur|j|jjS|jr0tt|jjSdd|jDSdS)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. NcSsi|]\}}||jqSr+)rrrr+r+r,r,sz.)rr`rrrrrrrrr+r+r,get_max_threads_per_blocks z(CUDADispatcher.get_max_threads_per_blockcCsH|dur|j|jjS|jr0tt|jjSdd|jDSdS)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. NcSsi|]\}}||jqSr+)rtrr+r+r,r?sz;CUDADispatcher.get_local_mem_per_thread..)rr`rtrrrrrrr+r+r,get_local_mem_per_thread/s 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}))keyZ signatures) _can_compilecompile_devicer_r2rformatrZmake_concrete_templateZnopython_signaturesrZ pysignature)rOr`r func_namerFZ call_templateZpysigr+r+r,get_call_templateBs  z CUDADispatcher.get_call_templatec Cs||jvr|j|jd}|jd}|jd}|jd}|jdrRdnd|d}tj}t|j||||||||d } | |j|<| j | j | j | j gWd q1s0Yn |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#rrrr)rrr N) rZ_compiling_counterrr rr5rr2r7Zinsert_user_functionr1r<r;) rOr` return_typer!r"r#rr$r%rRr+r+r,r]s2        $ zCUDADispatcher.compile_devicecCs,dd|D}|j||dd||j|<dS)NcSsg|] }|jqSr+)_coder r+r+r,r-rz/CUDADispatcher.add_overload..Tr)Z_insertr)rOrVr3Zc_sigr+r+r, add_overloadszCUDADispatcher.add_overloadcCst|\}}|dus$|tjks$J|jr.cSsi|]\}}||qSr+)rvrr+r+r,rs)rr rr;rurvrrOrHrQr+r+r,rvs zCUDADispatcher.inspect_llvmcs|tj|jd}|durD|r2|j|jS|j|Sn4|r`fdd|jDSfdd|jDSdS)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. rQNcsi|]\}}||jqSr+)r;r&rrwr+r,rsz.CUDADispatcher.inspect_asm..csi|]\}}||qSr+)rxrrwr+r,rs) rr5rr rr;r&rxrr&r+rwr,rxs    zCUDADispatcher.inspect_asmcCsB|jdrtd|dur*|j|Sdd|jDSdS)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. rQz'Cannot get the CFG of a device functionNcSsi|]\}}||qSr+)ryr'rdefnr+r+r,rsz3CUDADispatcher.inspect_sass_cfg..)rr r.rryrrr+r+r,rys zCUDADispatcher.inspect_sass_cfgcCsB|jdrtd|dur*|j|Sdd|jDSdS)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. rQz(Cannot inspect SASS of a device functionNcSsi|]\}}||qSr+)rzr'r+r+r,rsz/CUDADispatcher.inspect_sass..)rr r.rrzrrr+r+r,rzs zCUDADispatcher.inspect_sasscCs2|durtj}|jD]\}}|j|dqdS)r{Nr|)rrrrr)rOr}rr(r+r+r,rszCUDADispatcher.inspect_typescCs|||}|S)rbr+)rdr2rrfr+r+r,rgs zCUDADispatcher._rebuildcCst|j|jdS)zd Reduce the instance for serialization. Compiled definitions are discarded. )r2r)rhr2rr\r+r+r,riszCUDADispatcher._reduce_states)rr)rrr)N)N)N)N)N)N)N)N)N)N)N)*rrrrZ _fold_argsrZ targetdescrrr0rrrr lru_cacherrrr4rrrr rrrrrrrr!rr$r rvrxryrzrrrgrirr+r+rYr,r@sJ              '$      r)9numpyrr?rrrZ numba.corerrrrrrZnumba.core.cachingrr Znumba.core.compiler_lockr Znumba.core.dispatcherr Znumba.core.errorsr Znumba.core.typing.typeofr rZnumba.cuda.apirZnumba.cuda.argsrZnumba.cuda.compilerrrZnumba.cuda.cudadrvrZnumba.cuda.cudadrv.devicesrZnumba.cuda.descriptorrZnumba.cuda.errorsrrZ numba.cudarZnumbarrwarningsrr>Z ReduceMixinrobjectrrrrrr+r+r+r,s>              /.