+ :i^RIt^RIt^RIt^RIt^RIt^RIHtHtH t H t H t H t ^RI HtHt^RIHt^RIHt^RIHtHt^RIHtHt^RIHt^R IHt^R IHtH t ^R I!H"t"^R I#H$t$^R I%H&t&^RI'H(t(H)t)^RI*H t+^RI,H-t-^RI,H.t.^RI/H0t0.ROt1!RR]Pd4t3!RR]44t5!RR4t6!RR]4t7!RR]4t8!RR]]Pd4t9R#) 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)warnc8aa]tRt^(toRt]RV3Rll4t]R4t]R4t Rt ]R4t ]R4t ] V3Rl4tR tR t]R 4t]R 4t]R 4t]R4t]R4tRtRtRtRtRRltRRltRRltRtRtVtV;t #)_Kernelzx CUDA Kernel specialized for a given set of argument types. When called, this object launches the kernel on the device. c d<V 'd \R4h\SV` 4RVnRVnWnW nW@nWPnT;'g.Vn RTRV 'd^M^/p \4Pp \VP \PVP VPVVVV V R7 pVPpVP P pVP"pVP$pVP'VP(VP*WEV VVV 4wppV'g.pRVP-49VnVP.'dRVn\2Uu.uFpRV 2VP-49gKVNK! ppV'dr\4P6P9\4P6P;\<44p\4P6P?VR 4pVPAV4VFpVPCV4K VPDVn#VPHVn$VPJVn&VVn'VPPVn(WnVP*VnVPRVn).Vn*.Vn+.Vn,R#uupi) z,Cannot compile a device function as a kernelFNfastmathoptdebuglineinfoinliner! nvvm_optionscccudaCGGetIntrinsicHandleT__numba_wrapper_zcpp_function_wrappers.cu)- RuntimeErrorsuper__init__ objectmode entry_pointpy_funcargtypesr$r% 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)selfr0r1linkr$r%r&r!r2 max_registersr"devicer'r(crestgt_ctxcodefilenamelinenumlibkernelfnresbasedirfunctions_cu_pathfilepath __class__s&&&&&&&&&&&& S/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/dispatcher.pyr-_Kernel.__init__.sU MN N       $**  1  ! 4 4DLL%**dmm"&**%-#)%-)5! #%%||$$##%%11$,, 27<2:G2?A V D69JJ    "&C 0B0b%bT*coo.??r0B ggoobggooh&?@G " W-G!I  KK) *H   *!++ $ 4 4++ &kk ++(*% ;Bs )J-J-cVP#N)rMrSs&rdr:_Kernel.librarys   cVP#rg)rLrhs&rdrK_Kernel.type_annotations$$$rjcVP#rg)rPrhs&rd_find_referenced_environments%_Kernel._find_referenced_environmentss,,,rjc6VPP4#rg)r5codegenrhs&rdrq_Kernel.codegens""**,,rjc@\VPP4#rg)tuplerJargsrhs&rdargument_types_Kernel.argument_typessT^^(())rjc  <VPV4p \W ` 4RV nWnW)nW9nRV nWInWYn Win Wyn Wn V #) Rebuild an instance. N) __new__r,r-r/r=rIrJrLrMr$r%rNr2) clsr=rHrJ codelibraryr$r%rNr2instancercs &&&&&&&&& rd_rebuild_Kernel._rebuildsb ;;s# c%'#*"&$(! +$*(rjc  \VPVPVPVPVP VP VPVPR7#)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. )r=rHrJr|r$r%rNr2) dictr=rIrJrMr$r%rNr2rhs&rd_reduce_states_Kernel._reduce_statessL 0 0t"nn$:K:K**t}} $ 0 0T__N Nrjc :VPP4R#)z' Force binding to current CUDA context N)rM get_cufuncrhs&rdbind _Kernel.binds $$&rjc ^VPP4PP#)z> The number of registers used by each thread for this kernel. )rMrattrsregsrhs&rdregs_per_thread_Kernel.regs_per_threads%   ++-33888rjc ^VPP4PP#)z4 The amount of constant memory used by this kernel. )rMrrconstrhs&rdconst_mem_size_Kernel.const_mem_size%   ++-33999rjc ^VPP4PP#)z= The amount of shared memory used per block for this kernel. )rMrrsharedrhs&rdshared_mem_per_block_Kernel.shared_mem_per_blocks%   ++-33:::rjc ^VPP4PP#)z* The maximum allowable threads per block. )rMrr maxthreadsrhs&rdmax_threads_per_block_Kernel.max_threads_per_blocks%   ++-33>>>rjc ^VPP4PP#)z= The amount of local memory used per thread for this kernel. )rMrrlocalrhs&rdlocal_mem_per_thread_Kernel.local_mem_per_threadrrjc 6VPP4#)z& Returns the LLVM IR for this kernel. )rM get_llvm_strrhs&rd inspect_llvm_Kernel.inspect_llvms  --//rjc :VPPVR7#)z' Returns the PTX code for this kernel. )r()rMr<)rSr(s&&rd inspect_asm_Kernel.inspect_asms  ,,,33rjc 6VPP4#)z^ Returns the CFG of the SASS for this kernel. Requires nvdisasm to be available on the PATH. )rM get_sass_cfgrhs&rdinspect_sass_cfg_Kernel.inspect_sass_cfgs   --//rjc 6VPP4#)zX Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )rMget_sassrhs&rd inspect_sass_Kernel.inspect_sasss   ))++rjc VPf \R4hVf\Pp\ VP : RVP : 2VR7\ RVR7\ VPVR7\ RVR7R#) 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================================================================================)rL ValueErrorsysstdoutprintrIrv)rSrs&&rd inspect_types_Kernel.inspect_typessg  (?@ @ <::D $*=*=>TJ hT" d##$/ hT"rjc \4pVPP4p\V\4'd\ P !RV4pVPVVV4pVPPpWV,#)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. cW,#rg)xys&&rd5_Kernel.max_cooperative_grid_blocks..&sQUrj) rrMr isinstancert functoolsreduce$get_active_blocks_per_multiprocessorrVMULTIPROCESSOR_COUNT)rSblockdim dynsmemsizectxcufunc active_per_smsm_counts&&& rdmax_cooperative_grid_blocks#_Kernel.max_cooperative_grid_blockssrm""--/ h & & ''(:HEH@@AIALN ::22''rjcaVPP4oVP'dSPR,pSPP V4wrxV\ P!\ P48XgQh\ P!4p VP^VR7.p .p \VPV4FwrVPWWJV 4K \P'd!\PP!^4pMRpT;'dVP";'gTp\P$!SP".VOVOVNVNV N5RVP&/VP'Ed$\P(!\ P*!X 4XX4V P,^8wdV3RlpRUu.uFpV!RV,4NK ppRUu.uFpV!RV,4NK ppV P,pVP.P1V4wpppVfRpM4Vwppp\2P4P7V4pR V: R V: R V: R 2pV: R V: RV: 2pV'd!V: RV^,: 23VR,,pMV3pV!V!hV F pV!4K R#uupiuupi) __errcode__)streamNr=c<SPPSP: RV: R24wr\P!4p\ P !\P!V4W4VP#)__) moduleget_global_symbolrHctypesc_intrdevice_to_host addressofvalue)rHmemszvalrs& rd load_symbol#_Kernel.launch..load_symbolSs\$mm==?E{{?C?EFGC!,,.C))&*:*:3*?I99$rjzyxtidctaidz In function z, file z, line z, ztid=z ctaid=z: :NN)rMrr$rHrrrsizeofrmemsetziprv _prepare_argsrUSE_NV_BINDINGbindingCUstreamhandle launch_kernelr=rrrrN get_exceptionr@rArC)rSrugriddimrr sharedmemexcnameexcmemexcszexcvalretr kernelargstv zero_stream stream_handlerirrrYexcclsexc_argsloclocinfosymrblinenoprefixwbrs&&&&&& @rdlaunch_Kernel.launch-st""--/ :::kkM1G"MM;;GDMFFMM&,,77 77\\^F MM!FM + ++T2DA   qV: >3   ..11!4KK006==??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 J=-Kc \VP4FpVPVVVVR7wrK \V\P 4'Ed\ V4PWC4p\Pp\P!^4p \P!^4p V!VP4p V!VPP4p \P!V4p \P 'd \#V 4p \P!V 4pVP%V 4VP%V 4VP%V 4VP%V 4VP%V4\'VP(4F+pVP%V!VP*V,44K- \'VP(4F+pVP%V!VP,V,44K- R#\V\P.4'd1\1\RV,4!V4pVP%V4R#V\P28Xd[\P4!\6P2!V4P9\6P:44pVP%V4R#V\P<8Xd*\P>!V4pVP%V4R#V\P@8Xd*\PB!V4pVP%V4R#V\PD8Xd3\PF!\#V44pVP%V4R#V\PH8XdaVP%\PB!VPJ44VP%\PB!VPL44R#V\PN8XdaVP%\P>!VPJ44VP%\P>!VPL44R#\V\PP\PR34'dEVP%\PT!VP9\6PV444R#\V\PX4'do\ V4PWC4pVPZp \P 'd \P!\#V 44p VP%V 4R#\V\P\4'dG\_V4\_V48XgQh\aW4FwppVPcVVW4V4K R#\V\Pd4'd,VPcVPVPfW4V4R#\iW4h \hd \iY4hi;i)z6 Convert arguments to ctypes and append to kernelargs )rrzc_%sN)5reversedr2 prepare_argsrrArrayr to_devicer c_ssize_tc_void_psizedtypeitemsizerdevice_pointerrintrFrangendimshapestridesIntegergetattrfloat16c_uint16npviewuint16float64c_doublefloat32c_floatbooleanc_uint8 complex64realimag complex128 NPDatetime NPTimedeltac_int64int64Recorddevice_ctypes_pointer BaseTuplelenrr EnumMemberrNotImplementedError)rStyrrrr extensiondevaryc_intpmeminfoparentnitemsr ptrdataaxcvaldevrecrrs&&&&&& rdr_Kernel._prepare_argsusG "$//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 :(X//Y)rMrPrLr1rNr=r$rIr/rOr2r;rQr%r.r0rRrJr5) NFFFFNNTFrg)rrr)!__name__ __module__ __qualname____firstlineno____doc__r r-propertyr:rKrnrqrv classmethodr~rrrrrrrrrrrrrrr__static_attributes____classdictcell__ __classcell__rc __classdict__s@@rdrr(s* ZZx!!%%---*** N' 99 :: ;; ?? :: 0 4 0,#"(,FP\/\/rjrc2a]tRtRtoRtRtRtRtVtR#)ForAllictV^8d\RV,4hWnW nW0nW@nWPnR#)rz0Can't create ForAll with negative task count: %sN)r dispatcherntasksthread_per_blockrr)rSrKrLtpbrrs&&&&&&rdr-ForAll.__init__s; A:O%&' '$ # "rjcXVP^8XdR#VPP'dVPpMVPP!V!pVP V4pVPV,^, V,pW$W0P VP 3,!V!#)rN)rLrK specialized specialize_compute_thread_per_blockrr)rSrurQrrs&* rd__call__ForAll.__call__s ;;!   ?? & & &//K//44d;K11+>;;)A-(:Hkk>>*+,02 2rjc&VPpV^8wdV#\4p\\VPP 444p\ VPP4^VPRR7pVP!R/VBwrbV#)ri)funcb2d_funcmemsizeblocksizelimitr) rMrnextiter overloadsvaluesrrMrrget_max_potential_block_size)rSrKrNrr]kwargs_s&& rdrS ForAll._compute_thread_per_blocks## !8J-C$z33::<=>F((335# F 55??FAJrj)rKrLrrrMN) r<r=r>r?r-rTrSrCrDrGs@rdrIrIs# 2rjrIc,a]tRtRtoRtRtRtVtR#)_LaunchConfigurationicWnW nW0nW@nWPn\ P 'dL^pV^,V^,,V^,,pWv8dRV R2p\\V44R#R#R#)z Grid size zB will likely result in GPU under-utilization due to low occupancy.N) rKrrrrrCUDA_LOW_OCCUPANCY_WARNINGSrr ) rSrKrrrr min_grid_size grid_sizemsgs &&&&&& rdr-_LaunchConfiguration.__init__s|$   "  - - - M WQZ/'!*r?r-rTrCrDrcs@rdreres3.AArjrec2a]tRtRtoRtRtRtRtVtR#) CUDACacheImplic"VP4#rg)r)rSr]s&&rdrCUDACacheImpl.reduce s$$&&rjc.\P!R/VB#)Nr)rr~)rSr5payloads&&&rdrebuildCUDACacheImpl.rebuild#s*'**rjcR#)Tr)rSrWs&&rdcheck_cachableCUDACacheImpl.check_cachable&srjrN) r<r=r>r?rrwrzrCrDrcs@rdrrrrs'+rjrrc:aa]tRtRtoRt]tV3RltRtVt V;t #) CUDACachei1zK Implements a cache that saves and loads CUDA kernels and compile results. c<^RIHpV!R4;_uu_4\SV` W4uuRRR4# +'giR#;i)r)target_overriderN)numba.core.target_extensionrr, load_overload)rSsigr5rrcs&&& rdrCUDACache.load_overload7s0 @ V $ $7(=% $ $ $s 2 A r) r<r=r>r?r@rr _impl_classrrCrDrErFs@@rdr}r}1s K>>rjr}ctaa]tRtRtoRtRt]t]3V3Rllt ] R4t Rt ] P!^R7R#Rl4tR tR$R lt] R 4tR tR tRtRtRt] R4tR%RltR%RltR%RltR%RltR%RltRtR%RltRt Rt!R%Rlt"R%Rlt#R%Rlt$R%Rlt%R%Rlt&]'R 4t(R!t)R"t*Vt+V;t,#)&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. FcF<\SV`WVR7RVn/VnR#)) targetoptionspipeline_classFN)r,r- _specializedspecializations)rSr0rrrcs&&&&rdr-CUDADispatcher.__init__Rs. (6  8" "rjc.\P!V4#rg) cuda_typesrrhs&rd _numba_type_CUDADispatcher._numba_type_bs((..rjc:\VP4VnR#rg)r}r0_cacherhs&rdenable_cachingCUDADispatcher.enable_cachingfs - rj)maxsizec6\W4wr\WW#V4#rg)rre)rSrrrrs&&&&&rd configureCUDADispatcher.configureis7J#D8YOOrjcX\V4R9d \R4hVP!V!#)z.must specify at least the griddim and blockdim)r)r+rrros&&rd __getitem__CUDADispatcher.__getitem__ns) t9I %MN N~~t$$rjc \WW#VR7#)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.)rNrr)rI)rSrLrNrrs&&&&&rdforallCUDADispatcher.forallss(diPPrjc 8VPPR4#)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. r2)rgetrhs&rdr2CUDADispatcher.extensionss&!!%%l33rjc \\4hrg)rr)rSrur`s&*,rdrTCUDADispatcher.__call__s233rjc VP'd.\\VPP 444pM"\ P P!V.VO5!pVPWW4V4R#)z: Compile if necessary and invoke this kernel with *args*. N) rQr[r\r]r^rr _cuda_callr)rSrurrrrr]s&&&&&& rdrnCUDADispatcher.callsW    $t~~44678F ++66tCdCF dXyArjcV'dQhVUu.uFq0PV4NK ppVP\V44#uupirg) typeof_pyvalcompilert)rSrukwsar1s&*, rd_compile_for_args CUDADispatcher._compile_for_argss>w267$Q%%a($7||E(O,,8sAc\V\P4# \\3dQ\ P !T4'd3\\ P!TRR7\P4u#hi;i)F)sync)rrargumentrrr is_cuda_array as_cuda_array)rSrs&&rdrCUDADispatcher.typeof_pyvalsk #w//0 0, !!#&&d005A%..00 s,A= 0A=;A=c aSP'd \R4h\4Pp\;QJd.V3RlV4FNK 5M!V3RlV44pSP P W#34pV'dV#SPp\SPVR7pVPV4VP4RVn VSP W#3&V#)zL Create a new instance of this dispatcher specialized for the given *args*. zDispatcher already specializedc3F<"TFpSPV4xK R#5irg)r).0rrSs& rd ,CUDADispatcher.specialize..s>)..1AA A    T^^22456FF F*.)=)=)?A)? 111)?A AAB,c dVe(VPVP,P#VP'd7\ \ VPP 444P#VPP4UUu/uFwr#W#PbK upp#uuppi)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. )r]rurrQr[r\r^rrs&& rdget_const_mem_size!CUDADispatcher.get_const_mem_sizes  >>)..1@@ @    T^^22456EE E*.)=)=)?A)? 000)?A AArc dVe(VPVP,P#VP'd7\ \ VPP 444P#VPP4UUu/uFwr#W#PbK upp#uuppi)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. )r]rurrQr[r\r^rrs&& rdget_shared_mem_per_block'CUDADispatcher.get_shared_mem_per_block  >>)..1FF F    T^^22456KK K*.)=)=)?A)? 666)?A AArc dVe(VPVP,P#VP'd7\ \ VPP 444P#VPP4UUu/uFwr#W#PbK upp#uuppi)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. )r]rurrQr[r\r^rrs&& rdget_max_threads_per_block(CUDADispatcher.get_max_threads_per_blocks  >>)..1GG G    T^^22456LL L*.)=)=)?A)? 777)?A AArc dVe(VPVP,P#VP'd7\ \ VPP 444P#VPP4UUu/uFwr#W#PbK upp#uuppi)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. )r]rurrQr[r\r^rrs&& rdget_local_mem_per_thread'CUDADispatcher.get_local_mem_per_thread/rrc 4VP'dVP\V44VPPpRP V4p\ P!WCVPR7p\P!VP4pWVW3#)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_devicertr0r<formatrmake_concrete_templatenopython_signaturesr pysignature)rSrur func_namerH call_templatepysigs&&& rdget_call_template CUDADispatcher.get_call_templateBs~       d ,LL)) ")))455 D,D,DF !!$,,/T..rjc WP9Ed0VP;_uu_4VPPR4pVPPR4pVPPR4pVPPR4pRVPPR4'd^M^RV/p\ 4P p\ VPW!VVVVVVR7 p WPV&V PPV PV PV P.4RRR4V #VPV,p V # +'giX #;i)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&r!r"r#N) r]_compiling_counterrrrr3rr0r5insert_user_functionr/r;r:) rSru return_typer$r%r&r!r'r(rWs &&& rdrCUDADispatcher.compile_device]s= ~~ %(((**..w7--11*=++//9--11*= 2 2 6 6u = =11  ()<<#DLL+*/-5+1-51=') +(,t$##889I9I9=:>,,I-)8 >>$'D 9)(8 s DE E' cVUu.uFq3PNK ppVPWARR7WPV&R#uupi)TrN)_code_insertr])rSr]r1rc_sigs&&& rd add_overloadCUDADispatcher.add_overloads9"*+(Q(+ U .#)x ,s>c \P!V4wr#VeV\P8XgQhVP'd-\ \ VPP444#VPPV4pVeV#VPPWP4pVe!VPV;;,^, uu&MVPV;;,^, uu&VP'g \!R4h\#VP$V3/VP&BpVP)4VPP+W4VP-WB4V#)zg Compile and bind to the current context a version of this kernel specialized for the given signature. zCompilation disabled)rnormalize_signaturernonerQr[r\r]r^rrr targetctx _cache_hits _cache_missesrr+rr0rr save_overloadr)rSrr1rr]s&& rdrCUDADispatcher.compiles' !) < Dc D\4PpVPPR4pVeWV'd-VPV,P P V4#VPV,PV4#V'dHVPP4UUu/uF wrEWEP P V4bK" upp#VPP4UUu/uFwrEWEPV4bK upp#uuppiuuppi)z 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. rV) rr3rrr]r:r<rr)rSrJr(rVrrs&& rdrCUDADispatcher.inspect_asms ! 4 4##''1  ~~i088DDRHH~~i0<>),==? ?&*^^%9%9%;=%; ..00%;= ==/Bc (VPPR4'd \R4hVe"VPV,P 4#VPP 4UUu/uFwr#W#P 4bK upp#uuppi)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. rVz(Cannot inspect SASS of a device function)rrr+r]rrrs&& rdrCUDADispatcher.inspect_sasss    ! !( + +IJ J  >>),99; ;&*^^%9%9%;=%; **,,%;= ==rc Vf\PpVPP4Fwr#VP VR7K R#)rNr)rrr]rr)rSrrars&& rdrCUDADispatcher.inspect_typess< <::D~~++-GA   D  ).rjc V!W4pV#)ryr)r{r0rr}s&&& rdr~CUDADispatcher._rebuilds w.rjc D\VPVPR7#)zL Reduce the instance for serialization. Compiled definitions are discarded. )r0r)rr0rrhs&rdrCUDADispatcher._reduce_statess DLL"&"4"46 6rj)rrrr;)rrrrg)-r<r=r>r?r@ _fold_argsr targetdescrrr-rArrr lru_cacherrrr2rTrnrrrRrQrrrrrrrrrrrrrrrBr~rrCrDrErFs@@rdrr@sJK>J" //.%P&P% Q,44(4 B- 0!! A&A(A&A*A&/6%N* "HE.E0=*=, *66rjr)hsinhcoshloghlog10hlog2hexphexp10hexp2hsqrthrsqrthfloorhceilhrcphrinthtrunchdiv):numpyrr@rrr 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.cudarnumbarrwarningsrr? ReduceMixinrobjectrIrerrr}rrrjrdr-s HH/9,F4-$:%2-<**i/i##i/X +V+\AA:I$ > >a6Z!6!6a6rj