ó
    „ñ:iÅA  ã                   óf  • S SK r S SKJr  S SKJr  S SKJr  S SKJ	r	J
r
JrJrJrJrJr  S SKJr  S SKJr  S SKJrJr  S SKJr  S S	KJr  S
SKJr  S SKJrJrJr  S SK J!r!   " S S\R&                  5      r"\ RF                  " S\ RH                  5      r% " S S\5      r& " S S\5      r' " S S\5      r(g)é    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                   ó.   ^ • \ rS rSrS rU 4S jrSrU =r$ )ÚCUDATypingContexté   c                 ó®  • SSK JnJnJnJn  SSKJnJn  U R                  UR                  5        U R                  UR                  5        U R                  UR                  5        U R                  [        R                  5        U R                  UR                  5        U R                  UR                  5        U R                  UR                  5        g )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          ÚT/srv/projetos/modelo_ml_acdoc/venv/lib/python3.13/site-packages/numba/cuda/target.pyÚload_additional_registriesÚ,CUDATypingContext.load_additional_registries   s—   € ßEÓEß:à×Ñ˜h×/Ñ/Ô0Ø×Ñ˜j×1Ñ1Ô2Ø×Ñ˜h×/Ñ/Ô0Ø×Ñœi×0Ñ0Ô1Ø×Ñ˜m×4Ñ4Ô5Ø×Ñ˜h×/Ñ/Ô0Ø×Ñ˜l×:Ñ:Õ;ó    c                 óº  >• SSK Jn  [        U[        5      (       a  [        X5      (       d   UR                  n[        [        U ];  U5      $ ! [
         aƒ    UR                  (       d  [        S5      eUR                  R                  5       nSUS'   UR                  SS5      US'   UR                  SS5      US'   U" UR                  U5      nXAl        Un NŸf = f)	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-   r7   ÚdispÚ	__class__s        €r(   r<   Ú$CUDATypingContext.resolve_value_type#   sÛ   ø€ å8Ü�sœJ×'Ñ'Ü˜3×/Ñ/ðà×&Ñ&�ô  Ô&¨Ñ@ÀÓEÐEøô "ó Ø×'×'Ü$ð &Gó Hð Hà #× 1Ñ 1× 6Ñ 6Ó 8�Ø*.�˜hÑ'Ø)6×):Ñ):¸7ÀEÓ)J�˜gÑ&Ø'4×'8Ñ'8¸ÀÓ'E�˜eÑ$Ù% c§k¡k°=ÓA�ð $(Ô Ø’ðús   ®A ÁB
CÃC© )Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__r)   r<   Ú__static_attributes__Ú__classcell__©r?   s   @r(   r   r      s   ø† ò
<÷Fó Fr+   r   z	[^a-z0-9]c                   óÜ   ^ • \ rS rSrSrSrSU 4S jjr\S 5       r\S 5       r	S r
S rS rS	 r\S
 5       r\S 5       r\S 5       rSSS.S jr SS jrS rS rS rS rS rS rSrU =r$ )ÚCUDATargetContextéC   Tc                 óv   >• [         TU ]  X5        [        R                  " [        R
                  5      U l        g ©N)r;   Ú__init__r   Úchainr   Údefault_managerÚdata_model_manager)r'   Ú	typingctxÚtargetr?   s      €r(   rN   ÚCUDATargetContext.__init__G   s,   ø€ Ü‰Ñ˜Ô+Ü"3×"9Ò"9Ü×%Ñ%ó#
ˆÕr+   c                 ó"   • [         R                  $ rM   )r   Ú	DIBuilder©r'   s    r(   rV   ÚCUDATargetContext.DIBuilderM   s   € ä×"Ñ"Ð"r+   c                 ó   • g)NFrA   rW   s    r(   Úenable_boundscheckÚ$CUDATargetContext.enable_boundscheckQ   s   € ð r+   c                 ó8   • U R                   R                  U5      $ rM   )Ú_internal_codegenÚ_create_empty_module)r'   Únames     r(   Úcreate_moduleÚCUDATargetContext.create_moduleW   s   € Ø×%Ñ%×:Ñ:¸4Ó@Ð@r+   c                 óH   • [         R                  " S5      U l        S U l        g )Nznumba.cuda.jit)r   ÚJITCUDACodegenr]   Ú_target_datarW   s    r(   ÚinitÚCUDATargetContext.initZ   s   € Ü!(×!7Ò!7Ð8HÓ!IˆÔØ ˆÕr+   c                 ó  • SSK JnJnJn  SSK JnJnJn  SSK JnJn  SSK J	n	  SSK
Jn
  SSKJn  SSKJn  S	S
KJnJnJnJnJn  SSKJn  U R/                  UR0                  5        U R/                  U
R0                  5        U R/                  UR0                  5        U R/                  UR0                  5        U R/                  U	R0                  5        U R/                  UR0                  5        U R/                  UR2                  5        g )Nr   )ÚnumbersÚtupleobjÚslicing)ÚrangeobjÚ	iteratorsÚenumimpl)ÚunicodeÚcharseq)Ú	cmathimpl)Úcffiimpl)Úarrayobj)Ú
npdatetimer   )ÚcudaimplÚ	printimplÚlibdeviceimplÚmathimplr   )Úndarray)Únumba.cpythonrh   ri   rj   rk   rl   rm   rn   ro   rp   Ú
numba.miscrq   Únumba.nprr   rs   r"   rt   ru   rv   rw   r   Únumba.np.unsaferx   r$   r%   Úimpl_registry)r'   rh   ri   rj   rk   rl   rm   rn   ro   rp   rq   rr   rs   rt   ru   rv   rw   r   rx   s                      r(   r)   Ú,CUDATargetContext.load_additional_registries^   sº   € ÷ 	=Ñ<ß?Ñ?ß2Ý+Ý'Ý%Ý'÷	
õ 	
õ 	,à×Ñ˜h×/Ñ/Ô0Ø×Ñ˜h×/Ñ/Ô0Ø×Ñ˜i×0Ñ0Ô1Ø×Ñ˜m×4Ñ4Ô5Ø×Ñ˜i×0Ñ0Ô1Ø×Ñ˜h×/Ñ/Ô0Ø×Ñ˜l×8Ñ8Õ9r+   c                 ó   • U R                   $ rM   )r]   rW   s    r(   r   ÚCUDATargetContext.codegenv   s   € Ø×%Ñ%Ð%r+   c                 ó¤   • U R                   c8  [        R                  " [        R                  " 5       R
                  5      U l         U R                   $ rM   )rd   ÚllÚcreate_target_datar   ÚNVVMÚdata_layoutrW   s    r(   Útarget_dataÚCUDATargetContext.target_datay   s9   € à×ÑÑ$Ü "× 5Ò 5´d·i²i³k×6MÑ6MÓ NˆDÔØ× Ñ Ð r+   c                 ó„   • SSK Jn  Sn[        U Vs/ s H  n[        R                  " U5      U4PM     sn5      nU$ s  snf )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   ©Úcuda)Ú	threadIdxÚblockDimÚblockIdxÚgridDimÚlaneidÚwarpsize)ÚnumbarŠ   Útupler	   ÚModule)r'   rŠ   Ú	nonconstsÚncÚnonconsts_with_mods        r(   Únonconst_module_attrsÚ'CUDATargetContext.nonconst_module_attrs   sM   € õ 	ð!ˆ	ä"Ù.7ó$9Ú.7¨ô &+§\¢\°$Ó%7¸Ó$<Ù.7ñ$9ó :Ðà!Ð!ùò$9s   ’"=c                 ó   • [        U 5      $ rM   )ÚCUDACallConvrW   s    r(   Ú	call_convÚCUDATargetContext.call_conv�   s   € ä˜DÓ!Ð!r+   rA   N©Úabi_tagsÚuidc                ó.   • [         R                  " XUUS9$ )Nr�   )r   Úmangle)r'   r_   Úargtypesrž   rŸ   s        r(   ÚmanglerÚCUDATargetContext.mangler‘   s   € Ü×%Ò% d¸xØ*-ñ/ð 	/r+   c	           	      óì   • [         R                  " UR                  SS9n	U R                  5       R	                  UR
                   S3U	UUS9n
U
R                  U5        U R                  X¢U	X4UU5      nX«4$ )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_libraryr_   Úadd_linking_libraryÚgenerate_kernel_wrapper)r'   ÚcodelibÚfndescr/   Úlineinfor«   ÚfilenameÚlinenumr¬   Úkernel_nameÚlibraryÚwrappers               r(   Úprepare_cuda_kernelÚ%CUDATargetContext.prepare_cuda_kernel•   sŠ   € ô. &×7Ò7Ø×!Ñ! hñ
ˆð —,‘,“.×/Ñ/°7·<±<°.ÀÐ0IØ;FØ=IØ>Kð 0ð Mˆð 	×#Ñ# GÔ,Ø×.Ñ.¨wÀØ/4ÀØ/6ó8ˆð ÐÐr+   c                 ó"	  ^!^"• UR                   nU R                  U5      n	[        U	R                  5      n
[        R
                  " [        R                  " 5       U
5      nU R                  S5      m"[        R
                  " [        R                  " S5      U R                  R                  [        R                  5      /U
-   5      n[        R                  " T"XÂR                  5      n[        R                   " UR"                  SS9n[        R                  " T"X¾5      m![        R$                  " T!R'                  S5      5      nU(       d  U(       aR  U=(       a    U(       + nU R)                  T"UU US9nUR*                  " T!X2R,                  X‡5        UR.                  " X÷5        U!U"4S jnU" S5      n/ n/ nS	 H7  nUR1                  U" S
U-  5      5        UR1                  U" SU-  5      5        M9     U	R3                  UT!R,                  5      nU R                  R5                  Xý[        R6                  UU5      u  nnU(       Gak  [8        R:                  " UUR<                  5         UR?                  5         SSS5        URA                  URC                  URD                  5      5         [        RF                  " URH                  RJ                  S5      nURM                  UUURN                  SS5      nURQ                  US5      n[R        RT                  " U5      nURA                  U5         [W        S	U5       H)  u  nnURY                  U5      n UR[                  U U5        M+     [W        S	U5       H)  u  nnUR]                  U5      n UR[                  U U5        M+     SSS5        SSS5        UR?                  5         [^        R`                  " T!5        URc                  T"5        U(       d  U(       a  WRd                  " 5         URe                  5         [f        Rh                  (       a  [j        Rl                  " UT"5        URo                  T!R"                  5      $ ! , (       d  f       GNí= f! , (       d  f       NÕ= f! , (       d  f       NÞ= f)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                 óÞ   >• TR                   U -   n[        R                  " T[        R                  " S5      U5      n[        R
                  " UR                  R                  S 5      Ul        U$ )Nr½   )	r_   r   Úadd_global_variabler   ÚIntTypeÚConstantÚtypeÚpointeeÚinitializer)Úpostfixr_   ÚgvÚwrapfnÚwrapper_modules      €€r(   Údefine_error_gvÚBCUDATargetContext.generate_kernel_wrapper.<locals>.define_error_gvÛ   sO   ø€ Ø—;‘; Ñ(ˆDÜ×,Ò,¨^¼R¿ZºZÈ»^Ø-1ó3ˆBäŸ[š[¨¯©¯©¸$Ó?ˆBŒNØˆIr+   Ú__errcode__Úxyzz	__tid%s__z__ctaid%s__NÚ	monotonicr   )8r¢   Úget_arg_packerÚlistÚargument_typesr   ÚFunctionTypeÚVoidTyper`   rÄ   r›   Úget_return_typer	   ÚpyobjectÚFunctionr®   r   r­   r_   Ú	IRBuilderÚappend_basic_blockrV   Ú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±   Ú)CUDATargetContext.generate_kernel_wrapper¹   s{  ù€ ð —?‘?ˆØ×%Ñ% hÓ/ˆÜ�g×,Ñ,Ó-ˆÜ—?’?¤2§;¢;£=°&Ó9ˆØ×+Ñ+Ð,AÓBˆÜ�ŠœrŸzšz¨"›~Ø $§¡× >Ñ >¼u¿~¹~Ó NÐOØ!'ñ (ó)ˆô �{Š{˜>¨4×1FÑ1FÓGˆä"×4Ò4°T·Y±YÀ8ÑLˆÜ—’˜^¨XÓ@ˆÜ—,’,˜v×8Ñ8¸Ó<Ó=ˆæ–HØ&×4¨u¬9ˆOØŸ™¨nØ08Ø-1Ø7Fð 'ð HˆIð ×%Ò%Ø˜§[¡[°(ôð ×#Ò# GÔ5ö	ñ ! Ó/ˆØˆØˆÛˆAØ�M‰M™/¨+¸©/Ó:Ô;Ø�O‰O™O¨M¸AÑ,=Ó>Ö?ñ ð ×)Ñ)¨'°6·;±;Ó?ˆØ—N‘N×0Ñ0Øœ5Ÿ:™: x°ó;‰	ˆ�÷ ä×"Ò" 7¨F¯L©LÕ9Ø× Ñ Ô"÷ :ð —‘ §¡¨f×.BÑ.BÓ!CÕDä—k’k &§+¡+×"5Ñ"5°tÓ<�ð
 —‘ v¨s°F·K±KØ'2°KóA�à!×/Ñ/°°aÓ8�ô !×,Ò,¨WÓ5�Ø—_‘_ WÕ-Ü%(¨°Ö%7™	˜˜SØ"Ÿh™h s›m˜ØŸ™ c¨3Ö/ñ &8ô &)¨°Ö%9™	˜˜SØ"Ÿj™j¨›o˜ØŸ™ c¨3Ö/ñ &:÷ .÷ Eð, 	×ÑÔä×Ò˜VÔ$Ø×Ñ˜nÔ-Þ–HØ×ÒÔ Ø×ÑÔä××Ü�OŠO˜F NÔ3à×#Ñ# F§K¡KÓ0Ð0÷I :Ö9ú÷  .Õ-ú÷ EÕDús2   É&QÊ)BR Ì-A3Q/Î R Ñ
Q,Ñ/
Q=	Ñ9R Ò 
Rc           
      ó²  • UR                   n[        UR                  SS95       Vs/ s H#  nU R                  [        R
                  U5      PM%     nn[        R                  " [        R                  " S5      [        U5      5      n[        R                  " Xv5      n[        R                  n	[        R                  " XHR                  SU	S9n
SU
l        SU
l        XŠl        U R'                  UR(                  5      nU R+                  U5      nSUS	-
  R-                  5       -  U
l        [        R0                  " [        R                  " S5      5      nUR3                  X­S
5      nU R5                  U5      " X5      nUR6                   Vs/ s H#  nU R                  [        R8                  U5      PM%     nnUR:                   Vs/ s H#  nU R                  [        R8                  U5      PM%     nnU R=                  XñR?                  XïR@                  R                  5      UUURB                  URD                  SS9  URG                  5       $ s  snf s  snf s  snf )úQ
Unlike the parent version.  This returns a a pointer in the constant
addrspace.
ÚA)Úorderé   Ú_cudapy_cmem©Ú	addrspaceÚinternalTé   r   ÚgenericN)ÚdataÚshapeÚstridesÚitemsizeÚparentÚmeminfo)$r¾   ÚiterÚtobytesÚget_constantr	   Úbyter   Ú	ArrayTyperÄ   ÚlenrÅ   r   ÚADDRSPACE_CONSTANTr   rÃ   rÆ   ÚlinkageÚglobal_constantrÈ   Úget_data_typeÚdtypeÚget_abi_sizeofÚ
bit_lengthÚalignÚPointerTypeÚaddrspacecastÚ
make_arrayr  Úintpr  Úpopulate_arrayÚbitcastr  r  r  Ú	_getvalue)r'   rý   ÚarytyÚarrÚlmodr  Ú	constvalsÚ
constarytyÚconstaryr  rÊ   Úlldtyper*  ÚptrtyÚgenptrÚaryÚsÚkshapeÚkstridess                      r(   Úmake_constant_arrayÚ%CUDATargetContext.make_constant_array  sß  € ð �~‰~ˆô ˜#Ÿ+™+¨C˜+Ð0Ô1ó
â1�ð ×ÑœeŸj™j¨!Ö,Ù1ð 	ð 
ô —\’\¤"§*¢*¨Q£-´°Y³Ó@ˆ
Ü—;’;˜zÓ5ˆä×+Ñ+ˆ	Ü×(Ò(¨¯}©}¸nØ3<ñ>ˆàˆŒ
Ø!ˆÔØ!Œð ×$Ñ$ U§[¡[Ó1ˆØ×#Ñ# GÓ,ˆØ˜ ™×.Ñ.Ó0Ñ0ˆŒô —’œrŸzšz¨!›}Ó-ˆØ×&Ñ& r°)Ó<ˆð �o‰o˜eÔ$ TÓ3ˆØ<?¿IºIÓFºI°q�$×#Ñ#¤E§J¡J°Ö2¹IˆÐFØ>A¿kºkÓJºk¸�D×%Ñ%¤e§j¡j°!Ö4¹kˆÐJØ×Ñ˜C§o¡o°f¿h¹h¿m¹mÓ&LØ"(Ø$,Ø%(§\¡\¸#¿*¹*Ø$(ð	 	ñ 	*ð �}‰}‹ÐùòC
ùò2 GùÚJs   §*I
Æ	*IÇ*Ic                 óü  • [         R                  " UR                  S5      S-   5      nSR                  S[        R
                  " U5      /5      nUR                  R                  U5      nUcB  [         R                  " XR                  U[        R                  S9nSUl        SUl        X5l        UR                  R                  R                   nUR#                  UR%                  [        R                  5      5      $ )r  zutf-8ó    Ú$Ú__conststring__r  r  T)r   Úmake_bytearrayÚencodeÚjoinr   Úmangle_identifierÚglobalsr9   rÃ   rÆ   r   r#  r$  r%  rÈ   rÇ   Úelementr0  Ú
as_pointer)r'   ÚmodÚstringÚtextr_   rÊ   Úchartys          r(   Úinsert_const_stringÚ%CUDATargetContext.insert_const_string@  sÌ   € ô
 ×%Ò% f§m¡m°GÓ&<¸wÑ&FÓGˆØ�x‰xÐ*Ü(×:Ò:¸6ÓBðDó Eˆð �[‰[�_‰_˜TÓ"ˆØ‰:ä×,Ò,¨S·)±)¸TÜ7;×7NÑ7NñPˆBà#ˆBŒJØ!%ˆBÔØ!ŒNð —‘—‘×(Ñ(ˆØ�z‰z˜&×+Ñ+¬D×,CÑ,CÓDÓEÐEr+   c                 ó´   • UR                   nU R                  X25      n[        R                  " [        R                  " S5      5      nUR                  XES5      $ )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¾   rP  r   r+  rÄ   r,  )r'   rý   rM  r4  rÊ   Ú	charptrtys         r(   Úinsert_string_const_addrspaceÚ/CUDATargetContext.insert_string_const_addrspaceV  sF   € ð �~‰~ˆØ×%Ñ% dÓ3ˆÜ—N’N¤2§:¢:¨a£=Ó1ˆ	Ø×$Ñ$ R°IÓ>Ð>r+   c                 ó   • g)zRun O1 function passes
        NrA   ©r'   rû   s     r(   Úoptimize_functionÚ#CUDATargetContext.optimize_functionb  s   € ð 	r+   c                 ó.   • [         R                  " U5      $ rM   )r   Úget_ufunc_info)r'   Ú	ufunc_keys     r(   r[  Ú CUDATargetContext.get_ufunc_infoo  s   € Ü×$Ò$ YÓ/Ð/r+   )r]   rd   rQ   r‰   rM   )rB   rC   rD   rE   Úimplement_powi_as_math_callÚstrict_alignmentrN   ÚpropertyrV   rZ   r`   re   r)   r   r†   r   r—   r›   r£   rº   r±   r?  rP  rT  rX  r[  rF   rG   rH   s   @r(   rJ   rJ   C   sË   ø† Ø"&ÐØÐ÷
ð ñ#ó ð#ð ñó ðò
Aò!ò:ò0&ð ñ!ó ð!ð
 ñ"ó ð"ð ñ"ó ð"ð 35¸$õ /ð +/ô" òHZ1òx)òVFò,
?ò÷0ð 0r+   rJ   c                   ó   • \ rS rSrSrg)rš   is  rA   N)rB   rC   rD   rE   rF   rA   r+   r(   rš   rš   s  s   † Úr+   rš   c                   óZ   • \ rS rSrSrS rS r  SS jrS rS r	SS	 jr
S
 rS rS rSrg)ÚCUDACABICallConviw  zÀ
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                 ó   • g rM   rA   )r'   rý   s     r(   Ú_make_call_helperÚ"CUDACABICallConv._make_call_helper�  s   € ð r+   c                 ó$   • UR                  U5      $ rM   )Úret)r'   rý   Úretvals      r(   Úreturn_valueÚCUDACABICallConv.return_value†  s   € Ø�{‰{˜6Ó"Ð"r+   Nc                 ó   • Sn[        U5      e)Nz7Python exceptions are unsupported in the CUDA C/C++ ABI©ÚNotImplementedError)r'   rý   ÚexcÚexc_argsÚlocÚ	func_nameÚmsgs          r(   Úreturn_user_excÚ CUDACABICallConv.return_user_exc‰  s   € àGˆÜ! #Ó&Ð&r+   c                 ó   • Sn[        U5      e)Nz2Return status is unsupported in the CUDA C/C++ ABIrm  )r'   rý   r  rs  s       r(   Úreturn_status_propagateÚ(CUDACABICallConv.return_status_propagateŽ  s   € ØBˆÜ! #Ó&Ð&r+   c                 óž   • U R                  U5      n[        UR                  5      n[        R                  " U R                  U5      U5      nU$ )z=
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_typeÚ"CUDACABICallConv.get_function_type’  sD   € ð ×&Ñ& xÓ0ˆÜ˜×.Ñ.Ó/ˆÜ�Š˜t×3Ñ3°GÓ<¸hÓGˆØˆr+   c                 ó¬   • U(       a   eU R                  U5      nUR                  U R                  U5      U Vs/ s H  nSU-   PM
     sn5        gs  snf )z1
Set names and attributes of function arguments.
zarg.N)rz  Úassign_namesÚget_arguments)r'   ÚfnrÝ   Úfe_argtypesÚnoaliasr÷   Úas          r(   Údecorate_functionÚ"CUDACABICallConv.decorate_function›  sP   € ö Ðˆ{Ø×&Ñ& {Ó3ˆØ×Ñ˜T×/Ñ/°Ó3Ù26Ó7²$¨Q˜f qœj±$Ñ7õ	9ùÚ7s   ºA
c                 ó   • UR                   $ )z0
Get the Python-level arguments of LLVM *func*.
)rÝ   rW  s     r(   r€  ÚCUDACABICallConv.get_arguments¤  s   € ð �y‰yÐr+   c                 óª   • U R                  U5      nUR                  X5      nUR                  X'5      nSn	U R                  R	                  XU5      n
Xš4$ )z#
Call the Numba-compiled *callee*.
N)rz  Úas_argumentsÚcallÚcontextÚget_returned_value)r'   rý   ÚcalleeÚrestyrø   rÝ   r÷   Úrealargsrê   r  Úouts              r(   rá   ÚCUDACABICallConv.call_functionª  sW   € ð ×&Ñ& vÓ.ˆØ×'Ñ'¨Ó6ˆØ�|‰|˜FÓ-ˆð ˆØ�l‰l×-Ñ-¨g¸dÓCˆØˆ{Ðr+   c                 óP   • U R                   R                  U   R                  5       $ rM   )rŒ  rQ   r×   )r'   Útys     r(   r×   Ú CUDACABICallConv.get_return_type·  s    € Ø�|‰|×.Ñ.¨rÑ2×BÑBÓDÐDr+   rA   )NNN)F)rB   rC   rD   rE   Ú__doc__re  rj  rt  rw  r|  r…  r€  rá   r×   rF   rA   r+   r(   rc  rc  w  s?   † ñòò
#ð @DØ"&ô'ò
'òô9òòõEr+   rc  ))ÚreÚ	functoolsr   Úllvmlite.bindingÚbindingr‚   Ú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_CHARSrJ   rš   rc  rA   r+   r(   Ú<module>r¦     s‘   ðÛ 	Ý %Ý Ý ÷'÷ 'ñ 'å ,Ý 'ß =Ý 'Ý  å ß 1Ñ 1Ý /ô$F˜×*Ñ*ô $FðT �jŠj˜ r§t¡tÓ,€ôm0˜ô m0ô`		�?ô 	ôAE�|õ AEr+   