ZhݜSSKrSSKrSSKrSSKrSSKrSSKJrJrJ r J r J r J r SSK JrJr SSKJr SSKJr SSKJr SSKJrJr SSKJr SS KJr SS KJrJr SS K J!r! SS K"J#r# SS K$J%r% SSK&J'r'J(r( SSK)J r* SSK+J,r, SSK+J-r- SSK.J/r/ /SQr0"SS\Rb5r2"SS\35r4"SS5r5"SS\5r6"SS\5r7"SS\\Rb5r8g) 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)hsinhcoshloghlog10hlog2hexphexp10hexp2hsqrthrsqrthfloorhceilhrcphrinthtrunchdivc4^\rSrSrSr\SU4Sjj5r\S5r\S5r Sr \S5r \S5r \ U4S j5rS rS r\S 5r\S 5r\S5r\S5r\S5rSrSrSrSrSSjrSSjrSSjrSrSrU=r$)_Kernel(zx CUDA Kernel specialized for a given set of argument types. When called, this object launches the kernel on the device. c N>U (a [S5e[TU] 5 SUlSUlXlX lX@lXPlU=(d /Ul UU (aSOSS.n [5Rn [UR [RUR URUUUU U S9 nURnUR R nUR"nUR$nUR'UR(UR*XEU UUU 5unnU(d/nSUR-5;UlUR.(aSUl[2Vs/sHnS U3UR-5;dMUPM nnU(aq[4R6R9[4R6R;[<55n[4R6R?US 5nURAU5 UHnURCU5 M URDUl#URHUl$URJUl&UUl'URPUl(XlUR*UlURRUl)/Ul*/Ul+/Ul,gs snf) 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 d/home/gothic/public_html/Fooocus/fooocus_env/lib64/python3.13/site-packages/numba/cuda/dispatcher.pyr>_Kernel.__init__.sI 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 ++(*% ;Bs J"?J"cUR$N)r^rds rurK_Kernel.librarys   cUR$rx)r]rys rur\_Kernel.type_annotations$$$r{cUR$rx)rarys ru_find_referenced_environments%_Kernel._find_referenced_environmentss,,,r{c6URR5$rx)rFcodegenrys rur_Kernel.codegens""**,,r{c@[URR5$rx)tupler[argsrys ruargument_types_Kernel.argument_typessT^^(())r{c >URU5n [X ] 5 SU lXlX)lX9lSU lXIlXYl Xil Xyl Xl U $) Rebuild an instance. N) __new__r=r>r@rNrZr[r]r^r5r6r_rC) clsrNrYr[ codelibraryr5r6r_rCinstancerts ru_rebuild_Kernel._rebuildsb ;;s# c%'#*"&$(! +$*(r{c [URURURURUR UR URURS9$)z 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_rCrys ru_reduce_states_Kernel._reduce_statessL 0 0t"nn$:K:K**t}} $ 0 0T__N Nr{c8URR5 g)z' Force binding to current CUDA context N)r^ get_cufuncrys rubind _Kernel.binds $$&r{c^URR5RR$)z> The number of registers used by each thread for this kernel. )r^rattrsregsrys ruregs_per_thread_Kernel.regs_per_threads%   ++-33888r{c^URR5RR$)z4 The amount of constant memory used by this kernel. )r^rrconstrys ruconst_mem_size_Kernel.const_mem_size%   ++-33999r{c^URR5RR$)z= The amount of shared memory used per block for this kernel. )r^rrsharedrys rushared_mem_per_block_Kernel.shared_mem_per_blocks%   ++-33:::r{c^URR5RR$)z* The maximum allowable threads per block. )r^rr maxthreadsrys rumax_threads_per_block_Kernel.max_threads_per_blocks%   ++-33>>>r{c^URR5RR$)z= The amount of local memory used per thread for this kernel. )r^rrlocalrys rulocal_mem_per_thread_Kernel.local_mem_per_threadrr{c6URR5$)z& Returns the LLVM IR for this kernel. )r^ get_llvm_strrys ru inspect_llvm_Kernel.inspect_llvms  --//r{c4URRUS9$)z' Returns the PTX code for this kernel. )r9)r^rM)rdr9s ru inspect_asm_Kernel.inspect_asms  ,,,33r{c6URR5$)z^ Returns the CFG of the SASS for this kernel. Requires nvdisasm to be available on the PATH. )r^ get_sass_cfgrys ruinspect_sass_cfg_Kernel.inspect_sass_cfgs   --//r{c6URR5$)zX Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )r^get_sassrys ru inspect_sass_Kernel.inspect_sasss   ))++r{cURc [S5eUc[Rn[ UR <SUR <3US9 [ SUS9 [ URUS9 [ SUS9 g) 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_types_Kernel.inspect_typessg  (?@ @ <::D $*=*=>TJ hT" d##$/ hT"r{c[5nURR5n[U[5(a[ R "SU5nURUUU5nURRnXV-$)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 X-$rx)xys ru5_Kernel.max_cooperative_grid_blocks..&sQUr{) rr^r isinstancer functoolsreduce$get_active_blocks_per_multiprocessorrgMULTIPROCESSOR_COUNT)rdblockdim dynsmemsizectxcufunc active_per_smsm_counts rumax_cooperative_grid_blocks#_Kernel.max_cooperative_grid_blockssrm""--/ h & & ''(:HEH@@AIALN ::22''r{cF^URR5mUR(a{TRS-nTRR U5upxU[ R"[ R5:Xde[ R"5n URSUS9 /n /n [URU5HupURXXJU 5 M [R(a [RR!S5nOSnU=(a UR"=(d Un[R$"TR"/UQUQUPUPU P7SUR&06 UR(Ga [R("[ R*"W 5WW5 U R,S:waU4SjnSVs/sHnU"SU-5PM nnSVs/sHnU"SU-5PM nnU R,nUR.R1U5unnnUcS nO4Uunnn4R7U5nS U<S U<S U<S 3nU<SU<SU<3nU(aU<SUS<34USS-nOU4nU"U6eU H nU"5 M gs snfs snf)N __errcode__r)streamrNc>TRRTR<SU<S35up[R"5n[ R "[R"U5X5 UR$)N__) moduleget_global_symbolrYctypesc_intrdevice_to_host addressofvalue)rYmemszvalrs ru load_symbol#_Kernel.launch..load_symbolSs\$mm==?E{{?C?EFGC!,,.C))&*:*:3*?I99$r{zyxtidctaidz 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 @rulaunch_Kernel.launch-sk""--/ ::kkM1G"MM;;GDMFFMM&,,77 77\\^F MM!FM + ++T2DA   qV: >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 JJc< [UR5HnURUUUUS9upM [U[R 5(Ga[ U5RXC5n[Rn[R"S5n [R"S5n U"UR5n U"URR5n [R"U5n [R (a [#U 5n [R"U 5nUR%U 5 UR%U 5 UR%U 5 UR%U 5 UR%U5 ['UR(5H'nUR%U"UR*U55 M) ['UR(5H'nUR%U"UR,U55 M) g[U[R.5(a+[1[SU-5"U5nUR%U5 gU[R2:XaY[R4"[6R2"U5R9[6R:55nUR%U5 gU[R<:Xa([R>"U5nUR%U5 gU[R@:Xa([RB"U5nUR%U5 gU[RD:Xa1[RF"[#U55nUR%U5 gU[RH:Xa_UR%[RB"URJ55 UR%[RB"URL55 gU[RN:Xa_UR%[R>"URJ55 UR%[R>"URL55 g[U[RP[RR45(aCUR%[RT"UR9[6RV555 g[U[RX5(al[ U5RXC5nURZn [R (a[R"[#U 55n UR%U 5 g[U[R\5(aD[_U5[_U5:Xde[aX5HunnURcUUX4U5 M g[U[Rd5(a*URcURURfX4U5 g[iX5e![ha [iX5ef=f)z6 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 rur_Kernel._prepare_argsusA "$//2I,, -GB3 b%++ & &c],,T:F%%Fooa(G__Q'FFKK(Ffll334H''/C$$#h??3'D   g &   f %   f %   h '   d #FKK(!!&b)9":;)FKK(!!&);"<=)EMM * *66B;/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 B   fnnSXXbhh-?@ A ELL ) )c],,T:F..C$$ooc#h/   c " EOO , ,r7c#h& &&B 1""1azB%E,, - - 3""HHciiz&b. . ' 3)"22 3s (XX)r^rar]rBr_rNr5rZr@r`rCrLrbr6r?rArcr[rF) NFFFFNNTFrx)rrr) __name__ __module__ __qualname____firstlineno____doc__r r>propertyrKr\rrr classmethodrrrrrrrrrrrrrrrr__static_attributes__ __classcell__rts@rur.r.(s1 ;@JN6;ZZx!!%%---*** N' 99 :: ;; ?? :: 0 4 0,#"(,FP\/\/r{r.c&\rSrSrSrSrSrSrg)ForAllichUS:a[SU-5eXlX lX0lX@lXPlg)Nrz0Can't create ForAll with negative task count: %s)r dispatcherntasksthread_per_blockrr)rdrZr[tpbrrs rur>ForAll.__init__s; A:O%&' '$ # "r{c.URS:XagURR(a URnOURR"U6nUR U5nURU-S- U-nX$X0R UR 4"U6$)Nrr)r[rZ specialized specialize_compute_thread_per_blockrr)rdrr`rrs ru__call__ForAll.__call__s ;;!   ?? & &//K//44d;K11+>;;)A-(:Hkk>>*+,02 2r{cURnUS:waU$[5n[[URR 555n[ URR5SURSS9nUR"S0UD6upbU$)Nri)funcb2d_funcmemsizeblocksizelimitr) r\rnextiter overloadsvaluesrr^rrget_max_potential_block_size)rdrZr]rrnkwargs_s rurb ForAll._compute_thread_per_blocks## !8J-C$z33::<=>F((335# F 55??FAJr{)rZr[rrr\N)rMrNrOrPr>rcrbrTrr{rurXrXs# 2r{rXc \rSrSrSrSrSrg)_LaunchConfigurationicXlX lX0lX@lXPl[ R (a4SnUSUS-US-nXv:aSUS3n[[U55 ggg)Nrrz Grid size zB will likely result in GPU under-utilization due to low occupancy.) rZrrrrrCUDA_LOW_OCCUPANCY_WARNINGSrr ) rdrZrrrr min_grid_size grid_sizemsgs rur>_LaunchConfiguration.__init__sy$   "  - - M WQZ/'!*rcrTrr{rursrss 3.Ar{rsc&\rSrSrSrSrSrSrg) CUDACacheImplic"UR5$rx)r)rdrns rurCUDACacheImpl.reduce s$$&&r{c.[R"S0UD6$)Nr)r.r)rdrFpayloads rurebuildCUDACacheImpl.rebuild#s*'**r{cg)NTr)rdrhs rucheck_cachableCUDACacheImpl.check_cachable&sr{rN)rMrNrOrPrrrrTrr{rurrs'+r{rc0^\rSrSrSr\rU4SjrSrU=r $) CUDACachei1zK Implements a cache that saves and loads CUDA kernels and compile results. ct>SSKJn U"S5 [TU] X5sSSS5 $!,(df  g=f)Nr)target_overrider)numba.core.target_extensionrr= load_overload)rdsigrFrrts rurCUDACache.load_overload7s* @ V $7(=% $ $s) 7r) rMrNrOrPrQr _impl_classrrTrUrVs@rurr1s K>>r{rcd^\rSrSrSrSr\r\4U4Sjjr \ S5r Sr \ R"SS9S$S j5rS rS%S jr\ S 5rS rSrSrSrSr\ S5rS&SjrS&SjrS&SjrS&SjrS&SjrSrS&SjrSr Sr!S&Sjr"S&Sjr#S&Sjr$S&Sjr%S&S jr&\'S!5r(S"r)S#r*U=r+$)'CUDADispatcheri@az 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. Fc>>[TU]XUS9 SUl0Ulg)N) targetoptionspipeline_classF)r=r> _specializedspecializations)rdrArrrts rur>CUDADispatcher.__init__Rs. (6  8" "r{c.[R"U5$rx) cuda_typesrrys ru _numba_type_CUDADispatcher._numba_type_bs((..r{c8[UR5Ulgrx)rrA_cacherys ruenable_cachingCUDADispatcher.enable_cachingfs - r{ru)maxsizec6[X5up[XX#U5$rx)rrs)rdrrrrs ru configureCUDADispatcher.configureis7J#D8YOOr{cT[U5S;a [S5eUR"U6$)N)rvr1z.must specify at least the griddim and blockdim)r<rrr~s ru __getitem__CUDADispatcher.__getitem__ns) t9I %MN N~~t$$r{c[XX#US9$)aReturns 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.)r]rr)rX)rdr[r]rrs ruforallCUDADispatcher.forallss(diPPr{c8URRS5$)a 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)rgetrys rurCCUDADispatcher.extensionss&!!%%l33r{c [[5erx)rr)rdrros rurcCUDADispatcher.__call__s233r{cUR(a-[[URR 555nO![ R R"U/UQ76nURXX4U5 g)z: Compile if necessary and invoke this kernel with *args*. N) r`rjrkrlrmrr _cuda_callr)rdrrrrrrns rur}CUDADispatcher.callsT   $t~~44678F ++66tCdCF dXyAr{cU(aeUVs/sHo0RU5PM nnUR[U55$s snfrx) typeof_pyvalcompiler)rdrkwsarBs ru_compile_for_args CUDADispatcher._compile_for_argss>w267$Q%%a($7||E(O,,8sAc[U[R5$![aL [R "U5(a/[[R "USS9[R5s$ef=f)NF)sync)rrargumentrr is_cuda_array as_cuda_array)rdrs rurCUDADispatcher.typeof_pyvalsf #w//0 0 !!#&&d005A%..00 sAA20A2c^TR(a [S5e[5Rn[ U4SjU55nTR R X#45nU(aU$TRn[TRUS9nURU5 UR5 SUl UTR X#4'U$)zL Create a new instance of this dispatcher specialized for the given *args*. zDispatcher already specializedc3F># UHnTRU5v M g7frx)r).0rrds ru ,CUDADispatcher.specialize..s>)..1AA A   T^^22456FF F*.)=)=)?A)? 111)?A AA B'cZUb#URURR$UR(a6[ [ URR 555R$URR5VVs0sHup#X#R_M snn$s snnf)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. )rlrrr`rjrkrmrrs ruget_const_mem_size!CUDADispatcher.get_const_mem_sizes  >>)..1@@ @   T^^22456EE E*.)=)=)?A)? 000)?A AArcZUb#URURR$UR(a6[ [ URR 555R$URR5VVs0sHup#X#R_M snn$s snnf)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. )rlrrr`rjrkrmrrs ruget_shared_mem_per_block'CUDADispatcher.get_shared_mem_per_block  >>)..1FF F   T^^22456KK K*.)=)=)?A)? 666)?A AArcZUb#URURR$UR(a6[ [ URR 555R$URR5VVs0sHup#X#R_M snn$s snnf)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. )rlrrr`rjrkrmrrs ruget_max_threads_per_block(CUDADispatcher.get_max_threads_per_blocks  >>)..1GG G   T^^22456LL L*.)=)=)?A)? 777)?A AArcZUb#URURR$UR(a6[ [ URR 555R$URR5VVs0sHup#X#R_M snn$s snnf)ay 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. )rlrrr`rjrkrmrrs ruget_local_mem_per_thread'CUDADispatcher.get_local_mem_per_thread/rrc,UR(aUR[U55 URRnSR U5n[ R"XCURS9n[R"UR5nXVX4$)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_devicerrArMformatrmake_concrete_templatenopython_signaturesr pysignature)rdrr func_namerY call_templatepysigs ruget_call_template CUDADispatcher.get_call_templateBs{      d ,LL)) ")))455 D,D,DF !!$,,/T..r{c XR;Ga"UR URRS5nURRS5nURRS5nURRS5nURRS5(aSOSUS.n[ 5R n[ URX!UUUUUUS 9 n XRU'U RRU RU RU R/5 S S S 5 U $URUn U $!,(df  W $=f) 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) rl_compiling_counterrrrrDrrArFinsert_user_functionr@rLrK) rdr return_typer5r6r7r2r8r9rhs rurCUDADispatcher.compile_device]s6 ~~ %((**..w7--11*=++//9--11*=!% 2 2 6 6u = =11 (  ()<<#DLL+*/-5+1-51=') +(,t$##889I9I9=:>,,I-)8 >>$'D 9)(8 s D E EcUVs/sHo3RPM nnURXASS9 XRU'gs snf)NTr)_code_insertrl)rdrnrBrc_sigs ru add_overloadCUDADispatcher.add_overloads9"*+(Q(+ U .#)x ,s;c[R"U5up#UbU[R:XdeUR(a,[ [ URR555$URRU5nUbU$URRXR5nUbURU==S- ss'OURU==S- ss'UR(d [!S5e[#UR$U40UR&D6nUR)5 URR+X5 UR-XB5 U$)zg Compile and bind to the current context a version of this kernel specialized for the given signature. rzCompilation disabled)rnormalize_signaturernoner`rjrkrlrmrrr targetctx _cache_hits _cache_missesrr<r.rArr save_overloadr)rdrrBrrns rurCUDADispatcher.compiles! !) < >),==? ?&*^^%9%9%;=%; ..00%;= ==)B cURRS5(a [S5eUbURUR 5$URR 5VVs0sHup#X#R 5_M snn$s snnf)ag 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)rrr<rlrrrs rurCUDADispatcher.inspect_sasss    ! !( + +IJ J  >>),99; ;&*^^%9%9%;=%; **,,%;= ==r cUc[RnURR5Hup#UR US9 M g)rNr)rrrlrr)rdrrpr s rurCUDADispatcher.inspect_typess< <::D~~++-GA   D  ).r{cU"X5nU$)rr)rrArrs rurCUDADispatcher._rebuilds w.r{c>[URURS9$)zL Reduce the instance for serialization. Compiled definitions are discarded. )rAr)rrArrys rurCUDADispatcher._reduce_statess DLL"&"4"46 6r{)rrrrL)rrrrx),rMrNrOrPrQ _fold_argsr targetdescrrr>rRrrr lru_cacherrrrCrcr}rrrar`rrrrrrrrrrrrrrrSrrrTrUrVs@rurr@sJK>J" //.%P&P% Q,44(4 B- 0!! A&A(A&A*A&/6%N* "HE.E0=*=, *66r{r)9numpyr(rQrrr numba.corerrrrrrnumba.core.cachingr r numba.core.compiler_lockr numba.core.dispatcherr numba.core.errorsr numba.core.typing.typeofrrnumba.cuda.apirnumba.cuda.argsrnumba.cuda.compilerrrnumba.cuda.cudadrvrnumba.cuda.cudadrv.devicesrnumba.cuda.descriptorrnumba.cuda.errorsrr numba.cudarnumbarrwarningsrrP ReduceMixinr.objectrXrsrrrrr{rur*s HH/9,54-$:%2-<**i/i##i/X +V+\AA:I$ > >a6Z!6!6a6r{