§
    1@Äj}C  ã                   óh  — d dl Z d dlmZ d dlmZ d dlmZ d dlm	Z	m
Z
mZmZmZ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mZ d dl m!Z!  G d„ dej        ¦  «        Z" e j#        de j$        ¦  «        Z% G d„ de¦  «        Z& G d„ de¦  «        Z' G d„ de¦  «        Z(dS )é    N)Úcached_property)Úir)ÚcgutilsÚconfigÚ	debuginfoÚitanium_manglerÚtypesÚtypingÚutils)Ú
Dispatcher)ÚBaseContext)ÚBaseCallConvÚMinimalCallConv)Ú	cmathdecl)Ú	datamodelé   )Únvvm)ÚcodegenÚ	nvvmutilsÚufuncs)Úcuda_data_managerc                   ó$   ‡ — e Zd Zd„ Zˆ fd„Zˆ xZS )ÚCUDATypingContextc                 ó¤  — ddl m}m}m}m} ddlm}m} |                      |j	        ¦  «         |                      |j	        ¦  «         |                      |j	        ¦  «         |                      t          j	        ¦  «         |                      |j	        ¦  «         |                      |j	        ¦  «         |                      |j        ¦  «         d S )Nr   )ÚcudadeclÚcudamathÚlibdevicedeclÚvector_typesr   )ÚenumdeclÚ
cffi_utils)Ú r   r   r   r   Únumba.core.typingr   r    Úinstall_registryÚregistryr   Útyping_registry)Úselfr   r   r   r   r   r    s          úBC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/target.pyÚload_additional_registriesz,CUDATypingContext.load_additional_registries   så   € ØEÐEÐEÐEÐEÐEÐEÐEÐEÐEÐEÐEØ:Ð:Ð:Ð:Ð:Ð:Ð:Ð:à×Ò˜hÔ/Ñ0Ô0Ð0Ø×Ò˜jÔ1Ñ2Ô2Ð2Ø×Ò˜hÔ/Ñ0Ô0Ð0Ø×Ò�iÔ0Ñ1Ô1Ð1Ø×Ò˜mÔ4Ñ5Ô5Ð5Ø×Ò˜hÔ/Ñ0Ô0Ð0Ø×Ò˜lÔ:Ñ;Ô;Ð;Ð;Ð;ó    c                 óÜ  •— ddl m} t          |t          ¦  «        r©t          ||¦  «        s™	 |j        }n�# t
          $ rƒ |j        st          d¦  «        ‚|j         	                    ¦   «         }d|d<   | 
                    dd¦  «        |d<   | 
                    dd¦  «        |d<    ||j        |¦  «        }||_        |}Y nw xY wt          t          | ¦  «                             |¦  «        S )	Nr   )ÚCUDADispatcherz<using cpu function on device but its compilation is disabledTÚdeviceÚdebugFÚopt)Únumba.cuda.dispatcherr+   Ú
isinstancer   Ú_CUDATypingContext__dispatcherÚAttributeErrorÚ_can_compileÚ
ValueErrorÚtargetoptionsÚcopyÚgetÚpy_funcÚsuperr   Úresolve_value_type)r&   Úvalr+   r5   ÚdispÚ	__class__s        €r'   r:   z$CUDATypingContext.resolve_value_type#   s#  ø€ à8Ð8Ð8Ð8Ð8Ð8Ý�s�JÑ'Ô'ð 	Ý˜3 Ñ/Ô/ð	ðàÔ&��øÝ!ð ð ð ØÔ'ð HÝ$ð &Gñ Hô Hð Hà #Ô 1× 6Ò 6Ñ 8Ô 8�Ø*.�˜hÑ'Ø)6×):Ò):¸7ÀEÑ)JÔ)J�˜gÑ&Ø'4×'8Ò'8¸ÀÑ'EÔ'E�˜eÑ$Ø%�~ c¤k°=ÑAÔA�ð $(�Ô Ø���ðøøøõ Õ&¨Ñ-Ô-×@Ò@ÀÑEÔEÐEs   ®6 ¶B
CÃC)Ú__name__Ú
__module__Ú__qualname__r(   r:   Ú__classcell__©r=   s   @r'   r   r      sP   ø€ € € € € ð
<ð 
<ð 
<ðFð Fð Fð Fð Fð Fð Fð Fð Fr)   r   z	[^a-z0-9]c                   óê   ‡ — e Zd ZdZdZdˆ fd„	Zed„ ¦   «         Zed„ ¦   «         Zd„ Z	d„ Z
d„ Zd	„ Zed
„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zdddœd„Z	 dd„Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zˆ xZS )ÚCUDATargetContextTÚcudac                 ó’   •— t          ¦   «                              ||¦  «         t          j        t          j        ¦  «        | _        d S ©N)r9   Ú__init__r   Úchainr   Údefault_managerÚdata_model_manager)r&   Ú	typingctxÚtargetr=   s      €r'   rH   zCUDATargetContext.__init__G   s>   ø€ Ý‰Œ×Ò˜ FÑ+Ô+Ð+Ý"3Ô"9ÝÔ%ñ#
ô #
ˆÔÐÐr)   c                 ó   — t           j        S rG   )r   Ú	DIBuilder©r&   s    r'   rO   zCUDATargetContext.DIBuilderM   s   € åÔ"Ð"r)   c                 ó   — dS )NF© rP   s    r'   Úenable_boundscheckz$CUDATargetContext.enable_boundscheckQ   s	   € ð ˆur)   c                 ó6   — | j                              |¦  «        S rG   )Ú_internal_codegenÚ_create_empty_module)r&   Únames     r'   Úcreate_modulezCUDATargetContext.create_moduleW   s   € ØÔ%×:Ò:¸4Ñ@Ô@Ð@r)   c                 óF   — t          j        d¦  «        | _        d | _        d S )Nznumba.cuda.jit)r   ÚJITCUDACodegenrU   Ú_target_datarP   s    r'   ÚinitzCUDATargetContext.initZ   s$   € Ý!(Ô!7Ð8HÑ!IÔ!IˆÔØ ˆÔÐÐr)   c                 ó  — ddl m}m}m} ddl m}m}m} ddl m}m} ddl m	}	 ddl
m}
 ddlm} ddlm} d	d
lm}m}m}m}m} ddlm} |                      |j        ¦  «         |                      |
j        ¦  «         |                      |j        ¦  «         |                      |j        ¦  «         |                      |	j        ¦  «         |                      |j        ¦  «         |                      |j        ¦  «         d S )Nr   )ÚnumbersÚtupleobjÚslicing)ÚrangeobjÚ	iteratorsÚenumimpl)ÚunicodeÚcharseq)Ú	cmathimpl)Úcffiimpl)Úarrayobj)Ú
npdatetimer   )ÚcudaimplÚ	printimplÚlibdeviceimplÚmathimplr   )Úndarray)Únumba.cpythonr^   r_   r`   ra   rb   rc   rd   re   rf   Ú
numba.miscrg   Únumba.nprh   ri   r!   rj   rk   rl   rm   r   Únumba.np.unsafern   r#   r$   Úimpl_registry)r&   r^   r_   r`   ra   rb   rc   rd   re   rf   rg   rh   ri   rj   rk   rl   rm   r   rn   s                      r'   r(   z,CUDATargetContext.load_additional_registries^   s¡  € ð 	=Ð<Ð<Ð<Ð<Ð<Ð<Ð<Ð<Ð<Ø?Ð?Ð?Ð?Ð?Ð?Ð?Ð?Ð?Ð?Ø2Ð2Ð2Ð2Ð2Ð2Ð2Ð2Ø+Ð+Ð+Ð+Ð+Ð+Ø'Ð'Ð'Ð'Ð'Ð'Ø%Ð%Ð%Ð%Ð%Ð%Ø'Ð'Ð'Ð'Ð'Ð'ð	
ð 	
ð 	
ð 	
ð 	
ð 	
ð 	
ð 	
ð 	
ð 	
ð 	
ð 	
ð 	
ð 	
ð 	,Ð+Ð+Ð+Ð+Ð+à×Ò˜hÔ/Ñ0Ô0Ð0Ø×Ò˜hÔ/Ñ0Ô0Ð0Ø×Ò˜iÔ0Ñ1Ô1Ð1Ø×Ò˜mÔ4Ñ5Ô5Ð5Ø×Ò˜iÔ0Ñ1Ô1Ð1Ø×Ò˜hÔ/Ñ0Ô0Ð0Ø×Ò˜lÔ8Ñ9Ô9Ð9Ð9Ð9r)   c                 ó   — | j         S rG   )rU   rP   s    r'   r   zCUDATargetContext.codegenv   s   € ØÔ%Ð%r)   c                 ó|   — | j         €/t          j        t          j        ¦   «         j        ¦  «        | _         | j         S rG   )r[   ÚllÚcreate_target_datar   ÚNVVMÚdata_layoutrP   s    r'   Útarget_datazCUDATargetContext.target_datay   s0   € àÔÐ$Ý "Ô 5µd´i±k´kÔ6MÑ NÔ NˆDÔØÔ Ð r)   c                 óN   ‡— ddl mŠ d}t          ˆfd„|D ¦   «         ¦  «        }|S )z×
        Some CUDA intrinsics are at the module level, but cannot be treated as
        constants, because they are loaded from a special register in the PTX.
        These include threadIdx, blockDim, etc.
        r   ©rE   )Ú	threadIdxÚblockDimÚblockIdxÚgridDimÚlaneidÚwarpsizec                 ó<   •— g | ]}t          j        ‰¦  «        |f‘ŒS rR   )r	   ÚModule)Ú.0ÚncrE   s     €r'   ú
<listcomp>z;CUDATargetContext.nonconst_module_attrs.<locals>.<listcomp>‰   s8   ø€ ð $9ð $9ð $9Ø(*õ &+¤\°$Ñ%7Ô%7¸Ð$<ð $9ð $9ð $9r)   )ÚnumbarE   Útuple)r&   Ú	nonconstsÚnonconsts_with_modrE   s      @r'   Únonconst_module_attrsz'CUDATargetContext.nonconst_module_attrs   sZ   ø€ ð 	ÐÐÐÐÐð!ˆ	å"ð $9ð $9ð $9ð $9Ø.7ð$9ñ $9ô $9ñ :ô :Ðà!Ð!r)   c                 ó    — t          | ¦  «        S rG   )ÚCUDACallConvrP   s    r'   Ú	call_convzCUDATargetContext.call_conv�   s   € å˜DÑ!Ô!Ð!r)   rR   N©Úabi_tagsÚuidc                ó2   — t          j        ||||¬¦  «        S )Nr�   )r   Úmangle)r&   rW   Úargtypesr‘   r’   s        r'   ÚmanglerzCUDATargetContext.mangler‘   s%   € ÝÔ% d¨H¸xØ*-ð/ñ /ô /ð 	/r)   c	           	      ó  — t          j        |j        d¬¦  «        }	|                      ¦   «                              |j        › d�|	||¬¦  «        }
|
                     |¦  «         |                      |
||	||||¦  «        }|
|fS )a»  
        Adapt a code library ``codelib`` with the numba compiled CUDA kernel
        with name ``fname`` and arguments ``argtypes`` for NVVM.
        A new library is created with a wrapper function that can be used as
        the kernel entry point for the given kernel.

        Returns the new code library and the wrapper function.

        Parameters:

        codelib:       The CodeLibrary containing the device function to wrap
                       in a kernel call.
        fndesc:        The FunctionDescriptor of the source function.
        debug:         Whether to compile with debug.
        lineinfo:      Whether to emit line info.
        nvvm_options:  Dict of NVVM options used when compiling the new library.
        filename:      The source filename that the function is contained in.
        linenum:       The source line that the function is on.
        max_registers: The max_registers argument for the code library.
        Úcudapy©ÚnsÚ_kernel_)Ú
entry_nameÚnvvm_optionsÚmax_registers)r   Úprepend_namespaceÚllvm_func_namer   Úcreate_libraryrW   Úadd_linking_libraryÚgenerate_kernel_wrapper)r&   ÚcodelibÚfndescr-   Úlineinfor�   ÚfilenameÚlinenumrž   Úkernel_nameÚlibraryÚwrappers               r'   Úprepare_cuda_kernelz%CUDATargetContext.prepare_cuda_kernel•   s§   € õ. &Ô7ØÔ! hð
ñ 
ô 
ˆð —,’,‘.”.×/Ò/°7´<Ð0IÐ0IÐ0IØ;FØ=IØ>Kð 0ñ Mô Mˆð 	×#Ò# GÑ,Ô,Ð,Ø×.Ò.¨w¸ÀØ/4°hÀØ/6ñ8ô 8ˆð ˜ÐÐr)   c                 óÀ  ‡!‡"— |j         }|                      |¦  «        }	t          |	j        ¦  «        }
t	          j        t	          j        ¦   «         |
¦  «        }|                      d¦  «        Š"t	          j        t	          j        d¦  «        | j	         
                    t          j        ¦  «        g|
z   ¦  «        }t	          j        ‰"||j        ¦  «        }t          j        |j        d¬¦  «        }t	          j        ‰"||¦  «        Š!t	          j        ‰!                     d¦  «        ¦  «        }|s|rH|o| }|                      ‰"|| |¬¦  «        } |j        ‰!||j        ||¦  «          |j        ||¦  «         ˆ!ˆ"fd„} |d¦  «        }g }g }d	D ]D}|                      |d
|z  ¦  «        ¦  «         |                      |d|z  ¦  «        ¦  «         ŒE|	                     |‰!j        ¦  «        }| j	                             ||t          j        ||¦  «        \  }}|�r¢t9          j        ||j        ¦  «        5  |                     ¦   «          ddd¦  «         n# 1 swxY w Y   |                      | !                    |j"        ¦  «        ¦  «        5  t	          j#        |j$        j%        d¦  «        }| &                    |||j'        dd¦  «        }| (                    |d¦  «        }tS          j*        |¦  «        }|                      |¦  «        5  tW          d	|¦  «        D ]0\  }}| ,                    |¦  «        } | -                    | |¦  «         Œ1tW          d	|¦  «        D ]0\  }}| .                    |¦  «        } | -                    | |¦  «         Œ1	 ddd¦  «         n# 1 swxY w Y   ddd¦  «         n# 1 swxY w Y   |                     ¦   «          t_          j0        ‰!¦  «         | 1                    ‰"¦  «         |s|r |j2        ¦   «          | 2                    ¦   «          tf          j4        rtk          j6        |‰"¦  «         | 7                    ‰!j        ¦  «        S )z°
        Generate the kernel wrapper in the given ``library``.
        The function being wrapped is described by ``fndesc``.
        The wrapper function is returned.
        zcuda.kernel.wrapperé    r˜   r™   r!   )ÚmoduleÚfilepathÚcgctxÚdirectives_onlyc                 ó´   •— ‰j         | z   }t          j        ‰t          j        d¦  «        |¦  «        }t          j        |j        j        d ¦  «        |_        |S )Nr®   )	rW   r   Úadd_global_variabler   ÚIntTypeÚConstantÚtypeÚpointeeÚinitializer)ÚpostfixrW   ÚgvÚwrapfnÚwrapper_modules      €€r'   Údefine_error_gvzBCUDATargetContext.generate_kernel_wrapper.<locals>.define_error_gvÛ   sL   ø€ Ø”; Ñ(ˆDÝÔ,¨^½R¼ZÈ¹^¼^Ø-1ñ3ô 3ˆBåœ[¨¬¬¸$Ñ?Ô?ˆBŒNØˆIr)   Ú__errcode__Úxyzz	__tid%s__z__ctaid%s__NÚ	monotonicr   )8r•   Úget_arg_packerÚlistÚargument_typesr   ÚFunctionTypeÚVoidTyperX   rµ   r�   Úget_return_typer	   ÚpyobjectÚFunctionr    r   rŸ   rW   Ú	IRBuilderÚappend_basic_blockrO   Úmark_subprogramÚargsÚmark_locationÚappendÚfrom_argumentsÚcall_functionÚvoidr   Ú	if_likelyÚis_okÚret_voidÚif_thenÚnot_Úis_python_excr¶   r·   r¸   ÚcmpxchgÚcodeÚextract_valuer   ÚSRegBuilderÚzipÚtidÚstoreÚctaidr   Úset_cuda_kernelÚadd_ir_moduleÚfinalizer   Ú	DUMP_LLVMr   Ú	dump_llvmÚget_function)#r&   rª   r¥   r©   r-   r¦   r§   r¨   r•   ÚarginfoÚargtysÚwrapfntyÚfntyÚfuncÚprefixedÚbuilderr²   r   r¾   Úgv_excÚgv_tidÚgv_ctaidÚiÚcallargsÚstatusÚ_ÚoldÚxchgÚchangedÚsregÚdimÚptrr;   r¼   r½   s#                                    @@r'   r£   z)CUDATargetContext.generate_kernel_wrapper¹   s.  øø€ ð ”?ˆØ×%Ò% hÑ/Ô/ˆÝ�gÔ,Ñ-Ô-ˆÝ”?¥2¤;¡=¤=°&Ñ9Ô9ˆØ×+Ò+Ð,AÑBÔBˆÝŒ�rœz¨"™~œ~Ø $¤× >Ò >½u¼~Ñ NÔ NÐOØ!'ñ (ñ)ô )ˆõ Œ{˜>¨4°Ô1FÑGÔGˆå"Ô4°T´YÀ8ÐLÑLÔLˆÝ”˜^¨X°xÑ@Ô@ˆÝ”,˜v×8Ò8¸Ñ<Ô<Ñ=Ô=ˆàð 		6�Hð 		6Ø&Ð4¨u¨9ˆOØŸš¨nØ08Ø-1Ø7Fð 'ñ Hô HˆIð &ˆIÔ%Ø˜ V¤[°(¸Gñô ð ð $ˆIÔ# G¨WÑ5Ô5Ð5ð	ð 	ð 	ð 	ð 	ð 	ð !� Ñ/Ô/ˆØˆØˆØð 	@ð 	@ˆAØ�MŠM˜/˜/¨+¸©/Ñ:Ô:Ñ;Ô;Ð;Ø�OŠO˜O˜O¨M¸AÑ,=Ñ>Ô>Ñ?Ô?Ð?Ð?à×)Ò)¨'°6´;Ñ?Ô?ˆØ”N×0Ò0Ø�T�5œ: x°ñ;ô ;‰	ˆ�ð ñ 	0åÔ" 7¨F¬LÑ9Ô9ð #ð #Ø× Ò Ñ"Ô"Ð"ð#ð #ð #ñ #ô #ð #ð #ð #ð #ð #ð #øøøð #ð #ð #ð #ð —’ §¢¨fÔ.BÑ!CÔ!CÑDÔDð 0ð 0å”k &¤+Ô"5°tÑ<Ô<�ð
 —’ v¨s°F´KØ'2°KñAô A�à!×/Ò/°°aÑ8Ô8�õ !Ô,¨WÑ5Ô5�Ø—_’_ WÑ-Ô-ð 0ð 0Ý%(¨°Ñ%7Ô%7ð 0ð 0™	˜˜SØ"Ÿhšh s™mœm˜ØŸš c¨3Ñ/Ô/Ð/Ð/å%(¨°Ñ%9Ô%9ð 0ð 0™	˜˜SØ"Ÿjšj¨™oœo˜ØŸš c¨3Ñ/Ô/Ð/Ð/ð0ð0ð 0ð 0ñ 0ô 0ð 0ð 0ð 0ð 0ð 0ð 0øøøð 0ð 0ð 0ð 0ð0ð 0ð 0ñ 0ô 0ð 0ð 0ð 0ð 0ð 0ð 0øøøð 0ð 0ð 0ð 0ð, 	×ÒÑÔÐåÔ˜VÑ$Ô$Ð$Ø×Ò˜nÑ-Ô-Ð-Øð 	!�Hð 	!ØˆIÔÑ Ô Ð Ø×ÒÑÔÐåÔð 	4ÝŒO˜F NÑ3Ô3Ð3à×#Ò# F¤KÑ0Ô0Ð0sJ   È9IÉIÉ!IÊA=N6ÌBNÎN6ÎN#	Î#N6Î&N#	Î'N6Î6N:Î=N:c           	      óÎ  ‡ — |j         }ˆ fd„t          |                     d¬¦  «        ¦  «        D ¦   «         }t          j        t          j        d¦  «        t          |¦  «        ¦  «        }t          j        ||¦  «        }t          j	        }t          j        ||j        d|¬¦  «        }	d|	_        d|	_        ||	_        ‰                      |j        ¦  «        }
‰                      |
¦  «        }d	|d
z
                       ¦   «         z  |	_        t          j        t          j        d¦  «        ¦  «        }|                     |	|d¦  «        } ‰                      |¦  «        ‰ |¦  «        }ˆ fd„|j        D ¦   «         }ˆ fd„|j        D ¦   «         }‰                      ||                     ||j        j        ¦  «        |||j        |j        d¬¦  «         |                     ¦   «         S )úi
        Unlike the parent version.  This returns a a pointer in the constant
        addrspace.
        c                 óP   •— g | ]"}‰                      t          j        |¦  «        ‘Œ#S rR   )Úget_constantr	   Úbyte)r…   rñ   r&   s     €r'   r‡   z9CUDATargetContext.make_constant_array.<locals>.<listcomp>  s;   ø€ ð 
ð 
ð 
àð ×Ò�eœj¨!Ñ,Ô,ð
ð 
ð 
r)   ÚA)Úorderé   Ú_cudapy_cmem©Ú	addrspaceÚinternalTé   r   Úgenericc                 óP   •— g | ]"}‰                      t          j        |¦  «        ‘Œ#S rR   ©rþ   r	   Úintp©r…   Úsr&   s     €r'   r‡   z9CUDATargetContext.make_constant_array.<locals>.<listcomp>6  s+   ø€ ÐFÐFÐF°q�$×#Ò#¥E¤J°Ñ2Ô2ÐFÐFÐFr)   c                 óP   •— g | ]"}‰                      t          j        |¦  «        ‘Œ#S rR   r
  r  s     €r'   r‡   z9CUDATargetContext.make_constant_array.<locals>.<listcomp>7  s+   ø€ ÐJÐJÐJ¸�D×%Ò%¥e¤j°!Ñ4Ô4ÐJÐJÐJr)   N)ÚdataÚshapeÚstridesÚitemsizeÚparentÚmeminfo) r¯   ÚiterÚtobytesr   Ú	ArrayTyperµ   Úlenr¶   r   ÚADDRSPACE_CONSTANTr   r´   r·   ÚlinkageÚglobal_constantr¹   Úget_data_typeÚdtypeÚget_abi_sizeofÚ
bit_lengthÚalignÚPointerTypeÚaddrspacecastÚ
make_arrayr  r  Úpopulate_arrayÚbitcastr  r  r  Ú	_getvalue)r&   rí   ÚarytyÚarrÚlmodÚ	constvalsÚ
constarytyÚconstaryr  r»   Úlldtyper   ÚptrtyÚgenptrÚaryÚkshapeÚkstridess   `                r'   Úmake_constant_arrayz%CUDATargetContext.make_constant_array  sÜ  ø€ ð Œ~ˆð
ð 
ð 
ð 
å˜#Ÿ+š+¨C˜+Ñ0Ô0Ñ1Ô1ð
ñ 
ô 
ˆ	õ ”\¥"¤*¨Q¡-¤-µ°Y±´Ñ@Ô@ˆ
Ý”;˜z¨9Ñ5Ô5ˆåÔ+ˆ	ÝÔ(¨¨x¬}¸nØ3<ð>ñ >ô >ˆàˆŒ
Ø!ˆÔØ!ˆŒð ×$Ò$ U¤[Ñ1Ô1ˆØ×#Ò# GÑ,Ô,ˆØ˜ ™×.Ò.Ñ0Ô0Ñ0ˆŒõ ”�rœz¨!™}œ}Ñ-Ô-ˆØ×&Ò& r¨5°)Ñ<Ô<ˆð %ˆd�oŠo˜eÑ$Ô$ T¨7Ñ3Ô3ˆØFÐFÐFÐF¸C¼IÐFÑFÔFˆØJÐJÐJÐJ¸c¼kÐJÑJÔJˆØ×Ò˜C g§o¢o°f¸c¼h¼mÑ&LÔ&LØ"(Ø$,Ø%(¤\¸#¼*Ø$(ð	 	ñ 	*ô 	*ð 	*ð �}Š}‰ŒÐr)   c                 óÞ  — t          j        |                     d¦  «        dz   ¦  «        }d                     dt	          j        |¦  «        g¦  «        }|j                             |¦  «        }|€<t          j        ||j	        |t          j        ¬¦  «        }d|_        d|_        ||_        |j	        j        j        }|                     |                     t          j        ¦  «        ¦  «        S )	rü   zutf-8ó    ú$Ú__conststring__Nr  r  T)r   Úmake_bytearrayÚencodeÚjoinr   Úmangle_identifierÚglobalsr7   r´   r·   r   r  r  r  r¹   r¸   Úelementr%  Ú
as_pointer)r&   ÚmodÚstringÚtextrW   r»   Úchartys          r'   Úinsert_const_stringz%CUDATargetContext.insert_const_string@  sÚ   € õ
 Ô% f§m¢m°GÑ&<Ô&<¸wÑ&FÑGÔGˆØ�xŠxÐ*Ý(Ô:¸6ÑBÔBðDñ Eô Eˆð Œ[�_Š_˜TÑ"Ô"ˆØˆ:åÔ,¨S°$´)¸TÝ7;Ô7NðPñ Pô PˆBà#ˆBŒJØ!%ˆBÔØ!ˆBŒNð ””Ô(ˆØ�zŠz˜&×+Ò+­DÔ,CÑDÔDÑEÔEÐEr)   c                 ó¶   — |j         }|                      ||¦  «        }t          j        t          j        d¦  «        ¦  «        }|                     ||d¦  «        S )z«
        Insert a constant string in the constant addresspace and return a
        generic i8 pointer to the data.

        This function attempts to deduplicate.
        r  r  )r¯   rC  r   r!  rµ   r"  )r&   rí   r@  r)  r»   Ú	charptrtys         r'   Úinsert_string_const_addrspacez/CUDATargetContext.insert_string_const_addrspaceV  sO   € ð Œ~ˆØ×%Ò% d¨FÑ3Ô3ˆÝ”N¥2¤:¨a¡=¤=Ñ1Ô1ˆ	Ø×$Ò$ R¨°IÑ>Ô>Ð>r)   c                 ó   — dS )zRun O1 function passes
        NrR   ©r&   rë   s     r'   Úoptimize_functionz#CUDATargetContext.optimize_functionb  s	   € ð 	ˆr)   c                 ó*   — t          j        |¦  «        S rG   )r   Úget_ufunc_info)r&   Ú	ufunc_keys     r'   rK  z CUDATargetContext.get_ufunc_infoo  s   € ÝÔ$ YÑ/Ô/Ð/r)   r|   rG   )r>   r?   r@   Úimplement_powi_as_math_callÚstrict_alignmentrH   ÚpropertyrO   rS   rX   r\   r(   r   rz   r   rŒ   r�   r–   r¬   r£   r3  rC  rF  rI  rK  rA   rB   s   @r'   rD   rD   C   s«  ø€ € € € € Ø"&ÐØÐð
ð 
ð 
ð 
ð 
ð 
ð ð#ð #ñ „Xð#ð ðð ñ „Xðð
Að Að Að!ð !ð !ð:ð :ð :ð0&ð &ð &ð ð!ð !ñ „Xð!ð
 ð"ð "ñ „_ð"ð ð"ð "ñ „_ð"ð 35¸$ð /ð /ð /ð /ð /ð +/ð" ð " ð " ð " ðHZ1ð Z1ð Z1ðx)ð )ð )ðVFð Fð Fð,
?ð 
?ð 
?ðð ð ð0ð 0ð 0ð 0ð 0ð 0ð 0r)   rD   c                   ó   — e Zd ZdS )rŽ   N)r>   r?   r@   rR   r)   r'   rŽ   rŽ   s  s   € € € € € Ø€Dr)   rŽ   c                   óP   — e Zd ZdZd„ Zd„ Z	 	 dd„Zd„ Zd„ Zdd	„Z	d
„ Z
d„ Zd„ ZdS )ÚCUDACABICallConvzÔ
    Calling convention aimed at matching the CUDA C/C++ ABI. The implemented
    function signature is:

        <Python return type> (<Python arguments>)

    Exceptions are unsupported in this convention.
    c                 ó   — d S rG   rR   )r&   rí   s     r'   Ú_make_call_helperz"CUDACABICallConv._make_call_helper�  s	   € ð ˆtr)   c                 ó,   — |                      |¦  «        S rG   )Úret)r&   rí   Úretvals      r'   Úreturn_valuezCUDACABICallConv.return_value†  s   € Ø�{Š{˜6Ñ"Ô"Ð"r)   Nc                 ó$   — d}t          |¦  «        ‚)Nz7Python exceptions are unsupported in the CUDA C/C++ ABI©ÚNotImplementedError)r&   rí   ÚexcÚexc_argsÚlocÚ	func_nameÚmsgs          r'   Úreturn_user_excz CUDACABICallConv.return_user_exc‰  s   € àGˆÝ! #Ñ&Ô&Ð&r)   c                 ó$   — d}t          |¦  «        ‚)Nz2Return status is unsupported in the CUDA C/C++ ABIrZ  )r&   rí   ró   r`  s       r'   Úreturn_status_propagatez(CUDACABICallConv.return_status_propagateŽ  s   € ØBˆÝ! #Ñ&Ô&Ð&r)   c                 ó¨   — |                       |¦  «        }t          |j        ¦  «        }t          j        |                      |¦  «        |¦  «        }|S )zM
        Get the LLVM IR Function type for *restype* and *argtypes*.
        )Ú_get_arg_packerrÃ   rÄ   r   rÅ   rÇ   )r&   Úrestyper•   rç   rê   s        r'   Úget_function_typez"CUDACABICallConv.get_function_type’  sK   € ð ×&Ò& xÑ0Ô0ˆÝ˜Ô.Ñ/Ô/ˆÝŒ˜t×3Ò3°GÑ<Ô<¸hÑGÔGˆØˆr)   Fc                 óž   — |rJ ‚|                       |¦  «        }|                     |                      |¦  «        d„ |D ¦   «         ¦  «         dS )zA
        Set names and attributes of function arguments.
        c                 ó   — g | ]}d |z   ‘ŒS )zarg.rR   )r…   Úas     r'   r‡   z6CUDACABICallConv.decorate_function.<locals>.<listcomp>¢  s   € Ð7Ð7Ð7¨Q˜f q™jÐ7Ð7Ð7r)   N)re  Úassign_namesÚget_arguments)r&   ÚfnrÍ   Úfe_argtypesÚnoaliasrç   s         r'   Údecorate_functionz"CUDACABICallConv.decorate_function›  se   € ð ÐÐˆ{Ø×&Ò& {Ñ3Ô3ˆØ×Ò˜T×/Ò/°Ñ3Ô3Ø7Ð7°$Ð7Ñ7Ô7ñ	9ô 	9ð 	9ð 	9ð 	9r)   c                 ó   — |j         S )z@
        Get the Python-level arguments of LLVM *func*.
        )rÍ   rH  s     r'   rl  zCUDACABICallConv.get_arguments¤  s   € ð ŒyÐr)   c                 óÈ   — |                       |¦  «        }|                     ||¦  «        }|                     ||¦  «        }d}	| j                             |||¦  «        }
|	|
fS )z3
        Call the Numba-compiled *callee*.
        N)re  Úas_argumentsÚcallÚcontextÚget_returned_value)r&   rí   ÚcalleeÚrestyrè   rÍ   rç   ÚrealargsrÚ   ró   Úouts              r'   rÑ   zCUDACABICallConv.call_functionª  si   € ð ×&Ò& vÑ.Ô.ˆØ×'Ò'¨°Ñ6Ô6ˆØ�|Š|˜F HÑ-Ô-ˆð ˆØŒl×-Ò-¨g°u¸dÑCÔCˆØ�sˆ{Ðr)   c                 óJ   — | j         j        |                              ¦   «         S rG   )ru  rK   rÇ   )r&   Útys     r'   rÇ   z CUDACABICallConv.get_return_type·  s   € ØŒ|Ô.¨rÔ2×BÒBÑDÔDÐDr)   )NNN)F)r>   r?   r@   Ú__doc__rT  rX  ra  rc  rg  rp  rl  rÑ   rÇ   rR   r)   r'   rR  rR  w  s¾   € € € € € ðð ðð ð ð
#ð #ð #ð @DØ"&ð'ð 'ð 'ð 'ð
'ð 'ð 'ðð ð ð9ð 9ð 9ð 9ðð ð ðð ð ðEð Eð Eð Eð Er)   rR  ))ÚreÚ	functoolsr   Úllvmlite.bindingÚbindingrv   Úllvmliter   Ú
numba.corer   r   r   r   r	   r
   r   Únumba.core.dispatcherr   Únumba.core.baser   Únumba.core.callconvr   r   r"   r   r   Úcudadrvr   Ú
numba.cudar   r   r   Únumba.cuda.modelsr   r   ÚcompileÚIÚVALID_CHARSrD   rŽ   rR  rR   r)   r'   ú<module>r�     s2  ðØ 	€	€	€	Ø %Ð %Ð %Ð %Ð %Ð %Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð ð'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'ð 'à ,Ð ,Ð ,Ð ,Ð ,Ð ,Ø 'Ð 'Ð 'Ð 'Ð 'Ð 'Ø =Ð =Ð =Ð =Ð =Ð =Ð =Ð =Ø 'Ð 'Ð 'Ð 'Ð 'Ð 'Ø  Ð  Ð  Ð  Ð  Ð  à Ð Ð Ð Ð Ð Ø 1Ð 1Ð 1Ð 1Ð 1Ð 1Ð 1Ð 1Ð 1Ð 1Ø /Ð /Ð /Ð /Ð /Ð /ð$Fð $Fð $Fð $Fð $F˜Ô*ñ $Fô $Fð $FðT ˆbŒj˜ r¤tÑ,Ô,€ðm0ð m0ð m0ð m0ð m0˜ñ m0ô m0ð m0ð`		ð 	ð 	ð 	ð 	�?ñ 	ô 	ð 	ðAEð AEð AEð AEð AE�|ñ AEô AEð AEð AEð AEr)   