J/PhݜddlZddlZddlZddlZddlZddlmZmZm Z m Z m Z m Z ddl mZmZddlmZddlmZddlmZddlmZmZddlmZdd lmZdd lmZmZdd 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/Z/gdZ0Gddej1Z2Gdde3Z4GddZ5GddeZ6GddeZ7Gddeej1Z8dS) 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)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)fastmathoptdebuglineinfoinliner1 nvvm_optionscccudaCGGetIntrinsicHandleTcFg|]}d|v|S)__numba_wrapper_) get_asm_str).0fnlibs U/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/dispatcher.py z$_Kernel.__init__..lsBBBBb*b**coo.?.??????zcpp_function_wrappers.cu)- RuntimeErrorsuper__init__ objectmode entry_pointpy_funcargtypesr4r5 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)selfrHrIlinkr4r5r6r1rJ max_registersr2devicer7r8crestgt_ctxcodefilenamelinenumkernelresbasedirfunctions_cu_pathfilepathr? __class__s @r@rEz_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 +(*% rBc|jSN)rdrjs r@rRz_Kernel.librarys   rBc|jSrz)rcr{s r@rbz_Kernel.type_annotations $$rBc|jSrz)rgr{s r@_find_referenced_environmentsz%_Kernel._find_referenced_environmentss ,,rBc4|jSrz)rMcodegenr{s r@rz_Kernel.codegens"**,,,rBc4t|jjSrz)tupleraargsr{s r@argument_typesz_Kernel.argument_typessT^()))rBc ||} t|| d| _|| _|| _|| _d| _|| _|| _ || _ || _ || _ | S)& Rebuild an instance. N) __new__rDrErGrTr`rarcrdr4r5rerJ) clsrTr_ra codelibraryr4r5rerJinstancerxs r@_rebuildz_Kernel._rebuilds ;;s## c8%%'''#*"&$(! +$*(rBc 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. )rTr_rarr4r5rerJ) dictrTr`rardr4r5rerJr{s r@_reduce_statesz_Kernel._reduce_statessE 0t"n$:K*t} $ 0T_NNN NrBc8|jdS)z7 Force binding to current CUDA context N)rd get_cufuncr{s r@bindz _Kernel.binds $$&&&&&rBcH|jjjS)zN The number of registers used by each thread for this kernel. )rdrattrsregsr{s r@regs_per_threadz_Kernel.regs_per_threads  ++--388rBcH|jjjS)zD The amount of constant memory used by this kernel. )rdrrconstr{s r@const_mem_sizez_Kernel.const_mem_size  ++--399rBcH|jjjS)zM The amount of shared memory used per block for this kernel. )rdrrsharedr{s r@shared_mem_per_blockz_Kernel.shared_mem_per_blocks  ++--3::rBcH|jjjS)z: The maximum allowable threads per block. )rdrr maxthreadsr{s r@max_threads_per_blockz_Kernel.max_threads_per_blocks  ++--3>>rBcH|jjjS)zM The amount of local memory used per thread for this kernel. )rdrrlocalr{s r@local_mem_per_threadz_Kernel.local_mem_per_threadrrBc4|jS)z6 Returns the LLVM IR for this kernel. )rd get_llvm_strr{s r@ inspect_llvmz_Kernel.inspect_llvms --///rBc8|j|S)z7 Returns the PTX code for this kernel. )r8)rdr<)rjr8s r@ inspect_asmz_Kernel.inspect_asms ,,,333rBc4|jS)zv Returns the CFG of the SASS for this kernel. Requires nvdisasm to be available on the PATH. )rd get_sass_cfgr{s r@inspect_sass_cfgz_Kernel.inspect_sass_cfgs  --///rBc4|jS)zp Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )rdget_sassr{s r@ inspect_sassz_Kernel.inspect_sasss  ))+++rBc|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================================================================================)rc ValueErrorsysstdoutprintr`r)rjrs r@ inspect_typesz_Kernel.inspect_typess  (?@@ @ <:D $*=*=>TJJJJ hT"""" d#$//// hT""""""rBrct}|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 ||zSrz)xys r@z5_Kernel.max_cooperative_grid_blocks..&s QUrB) rrdr isinstancer functoolsreduce$get_active_blocks_per_multiprocessorrmMULTIPROCESSOR_COUNT)rjblockdim dynsmemsizectxcufunc active_per_smsm_counts r@max_cooperative_grid_blocksz#_Kernel.max_cooperative_grid_blockssmm"--// h & & F '(:(:HEEH@@AIALNN :2x''rBcJ|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)streamrTcjjd|d\}}tj}t jtj||||jS)N__) moduleget_global_symbolr_ctypesc_intrdevice_to_host addressofvalue)r_memszvalrs r@ load_symbolz#_Kernel.launch..load_symbolSsk$m==?E{{{?Ctt?EFFGC!,..C)&*:3*?*?bIII9$rBc,g|]}d|zS)tidrr=irs r@rAz"_Kernel.launch..[s'===!{{519--===rBzyxc,g|]}d|zS)ctaidrrs r@rAz"_Kernel.launch..\s'AAAaWq[11AAArBz In function z, file z, line z, ztid=z ctaid=z: )rdrr4r_rrrsizeofrmemsetzipr _prepare_argsrUSE_NV_BINDINGbindingCUstreamhandle launch_kernelrTrrrre get_exceptionrWrXrZ)rjrgriddimrr sharedmemexcnameexcmemexcszexcvalretr kernelargstv zero_stream stream_handlerrrpexcclsexc_argsloclocinfosymrwlinenoprefixwbrrs @@r@launchz_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  rBc, 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|tj(tj)frF|tj*|t6j+dSt|tj,rnt |||}|j-} tjr!tj t#| } || dSt|tj.rSt_|t_|ksJta||D]\}}|1|||||dSt|tj2rD |1|j |j3|||dS#th$rti||wxYwti||)zF Convert arguments to ctypes and append to kernelargs )rrrzc_%sN)5reversedrJ prepare_argsrrArrayr to_devicer c_ssize_tc_void_psizedtypeitemsizerdevice_pointerrintr]rangendimshapestridesIntegergetattrfloat16c_uint16npviewuint16float64c_doublefloat32c_floatbooleanc_uint8 complex64realimag complex128 NPDatetime NPTimedeltac_int64int64Recorddevice_ctypes_pointer BaseTuplelenrr EnumMemberrNotImplementedError)rjtyrrrr extensiondevaryc_intpmeminfoparentnitemsr ptrdataaxcvaldevrecrrs r@rz_Kernel._prepare_argsus "$/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 U-u/@A B B /   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) NFFFFNNTFrz)rrr)__name__ __module__ __qualname____doc__r rEpropertyrRrbr~rr classmethodrrrrrrrrrrrrrrrr __classcell__rxs@r@r.r.(s? ;@JN6;ZZZZZZx!!X!%%X%-----X-**X*[* N N N''' 99X9 ::X: ;;X; ??X? ::X: 000 444 000,,,####"((((,FFFFP\/\/\/\/\/\/\/rBr.c eZdZdZdZdZdS)ForAllc||dkrtd|z||_||_||_||_||_dS)Nrz0Can't create ForAll with negative task count: %s)r dispatcherntasksthread_per_blockrr)rjrErFtpbrrs r@rEzForAll.__init__sP A::O%&'' '$ # "rBc|jdkrdS|jjr|j}n|jj|}||}|j|zdz |z}||||j|jf|S)Nrr)rFrE specialized specialize_compute_thread_per_blockrr)rjrrJrrs r@__call__zForAll.__call__s ;!   F ? & </KK4$/4d;K11+>>;)A-(:+{7Hdk>*+,02 2rBc$|j}|dkr|St}tt|j}t |jd|j d}|j di|\}}|S)Nri)funcb2d_funcmemsizeblocksizelimitr) rGrnextiter overloadsvaluesrrdrrget_max_potential_block_size)rjrErHrrskwargs_s r@rLz ForAll._compute_thread_per_blocks# !88J--C$z3::<<==>>F(3355# F 6S5????FAsJrBN)r:r;r<rErMrLrrBr@rCrCsA### 2 2 2rBrCceZdZdZdZdS)_LaunchConfigurationc||_||_||_||_||_t jrFd}|d|dz|dz}||kr&d|d}tt|dSdSdS)Nrrz Grid size zB will likely result in GPU under-utilization due to low occupancy.) rErrrrrCUDA_LOW_OCCUPANCY_WARNINGSrr ) rjrErrrr min_grid_size grid_sizemsgs r@rEz_LaunchConfiguration.__init__s$    "  - 3 M WQZ/'!* >77((n== > > > > > > > > > > > > > > > > > >s"AA A)r:r;r<r=rg _impl_classrsr@rAs@r@roro1sI K>>>>>>>>>rBroc^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)rDrE _specializedspecializations)rjrHryrzrxs r@rEzCUDADispatcher.__init__RsE  (6  8 8 8" "rBc*tj|Srz) cuda_typesrwr{s r@ _numba_type_zCUDADispatcher._numba_type_bs(...rBc8t|j|_dSrz)rorH_cacher{s r@enable_cachingzCUDADispatcher.enable_cachingfs -- rBr])maxsizercNt||\}}t|||||Srz)rr[)rjrrrrs r@ configurezCUDADispatcher.configureis,7JJ#D'8VYOOOrBcVt|dvrtd|j|S)N)r^r0z.must specify at least the griddim and blockdim)r*rrres r@ __getitem__zCUDADispatcher.__getitem__ns1 t99I % %MNN Nt~t$$rBc*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.)rHrr)rC)rjrFrHrrs r@forallzCUDADispatcher.forallss(dFFiPPPPrBc6|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. rJ)rygetr{s r@rJzCUDADispatcher.extensionss&!%%l333rBc*ttrz)rr)rjrrXs r@rMzCUDADispatcher.__call__s2333rBc|jr4tt|j}nt jj|g|R}||||||dS)zJ Compile if necessary and invoke this kernel with *args*. N) rJrSrTrUrVrr _cuda_callr)rjrrrrrrss r@rdzCUDADispatcher.callsp   D$t~44667788FF +6tCdCCCF dGXvyAAAAArBcl|rJfd|D}t|S)Nc:g|]}|Sr typeof_pyvalr=arjs r@rAz4CUDADispatcher._compile_for_args..s'777QD%%a((777rB)compiler)rjrkwsrIs` r@_compile_for_argsz CUDADispatcher._compile_for_argss=7777$777||E(OO,,,rBc t|tjS#t$rFt j|r0tt j|dtjcYSwxYw)NF)sync)rrargumentrr is_cuda_array as_cuda_array)rjrs r@rzCUDADispatcher.typeof_pyvals #w/00 0   !#&& d05AAA%.00000 sA A,*A,cjrtdtj}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|]}|VdSrzrrs r@ z,CUDADispatcher.specialize..s1< True if the Dispatcher has been specialized. )r{r{s r@rJzCUDADispatcher.specializeds   rBNc||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=rtoverloads r@ z6CUDADispatcher.get_regs_per_thread..s7AAA%X1AAArB)rUrrrJrSrTrVitemsrjras r@get_regs_per_threadz"CUDADispatcher.get_regs_per_threads  >).1A A   AT^22445566F FAA)-)=)=)?)?AAA ArBc||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 r@rz5CUDADispatcher.get_const_mem_size..s7AAA%X0AAArB)rUrrrJrSrTrVrrs r@get_const_mem_sizez!CUDADispatcher.get_const_mem_sizes  >).1@ @   AT^22445566E EAA)-)=)=)?)?AAA ArBc||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 r@rz;CUDADispatcher.get_shared_mem_per_block..7AAA%X6AAArB)rUrrrJrSrTrVrrs r@get_shared_mem_per_blockz'CUDADispatcher.get_shared_mem_per_block  >).1F F   AT^22445566K KAA)-)=)=)?)?AAA ArBc||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 r@rz.,s7AAA%X7AAArB)rUrrrJrSrTrVrrs r@get_max_threads_per_blockz(CUDADispatcher.get_max_threads_per_blocks  >).1G G   AT^22445566L LAA)-)=)=)?)?AAA ArBc||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 r@rz;CUDADispatcher.get_local_mem_per_thread..?rrB)rUrrrJrSrTrVrrs r@get_local_mem_per_threadz'CUDADispatcher.get_local_mem_per_thread/rrBc |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_devicerrHr:formatrmake_concrete_templatenopython_signaturesr pysignature)rjrr func_namer_ call_templatepysigs r@get_call_templatez CUDADispatcher.get_call_templateBs   -   d , , ,L) ")))445 iD,DFFF !$,//eT3..rBc 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`. r4r5r6r1r2r0r)r2r1r3N) rU_compiling_counterryrrrKrrHrMinsert_user_functionrGrSrR) rjr return_typer4r5r6r1r7r8rns r@rzCUDADispatcher.compile_device]s 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 r@rAz/CUDADispatcher.add_overload..s+++Q+++rBTr)_insertrU)rjrsrIc_sigs r@ add_overloadzCUDADispatcher.add_overloads?++(+++ UF ...#)x   rBctj|\}}||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_signaturernonerJrSrTrUrVrrrs targetctx _cache_hits _cache_missesrrCr.rHryr save_overloadr)rjrtrIrrss r@rzCUDADispatcher.compiles\ !) .sCEEE)CX-::<<EEErBc>i|]\}}||Sr)rrs r@rz/CUDADispatcher.inspect_llvm..s@EEE)CX2244EEErB)ryrrUrRrrr)rjrarms r@rzCUDADispatcher.inspect_llvms#''11   @~i08EEGGG~i0==??? EEE-1^-A-A-C-CEEEEEE-1^-A-A-C-CEEEErBctj|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. rmNcLi|] \}}||j!Sr)rRr<r=rtrr8s r@rz.CUDADispatcher.inspect_asm..sFEEE)CX-99"==EEErBcBi|]\}}||Sr)rrs r@rz.CUDADispatcher.inspect_asm..sCEEE)CX11"55EEErB) rrKryrrUrRr<rr)rjrarmr8s @r@rzCUDADispatcher.inspect_asms ! ! 4#''11   A~i08DDRHHH~i0<i|]\}}||Sr)rr=rtdefns r@rz3CUDADispatcher.inspect_sass_cfg..s:===!T..00===rB)ryrrCrUrrrs r@rzCUDADispatcher.inspect_sass_cfgs{   ! !( + + JHII I  >),==?? ?==%)^%9%9%;%;=== =rBc|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. rmz(Cannot inspect SASS of a device functionNc>i|]\}}||Sr)rrs r@rz/CUDADispatcher.inspect_sass..s:===!T**,,===rB)ryrrCrUrrrs r@rzCUDADispatcher.inspect_sasss{   ! !( + + KIJJ J  >),99;; ;==%)^%9%9%;%;=== =rBc| tj}|jD]\}}||dS)rNr)rrrUrr)rjrrYrs r@rzCUDADispatcher.inspect_typessU <:D~++-- * *GAt   D  ) ) ) ) * *rBc|||}|S)rr)rrHryrs r@rzCUDADispatcher._rebuilds 3w ..rBc8t|j|jS)zd Reduce the instance for serialization. Compiled definitions are discarded. )rHry)rrHryr{s r@rzCUDADispatcher._reduce_statess% DL"&"4666 6rBr9)rrrrz)*r:r;r<r= _fold_argsr targetdescrrrEr>rrr lru_cacherrrrJrMrdrrrKrJrrrrrrrrrrrrrrr?rrr@rAs@r@rwrw@sJK>J"""""" //X/...Y%%%PPP&%P%%% QQQQ,44X4(444 B B B---   0!!X! AAAA&AAAA(AAAA&AAAA*AAAA&///6%%%%N*** """HEEEE.EEEE0====*====, * * * *[6666666rBrw)9numpyrrWrrr 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.cudar~numbarrwarningsrrV ReduceMixinr.objectrCr[rgrorwrrBr@rs HHHHHHHHHHHHHHHH////////999999,,,,,,55555544444444------$$$$$$::::::::%%%%%%222222------<<<<<<<<*********i/i/i/i/i/i#i/i/i/X +++++V+++\AAAAAAAA:I$ > > > > > > > >a6a6a6a6a6Z!6a6a6a6a6a6rB