§
    ƒÛg;8  ã                   ó@  — 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 d dlmZ d dlmZ d dl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$dS )é    N)Úcached_property)Úir)ÚtypingÚtypesÚ	debuginfoÚitanium_manglerÚcgutils)Ú
Dispatcher)ÚBaseContext)Ú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          úG/granary/edev/quantum/lib/python3.11/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(   r2   ÚdispÚ	__class__s        €r$   r7   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%   r7   Ú__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)r6   Ú__init__r   Úchainr   Údefault_managerÚdata_model_manager)r#   Ú	typingctxÚtargetr:   s      €r$   rE   zCUDATargetContext.__init__F   s>   ø€ Ý‰Œ×Ò˜ FÑ+Ô+Ð+Ý"3Ô"9ÝÔ%ñ#
ô #
ˆÔÐÐr&   c                 ó   — t           j        S rD   )r   Ú	DIBuilder©r#   s    r$   rL   zCUDATargetContext.DIBuilderL   s   € åÔ"Ð"r&   c                 ó   — dS )NF© rM   s    r$   Úenable_boundscheckz$CUDATargetContext.enable_boundscheckP   s	   € ð ˆur&   c                 ó6   — | j                              |¦  «        S rD   )Ú_internal_codegenÚ_create_empty_module)r#   Únames     r$   Úcreate_modulezCUDATargetContext.create_moduleV   s   € ØÔ%×:Ò:¸4Ñ@Ô@Ð@r&   c                 óF   — t          j        d¦  «        | _        d | _        d S )Nznumba.cuda.jit)r   ÚJITCUDACodegenrR   Ú_target_datarM   s    r$   ÚinitzCUDATargetContext.initY   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]   r^   r_   r`   ra   rb   rc   Ú
numba.miscrd   Únumba.npre   rf   r   rg   rh   ri   rj   r   Únumba.np.unsaferk   r    r!   Úimpl_registry)r#   r[   r\   r]   r^   r_   r`   ra   rb   rc   rd   re   rf   rg   rh   ri   rj   r   rk   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 rD   )rR   rM   s    r$   r   zCUDATargetContext.codegenu   s   € ØÔ%Ð%r&   c                 ó|   — | j         €/t          j        t          j        ¦   «         j        ¦  «        | _         | j         S rD   )rX   ÚllÚcreate_target_datar   ÚNVVMÚdata_layoutrM   s    r$   Útarget_datazCUDATargetContext.target_datax   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   ©rB   )Ú	threadIdxÚblockDimÚblockIdxÚgridDimÚlaneidÚwarpsizec                 ó<   •— g | ]}t          j        ‰¦  «        |f‘ŒS rO   )r   ÚModule)Ú.0ÚncrB   s     €r$   ú
<listcomp>z;CUDATargetContext.nonconst_module_attrs.<locals>.<listcomp>ˆ   s8   ø€ ð $9ð $9ð $9Ø(*õ &+¤\°$Ñ%7Ô%7¸Ð$<ð $9ð $9ð $9r&   )ÚnumbarB   Útuple)r#   Ú	nonconstsÚnonconsts_with_modrB   s      @r$   Únonconst_module_attrsz'CUDATargetContext.nonconst_module_attrs~   sZ   ø€ ð 	ÐÐÐÐÐð!ˆ	å"ð $9ð $9ð $9ð $9Ø.7ð$9ñ $9ô $9ñ :ô :Ðà!Ð!r&   c                 ó    — t          | ¦  «        S rD   )ÚCUDACallConvrM   s    r$   Ú	call_convzCUDATargetContext.call_convŒ   s   € å˜DÑ!Ô!Ð!r&   rO   N©Úabi_tagsÚuidc                ó2   — t          j        ||||¬¦  «        S )Nr�   )r   Úmangle)r#   rT   Ú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_libraryrT   Ú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                    ¦   «          | 3                    ‰!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«   )	rT   r	   Úadd_global_variabler   ÚIntTypeÚConstantÚtypeÚpointeeÚinitializer)ÚpostfixrT   Ú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   )4r’   Úget_arg_packerÚlistÚargument_typesr   ÚFunctionTypeÚVoidTyperU   r²   rŒ   Úget_return_typer   ÚpyobjectÚFunctionr�   r   rœ   rT   Ú	IRBuilderÚappend_basic_blockrL   Ú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ÚfinalizeÚ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Úptrr8   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ÔÑ Ô Ð Ø×ÒÑÔÐØ×%Ò% f¤kÑ2Ô2ˆØˆsJ   È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 rO   )Ú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 rO   ©rù   r   Úintp©r‚   Úsr#   s     €r$   r„   z9CUDATargetContext.make_constant_array.<locals>.<listcomp>2  s+   ø€ ÐFÐFÐF°q�$×#Ò#¥E¤J°Ñ2Ô2ÐFÐFÐFr&   c                 óP   •— g | ]"}‰                      t          j        |¦  «        ‘Œ#S rO   r  r  s     €r$   r„   z9CUDATargetContext.make_constant_array.<locals>.<listcomp>3  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Úglobalsr4   r±   r´   r   r  r  r  r¶   rµ   Úelementr   Ú
as_pointer)r#   ÚmodÚstringÚtextrT   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¬   r>  r   r  r²   r  )r#   rè   r;  r$  r¸   Ú	charptrtys         r$   Úinsert_string_const_addrspacez/CUDATargetContext.insert_string_const_addrspaceR  sO   € ð Œ~ˆØ×%Ò% d¨FÑ3Ô3ˆÝ”N¥2¤:¨a¡=¤=Ñ1Ô1ˆ	Ø×$Ò$ R¨°IÑ>Ô>Ð>r&   c                 ó   — dS )zRun O1 function passes
        NrO   )r#   ræ   s     r$   Úoptimize_functionz#CUDATargetContext.optimize_function^  s	   € ð 	ˆr&   c                 ó*   — t          j        |¦  «        S rD   )r   Úget_ufunc_info)r#   Ú	ufunc_keys     r$   rE  z CUDATargetContext.get_ufunc_infok  s   € ÝÔ$ YÑ/Ô/Ð/r&   ry   rD   )r;   r<   r=   Úimplement_powi_as_math_callÚstrict_alignmentrE   ÚpropertyrL   rP   rU   rY   r%   r   rw   r   r‰   rŒ   r“   r©   r    r.  r>  rA  rC  rE  r>   r?   s   @r$   rA   rA   B   s«  ø€ € € € € Ø"&ÐØÐð
ð 
ð 
ð 
ð 
ð 
ð ð#ð #ñ „Xð#ð ðð ñ „Xðð
Að Að Að!ð !ð !ð:ð :ð :ð0&ð &ð &ð ð!ð !ñ „Xð!ð
 ð"ð "ñ „_ð"ð ð"ð "ñ „_ð"ð 35¸$ð /ð /ð /ð /ð /ð +/ð" ð " ð " ð " ðHWð Wð Wðr)ð )ð )ðVFð Fð Fð,
?ð 
?ð 
?ðð ð ð0ð 0ð 0ð 0ð 0ð 0ð 0r&   rA   c                   ó   — e Zd ZdS )r‹   N)r;   r<   r=   rO   r&   r$   r‹   r‹   o  s   € € € € € Ø€Dr&   r‹   )%ÚreÚ	functoolsr   Úllvmlite.bindingÚbindingrs   Úllvmliter   Ú
numba.corer   r   r   r   r	   Únumba.core.dispatcherr
   Únumba.core.baser   Únumba.core.callconvr   r   r   r   Úcudadrvr   Ú
numba.cudar   r   r   Únumba.cuda.modelsr   r   ÚcompileÚIÚVALID_CHARSrA   r‹   rO   r&   r$   ú<module>rZ     s¾  ðØ 	€	€	€	Ø %Ð %Ð %Ð %Ð %Ð %Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð à IÐ IÐ IÐ IÐ IÐ IÐ IÐ IÐ IÐ IÐ IÐ IÐ IÐ IØ ,Ð ,Ð ,Ð ,Ð ,Ð ,Ø 'Ð 'Ð 'Ð 'Ð 'Ð 'Ø /Ð /Ð /Ð /Ð /Ð /Ø 'Ð 'Ð 'Ð 'Ð 'Ð 'Ø  Ð  Ð  Ð  Ð  Ð  à Ð Ð Ð Ð Ð Ø 1Ð 1Ð 1Ð 1Ð 1Ð 1Ð 1Ð 1Ð 1Ð 1Ø /Ð /Ð /Ð /Ð /Ð /ð$Fð $Fð $Fð $Fð $F˜Ô*ñ $Fô $Fð $FðT ˆbŒj˜ r¤tÑ,Ô,€ðj0ð j0ð j0ð j0ð j0˜ñ j0ô j0ð j0ðZ		ð 	ð 	ð 	ð 	�?ñ 	ô 	ð 	ð 	ð 	r&   