Ë
    èÿæiÅA  ã                   óf  — 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 jF                  de jH                  «      Z% G d„ de«      Z& G d„ de«      Z' G d„ de«      Z(y)é    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                  |j                  «       | j                  |j                  «       | j                  t        j                  «       | j                  |j                  «       | j                  |j                  «       | j                  |j                  «       y )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          úf/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/numba/cuda/target.pyÚload_additional_registriesz,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                 óž  •— ddl m} t        |t        «      rt        ||«      s	 |j                  }t        t        | �;  |«      S # t
        $ r |j                  st        d«      ‚|j                  j                  «       }d|d<   |j                  dd«      |d<   |j                  dd«      |d<    ||j                  |«      }||_        |}Y Œšw xY w)	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Ü�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__r(   r:   Ú__classcell__©r=   s   @r'   r   r      s   ø„ ò
<÷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 )ÚCUDATargetContextTc                 óv   •— t         ‰| �  ||«       t        j                  t        j
                  «      | _        y ©N)r9   Ú__init__r   Úchainr   Údefault_managerÚdata_model_manager)r&   Ú	typingctxÚtargetr=   s      €r'   rG   zCUDATargetContext.__init__G   s.   ø€ Ü‰Ñ˜ FÔ+Ü"3×"9Ñ"9Ü×%Ñ%ó#
ˆÕr)   c                 ó"   — t         j                  S rF   )r   Ú	DIBuilder©r&   s    r'   rN   zCUDATargetContext.DIBuilderM   s   € ä×"Ñ"Ð"r)   c                  ó   — y)NF© rO   s    r'   Úenable_boundscheckz$CUDATargetContext.enable_boundscheckQ   s   € ð r)   c                 ó8   — | j                   j                  |«      S rF   )Ú_internal_codegenÚ_create_empty_module)r&   Únames     r'   Úcreate_modulezCUDATargetContext.create_moduleW   s   € Ø×%Ñ%×:Ñ:¸4Ó@Ð@r)   c                 óF   — t        j                  d«      | _        d | _        y )Nznumba.cuda.jit)r   ÚJITCUDACodegenrT   Ú_target_datarO   s    r'   ÚinitzCUDATargetContext.initZ   s   € Ü!(×!7Ñ!7Ð8HÓ!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/                  |j0                  «       | j/                  |
j0                  «       | j/                  |j0                  «       | j/                  |j0                  «       | j/                  |	j0                  «       | j/                  |j0                  «       | j/                  |j2                  «       y )Nr   )ÚnumbersÚtupleobjÚslicing)ÚrangeobjÚ	iteratorsÚenumimpl)ÚunicodeÚcharseq)Ú	cmathimpl)Úcffiimpl)Úarrayobj)Ú
npdatetimer   )ÚcudaimplÚ	printimplÚlibdeviceimplÚmathimplr   )Úndarray)Únumba.cpythonr]   r^   r_   r`   ra   rb   rc   rd   re   Ú
numba.miscrf   Únumba.nprg   rh   r!   ri   rj   rk   rl   r   Únumba.np.unsaferm   r#   r$   Úimpl_registry)r&   r]   r^   r_   r`   ra   rb   rc   rd   re   rf   rg   rh   ri   rj   rk   rl   r   rm   s                      r'   r(   z,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                 ó   — | j                   S rF   )rT   rO   s    r'   r   zCUDATargetContext.codegenv   s   € Ø×%Ñ%Ð%r)   c                 óž   — | j                   €6t        j                  t        j                  «       j
                  «      | _         | j                   S rF   )rZ   ÚllÚcreate_target_datar   ÚNVVMÚdata_layoutrO   s    r'   Útarget_datazCUDATargetContext.target_datay   s9   € à×ÑÐ$Ü "× 5Ñ 5´d·i±i³k×6MÑ6MÓ NˆDÔØ× Ñ Ð r)   c                 ó|   — ddl m} d}t        |D �cg c]  }t        j                  |«      |f‘Œ c}«      }|S c c}w )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_attrsz'CUDATargetContext.nonconst_module_attrs   sM   € õ 	ð!ˆ	ä"Ù.7ó$9Ù.7¨ô &+§\¡\°$Ó%7¸Ò$<Ø.7ñ$9ó :Ðà!Ð!ùò$9s   ’9c                 ó   — t        | «      S rF   )ÚCUDACallConvrO   s    r'   Ú	call_convzCUDATargetContext.call_conv�   s   € ä˜DÓ!Ð!r)   rQ   N©Úabi_tagsÚuidc                ó4   — t        j                  ||||¬«      S )Nr�   )r   Úmangle)r&   rV   ÚargtypesrŽ   r�   s        r'   ÚmanglerzCUDATargetContext.mangler‘   s   € Ü×%Ñ% d¨H¸xØ*-ô/ð 	/r)   c	           	      óü   — t        j                  |j                  d¬«      }	| j                  «       j	                  |j
                  › d�|	||¬«      }
|
j                  |«       | j                  |
||	||||«      }|
|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_libraryrV   Ú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Ñ7Ø×!Ñ! hô
ˆð —,‘,“.×/Ñ/°7·<±<°.ÀÐ0IØ;FØ=IØ>Kð 0ó Mˆð 	×#Ñ# GÔ,Ø×.Ñ.¨w¸ÀØ/4°hÀØ/6ó8ˆð ˜ÐÐr)   c                 ó¨  ‡!‡"— |j                   }| j                  |«      }	t        |	j                  «      }
t	        j
                  t	        j                  «       |
«      }| j                  d«      Š"t	        j
                  t	        j                  d«      | j                  j                  t        j                  «      g|
z   «      }t	        j                  ‰"||j                  «      }t        j                   |j"                  d¬«      }t	        j                  ‰"||«      Š!t	        j$                  ‰!j'                  d«      «      }|s|rO|xr | }| j)                  ‰"|| |¬«      } |j*                  ‰!||j,                  ||«        |j.                  ||«       ˆ!ˆ"fd„} |d«      }g }g }d	D ]6  }|j1                   |d
|z  «      «       |j1                   |d|z  «      «       Œ8 |	j3                  |‰!j,                  «      }| j                  j5                  ||t        j6                  ||«      \  }}|�ret9        j:                  ||j<                  «      5  |j?                  «        ddd«       |jA                  |jC                  |jD                  «      «      5  t	        jF                  |jH                  jJ                  d«      }|jM                  |||jN                  dd«      }|jQ                  |d«      }tS        jT                  |«      }|jA                  |«      5  tW        d	|«      D ](  \  }}|jY                  |«      } |j[                  | |«       Œ* tW        d	|«      D ](  \  }}|j]                  |«      } |j[                  | |«       Œ* 	 ddd«       ddd«       |j?                  «        t_        j`                  ‰!«       |jc                  ‰"«       |s|r jd                  «        |je                  «        tf        jh                  rtk        jl                  |‰"«       |jo                  ‰!j"                  «      S # 1 sw Y   �ŒÒxY w# 1 sw Y   ŒºxY w# 1 sw Y   Œ¾xY w)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«   )	rV   r   Úadd_global_variabler   ÚIntTypeÚConstantÚtypeÚpointeeÚinitializer)ÚpostfixrV   ÚgvÚwrapfnÚwrapper_modules      €€r'   Údefine_error_gvzBCUDATargetContext.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ÚVoidTyperW   r²   rŒ   Úget_return_typer	   ÚpyobjectÚFunctionr�   r   rœ   rV   Ú	IRBuilderÚappend_basic_blockrN   Ú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ˆØ×+Ñ+Ð,AÓBˆÜ�‰œrŸz™z¨"›~Ø $§¡× >Ñ >¼u¿~¹~Ó NÐOØ!'ñ (ó)ˆô �{‰{˜>¨4°×1FÑ1FÓGˆä"×4Ñ4°T·Y±YÀ8ÔLˆÜ—‘˜^¨X°xÓ@ˆÜ—,‘,˜v×8Ñ8¸Ó<Ó=ˆá‘HØ&Ò4¨u¨9ˆOØŸ™¨nØ08Ø-1Ø7Fð 'ó HˆIð &ˆI×%Ñ%Ø˜ V§[¡[°(¸Gôð $ˆI×#Ñ# G¨WÔ5õ	ñ ! Ó/ˆØˆØˆÛˆAØ�M‰M™/¨+¸©/Ó:Ô;Ø�O‰O™O¨M¸AÑ,=Ó>Õ?ð ð ×)Ñ)¨'°6·;±;Ó?ˆØ—N‘N×0Ñ0Ø�Tœ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ØˆI×ÑÔ Ø×ÑÔä×ÒÜ�O‰O˜F NÔ3à×#Ñ# F§K¡KÓ0Ð0÷I :Ñ9ú÷  .Ð-ú÷ EÐDús2   ÉP/ÊBQÌA/P<ÎQÐ/P9Ð<Q	ÑQÑQc           	      ó´  — |j                   }t        |j                  d¬«      «      D �cg c]"  }| j                  t        j
                  |«      ‘Œ$ }}t        j                  t        j                  d«      t        |«      «      }t        j                  ||«      }t        j                  }	t        j                  ||j                  d|	¬«      }
d|
_        d|
_        ||
_        | j'                  |j(                  «      }| j+                  |«      }d|d	z
  j-                  «       z  |
_        t        j0                  t        j                  d«      «      }|j3                  |
|d
«      } | j5                  |«      | |«      }|j6                  D �cg c]"  }| j                  t        j8                  |«      ‘Œ$ }}|j:                  D �cg c]"  }| j                  t        j8                  |«      ‘Œ$ }}| j=                  ||j?                  ||j@                  j                  «      |||jB                  |jD                  d¬«       |jG                  «       S c c}w c c}w c c}w )úi
        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_arrayz%CUDATargetContext.make_constant_array  sï  € ð �~‰~ˆô ˜#Ÿ+™+¨C˜+Ó0Ô1ó
á1�ð ×ÑœeŸj™j¨!Õ,Ø1ð 	ð 
ô —\‘\¤"§*¡*¨Q£-´°Y³Ó@ˆ
Ü—;‘;˜z¨9Ó5ˆä×+Ñ+ˆ	Ü×(Ñ(¨¨x¯}©}¸nØ3<ô>ˆàˆŒ
Ø!ˆÔØ!ˆŒð ×$Ñ$ U§[¡[Ó1ˆØ×#Ñ# GÓ,ˆØ˜ ™×.Ñ.Ó0Ñ0ˆŒô —‘œrŸz™z¨!›}Ó-ˆØ×&Ñ& r¨5°)Ó<ˆð %ˆd�o‰o˜eÓ$ T¨7Ó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 g§o¡o°f¸c¿h¹h¿m¹mÓ&LØ"(Ø$,Ø%(§\¡\¸#¿*¹*Ø$(ð	 	ô 	*ð �}‰}‹ÐùòC
ùò2 GùÚJs   ª'IÆ'IÇ'Ic                 óþ  — t        j                  |j                  d«      dz   «      }dj                  dt	        j
                  |«      g«      }|j                  j                  |«      }|€Ft        j                  ||j                  |t        j                  ¬«      }d|_        d|_        ||_        |j                  j                  j                   }|j#                  |j%                  t        j                  «      «      S )rù   zutf-8ó    Ú$Ú__conststring__rþ   r   T)r   Úmake_bytearrayÚencodeÚjoinr   Úmangle_identifierÚglobalsr7   r±   r´   r   r  r  r  r¶   rµ   Úelementr  Ú
as_pointer)r&   ÚmodÚstringÚtextrV   r¸   Úchartys          r'   Úinsert_const_stringz%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ÔØ!ˆBŒNð —‘—‘×(Ñ(ˆØ�z‰z˜&×+Ñ+¬D×,CÑ,CÓDÓEÐEr)   c                 ó´   — |j                   }| j                  ||«      }t        j                  t        j                  d«      «      }|j                  ||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ê   r8  r   r¸   Ú	charptrtys         r'   Úinsert_string_const_addrspacez/CUDATargetContext.insert_string_const_addrspaceV  sJ   € ð �~‰~ˆØ×%Ñ% d¨FÓ3ˆÜ—N‘N¤2§:¡:¨a£=Ó1ˆ	Ø×$Ñ$ R¨°IÓ>Ð>r)   c                  ó   — y)zRun O1 function passes
        NrQ   ©r&   rè   s     r'   Úoptimize_functionz#CUDATargetContext.optimize_functionb  s   € ð 	r)   c                 ó,   — t        j                  |«      S rF   )r   Úget_ufunc_info)r&   Ú	ufunc_keys     r'   rC  z CUDATargetContext.get_ufunc_infoo  s   € Ü×$Ñ$ YÓ/Ð/r)   r{   rF   )r>   r?   r@   Úimplement_powi_as_math_callÚstrict_alignmentrG   ÚpropertyrN   rR   rW   r[   r(   r   ry   r   r‰   rŒ   r“   r©   r    r+  r;  r>  rA  rC  rA   rB   s   @r'   rD   rD   C   sÆ   ø„ Ø"&ÐØÐõ
ð ñ#ó ð#ð ñó ðò
Aò!ò:ò0&ð ñ!ó ð!ð
 ñ"ó ð"ð ñ"ó ð"ð 35¸$ô /ð +/ó" òHZ1òx)òVFò,
?òö0r)   rD   c                   ó   — e Zd Zy)r‹   N)r>   r?   r@   rQ   r)   r'   r‹   r‹   s  s   „ Ør)   r‹   c                   óN   — 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y)Ú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                  ó   — y rF   rQ   )r&   rê   s     r'   Ú_make_call_helperz"CUDACABICallConv._make_call_helper�  s   € ð r)   c                 ó$   — |j                  |«      S rF   )Ú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++ ABIrR  )r&   rê   rð   rX  s       r'   Úreturn_status_propagatez(CUDACABICallConv.return_status_propagateŽ  s   € ØBˆÜ! #Ó&Ð&r)   c                 óœ   — | j                  |«      }t        |j                  «      }t        j                  | 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’  sD   € ð ×&Ñ& xÓ0ˆÜ˜×.Ñ.Ó/ˆÜ�‰˜t×3Ñ3°GÓ<¸hÓGˆØˆr)   c                 óœ   — |rJ ‚| j                  |«      }|j                  | j                  |«      |D �cg c]  }d|z   ‘Œ	 c}«       yc c}w )zA
        Set names and attributes of function arguments.
        zarg.N)r]  Úassign_namesÚget_arguments)r&   ÚfnrÊ   Úfe_argtypesÚnoaliasrä   Úas          r'   Údecorate_functionz"CUDACABICallConv.decorate_function›  sP   € ñ Ðˆ{Ø×&Ñ& {Ó3ˆØ×Ñ˜T×/Ñ/°Ó3Ù26Ó7±$¨Q˜f q›j°$Ñ7õ	9ùÚ7s   µA	
c                 ó   — |j                   S )z@
        Get the Python-level arguments of LLVM *func*.
        )rÊ   r@  s     r'   rb  zCUDACABICallConv.get_arguments¤  s   € ð �y‰yÐr)   c                 ó²   — | j                  |«      }|j                  ||«      }|j                  ||«      }d}	| j                  j	                  |||«      }
|	|
fS )z3
        Call the Numba-compiled *callee*.
        N)r]  Úas_argumentsÚcallÚcontextÚget_returned_value)r&   rê   ÚcalleeÚrestyrå   rÊ   rä   Úrealargsr×   rð   Úouts              r'   rÎ   zCUDACABICallConv.call_functionª  s_   € ð ×&Ñ& vÓ.ˆØ×'Ñ'¨°Ó6ˆØ�|‰|˜F HÓ-ˆð ˆØ�l‰l×-Ñ-¨g°u¸dÓCˆØ�sˆ{Ðr)   c                 óP   — | j                   j                  |   j                  «       S rF   )rl  rJ   rÄ   )r&   Útys     r'   rÄ   z CUDACABICallConv.get_return_type·  s    € Ø�|‰|×.Ñ.¨rÑ2×BÑBÓDÐDr)   )NNN)F)r>   r?   r@   Ú__doc__rL  rP  rY  r[  r_  rg  rb  rÎ   rÄ   rQ   r)   r'   rJ  rJ  w  s?   „ ñòò
#ð @DØ"&ó'ò
'òó9òòóEr)   rJ  ))ÚreÚ	functoolsr   Úllvmlite.bindingÚbindingru   Ú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‹   rJ  rQ   r)   r'   Ú<module>r„     s“   ðÛ 	Ý %Ý Ý ÷'÷ 'ñ 'å ,Ý 'ß =Ý 'Ý  å ß 1Ñ 1Ý /ô$F˜×*Ñ*ô $FðT ˆb�j‰j˜ r§t¡tÓ,€ôm0˜ô m0ô`		�?ô 	ôAE�|õ AEr)   