xKgddlZddlZddlZddlZddlZddlmZmZm 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/gdZ0GddejbZ2Gdde3Z4GddZ5GddeZ6GddeZ7GddeejbZ8y) 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)hsinhcoshloghlog10hlog2hexphexp10hexp2hsqrthrsqrthfloorhceilhrcphrinthtrunchdivc eZdZdZe dfd ZedZedZdZ edZ edZ e fdZ d Zd Zed Zed Zed ZedZedZdZdZdZdZddZddZddZdZxZS)_Kernelz CUDA Kernel specialized for a given set of argument types. When called, this object launches the kernel on the device. c | r tdt| d|_d|_||_||_||_||_|xsg|_ || rdndd} tj} t|j tj|j |j|||| |  }|j}|j j }|j"}|j$}|j'|j(|j*||| ||| \}}|sg}d|j-v|_|j.rd|_t2Dcgc]}d ||j-vr|}}|rqt4j6j9t4j6j;t<}t4j6j?|d }|jA||D]}|jC||jD|_#|jH|_$|jJ|_&||_'|jP|_(||_|j*|_|jR|_)g|_*g|_+g|_,ycc}w) Nz,Cannot compile a device function as a kernelFr)fastmathoptdebuglineinfoinliner1 nvvm_optionscccudaCGGetIntrinsicHandleT__numba_wrapper_zcpp_function_wrappers.cu)- RuntimeErrorsuper__init__ objectmode entry_pointpy_funcargtypesr4r5 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)selfr@rAlinkr4r5r6r1rB max_registersr2devicer7r8crestgt_ctxcodefilenamelinenumlibkernelfnresbasedirfunctions_cu_pathfilepath __class__s Y/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/dispatcher.pyr=z_Kernel.__init__.sU 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 ++(*% ;BsJc|jSN)r]rcs rtrJz_Kernel.librarys   c|jSrv)r\rws rtr[z_Kernel.type_annotations$$$rxc|jSrv)r`rws rt_find_referenced_environmentsz%_Kernel._find_referenced_environmentss,,,rxc6|jjSrv)rEcodegenrws rtr}z_Kernel.codegens""**,,rxc@t|jjSrv)tuplerZargsrws rtargument_typesz_Kernel.argument_typessT^^(())rxc |j|} t|| d| _|| _|| _|| _d| _|| _|| _ || _ || _ || _ | S)& Rebuild an instance. N) __new__r<r=r?rMrYrZr\r]r4r5r^rB) clsrMrXrZ codelibraryr4r5r^rBinstancerss rt_rebuildz_Kernel._rebuildst ;;s# c8%'#*"&$(! +$*(rxc 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. )rMrXrZrr4r5r^rB) dictrMrYrZr]r4r5r^rBrws rt_reduce_statesz_Kernel._reduce_statessL 0 0t"nn$:K:K**t}} $ 0 0T__N Nrxc8|jjy)z7 Force binding to current CUDA context N)r] get_cufuncrws rtbindz _Kernel.binds $$&rxc^|jjjjS)zN The number of registers used by each thread for this kernel. )r]rattrsregsrws rtregs_per_threadz_Kernel.regs_per_threads%   ++-33888rxc^|jjjjS)zD The amount of constant memory used by this kernel. )r]rrconstrws rtconst_mem_sizez_Kernel.const_mem_size%   ++-33999rxc^|jjjjS)zM The amount of shared memory used per block for this kernel. )r]rrsharedrws rtshared_mem_per_blockz_Kernel.shared_mem_per_blocks%   ++-33:::rxc^|jjjjS)z: The maximum allowable threads per block. )r]rr maxthreadsrws rtmax_threads_per_blockz_Kernel.max_threads_per_blocks%   ++-33>>>rxc^|jjjjS)zM The amount of local memory used per thread for this kernel. )r]rrlocalrws rtlocal_mem_per_threadz_Kernel.local_mem_per_threadrrxc6|jjS)z6 Returns the LLVM IR for this kernel. )r] get_llvm_strrws rt inspect_llvmz_Kernel.inspect_llvms  --//rxc:|jj|S)z7 Returns the PTX code for this kernel. )r8)r]rL)rcr8s rt inspect_asmz_Kernel.inspect_asms  ,,,33rxc6|jjS)zv Returns the CFG of the SASS for this kernel. Requires nvdisasm to be available on the PATH. )r] get_sass_cfgrws rtinspect_sass_cfgz_Kernel.inspect_sass_cfgs   --//rxc6|jjS)zp Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )r]get_sassrws rt inspect_sassz_Kernel.inspect_sasss   ))++rxc|j td|tj}t |j d|j |t d|t |j|t d|y) 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\ ValueErrorsysstdoutprintrYr)rcrs rt inspect_typesz_Kernel.inspect_typessg  (?@ @ <::D $*=*=>TJ hT" d##$/ hT"rxct}|jj}t|trt j d|}|j|||}|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 ||zSrv)xys rtz5_Kernel.max_cooperative_grid_blocks..&sQUrx) rr]r isinstancer functoolsreduce$get_active_blocks_per_multiprocessorrfMULTIPROCESSOR_COUNT)rcblockdim dynsmemsizectxcufunc active_per_smsm_counts rtmax_cooperative_grid_blocksz#_Kernel.max_cooperative_grid_blockssqm""--/ h & ''(:HEH@@AIALN ::22x''rxc|jj|jr|jdz}jj |\}}|t jt jk(sJt j} |jd|g} g} t|j|D]\} } |j| | || | tjr tjj!d}nd}|xr |j"xs|}tj$j"g||||| d|j&i|jrtj(t j* | j,dk7rfd}dDcgc] }|d|z}}dDcgc] }|d|z}}| j,}|j.j1|\}}}|d }n1|\}}}t2j4j7|}d |d |d |d }|d|d|}|r|d|df|ddz}n|f}||| D] }| ycc}wcc}w)N __errcode__r)streamrMcjjjd|d\}}tj}t j tj||||jS)N__) moduleget_global_symbolrXctypesc_intrdevice_to_host addressofvalue)rXmemszvalrs rt load_symbolz#_Kernel.launch..load_symbolSs`$mm==?E{{?C?EFGC!,,.C))&*:*:3*?bI99$rxzyxtidctaidz In function z, file z, line z, ztid=z ctaid=z: )r]rr4rXrrrsizeofrmemsetzipr _prepare_argsrUSE_NV_BINDINGbindingCUstreamhandle launch_kernelrMrrrr^ get_exceptionrPrQrS)rcrgriddimrr sharedmemexcnameexcmemexcszexcvalretr kernelargstv zero_stream stream_handlerirrriexcclsexc_argsloclocinfosymrrlinenoprefixwbrs @rtlaunchz_Kernel.launch-sq""--/ ::kkM1G"MM;;GDMFEFMM&,,77 77\\^F MM!FM + ++T2DAq   q!VT: >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 $I&<I+c t|jD]}|j||||\}}t|tj rt |j||}tj}tjd} tjd} ||j} ||jj} tj|} tj r t#| } tj| }|j%| |j%| |j%| |j%| |j%|t'|j(D]&}|j%||j*|(t'|j(D]&}|j%||j,|(yt|tj.r+t1td|z|}|j%|y|tj2k(rWtj4t7j2|j9t6j:}|j%|y|tj<k(r'tj>|}|j%|y|tj@k(r'tjB|}|j%|y|tjDk(r0tjFt#|}|j%|y|tjHk(r]|j%tjB|jJ|j%tjB|jLy|tjNk(r]|j%tj>|jJ|j%tj>|jLyt|tjPtjRfrB|j%tjT|j9t6jVyt|tjXrgt |j||}|jZ} tj rtjt#| } |j%| yt|tj\rCt_|t_|k(sJta||D]\}}|jc|||||yt|tjdr+ |jc|j|jf|||yti||#th$r ti||wxYw)zF Convert arguments to ctypes and append to kernelargs )rrrzc_%sN)5reversedrB prepare_argsrrArrayr to_devicer c_ssize_tc_void_psizedtypeitemsizerdevice_pointerrintrVrangendimshapestridesIntegergetattrfloat16c_uint16npviewuint16float64c_doublefloat32c_floatbooleanc_uint8 complex64realimag complex128 NPDatetime NPTimedeltac_int64int64Recorddevice_ctypes_pointer BaseTuplelenrr EnumMemberrNotImplementedError)rctyrrrr extensiondevaryc_intpmeminfoparentnitemsrptrdataaxcvaldevrecrrs rtrz_Kernel._prepare_argsus? "$//2I,, -GB3 b%++ &c],,T6:F%%Fooa(G__Q'FFKK(Ffll334H''/C$$#h??3'D   g &   f %   f %   h '   d #FKK(!!&b)9":;)FKK(!!&);"<=)EMM */766B;/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   fnnSXXbhh-?@ A ELL )c],,T6:F..C$$ooc#h/   c " EOO ,r7c#h& &&B 1""1azB%E,, - 3""HHciiz&b#. . ' 3)"c22 3s )W W) NFFFFNNTFrv)rrr)__name__ __module__ __qualname____doc__r r=propertyrJr[r{r}r classmethodrrrrrrrrrrrrrrrr __classcell__rss@rtr.r.(s+ ;@JN6;ZZx!!%%---*** N' 99 :: ;; ?? :: 0 4 0,#"(,FP\/rxr.ceZdZdZdZdZy)ForAllcp|dkrtd|z||_||_||_||_||_y)Nrz0Can't create ForAll with negative task count: %s)r dispatcherntasksthread_per_blockrr)rcr?r@tpbrrs rtr=zForAll.__init__sE A:O%&' '$ # "rxc&|jdk(ry|jjr |j}n|jj|}|j |}|j|zdz |z}||||j |j f|S)Nrr)r@r? specialized specialize_compute_thread_per_blockrr)rcrrDrrs rt__call__zForAll.__call__s ;;!   ?? & &//K4$//44d;K11+>;;)A-(:+{7Hdkk>>*+,02 2rxc$|j}|dk7r|St}tt|jj }t |jjd|jd}|jdi|\}}|S)Nri)funcb2d_funcmemsizeblocksizelimitr) rArnextiter overloadsvaluesrr]rrget_max_potential_block_size)rcr?rBrrmkwargs_s rtrFz ForAll._compute_thread_per_blocks## !8J-C$z33::<=>F((335# F 6S55??FAsJrxN)r4r5r6r=rGrFrrxrtr=r=s# 2rxr=ceZdZdZdZy)_LaunchConfigurationc||_||_||_||_||_t j r4d}|d|dz|dz}||krd|d}tt|yyy)Nrrz Grid size zB will likely result in GPU under-utilization due to low occupancy.) r?rrrrrCUDA_LOW_OCCUPANCY_WARNINGSrr ) rcr?rrrr min_grid_size grid_sizemsgs rtr=z_LaunchConfiguration.__init__s$    "  - - M WQZ/'!*>rxricDeZdZdZdZeZeffd Ze dZ dZ e jdd"dZd Zd#d Ze d Zd Zd ZdZdZdZe dZd$dZd$dZd$dZd$dZd$dZdZd$dZdZdZ d$dZ!d$dZ"d$dZ#d$dZ$d$dZ%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. FcFt||||d|_i|_y)N) targetoptionspipeline_classF)r<r= _specializedspecializations)rcr@rsrtrss rtr=zCUDADispatcher.__init__Rs0  (6  8" "rxc,tj|Srv) cuda_typesrqrws rt _numba_type_zCUDADispatcher._numba_type_bs((..rxc8t|j|_yrv)rir@_cacherws rtenable_cachingzCUDADispatcher.enable_cachingfs - rxrW)maxsizec>t||\}}t|||||Srv)rrU)rcrrrrs rt configurezCUDADispatcher.configureis&7J#D'8VYOOrxcPt|dvr td|j|S)N)rXr0z.must specify at least the griddim and blockdim)r$rrr_s rt __getitem__zCUDADispatcher.__getitem__ns+ t9I %MN Nt~~t$$rxc"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.)rBrr)r=)rcr@rBrrs rtforallzCUDADispatcher.forallss(dFFiPPrxc8|jjdS)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. rB)rsgetrws rtrBzCUDADispatcher.extensionss&!!%%l33rxc ttrv)rr)rcrrRs rtrGzCUDADispatcher.__call__s233rxc|jr-tt|jj }n t j j|g|}|j|||||y)zJ Compile if necessary and invoke this kernel with *args*. N) rDrMrNrOrPrr _cuda_callr)rcrrrrrrms rtr^zCUDADispatcher.callsX   $t~~44678F ++66tCdCF dGXvyArxc|rJ|Dcgc]}|j|}}|jt|Scc}wrv) typeof_pyvalcompiler)rcrkwsarAs rt_compile_for_argsz CUDADispatcher._compile_for_argss@w267$QD%%a($7||E(O,,8s=c t|tjS#t$rHt j |r1tt j |dtjcYSwxYw)NF)sync)rrargumentrr is_cuda_array as_cuda_array)rcrs rtrzCUDADispatcher.typeof_pyvalsc #w//0 0 !!#&d005A%..00 sA A-+A-ctj}t|Dcgc]}|jj |c}}|j r t d|jj||f}|r|S|j}t|j|}|j||jd|_||j||f<|Scc}w)zd Create a new instance of this dispatcher specialized for the given *args*. zDispatcher already specialized)rsT)rrCr typingctxresolve_argument_typerDr;rvrrsrqr@rdisable_compileru)rcrr8rrAspecializationrss rtrEzCUDADispatcher.specializes  ! 4 4>B CdT^^ 1 1! 4d CE   ?@ @--112x.A ! !** ' 6CEx(&&(&*#-;R\* Ds"Cc|jS)z> True if the Dispatcher has been specialized. )rurws rtrDzCUDADispatcher.specializeds    rxcL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rOrrrDrMrNrPitemsrcrZrnoverloads rtget_regs_per_threadz"CUDADispatcher.get_regs_per_threads  >>)..1AA A   T^^22456FF F*.)=)=)?A)? X111)?A AAB cL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rOrrrDrMrNrPrrs rtget_const_mem_sizez!CUDADispatcher.get_const_mem_sizes  >>)..1@@ @   T^^22456EE E*.)=)=)?A)? X000)?A AArcL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rOrrrDrMrNrPrrs rtget_shared_mem_per_blockz'CUDADispatcher.get_shared_mem_per_block  >>)..1FF F   T^^22456KK K*.)=)=)?A)? X666)?A AArcL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rOrrrDrMrNrPrrs rtget_max_threads_per_blockz(CUDADispatcher.get_max_threads_per_blocks  >>)..1GG G   T^^22456LL L*.)=)=)?A)? X777)?A AArcL|#|j|jjS|jr6t t |jj jS|jjDcic]\}}||jc}}Scc}}w)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. )rOrrrDrMrNrPrrs rtget_local_mem_per_threadz'CUDADispatcher.get_local_mem_per_thread/rrc*|jr|jt||jj}dj |}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_devicerr@r4formatrmake_concrete_templatenopython_signaturesr pysignature)rcrr func_namerX call_templatepysigs rtget_call_templatez CUDADispatcher.get_call_templateBs      d ,LL)) ")))455 iD,D,DF !!$,,/eT3..rxc ||jvr"|j5|jjd}|jjd}|jjd}|jjd}|jjdrdnd|d}t j }t |j|||||||| } | |j|<| jj| j| j| jgd d d | S|j|} | S#1swY SxYw) 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) rO_compiling_counterrsrrrCrr@rEinsert_user_functionr?rKrJ) rcr return_typer4r5r6r1r7r8rgs rtrzCUDADispatcher.compile_device]s6 t~~ %((**..w7--11*=++//9--11*=!% 2 2 6 6u =11 (  ()<<#DLL+t*/-5+1-51=') +(,t$##889I9I9=:>,,I-)8 >>$'D 9)8 s D EE c|Dcgc]}|j}}|j||d||j|<ycc}w)NTr)_code_insertrO)rcrmrArc_sigs rt add_overloadzCUDADispatcher.add_overloads?"*+(Q(+ UF .#)x ,s>ctj|\}}||tjk(sJ|jr,t t |jjS|jj|}||S|jj||j}||j|xxdz cc<n{|j|xxdz cc<|js t!dt#|j$|fi|j&}|j)|jj+|||j-|||S)z Compile and bind to the current context a version of this kernel specialized for the given signature. rzCompilation disabled)rnormalize_signaturernonerDrMrNrOrPrr{rm targetctx _cache_hits _cache_missesrr;r.r@rsr save_overloadr)rcrnrArrms rtrzCUDADispatcher.compiles) !) < >),==? ?&*^^%9%9%;=%; T..00%;= ==#Bc|jjdr td||j|j S|jj Dcic]\}}||j c}}Scc}}w)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. rfz(Cannot inspect SASS of a device function)rsrr;rOrrrs rtrzCUDADispatcher.inspect_sasss    ! !( +IJ J  >>),99; ;&*^^%9%9%;=%; T**,,%;= ==rc|tj}|jjD]\}}|j |y)rNr)rrrOrr)rcrrSrs rtrzCUDADispatcher.inspect_typess> <::D~~++-GAt   D  ).rxc|||}|S)rr)rr@rsrs rtrzCUDADispatcher._rebuilds w .rxcDt|j|jS)zd Reduce the instance for serialization. Compiled definitions are discarded. )r@rs)rr@rsrws rtrzCUDADispatcher._reduce_statess DLL"&"4"46 6rxr3)rrrrv)*r4r5r6r7 _fold_argsr targetdescrrr=r8ryr|r lru_cacherrrrBrGr^rrrErDrrrrrrrrrrrrrrr9rrr:r;s@rtrqrq@s JK>J" //.Y%P&P% Q,44(4 B- 0!! A&A(A&A*A&/6%N* "HE.E0=*=, *6rxrq)9numpyrrPrrr 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.cudarxnumbarrwarningsrrO ReduceMixinr.objectr=rUrarirqrrxrtrs HH/9,54-$:%2-<**i/i##i/X +V+\AA:I$ > >a6Z!6!6a6rx