§
    ƒÛgOY  ã                   óz  — d Z ddlZddlZddlZddlZddlmZmZmZm	Z	m
Z
mZmZ ddlZddlmZ ddlmZmZmZ ddlmZmZmZ ddlmZmZ  ej        e¦  «        ZdZdZd	Z d
Z!dZ"eZ#eZ$d %                    ¦   «         Z& e'e&¦  «        D ]\  Z(Z) e*ej+        e         e)e(¦  «         ŒdZ,dZ-d„ Z. ej/        ¦   «         Z0 G d„ de1¦  «        Z2 G d„ de1¦  «        Z3dZ4dddddddddddœ
Z5d„ Z6d„ Z7d„ Z8d„ Z9dZ: G d„ d e1¦  «        Z;d!Z<d"Z=d#Z>d$Z?d%Z@d&„ ZAd'„ ZBd(„ ZCd)„ ZDd*„ ZEd+„ ZFd,„ ZG ejH        d-¦  «        ZId.„ ZJd/„ ZKd0„ ZLdS )1z(
This is a direct translation of nvvm.h
é    N)Úc_void_pÚc_intÚPOINTERÚc_char_pÚc_size_tÚbyrefÚc_char)Úiré   )Ú	NvvmErrorÚNvvmSupportErrorÚNvvmWarning)Úget_libdeviceÚopen_libdeviceÚopen_cudalib)ÚcgutilsÚconfigé   é   é   a  
NVVM_SUCCESS
NVVM_ERROR_OUT_OF_MEMORY
NVVM_ERROR_PROGRAM_CREATION_FAILURE
NVVM_ERROR_IR_VERSION_MISMATCH
NVVM_ERROR_INVALID_INPUT
NVVM_ERROR_INVALID_PROGRAM
NVVM_ERROR_INVALID_IR
NVVM_ERROR_INVALID_OPTION
NVVM_ERROR_NO_MODULE_IN_PROGRAM
NVVM_ERROR_COMPILATION
z�e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64zŽe-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64c                  óF   — 	 t          ¦   «          dS # t          $ r Y dS w xY w)z(
    Return if libNVVM is available
    TF)ÚNVVMr   © ó    úM/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/cudadrv/nvvm.pyÚis_availabler   <   s;   € ðÝ‰Œˆð ˆtøõ ð ð ð Øˆuˆuðøøøs   ‚ ’
 Ÿ c                   ó¬  — e Zd ZdZe ee¦  «         ee¦  «        fe ee¦  «        fe ee¦  «        feeee	efeeee	efeee ee¦  «        fee ee	¦  «        feeefee ee	¦  «        feeefe ee¦  «         ee¦  «         ee¦  «         ee¦  «        feee ee¦  «        fdœZ
dZd„ Zd„ Zed„ ¦   «         Zed„ ¦   «         Zd„ Zd	„ Zdd„ZdS )r   zProcess-wide singleton.
    )ÚnvvmVersionÚnvvmCreateProgramÚnvvmDestroyProgramÚnvvmAddModuleToProgramÚnvvmLazyAddModuleToProgramÚnvvmCompileProgramÚnvvmGetCompiledResultSizeÚnvvmGetCompiledResultÚnvvmGetProgramLogSizeÚnvvmGetProgramLogÚnvvmIRVersionÚnvvmVerifyProgramNc                 óæ  — t           5  | j        €Åt                               | ¦  «        x| _        }	 t	          d¦  «        |_        n-# t          $ r }d | _        d}t          ||z  ¦  «        ‚d }~ww xY w|j         	                    ¦   «         D ]G\  }}t          |j        |¦  «        }|d         |_        |dd …         |_        t          |||¦  «         ŒHd d d ¦  «         n# 1 swxY w Y   | j        S )NÚnvvmz;libNVVM cannot be found. Do `conda install cudatoolkit`:
%sr   r   )Ú
_nvvm_lockÚ_NVVM__INSTANCEÚobjectÚ__new__r   ÚdriverÚOSErrorr   Ú_PROTOTYPESÚitemsÚgetattrÚrestypeÚargtypesÚsetattr)ÚclsÚinstÚeÚerrmsgÚnameÚprotoÚfuncs          r   r/   zNVVM.__new__†   s@  € Ýð 	.ð 	.ØŒ~Ð%Ý(.¯ª°sÑ(;Ô(;Ð;�” ð7Ý".¨vÑ"6Ô"6�D”K�KøÝð 7ð 7ð 7Ø%)�C”Nð2�Få*¨6°A©:Ñ6Ô6Ð6øøøøð	7øøøð $(Ô#3×#9Ò#9Ñ#;Ô#;ð .ð .‘K�D˜%Ý" 4¤;°Ñ5Ô5�DØ#(¨¤8�D”LØ$)¨!¨"¨"¤I�D”MÝ˜D $¨Ñ-Ô-Ð-Ð-ð!	.ð 	.ð 	.ñ 	.ô 	.ð 	.ð 	.ð 	.ð 	.ð 	.ð 	.øøøð 	.ð 	.ð 	.ð 	.ð$ Œ~Ðs5   ˆ)C!²AÁC!Á
A1ÁA,Á,A1Á1A$C!Ã!C%Ã(C%c                 ó¼   — |                       ¦   «         }|d         | _        |d         | _        |d         | _        |d         | _        t          ¦   «         | _        d S )Nr   r   é   r   )Úget_ir_versionÚ_majorIRÚ_minorIRÚ	_majorDbgÚ	_minorDbgÚget_supported_ccsÚ_supported_ccs)ÚselfÚir_versionss     r   Ú__init__zNVVM.__init__›   sS   € Ø×)Ò)Ñ+Ô+ˆØ# AœˆŒØ# AœˆŒØ$ QœˆŒØ$ QœˆŒÝ/Ñ1Ô1ˆÔÐÐr   c                 óB   — | j         | j        fdk     rt          S t          S )N)r   é   )rB   rC   Ú_datalayout_originalÚ_datalayout_i128©rH   s    r   Údata_layoutzNVVM.data_layout£   s"   € àŒM˜4œ=Ð)¨FÒ2Ð2Ý'Ð'å#Ð#r   c                 ó   — | j         S ©N)rG   rO   s    r   Úsupported_ccszNVVM.supported_ccsª   s   € àÔ"Ð"r   c                 óâ   — t          ¦   «         }t          ¦   «         }|                      t          |¦  «        t          |¦  «        ¦  «        }|                      |d¦  «         |j        |j        fS )NzFailed to get version.)r   r   r   Úcheck_errorÚvalue)rH   ÚmajorÚminorÚerrs       r   Úget_versionzNVVM.get_version®   s[   € Ý‘”ˆÝ‘”ˆØ×Ò�u U™|œ|­U°5©\¬\Ñ:Ô:ˆØ×Ò˜Ð6Ñ7Ô7Ð7ØŒ{˜EœKÐ'Ð'r   c                 ój  — t          ¦   «         }t          ¦   «         }t          ¦   «         }t          ¦   «         }|                      t          |¦  «        t          |¦  «        t          |¦  «        t          |¦  «        ¦  «        }|                      |d¦  «         |j        |j        |j        |j        fS )NzFailed to get IR version.)r   r(   r   rU   rV   )rH   ÚmajorIRÚminorIRÚmajorDbgÚminorDbgrY   s         r   rA   zNVVM.get_ir_versionµ   sŒ   € Ý‘'”'ˆÝ‘'”'ˆÝ‘7”7ˆÝ‘7”7ˆØ× Ò ¥ w¡¤µ°w±´Ý!& x¡¤µ%¸±/´/ñCô Cˆà×Ò˜Ð9Ñ:Ô:Ð:ØŒ}˜gœm¨X¬^¸X¼^ÐKÐKr   Fc                 ó’   — |rDt          |t          |         ¦  «        }|r%t          |¦  «         t          j        d¦  «         d S |‚d S )Nr   )r   ÚRESULT_CODE_NAMESÚprintÚsysÚexit)rH   ÚerrorÚmsgrd   Úexcs        r   rU   zNVVM.check_error¿   sR   € Øð 	Ý˜CÕ!2°5Ô!9Ñ:Ô:ˆCØð Ý�c‘
”
�
Ý”˜‘”���à�	ð	ð 	r   )F)Ú__name__Ú
__module__Ú__qualname__Ú__doc__Únvvm_resultr   r   Únvvm_programr   r   r2   r-   r/   rJ   ÚpropertyrP   rS   rZ   rA   rU   r   r   r   r   r   K   sÕ  € € € € € ðð ð
 $ W W¨U¡^¤^°W°W¸U±^´^ÐDð *¨7¨7°<Ñ+@Ô+@ÐAð  +¨G¨G°LÑ,AÔ,AÐBð
 ˜ x°¸8ð#Eð ˜ x°¸8ð'Eð ˜ u¨g¨g°hÑ.?Ô.?ðAð ˜ w w¨xÑ'8Ô'8ð&:ð #.¨|¸XÐ!Fð #.¨|¸W¸WÀXÑ=NÔ=NÐ!Oð *¨<¸ÐBð & w w¨u¡~¤~°w°w¸u±~´~Ø!˜' %™.œ.¨'¨'°%©.¬.ð:ð *¨<¸Ø%˜g hÑ/Ô/ð1ðc3ð 3€Kðl €Jðð ð ð*2ð 2ð 2ð ð$ð $ñ „Xð$ð ð#ð #ñ „Xð#ð(ð (ð (ðLð Lð Lðð ð ð ð ð r   r   c                   ó8   — e Zd Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z	dS )	ÚCompilationUnitc                 óà   — t          ¦   «         | _        t          ¦   «         | _        | j                             t          | j        ¦  «        ¦  «        }| j                             |d¦  «         d S )NzFailed to create CU)r   r0   rm   Ú_handler   r   rU   )rH   rY   s     r   rJ   zCompilationUnit.__init__Ê   sV   € Ý‘f”fˆŒÝ#‘~”~ˆŒØŒk×+Ò+­E°$´,Ñ,?Ô,?Ñ@Ô@ˆØŒ×Ò Ð%:Ñ;Ô;Ð;Ð;Ð;r   c                 ó    — t          ¦   «         }|                     t          | j        ¦  «        ¦  «        }|                     |dd¬¦  «         d S )NzFailed to destroy CUT)rd   )r   r    r   rr   rU   )rH   r0   rY   s      r   Ú__del__zCompilationUnit.__del__Ð   sJ   € Ý‘”ˆØ×'Ò'­¨d¬lÑ(;Ô(;Ñ<Ô<ˆØ×Ò˜3Ð 6¸TÐÑBÔBÐBÐBÐBr   c                 óš   — | j                              | j        |t          |¦  «        d¦  «        }| j                              |d¦  «         dS )zÔ
         Add a module level NVVM IR to a compilation unit.
         - The buffer should contain an NVVM module IR either in the bitcode
           representation (LLVM3.0) or in the text representation.
        NúFailed to add module)r0   r!   rr   ÚlenrU   ©rH   ÚbufferrY   s      r   Ú
add_modulezCompilationUnit.add_moduleÕ   sM   € ð Œk×0Ò0°´¸vÝ14°V±´¸dñDô DˆàŒ×Ò Ð%;Ñ<Ô<Ð<Ð<Ð<r   c                 óš   — | j                              | j        |t          |¦  «        d¦  «        }| j                              |d¦  «         dS )zÂ
        Lazily add an NVVM IR module to a compilation unit.
        The buffer should contain NVVM module IR either in the bitcode
        representation or in the text representation.
        Nrv   )r0   r"   rr   rw   rU   rx   s      r   Úlazy_add_modulezCompilationUnit.lazy_add_moduleß   sM   € ð Œk×4Ò4°T´\À6Ý58¸±[´[À$ñHô HˆàŒ×Ò Ð%;Ñ<Ô<Ð<Ð<Ð<r   c                 ón  ‡— d„ Šˆfd„|                      ¦   «         D ¦   «         }t          t          |¦  «        z  d„ |D ¦   «         Ž }| j                             | j        t          |¦  «        |¦  «        }|                      |d¦  «         | j                             | j        t          |¦  «        |¦  «        }|                      |d¦  «         t          ¦   «         }| j         	                    | j        t          |¦  «        ¦  «        }|                      |d¦  «         t          |j        z  ¦   «         }| j                             | j        |¦  «        }|                      |d¦  «         |                      ¦   «         | _        | j        r t!          j        | j        t$          ¬¦  «         |d	d	…         S )
a  Perform Compilation.

        Compilation options are accepted as keyword arguments, with the
        following considerations:

        - Underscores (`_`) in option names are converted to dashes (`-`), to
          match NVVM's option name format.
        - Options that take a value will be emitted in the form
          "-<name>=<value>".
        - Booleans passed as option values will be converted to integers.
        - Options which take no value (such as `-gen-lto`) should have a value
          of `None` passed in and will be emitted in the form "-<name>".

        For documentation on NVVM compilation options, see the CUDA Toolkit
        Documentation:

        https://docs.nvidia.com/cuda/libnvvm-api/index.html#_CPPv418nvvmCompileProgram11nvvmProgramiPPKc
        c                 ó”   — |                       dd¦  «        } |€d| › �S t          |t          ¦  «        rt          |¦  «        }d| › d|› �S )NÚ_ú-ú=)ÚreplaceÚ
isinstanceÚboolÚint)ÚkÚvs     r   Ústringify_optionz1CompilationUnit.compile.<locals>.stringify_optioný   sT   € Ø—	’	˜#˜sÑ#Ô#ˆAàˆyØ˜1�w�w�å˜!�TÑ"Ô"ð Ý˜‘F”F�à�q�;�;˜1�;�;Ðr   c                 ó.   •— g | ]\  }} ‰||¦  «        ‘ŒS r   r   )Ú.0r†   r‡   rˆ   s      €r   ú
<listcomp>z+CompilationUnit.compile.<locals>.<listcomp>  s+   ø€ ÐFÐFÐF©d¨a°Ð#Ð# A qÑ)Ô)ÐFÐFÐFr   c                 óR   — g | ]$}t          |                     d ¦  «        ¦  «        ‘Œ%S )Úutf8)r   Úencode)rŠ   Úxs     r   r‹   z+CompilationUnit.compile.<locals>.<listcomp>
  s<   € ð -?ð -?ð -?Ø12õ .6°a·h²h¸vÑ6FÔ6FÑ-GÔ-Gð -?ð -?ð -?r   zFailed to verify
zFailed to compile
z&Failed to get size of compiled result.zFailed to get compiled result.)ÚcategoryN)r3   r   rw   r0   r)   rr   Ú
_try_errorr#   r   r$   r   r	   rV   r%   Úget_logÚlogÚwarningsÚwarnr   )rH   ÚoptionsÚc_optsrY   ÚreslenÚptxbufrˆ   s         @r   ÚcompilezCompilationUnit.compileé   s—  ø€ ð(		ð 		ð 		ð GÐFÐFÐF°g·m²m±o´oÐFÑFÔFˆå�S ™\œ\Ñ)ð -?ð -?Ø6=ð-?ñ -?ô -?ð @ˆð Œk×+Ò+¨D¬L½#¸g¹,¼,ÈÑOÔOˆØ�Š˜Ð1Ñ2Ô2Ð2ð Œk×,Ò,¨T¬\½3¸w¹<¼<ÈÑPÔPˆØ�Š˜Ð2Ñ3Ô3Ð3õ ‘”ˆØŒk×3Ò3°D´LÅ%ÈÁ-Ä-ÑPÔPˆà�Š˜ÐEÑFÔFÐFå˜6œ<Ñ'Ñ*Ô*ˆØŒk×/Ò/°´¸fÑEÔEˆØ�Š˜Ð=Ñ>Ô>Ð>ð —<’<‘>”>ˆŒØŒ8ð 	:ÝŒM˜$œ(­[Ð9Ñ9Ô9Ð9à�a�a�aŒyÐr   c                 ój   — | j                              ||›d|                      ¦   «         ›�¦  «         d S )Nú
)r0   rU   r’   )rH   rY   rf   s      r   r‘   zCompilationUnit._try_error%  s3   € ØŒ×Ò °°°°d·l²l±n´n°nÐ%EÑFÔFÐFÐFÐFr   c                 óž  — t          ¦   «         }| j                             | j        t	          |¦  «        ¦  «        }| j                             |d¦  «         |j        dk    rkt          |j        z  ¦   «         }| j                             | j        |¦  «        }| j                             |d¦  «         |j         	                    d¦  «        S dS )Nz#Failed to get compilation log size.r   zFailed to get compilation log.r�   Ú )
r   r0   r&   rr   r   rU   rV   r	   r'   Údecode)rH   r˜   rY   Úlogbufs       r   r’   zCompilationUnit.get_log(  s¯   € Ý‘”ˆØŒk×/Ò/°´½eÀF¹m¼mÑLÔLˆØŒ×Ò Ð%JÑKÔKÐKàŒ<˜!ÒÐÝ˜vœ|Ñ+Ñ.Ô.ˆFØ”+×/Ò/°´¸fÑEÔEˆCØŒK×#Ò# CÐ)IÑJÔJÐJà”<×&Ò& vÑ.Ô.Ð.àˆrr   N)
rh   ri   rj   rJ   rt   rz   r|   rš   r‘   r’   r   r   r   rp   rp   É   s„   € € € € € ð<ð <ð <ðCð Cð Cð
=ð =ð =ð=ð =ð =ð:ð :ð :ðxGð Gð Gðð ð ð ð r   rp   )©r   r   )r   é   ©r   r   )r   r@   )r   r   )é   r   )r¤   r   )r¤   r@   )r¢   r   )r¢   r@   )r¢   r   )rL   r   ©rL   r¤   ©rL   r¢   )rL   é	   ©r§   r   )r¡   r¥   )r¡   r¦   )r¡   r¨   )r£   r¨   )
)é   r@   )r©   r   )r©   r   )r©   r   )r©   r¤   )r©   r¢   )r©   rL   )é   r   )rª   r   )rª   r@   c                 óÆ   ‡‡— 	 t           |          \  ŠŠt          ˆˆfd„t          D ¦   «         ¦  «        S # t          $ r! t          d„ t          D ¦   «         ¦  «        cY S w xY w)Nc                 ó4   •— g | ]}‰|cxk    r‰k    ¯n n|‘ŒS r   r   )rŠ   ÚccÚmax_ccÚmin_ccs     €€r   r‹   z(ccs_supported_by_ctk.<locals>.<listcomp>S  sD   ø€ ð 1ð 1ð 1˜RØ 2Ð/Ð/Ò/Ð/¨Ò/Ð/Ð/Ð/Ð/ð Ø/Ð/Ð/r   c                 ó2   — g | ]}|t           j        k    ¯|‘ŒS r   )r   ÚCUDA_DEFAULT_PTX_CC)rŠ   r­   s     r   r‹   z(ccs_supported_by_ctk.<locals>.<listcomp>X  s/   € ð ;ð ;ð ;˜RØ�vÔ9Ò9Ð9ð Ø9Ð9Ð9r   )ÚCTK_SUPPORTEDÚtupleÚCOMPUTE_CAPABILITIESÚKeyError)Úctk_versionr®   r¯   s    @@r   Úccs_supported_by_ctkr·   O  s©   øø€ ð	<å& {Ô3‰ˆ�Ýð 1ð 1ð 1ð 1ð 1Õ#7ð 1ñ 1ô 1ñ 2ô 2ð 	2øåð <ð <ð <õ ð ;ð ;Õ#7ð ;ñ ;ô ;ñ <ô <ð 	<ð 	<ð 	<ð<øøøs   „05 µ(A ÁA c                  ó,  — 	 ddl m}  |                      ¦   «         }n#  d}|cY S xY wt          t          ¦  «        }||k     rCd}|d         › d|d         › �}d|› d|d         › d|d         › d�}t          j        |¦  «         |S t          |¦  «        }|S )	Nr   )Úruntimer   ú.r   zCUDA Toolkit z is unsupported by Numba - z! is the minimum required version.)Únumba.cuda.cudadrv.runtimer¹   rZ   Úminr²   r”   r•   r·   )r¹   Úcudart_versionÚ_supported_ccÚ
min_cudartÚctk_verÚunsupported_vers         r   rF   rF   \  sé   € ðØ6Ð6Ð6Ð6Ð6Ð6Ø ×,Ò,Ñ.Ô.ˆˆøðð ˆØÐÐÐøøøõ •]Ñ#Ô#€JØ˜
Ò"Ð"ØˆØ# AÔ&Ð<Ð<¨¸Ô):Ð<Ð<ˆð/¨7ð /ð /Ø(¨œmð/ð /Ø.8¸¬mð/ð /ð /ˆõ 	Œ�oÑ&Ô&Ð&ØÐå(¨Ñ8Ô8€MØÐs   ‚ �%c                 óü   — t          ¦   «         j        }|sd}t          |¦  «        ‚t          |¦  «        D ]?\  }}|| k    r|c S || k    r*|dk    rd| |z   z  }t          |¦  «        ‚||dz
           c S Œ@|d         S )zñ
    Given a compute capability, return the closest compute capability supported
    by the CUDA toolkit.

    :param mycc: Compute capability as a tuple ``(MAJOR, MINOR)``
    :return: Closest supported CC as a tuple ``(MAJOR, MINOR)``
    zmNo supported GPU compute capabilities found. Please check your cudatoolkit version matches your CUDA version.r   z?GPU compute capability %d.%d is not supported(requires >=%d.%d)r   éÿÿÿÿ)r   rS   r   Ú	enumerate)ÚmyccrS   rf   Úir­   s        r   Úfind_closest_archrÇ   u  s·   € õ ‘F”FÔ(€Màð $ðQˆå˜sÑ#Ô#Ð#å˜=Ñ)Ô)ð ,ð ,‰ˆˆ2Ø�Š:ˆ:àˆIˆIˆIØ�$ŠYˆYà�AŠvˆvð+Ø.2°R©iñ9�å& sÑ+Ô+Ð+ð % Q¨¡UÔ+Ð+Ð+Ð+ð ð ˜ÔÐr   c                 ó`   — t           j        rt           j        }nt          | |f¦  «        }d|z  S )z1Matches with the closest architecture option
    zcompute_%d%d)r   ÚFORCE_CUDA_CCrÇ   )rW   rX   Úarchs      r   Úget_arch_optionrË   —  s5   € õ Ôð 1ÝÔ#ˆˆå  %¨ Ñ0Ô0ˆØ˜DÑ Ð r   z~Missing libdevice file.
Please ensure you have package cudatoolkit >= 11.0
Install package by:

    conda install cudatoolkit
c                   ó   — e Zd ZdZd„ Zd„ ZdS )Ú	LibDeviceNc                 ó–   — | j         €5t          ¦   «         €t          t          ¦  «        ‚t	          ¦   «         | _         | j         | _        d S rR   )Ú_cache_r   ÚRuntimeErrorÚMISSING_LIBDEVICE_FILE_MSGr   ÚbcrO   s    r   rJ   zLibDevice.__init__¬  s=   € ØŒ<ÐÝ‰ŒÐ&Ý"Õ#=Ñ>Ô>Ð>Ý)Ñ+Ô+ˆDŒLà”,ˆŒˆˆr   c                 ó   — | j         S rR   )rÒ   rO   s    r   ÚgetzLibDevice.get´  s	   € ØŒwˆr   )rh   ri   rj   rÏ   rJ   rÔ   r   r   r   rÍ   rÍ   ©  s7   € € € € € Ø€Gðð ð ðð ð ð ð r   rÍ   z’
    %cas_success = cmpxchg volatile {Ti}* %iptr, {Ti} %old, {Ti} %new monotonic monotonic
    %cas = extractvalue {{ {Ti}, i1 }} %cas_success, 0
a  
define internal {T} @___numba_atomic_{T}_{FUNC}({T}* %ptr, {T} %val) alwaysinline {{
entry:
    %iptr = bitcast {T}* %ptr to {Ti}*
    %old2 = load volatile {Ti}, {Ti}* %iptr
    br label %attempt

attempt:
    %old = phi {Ti} [ %old2, %entry ], [ %cas, %attempt ]
    %dold = bitcast {Ti} %old to {T}
    %dnew = {OP} {T} %dold, %val
    %new = bitcast {T} %dnew to {Ti}
    {CAS}
    %repeat = icmp ne {Ti} %cas, %old
    br i1 %repeat, label %attempt, label %done

done:
    %result = bitcast {Ti} %old to {T}
    ret {T} %result
}}
aÇ  
define internal {T} @___numba_atomic_{Tu}_inc({T}* %iptr, {T} %val) alwaysinline {{
entry:
    %old2 = load volatile {T}, {T}* %iptr
    br label %attempt

attempt:
    %old = phi {T} [ %old2, %entry ], [ %cas, %attempt ]
    %bndchk = icmp ult {T} %old, %val
    %inc = add {T} %old, 1
    %new = select i1 %bndchk, {T} %inc, {T} 0
    {CAS}
    %repeat = icmp ne {T} %cas, %old
    br i1 %repeat, label %attempt, label %done

done:
    ret {T} %old
}}
aË  
define internal {T} @___numba_atomic_{Tu}_dec({T}* %iptr, {T} %val) alwaysinline {{
entry:
    %old2 = load volatile {T}, {T}* %iptr
    br label %attempt

attempt:
    %old = phi {T} [ %old2, %entry ], [ %cas, %attempt ]
    %dec = add {T} %old, -1
    %bndchk = icmp ult {T} %dec, %val
    %new = select i1 %bndchk, {T} %dec, {T} %val
    {CAS}
    %repeat = icmp ne {T} %cas, %old
    br i1 %repeat, label %attempt, label %done

done:
    ret {T} %old
}}
a”  
define internal {T} @___numba_atomic_{T}_{NAN}{FUNC}({T}* %ptr, {T} %val) alwaysinline {{
entry:
    %ptrval = load volatile {T}, {T}* %ptr
    ; Return early when:
    ; - For nanmin / nanmax when val is a NaN
    ; - For min / max when val or ptr is a NaN
    %early_return = fcmp uno {T} %val, %{PTR_OR_VAL}val
    br i1 %early_return, label %done, label %lt_check

lt_check:
    %dold = phi {T} [ %ptrval, %entry ], [ %dcas, %attempt ]
    ; Continue attempts if dold less or greater than val (depending on whether min or max)
    ; or if dold is NaN (for nanmin / nanmax)
    %cmp = fcmp {OP} {T} %dold, %val
    br i1 %cmp, label %attempt, label %done

attempt:
    ; Attempt to swap in the value
    %old = bitcast {T} %dold to {Ti}
    %iptr = bitcast {T}* %ptr to {Ti}*
    %new = bitcast {T} %val to {Ti}
    {CAS}
    %dcas = bitcast {Ti} %cas to {T}
    br label %lt_check

done:
    ret {T} %ptrval
}}
c                 ó8   — t                                | ¬¦  «        S )N©ÚTi)Úcas_nvvmÚformatrÖ   s    r   Úir_casrÚ     s   € Ý�?Š?˜bˆ?Ñ!Ô!Ð!r   c           	      óf   — t          | |||t          |¦  «        ¬¦  «        }t          j        di |¤ŽS )N)ÚTr×   ÚOPÚFUNCÚCASr   )ÚdictrÚ   Úir_numba_atomic_binary_templaterÙ   )rÜ   r×   rÝ   rÞ   Úparamss        r   Úir_numba_atomic_binaryrã   !  s7   € Ý�A˜" ¨$µF¸2±J´JÐ?Ñ?Ô?€FÝ*Ô1Ð;Ð;°FÐ;Ð;Ð;r   c                 ój   — t          | |||||t          |¦  «        ¬¦  «        }t          j        di |¤ŽS )N)rÜ   r×   ÚNANrÝ   Ú
PTR_OR_VALrÞ   rß   r   )rà   rÚ   Úir_numba_atomic_minmax_templaterÙ   )rÜ   r×   rå   rÝ   ræ   rÞ   râ   s          r   Úir_numba_atomic_minmaxrè   &  sD   € Ý�A˜" #¨"¸Ø¥¨¡¤ð-ñ -ô -€Fõ +Ô1Ð;Ð;°FÐ;Ð;Ð;r   c                 óV   — t                                | |t          | ¦  «        ¬¦  «        S ©N)rÜ   ÚTurß   )Úir_numba_atomic_inc_templaterÙ   rÚ   ©rÜ   rë   s     r   Úir_numba_atomic_incrî   -  ó#   € Ý'×.Ò.°°r½vÀa¹y¼yÐ.ÑIÔIÐIr   c                 óV   — t                                | |t          | ¦  «        ¬¦  «        S rê   )Úir_numba_atomic_dec_templaterÙ   rÚ   rí   s     r   Úir_numba_atomic_decrò   1  rï   r   c                 ó†  — dt          dddd¬¦  «        fdt          dd	d
d¬¦  «        fdt          ddd
d¬¦  «        fdt          dd¬¦  «        fdt          dd¬¦  «        fdt          dd	dddd¬¦  «        fdt          dddddd¬¦  «        fdt          dd	dddd¬¦  «        fdt          dddddd¬¦  «        fdt          dd	dddd¬¦  «        fdt          dddddd¬¦  «        fd t          dd	dd!dd¬¦  «        fd"t          dddd!dd¬¦  «        fd#g}|D ]\  }}|                      ||¦  «        } Œt          | ¦  «        } | S )$NzIdeclare double @"___numba_atomic_double_add"(double* %".1", double %".2")ÚdoubleÚi64ÚfaddÚadd)rÜ   r×   rÝ   rÞ   zEdeclare float @"___numba_atomic_float_sub"(float* %".1", float %".2")ÚfloatÚi32ÚfsubÚsubzIdeclare double @"___numba_atomic_double_sub"(double* %".1", double %".2")z=declare i64 @"___numba_atomic_u64_inc"(i64* %".1", i64 %".2")Úu64rí   z=declare i64 @"___numba_atomic_u64_dec"(i64* %".1", i64 %".2")zEdeclare float @"___numba_atomic_float_max"(float* %".1", float %".2")rž   znnan oltÚptrÚmax)rÜ   r×   rå   rÝ   ræ   rÞ   zIdeclare double @"___numba_atomic_double_max"(double* %".1", double %".2")zEdeclare float @"___numba_atomic_float_min"(float* %".1", float %".2")znnan ogtr¼   zIdeclare double @"___numba_atomic_double_min"(double* %".1", double %".2")zHdeclare float @"___numba_atomic_float_nanmax"(float* %".1", float %".2")ÚnanÚultzLdeclare double @"___numba_atomic_double_nanmax"(double* %".1", double %".2")zHdeclare float @"___numba_atomic_float_nanmin"(float* %".1", float %".2")ÚugtzLdeclare double @"___numba_atomic_double_nanmin"(double* %".1", double %".2"))Úimmargrž   )rã   rî   rò   rè   r‚   Úllvm140_to_70_ir)ÚllvmirÚreplacementsÚdeclÚfns       r   Úllvm_replacer  5  s   € à	TÝ	 (¨u¸ÀeÐ	LÑ	LÔ	Lð	Nà	PÝ	 '¨e¸ÀUÐ	KÑ	KÔ	Kð	Mà	TÝ	 (¨u¸ÀeÐ	LÑ	LÔ	Lð	Nà	HÝ	˜u¨Ð	/Ñ	/Ô	/ð	1à	HÝ	˜u¨Ð	/Ñ	/Ô	/ð	1à	PÝ	 '¨e¸À
Ø+0°uð
>ñ 
>ô 
>ð	?ð 
UÝ	 (¨u¸"ÀØ+0°uð
>ñ 
>ô 
>ð	?ð 
QÝ	 '¨e¸À
Ø+0°uð
>ñ 
>ô 
>ð	?ð 
UÝ	 (¨u¸"ÀØ+0°uð
>ñ 
>ô 
>ð	?ð 
TÝ	 '¨e¸À5Ø+-°Eð
;ñ 
;ô 
;ð	<ð 
XÝ	 (¨u¸%ÀEØ+-°Eð
;ñ 
;ô 
;ð	<ð 
TÝ	 '¨e¸À5Ø+-°Eð
;ñ 
;ô 
;ð	<ð 
XÝ	 (¨u¸%ÀEØ+-°Eð
;ñ 
;ô 
;ð	<ð 	ðG$€LðL !ð *ð *‰ˆˆbØ—’  bÑ)Ô)ˆˆå˜fÑ%Ô%€Fà€Mr   c                 óª  — t          | t          ¦  «        r| g} |                     dd¦  «        r|                     dddddœ¦  «         t	          ¦   «         }t          ¦   «         }| D ]9}t          |¦  «        }|                     |                     d¦  «        ¦  «         Œ:| 	                    | 
                    ¦   «         ¦  «          |j        di |¤ŽS )NÚfastmathFT)ÚftzÚfmaÚprec_divÚ	prec_sqrtr�   r   )rƒ   ÚstrÚpopÚupdaterp   rÍ   r  rz   rŽ   r|   rÔ   rš   )r  ÚoptsÚcuÚ	libdeviceÚmods        r   Úllvm_to_ptxr  d  sä   € Ý�&�#ÑÔð Ø�ˆà‡x‚x�
˜EÑ"Ô"ð Ø�ŠØØØØð	
ð 
ñ 	ô 	ð 	õ 
Ñ	Ô	€BÝ‘”€Iàð *ð *ˆÝ˜3ÑÔˆØ
�Š�c—j’j Ñ(Ô(Ñ)Ô)Ð)Ð)Ø×Ò�y—}’}‘”Ñ'Ô'Ð'àˆ2Œ:ÐÐ˜ÐÐÐr   z"^attributes #\d+ = \{ ([\w\s]+)\ }c                 óº  — g } | j         ¦   «         D ]µ}|                     d¦  «        r‰t                               |¦  «        }|                     d¦  «                             ¦   «         }d                     d„ |D ¦   «         ¦  «        }|                     |                     d¦  «        |¦  «        }|                     |¦  «         Œ¶d                     |¦  «        S )z,
    Convert LLVM 14.0 IR for LLVM 7.0.
    zattributes #r   ú c              3   ó&   K  — | ]}|d k    ¯|V — ŒdS )Ú
willreturnNr   )rŠ   Úas     r   ú	<genexpr>z#llvm140_to_70_ir.<locals>.<genexpr>ˆ  s,   è è € ÐCÐC 1°°lÒ1BÐ1B˜QÐ1BÐ1BÐ1BÐ1BÐCÐCr   rœ   )	Ú
splitlinesÚ
startswithÚre_attributes_defÚmatchÚgroupÚsplitÚjoinr‚   Úappend)r
   ÚbufÚlineÚmÚattrss        r   r  r  ~  sÄ   € ð €CØ�”‘”ð ð ˆØ�?Š?˜>Ñ*Ô*ð 	3å!×'Ò'¨Ñ-Ô-ˆAØ—G’G˜A‘J”J×$Ò$Ñ&Ô&ˆEØ—H’HÐCÐC¨ÐCÑCÔCÑCÔCˆEØ—<’< §¢¨¡
¤
¨EÑ2Ô2ˆDà�
Š
�4ÑÔÐÐà�9Š9�S‰>Œ>Ðr   c                 óD  — | j         }t          j        |d¦  «        }t          j        t          j        d¦  «        d¦  «        }|                     | ||f¦  «        }t          j        |d¦  «        }|                     |¦  «         | j	         
                    d¦  «         d S )NÚkernelé    r   znvvm.annotationsÚnoinline)Úmoduler
   ÚMetaDataStringÚConstantÚIntTypeÚadd_metadatar   Úget_or_insert_named_metadatar÷   Ú
attributesÚdiscard)Úlfuncr  ÚmdstrÚmdvalueÚmdÚnmds         r   Úset_cuda_kernelr:  �  s�   € Ø
Œ,€CåÔ˜c 8Ñ,Ô,€EÝŒk�"œ* R™.œ.¨!Ñ,Ô,€GØ	×	Ò	˜5 %¨Ð1Ñ	2Ô	2€Bå
Ô
.¨sÐ4FÑ
GÔ
G€CØ‡G‚GˆB�K„K€Kð 
Ô×Ò˜ZÑ(Ô(Ð(Ð(Ð(r   c                 óÞ   ‡— t          j        d¦  «        Šˆfd„t          ¦   «                              ¦   «         D ¦   «         }|                      |¦  «        }|                      d|¦  «         dS )zAdd NVVM IR version to moduler+  c                 ó&   •— g | ]} ‰|¦  «        ‘ŒS r   r   )rŠ   r‡   rù   s     €r   r‹   z"add_ir_version.<locals>.<listcomp>£  s!   ø€ Ð;Ð;Ð;˜a�3�3�q‘6”6Ð;Ð;Ð;r   znvvmir.versionN)r
   r0  r   rA   r1  Úadd_named_metadata)r  rI   Úmd_verrù   s      @r   Úadd_ir_versionr?  Ÿ  sk   ø€ õ Œ*�R‰.Œ.€CØ;Ð;Ð;Ð;¥4¡6¤6×#8Ò#8Ñ#:Ô#:Ð;Ñ;Ô;€KØ×Ò˜kÑ*Ô*€FØ×ÒÐ+¨VÑ4Ô4Ð4Ð4Ð4r   )Mrk   ÚloggingÚrerc   r”   Úctypesr   r   r   r   r   r   r	   Ú	threadingÚllvmliter
   re   r   r   r   Úlibsr   r   r   Ú
numba.corer   r   Ú	getLoggerrh   ÚloggerÚADDRSPACE_GENERICÚADDRSPACE_GLOBALÚADDRSPACE_SHAREDÚADDRSPACE_CONSTANTÚADDRSPACE_LOCALrm   rl   r"  ra   rÄ   rÆ   r†   r7   ÚmodulesrM   rN   r   ÚLockr,   r.   r   rp   r´   r²   r·   rF   rÇ   rË   rÑ   rÍ   rØ   rá   rì   rñ   rç   rÚ   rã   rè   rî   rò   r  r  rš   r  r  r:  r?  r   r   r   ú<module>rP     s²  ððð ð €€€Ø 	€	€	€	Ø 
€
€
€
Ø €€€ðð ð ð ð ð ð ð ð ð ð ð ð ð ð ð ð ð ð Ð Ð Ð à Ð Ð Ð Ð Ð à ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ø =Ð =Ð =Ð =Ð =Ð =Ð =Ð =Ð =Ð =Ø &Ð &Ð &Ð &Ð &Ð &Ð &Ð &ð 
ˆÔ	˜8Ñ	$Ô	$€àÐ ØÐ ØÐ ØÐ Ø€ð €ð €ð÷ 
‚E�G„Gð ð ˆIÐ'Ñ(Ô(ð )ð )�D€A€qØ€GˆCŒK˜Ô! 1 aÑ(Ô(Ð(Ð(ð;Ð ð7Ð ð
	ð 	ð 	ð ˆYŒ^ÑÔ€
ð{ð {ð {ð {ð {ˆ6ñ {ô {ð {ð|kð kð kð kð k�fñ kô kð kð\Ð ð ØØØØØØØØØðð €ð
<ð 
<ð 
<ðð ð ð2ð ð ðD!ð !ð !ðÐ ðð ð ð ð �ñ ô ð ð€ð#Ð ð, Ð ð( Ð ð(#Ð ð@"ð "ð "ð<ð <ð <ð
<ð <ð <ðJð Jð JðJð Jð Jð,ð ,ð ,ð^ð ð ð. �B”JÐDÑEÔEÐ ðð ð ð$)ð )ð )ð5ð 5ð 5ð 5ð 5r   