§
    ƒÛgU›  ã                   óâ  — d dl Zd dlZd dlZd dlZd dlZd dlmZmZm	Z	m
Z
mZmZ d dlmZmZ d dlmZ d dlmZ d dlmZ d dlmZmZ d dlmZ d d	lmZ d d
lmZmZ d dl m!Z! d dl"m#Z# d dl$m%Z% d dl&m'Z'm(Z( d dl)m
Z* d dl+m,Z, d dl+m-Z- d dl.m/Z/ g d¢Z0 G d„ dej1        ¦  «        Z2 G d„ de3¦  «        Z4 G d„ d¦  «        Z5 G d„ de¦  «        Z6 G d„ de¦  «        Z7 G d„ deej1        ¦  «        Z8dS ) é    N)ÚconfigÚ	serializeÚsigutilsÚtypesÚtypingÚutils)ÚCacheÚ	CacheImpl)Úglobal_compiler_lock)Ú
Dispatcher)ÚNumbaPerformanceWarning)ÚPurposeÚtypeof)Úget_current_device)Úwrap_arg)Úcompile_cudaÚCUDACompiler)Údriver)Úget_context)Úcuda_target)Úmissing_launch_config_msgÚnormalize_kernel_dimensions)r   ©Úcuda)Ú_dispatcher)Úwarn)ÚhsinÚhcosÚhlogÚhlog10Úhlog2ÚhexpÚhexp10Úhexp2ÚhsqrtÚhrsqrtÚhfloorÚhceilÚhrcpÚhrintÚhtruncÚhdivc                   ób  ‡ — e Zd ZdZe	 	 	 dˆ fd„	¦   «         Zed„ ¦   «         Zed„ ¦   «         Zd„ Z	ed	„ ¦   «         Z
ed
„ ¦   «         Zeˆ fd„¦   «         Zd„ Zd„ Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zd„ Zd„ Zd„ Zd„ Zdd„Zdd„Zdd„Zd„ Zˆ xZS ) Ú_Kernelz„
    CUDA Kernel specialized for a given set of argument types. When called, this
    object launches the kernel on the device.
    NFTc                 óz  •‡— |rt          d¦  «        ‚t          ¦   «                              ¦   «          d| _        d | _        || _        || _        || _        || _        |pg | _	        ||
rdnddœ}t          ¦   «         j        }t          | j        t          j        | j        | j        |||||¬¦	  «	        }|j        }| j        j        }|j        }|j        }|                     |j        |j        ||||||	¦  «        \  Š}|sg }d‰                     ¦   «         v | _        | j        rd‰_        ˆfd	„t2          D ¦   «         }|rvt4          j                             t4          j                             t<          ¦  «        ¦  «        }t4          j                             |d
¦  «        }|                      |¦  «         |D ]}‰ !                    |¦  «         Œ|j"        | _#        |j$        | _$        |j%        | _&        ‰| _'        |j(        | _(        || _        |j        | _        |j)        | _)        g | _*        g | _+        g | _,        d S )Nz,Cannot compile a device function as a kernelFé   r   )ÚfastmathÚopt©ÚdebugÚlineinfoÚinliner1   Únvvm_optionsÚccÚcudaCGGetIntrinsicHandleTc                 óF   •— g | ]}d |› �‰                      ¦   «         v ¯|‘ŒS )Ú__numba_wrapper_)Úget_asm_str)Ú.0ÚfnÚlibs     €úK/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/dispatcher.pyú
<listcomp>z$_Kernel.__init__.<locals>.<listcomp>l   sB   ø€ ð Bð Bð B�bØ* bÐ*Ð*¨c¯oªoÑ.?Ô.?Ð?Ð?ð Ø?Ð?Ð?ó    zcpp_function_wrappers.cu)-ÚRuntimeErrorÚsuperÚ__init__Ú
objectmodeÚentry_pointÚpy_funcÚargtypesr4   r5   Ú
extensionsr   Úcompute_capabilityr   r   ÚvoidÚtarget_contextÚ__code__Úco_filenameÚco_firstlinenoÚprepare_cuda_kernelÚlibraryÚfndescr<   ÚcooperativeÚneeds_cudadevrtÚcuda_fp16_math_funcsÚosÚpathÚdirnameÚabspathÚ__file__ÚjoinÚappendÚadd_linking_fileÚnameÚ
entry_nameÚ	signatureÚtype_annotationÚ_type_annotationÚ_codelibraryÚcall_helperÚenvironmentÚ_referenced_environmentsÚliftedÚreload_init)ÚselfrH   rI   Úlinkr4   r5   r6   r1   rJ   Úmax_registersr2   Údevicer7   r8   ÚcresÚtgt_ctxÚcodeÚfilenameÚlinenumÚkernelÚresÚbasedirÚfunctions_cu_pathÚfilepathr?   Ú	__class__s                           @€r@   rE   z_Kernel.__init__.   s}  øø€ ð
 ð 	OÝÐMÑNÔNÐNå‰Œ×ÒÑÔÐð  ˆŒð  ˆÔàˆŒØ ˆŒØˆŒ
Ø ˆŒØ$Ð*¨ˆŒð !Ø�?�1�1 ð
ð 
ˆõ
  Ñ!Ô!Ô4ˆÝ˜DœL­%¬*°d´mØ"&¤*Ø%-Ø#)Ø%-Ø)5Ø!ð#ñ #ô #ˆð Ô%ˆØŒ|Ô$ˆØÔ#ˆØÔ%ˆØ×1Ò1°$´,ÀÄØ27¸À<Ø2:¸GØ2?ñAô A‰ˆˆVð
 ð 	ØˆDð 6¸¿ºÑ9JÔ9JÐJˆÔàÔð 	'Ø"&ˆCÔðBð Bð Bð BÕ0ð Bñ Bô Bˆð ð 	+å”g—o’o¥b¤g§o¢oµhÑ&?Ô&?Ñ@Ô@ˆGÝ "¤§¢¨WØ-Gñ!Iô !IÐà�KŠKÐ)Ñ*Ô*Ð*àð 	+ð 	+ˆHØ× Ò  Ñ*Ô*Ð*Ð*ð !œ+ˆŒØœˆŒØ $Ô 4ˆÔØˆÔØÔ+ˆÔð &ˆÔØ”kˆŒØÔ+ˆÔØ(*ˆÔ%ØˆŒØˆÔÐÐrB   c                 ó   — | j         S ©N)rd   ©rj   s    r@   rR   z_Kernel.library‹   s   € àÔ Ð rB   c                 ó   — | j         S rz   )rc   r{   s    r@   rb   z_Kernel.type_annotation�   s   € àÔ$Ð$rB   c                 ó   — | j         S rz   )rg   r{   s    r@   Ú_find_referenced_environmentsz%_Kernel._find_referenced_environments“   s   € ØÔ,Ð,rB   c                 ó4   — | j                              ¦   «         S rz   )rM   Úcodegenr{   s    r@   r€   z_Kernel.codegen–   s   € àÔ"×*Ò*Ñ,Ô,Ð,rB   c                 ó4   — t          | j        j        ¦  «        S rz   )Útuplera   Úargsr{   s    r@   Úargument_typesz_Kernel.argument_typesš   s   € å�T”^Ô(Ñ)Ô)Ð)rB   c	                 ó  •— |                       | ¦  «        }	t          | |	¦  «                             ¦   «          d|	_        ||	_        ||	_        ||	_        d|	_        ||	_        ||	_	        ||	_
        ||	_        ||	_        |	S )ú&
        Rebuild an instance.
        N)Ú__new__rD   rE   rG   rT   r`   ra   rc   rd   r4   r5   re   rJ   )ÚclsrT   r_   ra   Úcodelibraryr4   r5   re   rJ   Úinstancerx   s             €r@   Ú_rebuildz_Kernel._rebuildž   s†   ø€ ð —;’;˜sÑ#Ô#ˆåˆc�8ÑÔ×%Ò%Ñ'Ô'Ð'à#ˆÔØ*ˆÔØ"ˆÔØ&ˆÔØ$(ˆÔ!Ø +ˆÔØˆŒØ$ˆÔØ*ˆÔØ(ˆÔØˆrB   c           
      ó€   — t          | j        | j        | j        | j        | j        | j        | j        | j        ¬¦  «        S )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.
        )rT   r_   ra   r‰   r4   r5   re   rJ   )	ÚdictrT   r`   ra   rd   r4   r5   re   rJ   r{   s    r@   Ú_reduce_statesz_Kernel._reduce_states´   sE   € õ  Ô 0°t´Ø"œn¸$Ô:KØœ*¨t¬}Ø $Ô 0¸T¼_ðNñ Nô Nð 	NrB   c                 ó8   — | j                              ¦   «          dS )z7
        Force binding to current CUDA context
        N)rd   Ú
get_cufuncr{   s    r@   Úbindz_Kernel.bindÁ   s   € ð 	Ô×$Ò$Ñ&Ô&Ð&Ð&Ð&rB   c                 óH   — | j                              ¦   «         j        j        S )zN
        The number of registers used by each thread for this kernel.
        )rd   r�   ÚattrsÚregsr{   s    r@   Úregs_per_threadz_Kernel.regs_per_threadÇ   s   € ð
 Ô ×+Ò+Ñ-Ô-Ô3Ô8Ð8rB   c                 óH   — | j                              ¦   «         j        j        S )zD
        The amount of constant memory used by this kernel.
        )rd   r�   r“   Úconstr{   s    r@   Úconst_mem_sizez_Kernel.const_mem_sizeÎ   ó   € ð
 Ô ×+Ò+Ñ-Ô-Ô3Ô9Ð9rB   c                 óH   — | j                              ¦   «         j        j        S )zM
        The amount of shared memory used per block for this kernel.
        )rd   r�   r“   Úsharedr{   s    r@   Úshared_mem_per_blockz_Kernel.shared_mem_per_blockÕ   s   € ð
 Ô ×+Ò+Ñ-Ô-Ô3Ô:Ð:rB   c                 óH   — | j                              ¦   «         j        j        S )z:
        The maximum allowable threads per block.
        )rd   r�   r“   Ú
maxthreadsr{   s    r@   Úmax_threads_per_blockz_Kernel.max_threads_per_blockÜ   s   € ð
 Ô ×+Ò+Ñ-Ô-Ô3Ô>Ð>rB   c                 óH   — | j                              ¦   «         j        j        S )zM
        The amount of local memory used per thread for this kernel.
        )rd   r�   r“   Úlocalr{   s    r@   Úlocal_mem_per_threadz_Kernel.local_mem_per_threadã   r™   rB   c                 ó4   — | j                              ¦   «         S )z6
        Returns the LLVM IR for this kernel.
        )rd   Úget_llvm_strr{   s    r@   Úinspect_llvmz_Kernel.inspect_llvmê   s   € ð Ô ×-Ò-Ñ/Ô/Ð/rB   c                 ó8   — | j                              |¬¦  «        S )z7
        Returns the PTX code for this kernel.
        )r8   )rd   r<   )rj   r8   s     r@   Úinspect_asmz_Kernel.inspect_asmð   s   € ð Ô ×,Ò,°Ð,Ñ3Ô3Ð3rB   c                 ó4   — | j                              ¦   «         S )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_cfgö   s   € ð Ô ×-Ò-Ñ/Ô/Ð/rB   c                 ó4   — | j                              ¦   «         S )zp
        Returns the SASS code for this kernel.

        Requires nvdisasm to be available on the PATH.
        )rd   Úget_sassr{   s    r@   Úinspect_sassz_Kernel.inspect_sassþ   s   € ð Ô ×)Ò)Ñ+Ô+Ð+rB   c                 óþ   — | j         €t          d¦  «        ‚|€t          j        }t	          | j        ›d| j        ›�|¬¦  «         t	          d|¬¦  «         t	          | j         |¬¦  «         t	          d|¬¦  «         dS )úÚ
        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   Ú
ValueErrorÚsysÚstdoutÚprintr`   r„   )rj   r²   s     r@   Úinspect_typesz_Kernel.inspect_types  sŽ   € ð Ô Ð(ÝÐ?Ñ@Ô@Ð@àˆ<Ý”:ˆDå˜œ˜˜¨$Ô*=Ð*=Ð>ÀTÐJÑJÔJÐJÝˆh˜TÐ"Ñ"Ô"Ð"ÝˆdÔ#¨$Ð/Ñ/Ô/Ð/Ýˆh˜TÐ"Ñ"Ô"Ð"Ð"Ð"rB   r   c                 óö   — t          ¦   «         }| j                             ¦   «         }t          |t          ¦  «        rt          j        d„ |¦  «        }|                     |||¦  «        }|j        j	        }||z  S )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                 ó   — | |z  S rz   © )ÚxÚys     r@   ú<lambda>z5_Kernel.max_cooperative_grid_blocks.<locals>.<lambda>&  s
   € °Q¸±U€ rB   )
r   rd   r�   Ú
isinstancer‚   Ú	functoolsÚreduceÚ$get_active_blocks_per_multiprocessorrm   ÚMULTIPROCESSOR_COUNT)rj   ÚblockdimÚdynsmemsizeÚctxÚcufuncÚactive_per_smÚsm_counts          r@   Úmax_cooperative_grid_blocksz#_Kernel.max_cooperative_grid_blocks  s€   € õ ‰mŒmˆØÔ"×-Ò-Ñ/Ô/ˆå�h¥Ñ&Ô&ð 	FÝ Ô'Ð(:Ð(:¸HÑEÔEˆHØ×@Ò@ÀØAIØALñNô Nˆð ”:Ô2ˆØ˜xÑ'Ð'rB   c                 óJ  ‡‡— | j                              ¦   «         Š| j        ru‰j        dz   }‰j                             |¦  «        \  }}|t          j        t          j        ¦  «        k    sJ ‚t          j        ¦   «         }	| 	                    d|¬¦  «         g }
g }t          | j        |¦  «        D ]\  }}|                      ||||
|¦  «         Œt          j        r t          j                             d¦  «        }nd }|r|j        p|}t          j        ‰j        g|¢|¢|‘|‘|‘R d| j        iŽ | j        rÝt          j        t          j        |	¦  «        ||¦  «         |	j        dk    rªˆfd„Šˆfd„dD ¦   «         }ˆfd„dD ¦   «         }|	j        }| j                             |¦  «        \  }}}|€d	}n1|\  }}}t2          j                             |¦  «        }d
|›d|›d|›d�}|›d|›d|›�}|r|›d|d         ›�f|dd …         z   }n|f} ||Ž ‚|
D ]} |¦   «          Œd S )NÚ__errcode__r   )ÚstreamrT   c                 óØ   •— ‰j                              ‰j        ›d| ›d�¦  «        \  }}t          j        ¦   «         }t          j        t          j        |¦  «        ||¦  «         |j        S )NÚ__)	ÚmoduleÚget_global_symbolr_   ÚctypesÚc_intr   Údevice_to_hostÚ	addressofÚvalue)r_   ÚmemÚszÚvalrÆ   s       €r@   Úload_symbolz#_Kernel.launch.<locals>.load_symbolS  sk   ø€ Ø$œm×=Ò=Ø?E¼{¸{¸{Ø?C¸t¸tð?Eñ Fô F‘G�C˜õ !œ,™.œ.�CÝÔ)­&Ô*:¸3Ñ*?Ô*?ÀÀbÑIÔIÐIØœ9Ð$rB   c                 ó,   •— g | ]} ‰d |z   ¦  «        ‘ŒS )Útidrº   ©r=   ÚirÙ   s     €r@   rA   z"_Kernel.launch.<locals>.<listcomp>[  s'   ø€ Ð=Ð=Ð=°!�{�{ 5¨1¡9Ñ-Ô-Ð=Ð=Ð=rB   Úzyxc                 ó,   •— g | ]} ‰d |z   ¦  «        ‘ŒS )Úctaidrº   rÜ   s     €r@   rA   z"_Kernel.launch.<locals>.<listcomp>\  s'   ø€ ÐAÐAÐA°a˜˜ W¨q¡[Ñ1Ô1ÐAÐAÐArB   Ú zIn function z, file z, line z, ztid=z ctaid=z: é   )rd   r�   r4   r_   rÏ   rÐ   rÑ   ÚsizeofrÒ   ÚmemsetÚzipr„   Ú_prepare_argsr   ÚUSE_NV_BINDINGÚbindingÚCUstreamÚhandleÚlaunch_kernelrT   rÓ   rÔ   rÕ   re   Úget_exceptionrW   rX   rZ   )rj   rƒ   ÚgriddimrÃ   rÌ   Ú	sharedmemÚexcnameÚexcmemÚexcszÚexcvalÚretrÚ
kernelargsÚtÚvÚzero_streamÚstream_handlerÛ   rà   rp   ÚexcclsÚexc_argsÚlocÚlocinfoÚsymrw   ÚlinenoÚprefixÚwbrÆ   rÙ   s                               @@r@   Úlaunchz_Kernel.launch-  sã  øø€ àÔ"×-Ò-Ñ/Ô/ˆàŒ:ð 	,Ø”k MÑ1ˆGØ"œM×;Ò;¸GÑDÔD‰MˆF�EØ�FœM­&¬,Ñ7Ô7Ò7Ð7Ð7Ð7Ý”\‘^”^ˆFØ�MŠM˜! FˆMÑ+Ô+Ð+ð ˆàˆ
Ý˜Ô+¨TÑ2Ô2ð 	?ð 	?‰DˆAˆqØ×Ò˜q ! V¨T°:Ñ>Ô>Ð>Ð>åÔ ð 	Ý œ.×1Ò1°!Ñ4Ô4ˆKˆKàˆKàÐ0 6¤=Ð?°Kˆõ 	Ô˜Vœ]ð 	;Ø%ð	;à&ð	;ð 'ð	;ð +ð		;ð
 (ð	;ð 	;ð 	;ð *.Ô)9ð	;ð 	;ð 	;ð Œ:ð  	(ÝÔ!¥&Ô"2°6Ñ":Ô":¸FÀEÑJÔJÐJØŒ|˜qÒ Ð ð%ð %ð %ð %ð %ð >Ð=Ð=Ð=°uÐ=Ñ=Ô=�ØAÐAÐAÐA¸5ÐAÑAÔA�Ø”|�Ø(,Ô(8×(FÒ(FÀtÑ(LÔ(LÑ%�˜ #à�;Ø �G�Gà,/Ñ)�C˜ 6Ý!œwŸš¨xÑ8Ô8�H�HØFIÀcÀcØFNÀhÀhØFLÀfÀfðO�Gð 18°°¸¸¸¸e¸eÐD�Øð 'Ø,2¨F¨F°H¸Q´K°KÐ @ÐBØ    œñ %�H�Hð  &˜w�HØ�f˜hÐ'Ð'ð ð 	ð 	ˆBØˆB‰DŒDˆDˆDð	ð 	rB   c                 ó,  — t          | j        ¦  «        D ]}|                     ||||¬¦  «        \  }}Œt          |t          j        ¦  «        �rªt          |¦  «                             ||¦  «        }t          j	        }t          j
        d¦  «        }	t          j
        d¦  «        }
 ||j        ¦  «        } ||j        j        ¦  «        }t          j        |¦  «        }t          j        rt#          |¦  «        }t          j
        |¦  «        }|                     |	¦  «         |                     |
¦  «         |                     |¦  «         |                     |¦  «         |                     |¦  «         t'          |j        ¦  «        D ]+}|                      ||j        |         ¦  «        ¦  «         Œ,t'          |j        ¦  «        D ]+}|                      ||j        |         ¦  «        ¦  «         Œ,dS t          |t          j        ¦  «        r8 t1          t          d|z  ¦  «        |¦  «        }|                     |¦  «         dS |t          j        k    rZt          j        t7          j        |¦  «                             t6          j        ¦  «        ¦  «        }|                     |¦  «         dS |t          j        k    r+t          j        |¦  «        }|                     |¦  «         dS |t          j         k    r+t          j!        |¦  «        }|                     |¦  «         dS |t          j"        k    r8t          j#        t#          |¦  «        ¦  «        }|                     |¦  «         dS |t          j$        k    rZ|                     t          j!        |j%        ¦  «        ¦  «         |                     t          j!        |j&        ¦  «        ¦  «         dS |t          j'        k    rZ|                     t          j        |j%        ¦  «        ¦  «         |                     t          j        |j&        ¦  «        ¦  «         dS t          |t          j(        t          j)        f¦  «        rF|                     t          j*        |                     t6          j+        ¦  «        ¦  «        ¦  «         dS t          |t          j,        ¦  «        rnt          |¦  «                             ||¦  «        }|j-        }t          j        r!t          j
        t#          |¦  «        ¦  «        }|                     |¦  «         dS t          |t          j.        ¦  «        rSt_          |¦  «        t_          |¦  «        k    sJ ‚ta          ||¦  «        D ]\  }}|  1                    |||||¦  «         ŒdS t          |t          j2        ¦  «        rD	 |  1                    |j        |j3        |||¦  «         dS # th          $ r ti          ||¦  «        ‚w xY wti          ||¦  «        ‚)zF
        Convert arguments to ctypes and append to kernelargs
        )rÌ   ró   r   zc_%sN)5ÚreversedrJ   Úprepare_argsr¾   r   ÚArrayr   Ú	to_devicerÑ   Ú	c_ssize_tÚc_void_pÚsizeÚdtypeÚitemsizer   Údevice_pointerrç   Úintr]   ÚrangeÚndimÚshapeÚstridesÚIntegerÚgetattrÚfloat16Úc_uint16ÚnpÚviewÚuint16Úfloat64Úc_doubleÚfloat32Úc_floatÚbooleanÚc_uint8Ú	complex64ÚrealÚimagÚ
complex128Ú
NPDatetimeÚNPTimedeltaÚc_int64Úint64ÚRecordÚdevice_ctypes_pointerÚ	BaseTupleÚlenrå   ræ   Ú
EnumMemberrÕ   ÚNotImplementedError)rj   ÚtyrØ   rÌ   ró   rô   Ú	extensionÚdevaryÚc_intpÚmeminfoÚparentÚnitemsr  ÚptrÚdataÚaxÚcvalÚdevrecrõ   rö   s                       r@   ræ   z_Kernel._prepare_argsu  sš  € õ " $¤/Ñ2Ô2ð 	ð 	ˆIØ×,Ò,ØØØØð	 -ñ ô ‰GˆB��õ �b�%œ+Ñ&Ô&ñ O	/Ý˜c‘]”]×,Ò,¨T°6Ñ:Ô:ˆFåÔ%ˆFå”o aÑ(Ô(ˆGÝ”_ QÑ'Ô'ˆFØ�V˜FœKÑ(Ô(ˆFØ�v˜fœlÔ3Ñ4Ô4ˆHåÔ'¨Ñ/Ô/ˆCåÔ$ð Ý˜#‘h”h�å”? 3Ñ'Ô'ˆDà×Ò˜gÑ&Ô&Ð&Ø×Ò˜fÑ%Ô%Ð%Ø×Ò˜fÑ%Ô%Ð%Ø×Ò˜hÑ'Ô'Ð'Ø×Ò˜dÑ#Ô#Ð#Ý˜FœKÑ(Ô(ð <ð <�Ø×!Ò! & &¨¬°bÔ)9Ñ":Ô":Ñ;Ô;Ð;Ð;Ý˜FœKÑ(Ô(ð >ð >�Ø×!Ò! & &¨¬¸Ô);Ñ"<Ô"<Ñ=Ô=Ð=Ð=ð>ð >õ ˜�EœMÑ*Ô*ð 4	/Ø/•7�6 6¨B¡;Ñ/Ô/°Ñ4Ô4ˆDØ×Ò˜dÑ#Ô#Ð#Ð#Ð#à•5”=Ò Ð Ý”?¥2¤:¨c¡?¤?×#7Ò#7½¼	Ñ#BÔ#BÑCÔCˆDØ×Ò˜dÑ#Ô#Ð#Ð#Ð#à•5”=Ò Ð Ý”? 3Ñ'Ô'ˆDØ×Ò˜dÑ#Ô#Ð#Ð#Ð#à•5”=Ò Ð Ý”> #Ñ&Ô&ˆDØ×Ò˜dÑ#Ô#Ð#Ð#Ð#à•5”=Ò Ð Ý”>¥# c¡(¤(Ñ+Ô+ˆDØ×Ò˜dÑ#Ô#Ð#Ð#Ð#à•5”?Ò"Ð"Ø×Ò�fœn¨S¬XÑ6Ô6Ñ7Ô7Ð7Ø×Ò�fœn¨S¬XÑ6Ô6Ñ7Ô7Ð7Ð7Ð7à•5Ô#Ò#Ð#Ø×Ò�fœo¨c¬hÑ7Ô7Ñ8Ô8Ð8Ø×Ò�fœo¨c¬hÑ7Ô7Ñ8Ô8Ð8Ð8Ð8å˜�UÔ-­uÔ/@ÐAÑBÔBð 	/Ø×Ò�fœn¨S¯XªXµb´hÑ-?Ô-?Ñ@Ô@ÑAÔAÐAÐAÐAå˜�EœLÑ)Ô)ð 	/Ý˜c‘]”]×,Ò,¨T°6Ñ:Ô:ˆFØÔ.ˆCÝÔ$ð 0Ý”o¥c¨#¡h¤hÑ/Ô/�Ø×Ò˜cÑ"Ô"Ð"Ð"Ð"å˜�EœOÑ,Ô,ð 	/Ý�r‘7”7�c #™hœhÒ&Ð&Ð&Ð&Ý˜B ™œð Cð C‘��1Ø×"Ò" 1 a¨°°zÑBÔBÐBÐBðCð Cõ ˜�EÔ,Ñ-Ô-ð 		/ð3Ø×"Ò"Ø”H˜cœi¨°°zñô ð ð ð øõ 'ð 3ð 3ð 3Ý)¨"¨cÑ2Ô2Ð2ð3øøøõ & b¨#Ñ.Ô.Ð.s   ×#W( ×(X)	NFFFFNNTFrz   )r   ©r   r   )Ú__name__Ú
__module__Ú__qualname__Ú__doc__r   rE   ÚpropertyrR   rb   r~   r€   r„   Úclassmethodr‹   rŽ   r‘   r•   r˜   rœ   rŸ   r¢   r¥   r§   rª   r­   r·   rÉ   r  ræ   Ú__classcell__©rx   s   @r@   r.   r.   (   s?  ø€ € € € € ðð ð
 Ø;@ØJNØ6;ðZð Zð Zð Zð Zñ ÔðZðx ð!ð !ñ „Xð!ð ð%ð %ñ „Xð%ð-ð -ð -ð ð-ð -ñ „Xð-ð ð*ð *ñ „Xð*ð ðð ð ð ñ „[ðð*Nð Nð Nð'ð 'ð 'ð ð9ð 9ñ „Xð9ð ð:ð :ñ „Xð:ð ð;ð ;ñ „Xð;ð ð?ð ?ñ „Xð?ð ð:ð :ñ „Xð:ð0ð 0ð 0ð4ð 4ð 4ð0ð 0ð 0ð,ð ,ð ,ð#ð #ð #ð #ð"(ð (ð (ð (ð,Fð Fð Fð FðP\/ð \/ð \/ð \/ð \/ð \/ð \/rB   r.   c                   ó    — e Zd Zd„ Zd„ Zd„ ZdS )ÚForAllc                 ó|   — |dk     rt          d|z  ¦  «        ‚|| _        || _        || _        || _        || _        d S )Nr   z0Can't create ForAll with negative task count: %s)r³   Ú
dispatcherÚntasksÚthread_per_blockrÌ   rî   )rj   rE  rF  ÚtpbrÌ   rî   s         r@   rE   zForAll.__init__Õ  sP   € Ø�AŠ:ˆ:ÝÐOØ%ñ&ñ 'ô 'ð 'à$ˆŒØˆŒØ #ˆÔØˆŒØ"ˆŒˆˆrB   c                 óÞ   — | j         dk    rd S | j        j        r| j        }n | j        j        |Ž }|                      |¦  «        }| j         |z   dz
  |z  } |||| j        | j        f         |Ž S )Nr   râ   )rF  rE  ÚspecializedÚ
specializeÚ_compute_thread_per_blockrÌ   rî   )rj   rƒ   rJ  rÃ   rí   s        r@   Ú__call__zForAll.__call__ß  s�   € ØŒ;˜!ÒÐØˆFàŒ?Ô&ð 	<Øœ/ˆKˆKà4˜$œ/Ô4°dÐ;ˆKØ×1Ò1°+Ñ>Ô>ˆØ”; Ñ)¨AÑ-°(Ñ:ˆð+ˆ{˜7 H¨d¬kØœ>ð*ô +Ø,0ð2ð 	2rB   c                 ó$  — | j         }|dk    r|S t          ¦   «         }t          t          |j                             ¦   «         ¦  «        ¦  «        }t          |j                             ¦   «         d| j	        d¬¦  «        } |j
        di |¤Ž\  }}|S )Nr   i   )ÚfuncÚb2d_funcÚmemsizeÚblocksizelimitrº   )rG  r   ÚnextÚiterÚ	overloadsÚvaluesr�   rd   r�   rî   Úget_max_potential_block_size)rj   rE  rH  rÅ   rs   ÚkwargsÚ_s          r@   rL  z ForAll._compute_thread_per_blockí  s�   € ØÔ#ˆà�!Š8ˆ8ØˆJõ ‘-”-ˆCõ �$˜zÔ3×:Ò:Ñ<Ô<Ñ=Ô=Ñ>Ô>ˆFÝØÔ(×3Ò3Ñ5Ô5ØØœØ#ð	ñ ô ˆFð 6�SÔ5Ð?Ð?¸Ð?Ð?‰FˆAˆsØˆJrB   N)r:  r;  r<  rE   rM  rL  rº   rB   r@   rC  rC  Ô  sA   € € € € € ð#ð #ð #ð2ð 2ð 2ðð ð ð ð rB   rC  c                   ó   — e Zd Zd„ Zd„ ZdS )Ú_LaunchConfigurationc                 óô   — || _         || _        || _        || _        || _        t
          j        rFd}|d         |d         z  |d         z  }||k     r&d|› d�}t          t          |¦  «        ¦  «         d S d S d S )Né€   r   râ   é   z
Grid size zB will likely result in GPU under-utilization due to low occupancy.)	rE  rí   rÃ   rÌ   rî   r   ÚCUDA_LOW_OCCUPANCY_WARNINGSr   r   )	rj   rE  rí   rÃ   rÌ   rî   Úmin_grid_sizeÚ	grid_sizeÚmsgs	            r@   rE   z_LaunchConfiguration.__init__  s¨   € Ø$ˆŒØˆŒØ ˆŒØˆŒØ"ˆŒåÔ-ð 	3ð  ˆMØ œ
 W¨Q¤ZÑ/°'¸!´*Ñ<ˆIØ˜=Ò(Ð(ðA Ið Að Að A�åÕ,¨SÑ1Ô1Ñ2Ô2Ð2Ð2Ð2ð	3ð 	3ð )Ð(rB   c                 óf   — | j                              || j        | j        | j        | j        ¦  «        S rz   )rE  Úcallrí   rÃ   rÌ   rî   ©rj   rƒ   s     r@   rM  z_LaunchConfiguration.__call__  s2   € ØŒ×#Ò# D¨$¬,¸¼Ø$(¤K°´ñAô Að 	ArB   N)r:  r;  r<  rE   rM  rº   rB   r@   r[  r[    s7   € € € € € ð3ð 3ð 3ð.Að Að Að Að ArB   r[  c                   ó    — e Zd Zd„ Zd„ Zd„ ZdS )ÚCUDACacheImplc                 ó*   — |                      ¦   «         S rz   )rŽ   )rj   rs   s     r@   rÀ   zCUDACacheImpl.reduce   s   € Ø×$Ò$Ñ&Ô&Ð&rB   c                 ó$   — t          j        di |¤ŽS )Nrº   )r.   r‹   )rj   rM   Úpayloads      r@   ÚrebuildzCUDACacheImpl.rebuild#  s   € ÝÔÐ*Ð* 'Ð*Ð*Ð*rB   c                 ó   — dS )NTrº   )rj   rn   s     r@   Úcheck_cachablezCUDACacheImpl.check_cachable&  s	   € ð ˆtrB   N)r:  r;  r<  rÀ   rk  rm  rº   rB   r@   rg  rg    sA   € € € € € ð'ð 'ð 'ð+ð +ð +ðð ð ð ð rB   rg  c                   ó   — e Zd ZdZeZdS )Ú	CUDACachezS
    Implements a cache that saves and loads CUDA kernels and compile results.
    N)r:  r;  r<  r=  rg  Ú_impl_classrº   rB   r@   ro  ro  1  s   € € € € € ðð ð  €K€K€KrB   ro  c                   ó^  ‡ — e Zd ZdZdZeZefˆ fd„	Ze	d„ ¦   «         Z
d„ Z ej        d¬¦  «        d$d	„¦   «         Zd
„ Zd%d„Ze	d„ ¦   «         Zd„ Zd„ Zd„ Zd„ Zd„ Ze	d„ ¦   «         Zd&d„Zd&d„Zd&d„Zd&d„Zd&d„Zd„ Zd&d„Zd„ Zd„ Z d&d„Z!d&d„Z"d&d„Z#d&d „Z$d&d!„Z%e&d"„ ¦   «         Z'd#„ Z(ˆ xZ)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.
    Fc                 ól   •— t          ¦   «                              |||¬¦  «         d| _        i | _        d S )N)ÚtargetoptionsÚpipeline_classF)rD   rE   Ú_specializedÚspecializations)rj   rH   rt  ru  rx   s       €r@   rE   zCUDADispatcher.__init__J  sE   ø€ Ý‰Œ×Ò˜°Ø(6ð 	ñ 	8ô 	8ð 	8ð "ˆÔð  "ˆÔÐÐrB   c                 ó*   — t          j        | ¦  «        S rz   )Ú
cuda_typesrr  r{   s    r@   Ú_numba_type_zCUDADispatcher._numba_type_Z  s   € åÔ(¨Ñ.Ô.Ð.rB   c                 ó8   — t          | j        ¦  «        | _        d S rz   )ro  rH   Ú_cacher{   s    r@   Úenable_cachingzCUDADispatcher.enable_caching^  s   € Ý ¤Ñ-Ô-ˆŒˆˆrB   r]  )Úmaxsizer   c                 óN   — t          ||¦  «        \  }}t          | ||||¦  «        S rz   )r   r[  )rj   rí   rÃ   rÌ   rî   s        r@   Ú	configurezCUDADispatcher.configurea  s,   € å7¸ÀÑJÔJÑˆ�Ý# D¨'°8¸VÀYÑOÔOÐOrB   c                 óV   — t          |¦  «        dvrt          d¦  «        ‚ | j        |Ž S )N)r^  r0   é   z.must specify at least the griddim and blockdim)r*  r³   r€  re  s     r@   Ú__getitem__zCUDADispatcher.__getitem__f  s1   € Ýˆt‰9Œ9˜IÐ%Ð%ÝÐMÑNÔNÐNØˆtŒ~˜tÐ$Ð$rB   c                 ó*   — t          | ||||¬¦  «        S )a3  Returns 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.)rH  rÌ   rî   )rC  )rj   rF  rH  rÌ   rî   s        r@   ÚforallzCUDADispatcher.forallk  s   € õ( �d˜F¨°FÀiÐPÑPÔPÐPrB   c                 ó6   — | j                              d¦  «        S )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   )rt  Úgetr{   s    r@   rJ   zCUDADispatcher.extensions�  s   € ð& Ô!×%Ò% lÑ3Ô3Ð3rB   c                 ó*   — t          t          ¦  «        ‚rz   )r³   r   )rj   rƒ   rX  s      r@   rM  zCUDADispatcher.__call__–  s   € åÕ2Ñ3Ô3Ð3rB   c                 óÜ   — | j         r4t          t          | j                             ¦   «         ¦  «        ¦  «        }nt          j        j        | g|¢R Ž }|                     |||||¦  «         dS )zJ
        Compile if necessary and invoke this kernel with *args*.
        N)	rJ  rS  rT  rU  rV  r   r   Ú
_cuda_callr  )rj   rƒ   rí   rÃ   rÌ   rî   rs   s          r@   rd  zCUDADispatcher.callš  sp   € ð Ôð 	DÝ�$˜tœ~×4Ò4Ñ6Ô6Ñ7Ô7Ñ8Ô8ˆFˆFå Ô+Ô6°tÐC¸dÐCÐCÐCˆFà�Š�d˜G X¨v°yÑAÔAÐAÐAÐArB   c                 ól   ‡ — |rJ ‚ˆ fd„|D ¦   «         }‰                       t          |¦  «        ¦  «        S )Nc                 ó:   •— g | ]}‰                      |¦  «        ‘ŒS rº   )Útypeof_pyval©r=   Úarj   s     €r@   rA   z4CUDADispatcher._compile_for_args.<locals>.<listcomp>¨  s'   ø€ Ð7Ð7Ð7¨Q�D×%Ò% aÑ(Ô(Ð7Ð7Ð7rB   )Úcompiler‚   )rj   rƒ   ÚkwsrI   s   `   r@   Ú_compile_for_argsz CUDADispatcher._compile_for_args¥  s=   ø€ àˆˆˆØ7Ð7Ð7Ð7°$Ð7Ñ7Ô7ˆØ�|Š|�E (™OœOÑ,Ô,Ð,rB   c                 óÞ   — 	 t          |t          j        ¦  «        S # t          $ rF t	          j        |¦  «        r0t          t	          j        |d¬¦  «        t          j        ¦  «        cY S ‚ w xY w)NF)Úsync)r   r   Úargumentr³   r   Úis_cuda_arrayÚas_cuda_array)rj   rØ   s     r@   r�  zCUDADispatcher.typeof_pyval«  sƒ   € ð		Ý˜#�wÔ/Ñ0Ô0Ð0øÝð 	ð 	ð 	ÝÔ! #Ñ&Ô&ð õ �dÔ0°¸5ÐAÑAÔAÝ%Ô.ñ0ô 0ð 0ð 0ð 0ð ð	øøøs   ‚ œAA,Á*A,c                 ó‚  ‡ — t          ¦   «         j        }t          ˆ fd„|D ¦   «         ¦  «        }‰ j        rt	          d¦  «        ‚‰ j                             ||f¦  «        }|r|S ‰ j        }t          ‰ j	        |¬¦  «        }| 
                    |¦  «         |                     ¦   «          d|_        |‰ j        ||f<   |S )zd
        Create a new instance of this dispatcher specialized for the given
        *args*.
        c                 óD   •— g | ]}‰j                              |¦  «        ‘ŒS rº   )Ú	typingctxÚresolve_argument_typerŽ  s     €r@   rA   z-CUDADispatcher.specialize.<locals>.<listcomp>À  s)   ø€ ÐCÐCÐC¸ˆTŒ^×1Ò1°!Ñ4Ô4ÐCÐCÐCrB   zDispatcher already specialized)rt  T)r   rK   r‚   rJ  rC   rw  r‡  rt  rr  rH   r�  Údisable_compilerv  )rj   rƒ   r8   rI   Úspecializationrt  s   `     r@   rK  zCUDADispatcher.specialize¹  sé   ø€ õ
  Ñ!Ô!Ô4ˆÝØCÐCÐCÐC¸dÐCÑCÔCñEô EˆàÔð 	AÝÐ?Ñ@Ô@Ð@àÔ-×1Ò1°2°x°.ÑAÔAˆØð 	"Ø!Ð!àÔ*ˆÝ'¨¬Ø6CðEñ Eô Eˆà×Ò˜xÑ(Ô(Ð(Ø×&Ò&Ñ(Ô(Ð(Ø&*ˆÔ#Ø-;ˆÔ˜R ˜\Ñ*ØÐrB   c                 ó   — | j         S )z>
        True if the Dispatcher has been specialized.
        )rv  r{   s    r@   rJ  zCUDADispatcher.specializedÑ  s   € ð
 Ô Ð rB   Nc                 óø   — |�| j         |j                 j        S | j        r8t	          t          | j                              ¦   «         ¦  «        ¦  «        j        S d„ | j                              ¦   «         D ¦   «         S )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 | ]\  }}||j         “ŒS rº   )r•   ©r=   ÚsigÚoverloads      r@   ú
<dictcomp>z6CUDADispatcher.get_regs_per_thread.<locals>.<dictcomp>è  s7   € ð Að Að AÙ%˜˜Xð ˜Ô1ð Að Að ArB   )rU  rƒ   r•   rJ  rS  rT  rV  Úitems©rj   ra   s     r@   Úget_regs_per_threadz"CUDADispatcher.get_regs_per_threadØ  s„   € ð Ð Ø”> )¤.Ô1ÔAÐAØÔð 	AÝ�˜Tœ^×2Ò2Ñ4Ô4Ñ5Ô5Ñ6Ô6ÔFÐFðAð AØ)-¬×)=Ò)=Ñ)?Ô)?ðAñ Aô Að ArB   c                 óø   — |�| j         |j                 j        S | j        r8t	          t          | j                              ¦   «         ¦  «        ¦  «        j        S d„ | j                              ¦   «         D ¦   «         S )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 | ]\  }}||j         “ŒS rº   )r˜   r¡  s      r@   r¤  z5CUDADispatcher.get_const_mem_size.<locals>.<dictcomp>ü  s7   € ð Að Að AÙ%˜˜Xð ˜Ô0ð Að Að ArB   )rU  rƒ   r˜   rJ  rS  rT  rV  r¥  r¦  s     r@   Úget_const_mem_sizez!CUDADispatcher.get_const_mem_sizeë  s„   € ð Ð Ø”> )¤.Ô1Ô@Ð@ØÔð 	AÝ�˜Tœ^×2Ò2Ñ4Ô4Ñ5Ô5Ñ6Ô6ÔEÐEðAð AØ)-¬×)=Ò)=Ñ)?Ô)?ðAñ Aô Að ArB   c                 óø   — |�| j         |j                 j        S | j        r8t	          t          | j                              ¦   «         ¦  «        ¦  «        j        S d„ | j                              ¦   «         D ¦   «         S )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 | ]\  }}||j         “ŒS rº   )rœ   r¡  s      r@   r¤  z;CUDADispatcher.get_shared_mem_per_block.<locals>.<dictcomp>  ó7   € ð Að Að AÙ%˜˜Xð ˜Ô6ð Að Að ArB   )rU  rƒ   rœ   rJ  rS  rT  rV  r¥  r¦  s     r@   Úget_shared_mem_per_blockz'CUDADispatcher.get_shared_mem_per_blockÿ  ó„   € ð Ð Ø”> )¤.Ô1ÔFÐFØÔð 	AÝ�˜Tœ^×2Ò2Ñ4Ô4Ñ5Ô5Ñ6Ô6ÔKÐKðAð AØ)-¬×)=Ò)=Ñ)?Ô)?ðAñ Aô Að ArB   c                 óø   — |�| j         |j                 j        S | j        r8t	          t          | j                              ¦   «         ¦  «        ¦  «        j        S d„ | j                              ¦   «         D ¦   «         S )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 | ]\  }}||j         “ŒS rº   )rŸ   r¡  s      r@   r¤  z<CUDADispatcher.get_max_threads_per_block.<locals>.<dictcomp>$  s7   € ð Að Að AÙ%˜˜Xð ˜Ô7ð Að Að ArB   )rU  rƒ   rŸ   rJ  rS  rT  rV  r¥  r¦  s     r@   Úget_max_threads_per_blockz(CUDADispatcher.get_max_threads_per_block  s„   € ð Ð Ø”> )¤.Ô1ÔGÐGØÔð 	AÝ�˜Tœ^×2Ò2Ñ4Ô4Ñ5Ô5Ñ6Ô6ÔLÐLðAð AØ)-¬×)=Ò)=Ñ)?Ô)?ðAñ Aô Að ArB   c                 óø   — |�| j         |j                 j        S | j        r8t	          t          | j                              ¦   «         ¦  «        ¦  «        j        S d„ | j                              ¦   «         D ¦   «         S )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 | ]\  }}||j         “ŒS rº   )r¢   r¡  s      r@   r¤  z;CUDADispatcher.get_local_mem_per_thread.<locals>.<dictcomp>7  r­  rB   )rU  rƒ   r¢   rJ  rS  rT  rV  r¥  r¦  s     r@   Úget_local_mem_per_threadz'CUDADispatcher.get_local_mem_per_thread'  r¯  rB   c                 ó  — | j         r"|                      t          |¦  «        ¦  «         | j        j        }d                     |¦  «        }t          j        ||| j        ¬¦  «        }t          j
        | j        ¦  «        }||||fS )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.
        zCallTemplate({0}))ÚkeyÚ
signatures)Ú_can_compileÚcompile_devicer‚   rH   r:  Úformatr   Úmake_concrete_templateÚnopython_signaturesr   Úpysignature)rj   rƒ   r‘  Ú	func_namer_   Úcall_templateÚpysigs          r@   Úget_call_templatez CUDADispatcher.get_call_template:  s�   € ð Ôð 	-Ø×Ò¥ d¡¤Ñ,Ô,Ð,ð ”LÔ)ˆ	Ø"×)Ò)¨)Ñ4Ô4ˆåÔ5Ø�i¨DÔ,DðFñ Fô FˆåÔ! $¤,Ñ/Ô/ˆà˜e T¨3Ð.Ð.rB   c                 óR  — || j         v�r| j        5  | j                             d¦  «        }| j                             d¦  «        }| j                             d¦  «        }| j                             d¦  «        }| j                             d¦  «        rdnd|dœ}t	          ¦   «         j        }t          | j        ||||||||¬	¦	  «	        }	|	| j         |<   |	j         	                    |	j
        |	j        |	j        g¦  «         d
d
d
¦  «         n# 1 swxY w Y   n| j         |         }	|	S )zÎCompile the device function for the given argument types.

        Each signature is compiled once by caching the compiled function inside
        this object.

        Returns the `CompileResult`.
        r4   r5   r6   r1   r2   r0   r   )r2   r1   r3   N)rU  Ú_compiling_counterrt  r‡  r   rK   r   rH   rM   Úinsert_user_functionrG   rS   rR   )
rj   rƒ   Úreturn_typer4   r5   r6   r1   r7   r8   rn   s
             r@   rº  zCUDADispatcher.compile_deviceU  s•  € ð �t”~Ð%Ñ%ØÔ(ð Ið IàÔ*×.Ò.¨wÑ7Ô7�ØÔ-×1Ò1°*Ñ=Ô=�ØÔ+×/Ò/°Ñ9Ô9�ØÔ-×1Ò1°*Ñ=Ô=�ð !%Ô 2× 6Ò 6°uÑ =Ô =ÐD˜1˜1À1Ø (ð ð  �õ
 (Ñ)Ô)Ô<�Ý# D¤L°+¸tØ*/Ø-5Ø+1Ø-5Ø1=Ø')ð+ñ +ô +�ð (,�”˜tÑ$àÔ#×8Ò8¸Ô9IØ9=¼Ø:>¼,¸ñIô Ið Ið-Ið Ið Iñ Iô Ið Ið Ið Ið Ið Ið Iøøøð Ið Ið Ið Iøð4 ”> $Ô'ˆDàˆs   ’C0DÄDÄDc                 ób   — d„ |D ¦   «         }|                       ||d¬¦  «         || j        |<   d S )Nc                 ó   — g | ]	}|j         ‘Œ
S rº   )Ú_code)r=   r�  s     r@   rA   z/CUDADispatcher.add_overload.<locals>.<listcomp>}  s   € Ð+Ð+Ð+˜Q�”Ð+Ð+Ð+rB   Tr   )Ú_insertrU  )rj   rs   rI   Úc_sigs       r@   Úadd_overloadzCUDADispatcher.add_overload|  s?   € Ø+Ð+ (Ð+Ñ+Ô+ˆØ�Š�U˜F¨ˆÑ.Ô.Ð.Ø#)ˆŒ�xÑ Ð Ð rB   c                 óŒ  — t          j        |¦  «        \  }}|�|t          j        k    sJ ‚| j        r3t          t          | j                             ¦   «         ¦  «        ¦  «        S | j         	                    |¦  «        }|�|S | j
                             || j        ¦  «        }|�| j        |xx         dz  cc<   nr| j        |xx         dz  cc<   | j        st!          d¦  «        ‚t#          | j        |fi | j        ¤Ž}|                     ¦   «          | j
                             ||¦  «         |                      ||¦  «         |S )z
        Compile and bind to the current context a version of this kernel
        specialized for the given signature.
        Nrâ   zCompilation disabled)r   Únormalize_signaturer   ÚnonerJ  rS  rT  rU  rV  r‡  r|  Úload_overloadÚ	targetctxÚ_cache_hitsÚ_cache_missesr¹  rC   r.   rH   rt  r‘   Úsave_overloadrÌ  )rj   r¢  rI   rÆ  rs   s        r@   r�  zCUDADispatcher.compile�  s\  € õ
 !)Ô <¸SÑ AÔ AÑˆ�+ØÐ" kµU´ZÒ&?Ð&?Ð&?Ð&?ð Ôð 	Ý�˜Tœ^×2Ò2Ñ4Ô4Ñ5Ô5Ñ6Ô6Ð6à”^×'Ò'¨Ñ1Ô1ˆFØÐ!Ø�ð ”×*Ò*¨3°´Ñ?Ô?ˆàÐØÔ˜SÐ!Ð!Ô! QÑ&Ð!Ð!Ñ!Ð!ð Ô˜sÐ#Ð#Ô# qÑ(Ð#Ð#Ñ#ØÔ$ð ;Ý"Ð#9Ñ:Ô:Ð:å˜Tœ\¨8ÐJÐJ°tÔ7IÐJÐJˆFà�KŠK‰MŒMˆMØŒK×%Ò% c¨6Ñ2Ô2Ð2à×Ò˜& (Ñ+Ô+Ð+àˆrB   c                 óT  — | j                              d¦  «        }|�E|r$| j        |         j                             ¦   «         S | j        |                              ¦   «         S |r#d„ | j                             ¦   «         D ¦   «         S d„ | j                             ¦   «         D ¦   «         S )zó
        Return the LLVM IR for this kernel.

        :param signature: A tuple of argument types.
        :return: The LLVM IR for the given signature, or a dict of LLVM IR
                 for all previously-encountered signatures.

        rm   Nc                 óH   — i | ]\  }}||j                              ¦   «         “Œ S rº   )rR   r¤   r¡  s      r@   r¤  z/CUDADispatcher.inspect_llvm.<locals>.<dictcomp>¶  sC   € ð Eð Eð EÙ)˜C ð ˜XÔ-×:Ò:Ñ<Ô<ð Eð Eð ErB   c                 ó>   — i | ]\  }}||                      ¦   «         “ŒS rº   )r¥   r¡  s      r@   r¤  z/CUDADispatcher.inspect_llvm.<locals>.<dictcomp>¹  s@   € ð Eð Eð EÙ)˜C ð ˜X×2Ò2Ñ4Ô4ð Eð Eð ErB   )rt  r‡  rU  rR   r¤   r¥   r¥  )rj   ra   rm   s      r@   r¥   zCUDADispatcher.inspect_llvm¥  sÑ   € ð Ô#×'Ò'¨Ñ1Ô1ˆØÐ Øð @Ø”~ iÔ0Ô8×EÒEÑGÔGÐGà”~ iÔ0×=Ò=Ñ?Ô?Ð?àð EðEð EØ-1¬^×-AÒ-AÑ-CÔ-CðEñ Eô Eð EðEð EØ-1¬^×-AÒ-AÑ-CÔ-CðEñ Eô Eð ErB   c                 óˆ  ‡— t          ¦   «         j        Š| j                             d¦  «        }|�G|r%| j        |         j                             ‰¦  «        S | j        |                              ‰¦  «        S |r%ˆfd„| j                             ¦   «         D ¦   «         S ˆfd„| j                             ¦   «         D ¦   «         S )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.
        rm   Nc                 óL   •— i | ] \  }}||j                              ‰¦  «        “Œ!S rº   )rR   r<   ©r=   r¢  r£  r8   s      €r@   r¤  z.CUDADispatcher.inspect_asm.<locals>.<dictcomp>Î  sF   ø€ ð Eð Eð EÙ)˜C ð ˜XÔ-×9Ò9¸"Ñ=Ô=ð Eð Eð ErB   c                 óB   •— i | ]\  }}||                      ‰¦  «        “ŒS rº   )r§   rÚ  s      €r@   r¤  z.CUDADispatcher.inspect_asm.<locals>.<dictcomp>Ñ  sC   ø€ ð Eð Eð EÙ)˜C ð ˜X×1Ò1°"Ñ5Ô5ð Eð Eð ErB   )	r   rK   rt  r‡  rU  rR   r<   r§   r¥  )rj   ra   rm   r8   s      @r@   r§   zCUDADispatcher.inspect_asm¼  sü   ø€ õ  Ñ!Ô!Ô4ˆØÔ#×'Ò'¨Ñ1Ô1ˆØÐ Øð AØ”~ iÔ0Ô8×DÒDÀRÑHÔHÐHà”~ iÔ0×<Ò<¸RÑ@Ô@Ð@àð EðEð Eð Eð EØ-1¬^×-AÒ-AÑ-CÔ-CðEñ Eô Eð EðEð Eð Eð EØ-1¬^×-AÒ-AÑ-CÔ-CðEñ Eô Eð ErB   c                 óÜ   — | j                              d¦  «        rt          d¦  «        ‚|�| j        |                              ¦   «         S d„ | j                             ¦   «         D ¦   «         S )aƒ  
        Return this kernel's CFG for the device in the current context.

        :param signature: A tuple of argument types.
        :return: The CFG for the given signature, or a dict of CFGs
                 for all previously-encountered signatures.

        The CFG for the device in the current context is returned.

        Requires nvdisasm to be available on the PATH.
        rm   z'Cannot get the CFG of a device functionNc                 ó>   — i | ]\  }}||                      ¦   «         “ŒS rº   )rª   ©r=   r¢  Údefns      r@   r¤  z3CUDADispatcher.inspect_sass_cfg.<locals>.<dictcomp>æ  s:   € ð =ð =ð =Ù!˜˜Tð ˜×.Ò.Ñ0Ô0ð =ð =ð =rB   )rt  r‡  rC   rU  rª   r¥  r¦  s     r@   rª   zCUDADispatcher.inspect_sass_cfgÔ  s{   € ð Ô×!Ò! (Ñ+Ô+ð 	JÝÐHÑIÔIÐIàÐ Ø”> )Ô,×=Ò=Ñ?Ô?Ð?ð=ð =Ø%)¤^×%9Ò%9Ñ%;Ô%;ð=ñ =ô =ð =rB   c                 óÜ   — | j                              d¦  «        rt          d¦  «        ‚|�| j        |                              ¦   «         S d„ | j                             ¦   «         D ¦   «         S )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.
        rm   z(Cannot inspect SASS of a device functionNc                 ó>   — i | ]\  }}||                      ¦   «         “ŒS rº   )r­   rÞ  s      r@   r¤  z/CUDADispatcher.inspect_sass.<locals>.<dictcomp>ü  s:   € ð =ð =ð =Ù!˜˜Tð ˜×*Ò*Ñ,Ô,ð =ð =ð =rB   )rt  r‡  rC   rU  r­   r¥  r¦  s     r@   r­   zCUDADispatcher.inspect_sassé  s{   € ð Ô×!Ò! (Ñ+Ô+ð 	KÝÐIÑJÔJÐJàÐ Ø”> )Ô,×9Ò9Ñ;Ô;Ð;ð=ð =Ø%)¤^×%9Ò%9Ñ%;Ô%;ð=ñ =ô =ð =rB   c                 óŒ   — |€t           j        }| j                             ¦   «         D ]\  }}|                     |¬¦  «         ŒdS )r¯   Nr±   )r´   rµ   rU  r¥  r·   )rj   r²   rY  rß  s       r@   r·   zCUDADispatcher.inspect_typesÿ  sU   € ð ˆ<Ý”:ˆDà”~×+Ò+Ñ-Ô-ð 	*ð 	*‰GˆAˆtØ×Ò DÐÑ)Ô)Ð)Ð)ð	*ð 	*rB   c                 ó   —  | ||¦  «        }|S )r†   rº   )rˆ   rH   rt  rŠ   s       r@   r‹   zCUDADispatcher._rebuild  s   € ð
 �3�w Ñ.Ô.ˆØˆrB   c                 ó8   — t          | j        | j        ¬¦  «        S )zd
        Reduce the instance for serialization.
        Compiled definitions are discarded.
        )rH   rt  )r�   rH   rt  r{   s    r@   rŽ   zCUDADispatcher._reduce_states  s%   € õ
 ˜DœLØ"&Ô"4ð6ñ 6ô 6ð 	6rB   r9  )r   r   r   rz   )*r:  r;  r<  r=  Ú
_fold_argsr   Útargetdescrr   rE   r>  rz  r}  r¿   Ú	lru_cacher€  rƒ  r…  rJ   rM  rd  r’  r�  rK  rJ  r§  rª  r®  r²  rµ  rÂ  rº  rÌ  r�  r¥   r§   rª   r­   r·   r?  r‹   rŽ   r@  rA  s   @r@   rr  rr  8  s±  ø€ € € € € ðð ð €Jà€Kà>Jð "ð "ð "ð "ð "ð "ð  ð/ð /ñ „Xð/ð.ð .ð .ð €YÔ Ð%Ñ%Ô%ðPð Pð Pñ &Ô%ðPð%ð %ð %ð
Qð Qð Qð Qð, ð4ð 4ñ „Xð4ð(4ð 4ð 4ð	Bð 	Bð 	Bð-ð -ð -ðð ð ðð ð ð0 ð!ð !ñ „Xð!ðAð Að Að Að&Að Að Að Að(Að Að Að Að&Að Að Að Að*Að Að Að Að&/ð /ð /ð6%ð %ð %ð %ðN*ð *ð *ð
"ð "ð "ðHEð Eð Eð Eð.Eð Eð Eð Eð0=ð =ð =ð =ð*=ð =ð =ð =ð,
*ð 
*ð 
*ð 
*ð ðð ñ „[ðð6ð 6ð 6ð 6ð 6ð 6ð 6rB   rr  )9Únumpyr  rW   r´   rÑ   r¿   Ú
numba.corer   r   r   r   r   r   Únumba.core.cachingr	   r
   Únumba.core.compiler_lockr   Únumba.core.dispatcherr   Únumba.core.errorsr   Únumba.core.typing.typeofr   r   Únumba.cuda.apir   Únumba.cuda.argsr   Únumba.cuda.compilerr   r   Únumba.cuda.cudadrvr   Únumba.cuda.cudadrv.devicesr   Únumba.cuda.descriptorr   Únumba.cuda.errorsr   r   Ú
numba.cudary  Únumbar   r   Úwarningsr   rV   ÚReduceMixinr.   ÚobjectrC  r[  rg  ro  rr  rº   rB   r@   ú<module>rû     sê  ðØ Ð Ð Ð Ø 	€	€	€	Ø 
€
€
€
Ø €€€Ø Ð Ð Ð à HÐ HÐ HÐ HÐ HÐ HÐ HÐ HÐ HÐ HÐ HÐ HÐ HÐ HÐ HÐ HØ /Ð /Ð /Ð /Ð /Ð /Ð /Ð /Ø 9Ð 9Ð 9Ð 9Ð 9Ð 9Ø ,Ð ,Ð ,Ð ,Ð ,Ð ,Ø 5Ð 5Ð 5Ð 5Ð 5Ð 5Ø 4Ð 4Ð 4Ð 4Ð 4Ð 4Ð 4Ð 4à -Ð -Ð -Ð -Ð -Ð -Ø $Ð $Ð $Ð $Ð $Ð $Ø :Ð :Ð :Ð :Ð :Ð :Ð :Ð :Ø %Ð %Ð %Ð %Ð %Ð %Ø 2Ð 2Ð 2Ð 2Ð 2Ð 2Ø -Ð -Ð -Ð -Ð -Ð -ð<ð <ð <ð <ð <ð <ð <ð <à *Ð *Ð *Ð *Ð *Ð *à Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð à Ð Ð Ð Ð Ð ð*ð *ð *Ð ði/ð i/ð i/ð i/ð i/ˆiÔ#ñ i/ô i/ð i/ðX+ð +ð +ð +ð +ˆVñ +ô +ð +ð\Að Að Að Að Añ Aô Að Að:ð ð ð ð �Iñ ô ð ð$ ð  ð  ð  ð  �ñ  ô  ð  ða6ð a6ð a6ð a6ð a6�Z Ô!6ñ a6ô a6ð a6ð a6ð a6rB   