1@jddlZddlZddlZddlZddlZddlmZmZm 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/m Z0ddl1m2Z2gdZ3Gddej4Z5Gdde6Z7GddZ8GddeZ9GddeZ:Gddeej4Z;dS) 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)hsinhcoshloghlog10hlog2hexphexp10hexp2hsqrthrsqrthfloorhceilhrcphrinthtrunchdivcbeZdZdZe dfd ZedZedZdZ ed Z ed Z e fd Z d Zd ZedZedZedZedZedZdZdZdZdZddZddZddZdZxZS) _Kernelz CUDA Kernel specialized for a given set of argument types. When called, this object launches the kernel on the device. NFTc z| rtdtd|_d|_||_||_||_||_|pg|_ || rdndd} tj } t|jtj|j|j|||| |  }|j}|jj}|j}|j}||j|j||| ||| \}|sg}dv|_|jrd_fd t2D}|rvt4jt4jt<}t4j|d }| ||D]}!||j"|_#|j$|_$|j%|_&|_'|j(|_(||_|j|_|j)|_)g|_*g|_+g|_,dS) Nz,Cannot compile a device function as a kernelFr)fastmathoptdebuglineinfoinliner2 nvvm_optionscccudaCGGetIntrinsicHandleTcFg|]}d|v|S)__numba_wrapper_) get_asm_str).0fnlibs FC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/dispatcher.py z$_Kernel.__init__..msBBBBb*b**coo.?.??????zcpp_function_wrappers.cu)- RuntimeErrorsuper__init__ objectmode entry_pointpy_funcargtypesr5r6 extensionsrcompute_capabilityrrvoidtarget_context__code__ co_filenameco_firstlinenoprepare_cuda_kernellibraryfndescr= cooperativeneeds_cudadevrtcuda_fp16_math_funcsospathdirnameabspath__file__joinappendadd_linking_filename entry_name signaturetype_annotation_type_annotation _codelibrary call_helper environment_referenced_environmentslifted reload_init)selfrIrJlinkr5r6r7r2rK max_registersr3devicer8r9crestgt_ctxcodefilenamelinenumkernelresbasedirfunctions_cu_pathfilepathr@ __class__s @rArFz_Kernel.__init__/s}  OMNN N         $*!?11   ! ! 4DL%*dm"&*%-#)%-)5! ###%|$#%11$, 27<2:G2?AA V  D69J9JJ   '"&C BBBB0BBB  +goobgooh&?&?@@G " W-G!I!I  KK) * * * + +H   * * * *!+ $ 4+ &k +(*% rCc|jSN)rerks rArSz_Kernel.librarys   rCc|jSr{)rdr|s rArcz_Kernel.type_annotations $$rCc|jSr{)rhr|s rA_find_referenced_environmentsz%_Kernel._find_referenced_environmentss ,,rCc4|jSr{)rNcodegenr|s rArz_Kernel.codegens"**,,,rCc4t|jjSr{)tuplerbargsr|s rAargument_typesz_Kernel.argument_typessT^()))rCc ||} t|| d| _|| _|| _|| _d| _|| _|| _ || _ || _ || _ | S)& Rebuild an instance. N) __new__rErFrHrUrarbrdrer5r6rfrK) clsrUr`rb codelibraryr5r6rfrKinstancerys rA_rebuildz_Kernel._rebuilds ;;s## c8%%'''#*"&$(! +$*(rCc 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. )rUr`rbrr5r6rfrK) dictrUrarbrer5r6rfrKr|s rA_reduce_statesz_Kernel._reduce_statessE 0t"n$:K*t} $ 0T_NNN NrCc8|jdS)z7 Force binding to current CUDA context N)re get_cufuncr|s rAbindz _Kernel.binds $$&&&&&rCcH|jjjS)zN The number of registers used by each thread for this kernel. )rerattrsregsr|s rAregs_per_threadz_Kernel.regs_per_threads  ++--388rCcH|jjjS)zD The amount of constant memory used by this kernel. )rerrconstr|s rAconst_mem_sizez_Kernel.const_mem_size  ++--399rCcH|jjjS)zM The amount of shared memory used per block for this kernel. )rerrsharedr|s rAshared_mem_per_blockz_Kernel.shared_mem_per_blocks  ++--3::rCcH|jjjS)z: The maximum allowable threads per block. )rerr maxthreadsr|s rAmax_threads_per_blockz_Kernel.max_threads_per_blocks  ++--3>>rCcH|jjjS)zM The amount of local memory used per thread for this kernel. )rerrlocalr|s rAlocal_mem_per_threadz_Kernel.local_mem_per_threadrrCc4|jS)z6 Returns the LLVM IR for this kernel. )re get_llvm_strr|s rA inspect_llvmz_Kernel.inspect_llvms --///rCc8|j|S)z7 Returns the PTX code for this kernel. )r9)rer=)rkr9s rA inspect_asmz_Kernel.inspect_asms ,,,333rCc4|jS)zv Returns the CFG of the SASS for this kernel. Requires nvdisasm to be available on the PATH. )re get_sass_cfgr|s rAinspect_sass_cfgz_Kernel.inspect_sass_cfgs  --///rCc4|jS)zp Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )reget_sassr|s rA inspect_sassz_Kernel.inspect_sasss  ))+++rCc|jtd| tj}t |jd|j|t d|t |j|t d|dS) 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================================================================================)rd ValueErrorsysstdoutprintrar)rkrs rA inspect_typesz_Kernel.inspect_typess  (?@@ @ <:D $*=*=>TJJJJ hT"""" d#$//// hT""""""rCrct}|j}t|trt jd|}||||}|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 ||zSr{)xys rAz5_Kernel.max_cooperative_grid_blocks..'s QUrC) rrer isinstancer functoolsreduce$get_active_blocks_per_multiprocessorrnMULTIPROCESSOR_COUNT)rkblockdim dynsmemsizectxcufunc active_per_smsm_counts rAmax_cooperative_grid_blocksz#_Kernel.max_cooperative_grid_blockssmm"--// h & & F '(:(:HEEH@@AIALNN :2x''rCcJ|j|jrujdz}j|\}}|t jt jksJt j} | d|g} g} t|j |D]\} } | | | || | tjr tjd}nd}|r|jp|}tjjg||||| Rd|ji|jrtjt j| ||| jdkrfdfddD}fddD}| j}|j|\}}}|d }n1|\}}}t2j|}d |d |d |d }|d|d|}|r|d|df|ddz}n|f}||| D] }| dS)N __errcode__r)streamrUcjjd|d\}}tj}t jtj||||jS)N__) moduleget_global_symbolr`ctypesc_intrdevice_to_host addressofvalue)r`memszvalrs rA load_symbolz#_Kernel.launch..load_symbolTsk$m==?E{{{?Ctt?EFFGC!,..C)&*:3*?*?bIII9$rCc,g|]}d|zS)tidrr>irs rArBz"_Kernel.launch..\s'===!{{519--===rCzyxc,g|]}d|zS)ctaidrrs rArBz"_Kernel.launch..]s'AAAaWq[11AAArCz In function z, file z, line z, ztid=z ctaid=z: )rerr5r`rrrsizeofrmemsetzipr _prepare_argsrUSE_NV_BINDINGbindingCUstreamhandle launch_kernelrUrrrrf get_exceptionrXrYr[)rkrgriddimrr sharedmemexcnameexcmemexcszexcvalretr kernelargstv zero_stream stream_handlerrrqexcclsexc_argsloclocinfosymrxlinenoprefixwbrrs @@rAlaunchz_Kernel.launch.s"--// : ,kM1G"M;;GDDMFEFM&,777777\^^F MM!FM + + + +T22 ? ?DAq   q!VT: > > > >    .11!44KKK06=?K  V] ;% ;& ;' ;+  ; (  ; ; ; *.)9  ; ; ; : (  !&"26":":FE J J J|q  %%%%%>===u===AAAA5AAA|(,(8(F(Ft(L(L%#; GG,/)C6!wx88HHFIccFNhhFLffOG18eeD',2FFHQKK @B   %HH &wHfh''  B BDDDD  rCc, t|jD]}|||||\}}t|tjrt |||}tj }tj d} tj d} ||j } ||j j } tj|} tjrt#| } tj | }|| || || || ||t'|jD]+}|||j|,t'|jD]+}|||j|,dSt|tjr8t1td|z|}||dS|tjkrZtjt7j|t6j}||dS|tjkr+tj|}||dS|tj kr+tj!|}||dS|tj"kr8tj#t#|}||dS|tj$krZ|tj!|j%|tj!|j&dS|tj'krZ|tj|j%|tj|j&dSt|tPj)tPj*frF|tj+|t6j,dSt|tj-rnt |||}|j.} tjr!tj t#| } || dSt|tj/rSta|ta|ksJtc||D]\}}|2|||||dSt|tj3rD |2|j |j4|||dS#tj$rtk||wxYwtk||)zF Convert arguments to ctypes and append to kernelargs )rrrzc_%sN)6reversedrK prepare_argsrrArrayr to_devicer c_ssize_tc_void_psizedtypeitemsizerdevice_pointerrintr^rangendimshapestridesIntegergetattrfloat16c_uint16npviewuint16float64c_doublefloat32c_floatbooleanc_uint8 complex64realimag complex128 npy_types NPDatetime NPTimedeltac_int64int64Recorddevice_ctypes_pointer BaseTuplelenrr EnumMemberrNotImplementedError)rktyrrrr extensiondevaryc_intpmeminfoparentnitemsr ptrdataaxcvaldevrecrrs rArz_Kernel._prepare_argsvs "$/22  I,, -GB b%+ & &O /c]],,T6::F%Foa((G_Q''FVFK((Fvfl344H'//C$ #hh?3''D   g & & &   f % % %   f % % %   h ' ' '   d # # #FK(( < <!!&&b)9":":;;;;FK(( > >!!&&);"<"<==== > >EM * *4 //766B;//44D   d # # # # # 5= ?2:c??#7#7 #B#BCCD   d # # # # # 5= ?3''D   d # # # # # 5= >#&&D   d # # # # # 5= >#c((++D   d # # # # # 5? " "   fnSX66 7 7 7   fnSX66 7 7 7 7 7 5# # #   foch77 8 8 8   foch77 8 8 8 8 8 Y193HI J J /   fnSXXbh-?-?@@ A A A A A EL ) ) /c]],,T6::F.C$ 0oc#hh//   c " " " " " EO , , /r77c#hh&&&&B  C C1""1azBBBB C CE, - - / 3""Hciz' 3 3 3)"c222 3&b#.. .s #W((X) NFFFFNNTFr{)rrr)__name__ __module__ __qualname____doc__r rFpropertyrSrcrrr classmethodrrrrrrrrrrrrrrrr __classcell__rys@rAr/r/)s? ;@JN6;ZZZZZZx!!X!%%X%-----X-**X*[* N N N''' 99X9 ::X: ;;X; ??X? ::X: 000 444 000,,,####"((((,FFFFP\/\/\/\/\/\/\/rCr/c eZdZdZdZdZdS)ForAllc||dkrtd|z||_||_||_||_||_dS)Nrz0Can't create ForAll with negative task count: %s)r dispatcherntasksthread_per_blockrr)rkrGrHtpbrrs rArFzForAll.__init__sP A::O%&'' '$ # "rCc|jdkrdS|jjr|j}n|jj|}||}|j|zdz |z}||||j|jf|S)Nrr)rHrG specialized specialize_compute_thread_per_blockrr)rkrrLrrs rA__call__zForAll.__call__s ;!   F ? & </KK4$/4d;K11+>>;)A-(:+{7Hdk>*+,02 2rCc$|j}|dkr|St}tt|j}t |jd|j d}|j di|\}}|S)Nri)funcb2d_funcmemsizeblocksizelimitr) rIrnextiter overloadsvaluesrrerrget_max_potential_block_size)rkrGrJrrtkwargs_s rArNz ForAll._compute_thread_per_blocks# !88J--C$z3::<<==>>F(3355# F 6S5????FAsJrCN)r<r=r>rFrOrNrrCrArErEsA### 2 2 2rCrEceZdZdZdZdS)_LaunchConfigurationc||_||_||_||_||_t jrRt jsHd}|d|dz|dz}||kr(d|d}tt|dSdSdSdS)Nrrz Grid size zB will likely result in GPU under-utilization due to low occupancy.) rGrrrrrCUDA_LOW_OCCUPANCY_WARNINGSDISABLE_PERFORMANCE_WARNINGSrr ) rkrGrrrr min_grid_size grid_sizemsgs rArFz_LaunchConfiguration.__init__s$    "  . 37 3 M WQZ/'!*rFrOrrCrAr]r]s73334AAAAArCr]c eZdZdZdZdZdS) CUDACacheImplc*|Sr{)r)rkrts rArzCUDACacheImpl.reduce$s$$&&&rCc$tjdi|S)Nr)r/r)rkrNpayloads rArebuildzCUDACacheImpl.rebuild's**'***rCcdS)NTr)rkros rAcheck_cachablezCUDACacheImpl.check_cachable*s trCN)r<r=r>rrnrprrCrArjrj#sA'''+++rCrjc&eZdZdZeZfdZxZS) CUDACachezS Implements a cache that saves and loads CUDA kernels and compile results. cddlm}|d5t||cdddS#1swxYwYdS)Nr)target_overrider)numba.core.target_extensionrtrE load_overload)rksigrNrtrys rArvzCUDACache.load_overload;s @????? _V $ $ > >77((n== > > > > > > > > > > > > > > > > > >s"AA A)r<r=r>r?rj _impl_classrvrBrCs@rArrrr5sI K>>>>>>>>>rCrrc^eZdZdZdZeZeffd Ze dZ dZ e j dd$d Zd Zd%d Ze d Zd ZdZdZdZdZe dZd&dZd&dZd&dZd&dZd&dZdZd&dZdZdZ d&dZ!d&dZ"d&dZ#d&d Z$d&d!Z%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. Fclt|||d|_i|_dS)N) targetoptionspipeline_classF)rErF _specializedspecializations)rkrIr|r}rys rArFzCUDADispatcher.__init__VsE  (6  8 8 8" "rCc*tj|Sr{) cuda_typesrzr|s rA _numba_type_zCUDADispatcher._numba_type_fs(...rCc8t|j|_dSr{)rrrI_cacher|s rAenable_cachingzCUDADispatcher.enable_cachingjs -- rCr_)maxsizercNt||\}}t|||||Sr{)rr])rkrrrrs rA configurezCUDADispatcher.configurems,7JJ#D'8VYOOOrCcVt|dvrtd|j|S)N)r`r1z.must specify at least the griddim and blockdim)r,rrrhs rA __getitem__zCUDADispatcher.__getitem__rs1 t99I % %MNN Nt~t$$rCc*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.)rJrr)rE)rkrHrJrrs rAforallzCUDADispatcher.forallws(dFFiPPPPrCc6|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. rK)r|getr|s rArKzCUDADispatcher.extensionss&!%%l333rCc*ttr{)rr)rkrrZs rArOzCUDADispatcher.__call__s2333rCc|jr4tt|j}nt jj|g|R}||||||dS)zJ Compile if necessary and invoke this kernel with *args*. N) rLrUrVrWrXrr _cuda_callr)rkrrrrrrts rArgzCUDADispatcher.callsp   D$t~44667788FF +6tCdCCCF dGXvyAAAAArCcl|rJfd|D}t|S)Nc:g|]}|Sr typeof_pyvalr>arks rArBz4CUDADispatcher._compile_for_args..s'777QD%%a((777rC)compiler)rkrkwsrJs` rA_compile_for_argsz CUDADispatcher._compile_for_argss=w7777$777||E(OO,,,rCc t|tjS#ttf$rFt j|r0tt j|dtjcYSwxYw)NF)sync)rrargumentrrr is_cuda_array as_cuda_array)rkrs rArzCUDADispatcher.typeof_pyvals #w/00 0,   !#&& d05AAA%.00000 sAA31A3cjrtdtj}t fd|D}j||f}|r|Sj}tj |}| || d|_ |j||f<|S)zd Create a new instance of this dispatcher specialized for the given *args*. zDispatcher already specializedc3BK|]}|VdSr{rrs rA z,CUDADispatcher.specialize..s1< True if the Dispatcher has been specialized. )r~r|s rArLzCUDADispatcher.specializeds   rCNc||j|jjS|jr8t t |jjSd|jDS)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. Nc$i|] \}}||jSr)rr>rwoverloads rA z6CUDADispatcher.get_regs_per_thread..s7AAA%X1AAArC)rWrrrLrUrVrXitemsrkrbs rAget_regs_per_threadz"CUDADispatcher.get_regs_per_threads  >).1A A   AT^22445566F FAA)-)=)=)?)?AAA ArCc||j|jjS|jr8t t |jjSd|jDS)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. Nc$i|] \}}||jSr)rrs rArz5CUDADispatcher.get_const_mem_size..s7AAA%X0AAArC)rWrrrLrUrVrXrrs rAget_const_mem_sizez!CUDADispatcher.get_const_mem_sizes  >).1@ @   AT^22445566E EAA)-)=)=)?)?AAA ArCc||j|jjS|jr8t t |jjSd|jDS)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. Nc$i|] \}}||jSr)rrs rArz;CUDADispatcher.get_shared_mem_per_block..7AAA%X6AAArC)rWrrrLrUrVrXrrs rAget_shared_mem_per_blockz'CUDADispatcher.get_shared_mem_per_block   >).1F F   AT^22445566K KAA)-)=)=)?)?AAA ArCc||j|jjS|jr8t t |jjSd|jDS)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. Nc$i|] \}}||jSr)rrs rArz.0s7AAA%X7AAArC)rWrrrLrUrVrXrrs rAget_max_threads_per_blockz(CUDADispatcher.get_max_threads_per_blocks  >).1G G   AT^22445566L LAA)-)=)=)?)?AAA ArCc||j|jjS|jr8t t |jjSd|jDS)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. Nc$i|] \}}||jSr)rrs rArz;CUDADispatcher.get_local_mem_per_thread..CrrC)rWrrrLrUrVrXrrs rAget_local_mem_per_threadz'CUDADispatcher.get_local_mem_per_thread3rrCc |jr"|t||jj}d|}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_devicerrIr<formatrmake_concrete_templatenopython_signaturesr pysignature)rkrr func_namer` call_templatepysigs rAget_call_templatez CUDADispatcher.get_call_templateFs   -   d , , ,L) ")))445 iD,DFFF !$,//eT3..rCc R||jvr|j5|jd}|jd}|jd}|jd}|jdrdnd|d}t j}t |j|||||||| } | |j|<| j | j | j | j gd d d n #1swxYwYn |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`. r5r6r7r2r3r1r)r3r2r4N) rW_compiling_counterr|rrrLrrIrNinsert_user_functionrHrTrS) rkr return_typer5r6r7r2r8r9ros rArzCUDADispatcher.compile_deviceas t~ % %( I I*..w77-11*==+//99-11*==!% 2 6 6u = =D111 (   ())<#DL+t*/-5+1-51=') +++(,t$#889I9=:>,III- I I I I I I I I I I I I I I I4>$'D sC0DDDcbd|D}|||d||j|<dS)Ncg|] }|j Sr)_code)r>rs rArBz/CUDADispatcher.add_overload..s+++Q+++rCTr)_insertrW)rkrtrJc_sigs rA add_overloadzCUDADispatcher.add_overloads?++(+++ UF ...#)x   rCctj|\}}||tjksJ|jr3t t |jS|j |}||S|j ||j }||j |xxdz cc<nr|j|xxdz cc<|jst!dt#|j|fi|j}||j ||||||S)z Compile and bind to the current context a version of this kernel specialized for the given signature. NrzCompilation disabled)rnormalize_signaturernonerLrUrVrWrXrrrv targetctx _cache_hits _cache_missesrrDr/rIr|r save_overloadr)rkrwrJrrts rArzCUDADispatcher.compiles\ !) .sCEEE)CX-::<<EEErCc>i|]\}}||Sr)rrs rArz/CUDADispatcher.inspect_llvm..s@EEE)CX2244EEErC)r|rrWrSrrr)rkrbrns rArzCUDADispatcher.inspect_llvms#''11   @~i08EEGGG~i0==??? EEE-1^-A-A-C-CEEEEEE-1^-A-A-C-CEEEErCctj|jd}|G|r%|j|jS|j|S|r%fd|jDSfd|jDS)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. rnNcLi|] \}}||j!Sr)rSr=r>rwrr9s rArz.CUDADispatcher.inspect_asm..sFEEE)CX-99"==EEErCcBi|]\}}||Sr)rrs rArz.CUDADispatcher.inspect_asm..sCEEE)CX11"55EEErC) rrLr|rrWrSr=rr)rkrbrnr9s @rArzCUDADispatcher.inspect_asms ! ! 4#''11   A~i08DDRHHH~i0<i|]\}}||Sr)rr>rwdefns rArz3CUDADispatcher.inspect_sass_cfg..s:===!T..00===rC)r|rrDrWrrrs rArzCUDADispatcher.inspect_sass_cfgs{   ! !( + + JHII I  >),==?? ?==%)^%9%9%;%;=== =rCc|jdrtd||j|Sd|jDS)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. rnz(Cannot inspect SASS of a device functionNc>i|]\}}||Sr)rrs rArz/CUDADispatcher.inspect_sass..s:===!T**,,===rC)r|rrDrWrrrs rArzCUDADispatcher.inspect_sasss{   ! !( + + KIJJ J  >),99;; ;==%)^%9%9%;%;=== =rCc| tj}|jD]\}}||dS)rNr)rrrWrr)rkrr[rs rArzCUDADispatcher.inspect_types sU <:D~++-- * *GAt   D  ) ) ) ) * *rCc|||}|S)rr)rrIr|rs rArzCUDADispatcher._rebuilds 3w ..rCc8t|j|jS)zd Reduce the instance for serialization. Compiled definitions are discarded. )rIr|)rrIr|r|s rArzCUDADispatcher._reduce_statess% DL"&"4666 6rCr;)rrrr{)*r<r=r>r? _fold_argsr targetdescrrrFr@rrr lru_cacherrrrKrOrgrrrMrLrrrrrrrrrrrrrrrArrrBrCs@rArzrzDsJK>J"""""" //X/...Y%%%PPP&%P%%% QQQQ,44X4(444 B B B---   0!!X! AAAA&AAAA(AAAA&AAAA*AAAA&///6%%%%N*** """HEEEE.EEEE0====*====, * * * *[6666666rCrz)<numpyrrXrrr 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.cudarnumbarrnumba.npr$warningsrrW ReduceMixinr/objectrEr]rjrrrzrrCrArs HHHHHHHHHHHHHHHH////////999999,,,,,,FFFFFFFF44444444------$$$$$$::::::::%%%%%%222222------<<<<<<<<******''''''***i/i/i/i/i/i#i/i/i/X +++++V+++\AAAAAAAA@I$ > > > > > > > >a6a6a6a6a6Z!6a6a6a6a6a6rC