iddlZddlZddlZddlZddlZddlmZmZm Z m Z m Z m Z ddl mZmZddlmZddlmZddlmZmZddlmZmZddlmZdd lmZdd lmZm Z dd 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/m0Z0gdZ1GddejdZ3Gdde4Z5GddZ6GddeZ7GddeZ8GddeejdZ9y) N)config serializesigutilstypestypingutils)Cache CacheImpl)global_compiler_lock) Dispatcher)NumbaPerformanceWarningNumbaValueError)Purposetypeof)get_current_device)wrap_arg) compile_cuda CUDACompiler)driver) get_context) cuda_target)missing_launch_config_msgnormalize_kernel_dimensions)rcuda) _dispatcher)warn)hsinhcoshloghlog10hlog2hexphexp10hexp2hsqrthrsqrthfloorhceilhrcphrinthtrunchdivc eZdZdZe dfd ZedZedZdZ edZ edZ e fdZ d Zd Zed Zed Zed ZedZedZdZdZdZdZddZddZddZdZxZS)_Kernelz CUDA Kernel specialized for a given set of argument types. When called, this object launches the kernel on the device. c | r tdt| d|_d|_||_||_||_||_|xsg|_ || rdndd} tj} t|j tj|j |j|||| |  }|j}|j j }|j"}|j$}|j'|j(|j*||| ||| \}}|sg}d|j-v|_|j.rd|_t2Dcgc]}d ||j-vr|}}|rqt4j6j9t4j6j;t<}t4j6j?|d }|jA||D]}|jC||jD|_#|jH|_$|jJ|_&||_'|jP|_(||_|j*|_|jR|_)g|_*g|_+g|_,ycc}w) Nz,Cannot compile a device function as a kernelFr)fastmathoptdebuglineinfoinliner2 nvvm_optionscccudaCGGetIntrinsicHandleT__numba_wrapper_zcpp_function_wrappers.cu)- RuntimeErrorsuper__init__ objectmode entry_pointpy_funcargtypesr5r6 extensionsrcompute_capabilityrrvoidtarget_context__code__ co_filenameco_firstlinenoprepare_cuda_kernellibraryfndesc get_asm_str cooperativeneeds_cudadevrtcuda_fp16_math_funcsospathdirnameabspath__file__joinappendadd_linking_filename entry_name signaturetype_annotation_type_annotation _codelibrary call_helper environment_referenced_environmentslifted reload_init)selfrArBlinkr5r6r7r2rC max_registersr3devicer8r9crestgt_ctxcodefilenamelinenumlibkernelfnresbasedirfunctions_cu_pathfilepath __class__s j/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/numba/cuda/dispatcher.pyr>z_Kernel.__init__.sU MN N         $*!1  ! 4 4DLL%**dmm"&**%-#)%-)5! #%%||$$##%%11$,, 27<2:G2?A V D69JJ   "&C 0B0b%bT*coo.??0B ggoobggooh&?@G " W-G!I  KK) *H   *!++ $ 4 4++ &kk ++(*% ;BsJc|jSN)r^rds rurKz_Kernel.librarys   c|jSrw)r]rxs rur\z_Kernel.type_annotations$$$ryc|jSrw)rarxs ru_find_referenced_environmentsz%_Kernel._find_referenced_environmentss,,,ryc6|jjSrw)rFcodegenrxs rur~z_Kernel.codegens""**,,ryc@t|jjSrw)tupler[argsrxs ruargument_typesz_Kernel.argument_typessT^^(())ryc |j|} t|| d| _|| _|| _|| _d| _|| _|| _ || _ || _ || _ | S)& Rebuild an instance. N) __new__r=r>r@rNrZr[r]r^r5r6r_rC) clsrNrYr[ codelibraryr5r6r_rCinstancerts ru_rebuildz_Kernel._rebuildst ;;s# c8%'#*"&$(! +$*(ryc t|j|j|j|j|j |j |j|jS)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. )rNrYr[rr5r6r_rC) dictrNrZr[r^r5r6r_rCrxs ru_reduce_statesz_Kernel._reduce_statessL 0 0t"nn$:K:K**t}} $ 0 0T__N Nryc8|jjy)z7 Force binding to current CUDA context N)r^ get_cufuncrxs rubindz _Kernel.binds $$&ryc^|jjjjS)zN The number of registers used by each thread for this kernel. )r^rattrsregsrxs ruregs_per_threadz_Kernel.regs_per_threads%   ++-33888ryc^|jjjjS)zD The amount of constant memory used by this kernel. )r^rrconstrxs ruconst_mem_sizez_Kernel.const_mem_size%   ++-33999ryc^|jjjjS)zM The amount of shared memory used per block for this kernel. )r^rrsharedrxs rushared_mem_per_blockz_Kernel.shared_mem_per_blocks%   ++-33:::ryc^|jjjjS)z: The maximum allowable threads per block. )r^rr maxthreadsrxs rumax_threads_per_blockz_Kernel.max_threads_per_blocks%   ++-33>>>ryc^|jjjjS)zM The amount of local memory used per thread for this kernel. )r^rrlocalrxs rulocal_mem_per_threadz_Kernel.local_mem_per_threadrryc6|jjS)z6 Returns the LLVM IR for this kernel. )r^ get_llvm_strrxs ru inspect_llvmz_Kernel.inspect_llvms  --//ryc:|jj|S)z7 Returns the PTX code for this kernel. )r9)r^rM)rdr9s ru inspect_asmz_Kernel.inspect_asms  ,,,33ryc6|jjS)zv Returns the CFG of the SASS for this kernel. Requires nvdisasm to be available on the PATH. )r^ get_sass_cfgrxs ruinspect_sass_cfgz_Kernel.inspect_sass_cfgs   --//ryc6|jjS)zp Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )r^get_sassrxs ru inspect_sassz_Kernel.inspect_sasss   ))++ryc|j td|tj}t |j d|j |t d|t |j|t d|y) 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 available filezP--------------------------------------------------------------------------------zP================================================================================)r] ValueErrorsysstdoutprintrZr)rdrs ru inspect_typesz_Kernel.inspect_typessg  (?@ @ <::D $*=*=>TJ hT" d##$/ hT"ryct}|jj}t|trt j d|}|j|||}|jj}||zS)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. c ||zSrw)xys ruz5_Kernel.max_cooperative_grid_blocks..&sQUry) rr^r isinstancer functoolsreduce$get_active_blocks_per_multiprocessorrgMULTIPROCESSOR_COUNT)rdblockdim dynsmemsizectxcufunc active_per_smsm_counts rumax_cooperative_grid_blocksz#_Kernel.max_cooperative_grid_blockssqm""--/ h & ''(:HEH@@AIALN ::22x''ryc|jj|jr|jdz}jj |\}}|t jt jk(sJt j} |jd|g} g} t|j|D]\} } |j| | || | tjr tjj!d}nd}|xr |j"xs|}tj$j"g||||| d|j&i|jrtj(t j* | j,dk7rfd}dDcgc] }|d|z}}dDcgc] }|d|z}}| j,}|j.j1|\}}}|d }n1|\}}}t2j4j7|}d |d |d |d }|d|d|}|r|d|df|ddz}n|f}||| D] }| ycc}wcc}w)N __errcode__r)streamrNcjjjd|d\}}tj}t j tj||||jS)N__) moduleget_global_symbolrYctypesc_intrdevice_to_host addressofvalue)rYmemszvalrs ru load_symbolz#_Kernel.launch..load_symbolSs`$mm==?E{{?C?EFGC!,,.C))&*:*:3*?bI99$ryzyxtidctaidz In function z, file z, line z, ztid=z ctaid=z: )r^rr5rYrrrsizeofrmemsetzipr _prepare_argsrUSE_NV_BINDINGbindingCUstreamhandle launch_kernelrNrrrr_ get_exceptionrQrRrT)rdrgriddimrr sharedmemexcnameexcmemexcszexcvalretr kernelargstv zero_stream stream_handlerirrrjexcclsexc_argsloclocinfosymrslinenoprefixwbrs @rulaunchz_Kernel.launch-sq""--/ ::kkM1G"MM;;GDMFEFMM&,,77 77\\^F MM!FM + ++T2DAq   q!VT: >3   ..11!4KK06==?K  V]] ;% ;& ;' ;+  ; (  ; *.)9)9  ; ::  ! !&"2"26":FE J||q %8==u!{519-u=;@A5aWq[15A||(,(8(8(F(Ft(L%#; G,/)C6!wwx8HFIFNFLOG18eD,2HQK @B   %H &wHh''B D/>As $I&<I+c t|jD]}|j||||\}}t|tj rt |j||}tj}tjd} tjd} ||j} ||jj} tj|} tj r t#| } tj| }|j%| |j%| |j%| |j%| |j%|t'|j(D]&}|j%||j*|(t'|j(D]&}|j%||j,|(yt|tj.r+t1td|z|}|j%|y|tj2k(rWtj4t7j2|j9t6j:}|j%|y|tj<k(r'tj>|}|j%|y|tj@k(r'tjB|}|j%|y|tjDk(r0tjFt#|}|j%|y|tjHk(r]|j%tjB|jJ|j%tjB|jLy|tjNk(r]|j%tj>|jJ|j%tj>|jLyt|tjPtjRfrB|j%tjT|j9t6jVyt|tjXrgt |j||}|jZ} tj rtjt#| } |j%| yt|tj\rCt_|t_|k(sJta||D]\}}|jc|||||yt|tjdr+ |jc|j|jf|||yti||#th$r ti||wxYw)zF Convert arguments to ctypes and append to kernelargs )rrrzc_%sN)5reversedrC prepare_argsrrArrayr to_devicer c_ssize_tc_void_psizedtypeitemsizerdevice_pointerrintrWrangendimshapestridesIntegergetattrfloat16c_uint16npviewuint16float64c_doublefloat32c_floatbooleanc_uint8 complex64realimag complex128 NPDatetime NPTimedeltac_int64int64Recorddevice_ctypes_pointer BaseTuplelenrr EnumMemberrNotImplementedError)rdtyrrrr extensiondevaryc_intpmeminfoparentnitemsrptrdataaxcvaldevrecrrs rurz_Kernel._prepare_argsus? "$//2I,, -GB3 b%++ &c],,T6:F%%Fooa(G__Q'FFKK(Ffll334H''/C$$#h??3'D   g &   f %   f %   h '   d #FKK(!!&b)9":;)FKK(!!&);"<=)EMM */766B;/4D   d # 5== ??2::c?#7#7 #BCD   d # 5== ??3'D   d # 5== >>#&D   d # 5== >>#c(+D   d # 5?? "   fnnSXX6 7   fnnSXX6 7 5## #   foochh7 8   foochh7 8 U--u/@/@A B   fnnSXXbhh-?@ A ELL )c],,T6:F..C$$ooc#h/   c " EOO ,r7c#h& &&B 1""1azB%E,, - 3""HHciiz&b#. . ' 3)"c22 3s )W W) NFFFFNNTFrw)rrr)__name__ __module__ __qualname____doc__r r>propertyrKr\r|r~r classmethodrrrrrrrrrrrrrrrr __classcell__rts@rur/r/(s+ ;@JN6;ZZx!!%%---*** N' 99 :: ;; ?? :: 0 4 0,#"(,FP\/ryr/ceZdZdZdZdZy)ForAllcp|dkrtd|z||_||_||_||_||_y)Nrz0Can't create ForAll with negative task count: %s)r dispatcherntasksthread_per_blockrr)rdr@rAtpbrrs rur>zForAll.__init__sE A:O%&' '$ # "ryc&|jdk(ry|jjr |j}n|jj|}|j |}|j|zdz |z}||||j |j f|S)Nrr)rAr@ specialized specialize_compute_thread_per_blockrr)rdrrErrs ru__call__zForAll.__call__s ;;!   ?? & &//K4$//44d;K11+>;;)A-(:+{7Hdkk>>*+,02 2ryc$|j}|dk7r|St}tt|jj }t |jjd|jd}|jdi|\}}|S)Nri)funcb2d_funcmemsizeblocksizelimitr) rBrnextiter overloadsvaluesrr^rrget_max_potential_block_size)rdr@rCrrnkwargs_s rurGz ForAll._compute_thread_per_blocks## !8J-C$z33::<=>F((335# F 6S55??FAsJryN)r5r6r7r>rHrGrryrur>r>s# 2ryr>ceZdZdZdZy)_LaunchConfigurationc||_||_||_||_||_t j r4d}|d|dz|dz}||krd|d}tt|yyy)Nrrz Grid size zB will likely result in GPU under-utilization due to low occupancy.) r@rrrrrCUDA_LOW_OCCUPANCY_WARNINGSrr ) rdr@rrrr min_grid_size grid_sizemsgs rur>z_LaunchConfiguration.__init__s$    "  - - M WQZ/'!*rHrryrurVrVs 3.AryrVceZdZdZdZdZy) CUDACacheImplc"|jSrw)r)rdrns rurzCUDACacheImpl.reduce s$$&&ryc,tjdi|S)Nr)r/r)rdrFpayloads rurebuildzCUDACacheImpl.rebuild#s*'**rycy)NTr)rdrhs rucheck_cachablezCUDACacheImpl.check_cachable&sryN)r5r6r7rrfrhrryrurbrbs'+ryrbc&eZdZdZeZfdZxZS) CUDACachezS Implements a cache that saves and loads CUDA kernels and compile results. clddlm}|d5t| ||cdddS#1swYyxYw)Nr)target_overrider)numba.core.target_extensionrlr= load_overload)rdsigrFrlrts rurnzCUDACache.load_overload7s, @ V $7(n=% $ $s*3)r5r6r7r8rb _impl_classrnr;r<s@rurjrj1s K>>ryrjcDeZdZdZdZeZeffd Ze dZ dZ e jdd"dZd Zd#d Ze d Zd Zd ZdZdZdZe dZd$dZd$dZd$dZd$dZd$dZdZd$dZdZdZ d$dZ!d$dZ"d$dZ#d$dZ$d$dZ%e&d Z'd!Z(xZ)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. FcFt||||d|_i|_y)N) targetoptionspipeline_classF)r=r> _specializedspecializations)rdrArtrurts rur>zCUDADispatcher.__init__Rs0  (6  8" "ryc,tj|Srw) cuda_typesrrrxs ru _numba_type_zCUDADispatcher._numba_type_bs((..ryc8t|j|_yrw)rjrA_cacherxs ruenable_cachingzCUDADispatcher.enable_cachingfs - ryrX)maxsizec>t||\}}t|||||Srw)rrV)rdrrrrs ru configurezCUDADispatcher.configureis&7J#D'8VYOOrycPt|dvr td|j|S)N)rYr1z.must specify at least the griddim and blockdim)r%rrr`s ru __getitem__zCUDADispatcher.__getitem__ns+ t9I %MN Nt~~t$$ryc"t|||||S)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.)rCrr)r>)rdrArCrrs ruforallzCUDADispatcher.forallss(dFFiPPryc8|jjdS)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. rC)rtgetrxs rurCzCUDADispatcher.extensionss&!!%%l33ryc ttrw)rr)rdrrSs rurHzCUDADispatcher.__call__s233ryc|jr-tt|jj }n t j j|g|}|j|||||y)zJ Compile if necessary and invoke this kernel with *args*. N) rErNrOrPrQrr _cuda_callr)rdrrrrrrns rur_zCUDADispatcher.callsX   $t~~44678F ++66tCdCF dGXvyAryc|rJ|Dcgc]}|j|}}|jt|Scc}wrw) typeof_pyvalcompiler)rdrkwsarBs ru_compile_for_argsz CUDADispatcher._compile_for_argss@w267$QD%%a($7||E(O,,8s=c t|tjS#ttf$rHt j |r1tt j|dtjcYSwxYw)NF)sync)rrargumentrrr is_cuda_array as_cuda_array)rdrs rurzCUDADispatcher.typeof_pyvalsh #w//0 0, !!#&d005A%..00 sAA31A3cjr tdtj}t fd|D}j j ||f}|r|Sj}tj|}|j||jd|_ |j ||f<|S)zd Create a new instance of this dispatcher specialized for the given *args*. zDispatcher already specializedc3@K|]}j|ywrw)r).0rrds ru z,CUDADispatcher.specialize..s True if the Dispatcher has been specialized. )rvrxs rurEzCUDADispatcher.specializeds    rycL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rPrrrErNrOrQitemsrdr[rooverloads ruget_regs_per_threadz"CUDADispatcher.get_regs_per_threads  >>)..1AA A   T^^22456FF F*.)=)=)?A)? X111)?A AAB cL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rPrrrErNrOrQrrs ruget_const_mem_sizez!CUDADispatcher.get_const_mem_sizes  >>)..1@@ @   T^^22456EE E*.)=)=)?A)? X000)?A AArcL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rPrrrErNrOrQrrs ruget_shared_mem_per_blockz'CUDADispatcher.get_shared_mem_per_block  >>)..1FF F   T^^22456KK K*.)=)=)?A)? X666)?A AArcL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rPrrrErNrOrQrrs ruget_max_threads_per_blockz(CUDADispatcher.get_max_threads_per_blocks  >>)..1GG G   T^^22456LL L*.)=)=)?A)? X777)?A AArcL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rPrrrErNrOrQrrs ruget_local_mem_per_threadz'CUDADispatcher.get_local_mem_per_thread/rrc*|jr|jt||jj}dj |}t j|||j}tj|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_devicerrAr5formatrmake_concrete_templatenopython_signaturesr pysignature)rdrr func_namerY call_templatepysigs ruget_call_templatez CUDADispatcher.get_call_templateBs      d ,LL)) ")))455 iD,D,DF !!$,,/eT3..ryc ||jvr"|j5|jjd}|jjd}|jjd}|jjd}|jjdrdnd|d}t j }t |j|||||||| } | |j|<| jj| j| j| jgd d d | S|j|} | S#1swY SxYw) 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`. r5r6r7r2r3r1r)r3r2r4N) rP_compiling_counterrtrrrDrrArFinsert_user_functionr@rLrK) rdr return_typer5r6r7r2r8r9rhs rurzCUDADispatcher.compile_device]s6 t~~ %((**..w7--11*=++//9--11*=!% 2 2 6 6u =11 (  ()<<#DLL+t*/-5+1-51=') +(,t$##889I9I9=:>,,I-)8 >>$'D 9)8 s D EE c|Dcgc]}|j}}|j||d||j|<ycc}w)NTr)_code_insertrP)rdrnrBrc_sigs ru add_overloadzCUDADispatcher.add_overloads?"*+(Q(+ UF .#)x ,s>ctj|\}}||tjk(sJ|jr,t t |jjS|jj|}||S|jj||j}||j|xxdz cc<n{|j|xxdz cc<|js t!dt#|j$|fi|j&}|j)|jj+|||j-|||S)z Compile and bind to the current context a version of this kernel specialized for the given signature. rzCompilation disabled)rnormalize_signaturernonerErNrOrPrQrr|rn targetctx _cache_hits _cache_missesrr<r/rArtr save_overloadr)rdrorBrrns rurzCUDADispatcher.compiles) !) < >),==? ?&*^^%9%9%;=%; T..00%;= ==#Bc|jjdr td||j|j S|jj Dcic]\}}||j c}}Scc}}w)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. rgz(Cannot inspect SASS of a device function)rtrr<rPrrrs rurzCUDADispatcher.inspect_sasss    ! !( +IJ J  >>),99; ;&*^^%9%9%;=%; T**,,%;= ==rc|tj}|jjD]\}}|j |y)rNr)rrrPrr)rdrrTrs rurzCUDADispatcher.inspect_typess> <::D~~++-GAt   D  ).ryc|||}|S)rr)rrArtrs rurzCUDADispatcher._rebuilds w .rycDt|j|jS)zd Reduce the instance for serialization. Compiled definitions are discarded. )rArt)rrArtrxs rurzCUDADispatcher._reduce_statess DLL"&"4"46 6ryr4)rrrrw)*r5r6r7r8 _fold_argsr targetdescrrr>r9rzr}r lru_cacherrrrCrHr_rrrFrErrrrrrrrrrrrrrr:rrr;r<s@rurrrr@s JK>J" //.Y%P&P% Q,44(4 B- 0!! A&A(A&A*A&/6%N* "HE.E0=*=, *6ryrr):numpyrrQrrr numba.corerrrrrrnumba.core.cachingr r numba.core.compiler_lockr numba.core.dispatcherr numba.core.errorsr rnumba.core.typing.typeofrrnumba.cuda.apirnumba.cuda.argsrnumba.cuda.compilerrrnumba.cuda.cudadrvrnumba.cuda.cudadrv.devicesrnumba.cuda.descriptorrnumba.cuda.errorsrr numba.cudarynumbarrwarningsrrP ReduceMixinr/objectr>rVrbrjrrrryrurs HH/9,F4-$:%2-<**i/i##i/X +V+\AA:I$ > >a6Z!6!6a6ry