|hSSKrSSKrSSKrSSKrSSKrSSKJrJrJ r J r J r J r SSK JrJr SSKJr SSKJr SSKJrJr SSKJrJr SSKJr SS KJr SS KJrJ r 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/J0r0 /SQr1"SS\Rd5r3"SS\45r5"SS5r6"SS\5r7"SS\5r8"SS\\Rd5r9g) 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)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)fastmathoptdebuglineinfoinliner3 nvvm_optionscccudaCGGetIntrinsicHandleT__numba_wrapper_zcpp_function_wrappers.cu)- RuntimeErrorsuper__init__ objectmode entry_pointpy_funcargtypesr6r7 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)selfrBrClinkr6r7r8r3rD max_registersr4devicer9r:crestgt_ctxcodefilenamelinenumlibkernelfnresbasedirfunctions_cu_pathfilepath __class__s O/home/james-whalen/.local/lib/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_res rvrL_Kernel.librarys   cUR$ry)r^rzs rvr]_Kernel.type_annotations$$$r|cUR$ry)rbrzs rv_find_referenced_environments%_Kernel._find_referenced_environmentss,,,r|c6URR5$ry)rGcodegenrzs rvr_Kernel.codegens""**,,r|c@[URR5$ry)tupler\argsrzs rvargument_types_Kernel.argument_typessT^^(())r|c >URU5n [X ] 5 SU lXlX)lX9lSU lXIlXYl Xil Xyl Xl U $) Rebuild an instance. N) __new__r>r?rArOr[r\r^r_r6r7r`rD) clsrOrZr\ codelibraryr6r7r`rDinstancerus rv_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. )rOrZr\rr6r7r`rD) dictrOr[r\r_r6r7r`rDrzs rv_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_cufuncrzs rvbind _Kernel.binds $$&r|c^URR5RR$)z> The number of registers used by each thread for this kernel. )r_rattrsregsrzs rvregs_per_thread_Kernel.regs_per_threads%   ++-33888r|c^URR5RR$)z4 The amount of constant memory used by this kernel. )r_rrconstrzs rvconst_mem_size_Kernel.const_mem_size%   ++-33999r|c^URR5RR$)z= The amount of shared memory used per block for this kernel. )r_rrsharedrzs rvshared_mem_per_block_Kernel.shared_mem_per_blocks%   ++-33:::r|c^URR5RR$)z* The maximum allowable threads per block. )r_rr maxthreadsrzs rvmax_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_rrlocalrzs rvlocal_mem_per_thread_Kernel.local_mem_per_threadrr|c6URR5$)z& Returns the LLVM IR for this kernel. )r_ get_llvm_strrzs rv inspect_llvm_Kernel.inspect_llvms  --//r|c4URRUS9$)z' Returns the PTX code for this kernel. )r:)r_rN)rer:s rv 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_cfgrzs rvinspect_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_sassrzs rv 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^ ValueErrorsysstdoutprintr[r)rers rv 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-$ry)xys rv5_Kernel.max_cooperative_grid_blocks..&sQUr|) rr_r isinstancer functoolsreduce$get_active_blocks_per_multiprocessorrhMULTIPROCESSOR_COUNT)reblockdim dynsmemsizectxcufunc active_per_smsm_counts rvmax_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)streamrOc>TRRTR<SU<S35up[R"5n[ R "[R"U5X5 UR$)N__) moduleget_global_symbolrZctypesc_intrdevice_to_host addressofvalue)rZmemszvalrs rv 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_rr6rZrrrsizeofrmemsetzipr _prepare_argsrUSE_NV_BINDINGbindingCUstreamhandle launch_kernelrOrrrr` get_exceptionrRrSrU)rergriddimrr sharedmemexcnameexcmemexcszexcvalretr kernelargstv zero_stream stream_handlerirrrkexcclsexc_argsloclocinfosymrtlinenoprefixwbrs @rvlaunch_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)5reversedrD prepare_argsrrArrayr to_devicer c_ssize_tc_void_psizedtypeitemsizerdevice_pointerrintrXrangendimshapestridesIntegergetattrfloat16c_uint16npviewuint16float64c_doublefloat32c_floatbooleanc_uint8 complex64realimag complex128 NPDatetime NPTimedeltac_int64int64Recorddevice_ctypes_pointer BaseTuplelenrr EnumMemberrNotImplementedError)retyrrrr extensiondevaryc_intpmeminfoparentnitemsrptrdataaxcvaldevrecrrs rvr_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_rbr^rCr`rOr6r[rArarDrMrcr7r@rBrdr\rG) NFFFFNNTFry)rrr) __name__ __module__ __qualname____firstlineno____doc__r r?propertyrLr]rrr classmethodrrrrrrrrrrrrrrrr__static_attributes__ __classcell__rus@rvr/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)rer[r\tpbrrs rvr?ForAll.__init__s; A:O%&' '$ # "r|c.URS:XagURR(a URnOURR"U6nUR U5nURU-S- U-nX$X0R UR 4"U6$)Nrr)r\r[ specialized specialize_compute_thread_per_blockrr)rerrarrs rv__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)rer[r^rrokwargs_s rvrc ForAll._compute_thread_per_blocks## !8J-C$z33::<=>F((335# F 55??FAJr|)r[r\rrr]N)rNrOrPrQr?rdrcrUrr|rvrYrYs# 2r|rYc \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.) r[rrrrrCUDA_LOW_OCCUPANCY_WARNINGSrr ) rer[rrrr min_grid_size grid_sizemsgs rvr?_LaunchConfiguration.__init__sy$   "  - - M WQZ/'!*SSKJn U"S5 [TU] X5sSSS5 $!,(df  g=f)Nr)target_overrider)numba.core.target_extensionrr> load_overload)resigrGrrus rvrCUDACache.load_overload7s* @ V $7(=% $ $s) 7r) rNrOrPrQrRr _impl_classrrUrVrWs@rvrr1s 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)rerBrrrus rvr?CUDADispatcher.__init__Rs. (6  8" "r|c.[R"U5$ry) cuda_typesrrzs rv _numba_type_CUDADispatcher._numba_type_bs((..r|c8[UR5Ulgry)rrB_cacherzs rvenable_cachingCUDADispatcher.enable_cachingfs - r|rv)maxsizec6[X5up[XX#U5$ry)rrt)rerrrrs rv configureCUDADispatcher.configureis7J#D8YOOr|cT[U5S;a [S5eUR"U6$)N)rwr2z.must specify at least the griddim and blockdim)r=rrrs rv __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)rY)rer\r^rrs rvforallCUDADispatcher.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. rD)rgetrzs rvrDCUDADispatcher.extensionss&!!%%l33r|c [[5ery)rr)rerrps rvrdCUDADispatcher.__call__s233r|cUR(a-[[URR 555nO![ R R"U/UQ76nURXX4U5 g)z: Compile if necessary and invoke this kernel with *args*. N) rarkrlrmrnrr _cuda_callr)rerrrrrros rvr~CUDADispatcher.callsT   $t~~44678F ++66tCdCF dXyAr|cU(aeUVs/sHo0RU5PM nnUR[U55$s snfry) typeof_pyvalcompiler)rerkwsarCs rv_compile_for_args CUDADispatcher._compile_for_argss>w267$Q%%a($7||E(O,,8sAc[U[R5$![[4aL [ R "U5(a/[[ R"USS9[R5s$ef=f)NF)sync)rrargumentrrr is_cuda_array as_cuda_array)rers rvrCUDADispatcher.typeof_pyvalsk #w//0 0, !!#&&d005A%..00 sAA86A8c^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 g7fry)r).0rres rv ,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. )rmrrrarkrlrnrrs rvget_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. )rmrrrarkrlrnrrs rvget_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. )rmrrrarkrlrnrrs rvget_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. )rmrrrarkrlrnrrs rvget_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_devicerrBrNformatrmake_concrete_templatenopython_signaturesr pysignature)rerr func_namerZ call_templatepysigs rvget_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`. r6r7r8r3r4r2r)r4r3r5N) rm_compiling_counterrrrrErrBrGinsert_user_functionrArMrL) rer return_typer6r7r8r3r9r:ris rvrCUDADispatcher.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_insertrm)rerorCrc_sigs rv 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_signaturernonerarkrlrmrnrrr targetctx _cache_hits _cache_missesrr=r/rBrr save_overloadr)rerrCrros rvrCUDADispatcher.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. rhz(Cannot inspect SASS of a device function)rrr=rmrrr s rvrCUDADispatcher.inspect_sasss    ! !( + +IJ J  >>),99; ;&*^^%9%9%;=%; **,,%;= ==r cUc[RnURR5Hup#UR US9 M g)rNr)rrrmrr)rerrqr s rvrCUDADispatcher.inspect_typess< <::D~~++-GA   D  ).r|cU"X5nU$)rr)rrBrrs rvrCUDADispatcher._rebuilds w.r|c>[URURS9$)zL Reduce the instance for serialization. Compiled definitions are discarded. )rBr)rrBrrzs rvrCUDADispatcher._reduce_statess DLL"&"4"46 6r|)rrrrM)rrrry),rNrOrPrQrR _fold_argsr targetdescrrr?rSrrr lru_cacherrrrDrdr~rrrbrarrrrrrrrrrrrrrrTrrrUrVrWs@rvrr@sJK>J" //.%P&P% Q,44(4 B- 0!! A&A(A&A*A&/6%N* "HE.E0=*=, *66r|r):numpyr)rRrrr 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.cudarnumbarrwarningsrrQ ReduceMixinr/objectrYrtrrrrr|rvr+s HH/9,F4-$:%2-<**i/i##i/X +V+\AA:I$ > >a6Z!6!6a6r|