U
    hâËd—z  ã                   @   sÀ  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 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(ƒ qÊdZ,dZ-dd„ Z.e /¡ Z0G dd„ de1ƒZ2G dd„ de1ƒZ3dZ4ddddddddddddœZ5dd„ Z6dd„ Z7d d!„ Z8d"d#„ Z9d$Z:G d%d&„ d&e1ƒZ;d'Z<d(Z=d)Z>d*Z?d+Z@d,ZAd-ZBd.d/„ ZCd0d1„ ZDd2d3„ ZEd4d5„ ZFd6d7„ ZGd8d9„ ZHd:d;„ ZId<d=„ ZJe Kd>¡ZLe Kd?¡ZMe Kd@¡ZNdAZOe KeO PdBdC¡¡ZQe KdD¡ZRdEdFdGdHdIdJdKdLdMdNdOdPhZSe KdQ¡ZTe KdR¡ZUe KdS¡ZVe KdT¡ZWe KdU¡ZXe KdV¡ZYe KdW¡ZZe KdX¡Z[e KdY¡Z\dZd[d\œZ]d]d^„ Z^d_d`„ Z_dadb„ Z`dcdd„ Zadedf„ Zbdgdh„ ZcdS )iz(
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                   C   s*   z
t ƒ  W n tk
r    Y dS X dS dS )z(
    Return if libNVVM is available
    FTN)ÚNVVMr   © r   r   úP/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/cudadrv/nvvm.pyÚis_available<   s
    
r   c                   @   s  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d„ Zdd„ Zedd	„ ƒZed
d„ ƒZedd„ ƒZdd„ Zdd„ Zddd„ZdS )r   zProcess-wide singleton.
    )ÚnvvmVersionÚnvvmCreateProgramÚnvvmDestroyProgramÚnvvmAddModuleToProgramÚnvvmLazyAddModuleToProgramÚnvvmCompileProgramÚnvvmGetCompiledResultSizeÚnvvmGetCompiledResultÚnvvmGetProgramLogSizeÚnvvmGetProgramLogÚnvvmIRVersionÚnvvmVerifyProgramNc                 C   s¼   t �¬ | jd kr¬t | ¡ | _}ztdƒ|_W n8 tk
rf } zd | _d}t|| ƒ‚W 5 d }~X Y nX |j 	¡ D ]8\}}t
|j|ƒ}|d |_|dd … |_t|||ƒ qrW 5 Q R X | jS )NZnvvmz;libNVVM cannot be found. Do `conda install cudatoolkit`:
%sr   r
   )Ú
_nvvm_lockÚ_NVVM__INSTANCEÚobjectÚ__new__r   ÚdriverÚOSErrorr   Ú_PROTOTYPESÚitemsÚgetattrÚrestypeZargtypesÚsetattr)ÚclsÚinstÚeÚerrmsgÚnameÚprotoÚfuncr   r   r   r)   †   s    

zNVVM.__new__c                 C   s<   |   ¡ }|d | _|d | _|d | _|d | _tƒ | _d S )Nr   r
   é   r   )Úget_ir_versionÚ_majorIRÚ_minorIRZ	_majorDbgZ	_minorDbgÚget_supported_ccsÚ_supported_ccs)ÚselfÚir_versionsr   r   r   Ú__init__›   s    



zNVVM.__init__c                 C   s   | j | jfdkS )N)r
   é   )r:   r;   ©r>   r   r   r   Ú	is_nvvm70£   s    zNVVM.is_nvvm70c                 C   s   | j | jfdk rtS tS d S )N)r
   é   )r:   r;   Ú_datalayout_originalÚ_datalayout_i128rB   r   r   r   Údata_layoutª   s    zNVVM.data_layoutc                 C   s   | j S ©N)r=   rB   r   r   r   Úsupported_ccs±   s    zNVVM.supported_ccsc                 C   s8   t ƒ }t ƒ }|  t|ƒt|ƒ¡}|  |d¡ |j|jfS )NzFailed to get version.)r   r   r   Úcheck_errorÚvalue)r>   ÚmajorÚminorÚerrr   r   r   Úget_versionµ   s
    zNVVM.get_versionc                 C   sX   t ƒ }t ƒ }t ƒ }t ƒ }|  t|ƒt|ƒt|ƒt|ƒ¡}|  |d¡ |j|j|j|jfS )NzFailed to get IR version.)r   r$   r   rJ   rK   )r>   ZmajorIRZminorIRZmajorDbgZminorDbgrN   r   r   r   r9   ¼   s     ÿzNVVM.get_ir_versionFc                 C   s2   |r.t |t| ƒ}|r*t|ƒ t d¡ n|‚d S )Nr
   )r   ÚRESULT_CODE_NAMESÚprintÚsysÚexit)r>   ÚerrorÚmsgrS   Úexcr   r   r   rJ   Æ   s    zNVVM.check_error)F)Ú__name__Ú
__module__Ú__qualname__Ú__doc__Únvvm_resultr   r   Únvvm_programr   r   r,   r'   r)   r@   ÚpropertyrC   rG   rI   rO   r9   rJ   r   r   r   r   r   K   s`   

    ÿ    ÿ   ÿ  ÿ ÿÿÏ6



r   c                   @   sD   e Zd Zdd„ Zdd„ Zdd„ Zdd„ Zd	d
„ Zdd„ Zdd„ Z	dS )ÚCompilationUnitc                 C   s4   t ƒ | _tƒ | _| j t| jƒ¡}| j |d¡ d S )NzFailed to create CU)r   r*   r\   Ú_handler   r   rJ   )r>   rN   r   r   r   r@   Ñ   s    zCompilationUnit.__init__c                 C   s*   t ƒ }| t| jƒ¡}|j|ddd� d S )NzFailed to destroy CUT)rS   )r   r   r   r_   rJ   )r>   r*   rN   r   r   r   Ú__del__×   s    zCompilationUnit.__del__c                 C   s*   | 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)r*   r   r_   ÚlenrJ   ©r>   ÚbufferrN   r   r   r   Ú
add_moduleÜ   s
     ÿzCompilationUnit.add_modulec                 C   s*   | 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.
        Nra   )r*   r   r_   rb   rJ   rc   r   r   r   Úlazy_add_moduleæ   s
     ÿzCompilationUnit.lazy_add_modulec                 K   s€  g }d|kr |  d| d¡ ¡ | d¡r>|  d| d¡ ¡ d}|D ]8}||krFtt| |¡ƒƒ}|  d| dd¡|f ¡ qF|r¦d	 tt| 	¡ ƒ¡}t
d
 |¡ƒ‚tt|ƒ dd„ |D ƒŽ }| j | jt|ƒ|¡}|  |d¡ | j | jt|ƒ|¡}|  |d¡ tƒ }	| j | jt|	ƒ¡}|  |d¡ t|	j ƒ }
| j | j|
¡}|  |d¡ |  ¡ | _| j�rttj| jtd� |
dd… S )aj  Perform Compilation

        The valid compiler options are

         *   - -opt=
         *     - 0 (disable optimizations)
         *     - 3 (default, enable optimizations)
         *   - -arch=
         *     - compute_XX where XX is in (35, 37, 50, 52, 53, 60, 61, 62, 70,
         *                                  72, 75, 80, 86, 89, 90).
         *       The default is compute_52.
         *   - -ftz=
         *     - 0 (default, preserve denormal values, when performing
         *          single-precision floating-point operations)
         *     - 1 (flush denormal values to zero, when performing
         *          single-precision floating-point operations)
         *   - -prec-sqrt=
         *     - 0 (use a faster approximation for single-precision
         *          floating-point square root)
         *     - 1 (default, use IEEE round-to-nearest mode for
         *          single-precision floating-point square root)
         *   - -prec-div=
         *     - 0 (use a faster approximation for single-precision
         *          floating-point division and reciprocals)
         *     - 1 (default, use IEEE round-to-nearest mode for
         *          single-precision floating-point division and reciprocals)
         *   - -fma=
         *     - 0 (disable FMA contraction)
         *     - 1 (default, enable FMA contraction)
         *
         Úoptz-opt=%dÚarchz-arch=%s)ÚftzÚ	prec_sqrtÚprec_divÚfmaz-%s=%dÚ_ú-ú, zunsupported option {0}c                 S   s   g | ]}t | d ¡ƒ‘qS )Úutf8)r   Úencode)Ú.0Úxr   r   r   Ú
<listcomp>+  s   ÿz+CompilationUnit.compile.<locals>.<listcomp>zFailed to verify
zFailed to compile
z&Failed to get size of compiled result.zFailed to get compiled result.)ÚcategoryN)ÚappendÚpopÚgetÚintÚboolÚreplaceÚjoinÚmapÚreprÚkeysr   Úformatr   rb   r*   r%   r_   Ú
_try_errorr   r   r    r   r   rK   r!   Úget_logÚlogÚwarningsÚwarnr   )r>   ÚoptionsÚoptsZother_optionsÚkÚvZoptstrZc_optsrN   ÚreslenZptxbufr   r   r   Úcompileð   s<    "
ÿ
zCompilationUnit.compilec                 C   s   | j  |d||  ¡ f ¡ d S )Nz%s
%s)r*   rJ   r‚   )r>   rN   rU   r   r   r   r�   F  s    zCompilationUnit._try_errorc                 C   sl   t ƒ }| j | jt|ƒ¡}| j |d¡ |jdkrht|j ƒ }| j | j|¡}| j |d¡ |j 	d¡S dS )Nz#Failed to get compilation log size.r
   zFailed to get compilation log.rp   Ú )
r   r*   r"   r_   r   rJ   rK   r   r#   Údecode)r>   rŠ   rN   Zlogbufr   r   r   r‚   I  s    
zCompilationUnit.get_logN)
rW   rX   rY   r@   r`   re   rf   r‹   r�   r‚   r   r   r   r   r^   Ð   s   

Vr^   )©r   r   )r   é   ©r   r   )r   r8   )r   r   )rA   r   )rA   r
   )rA   r8   )r�   r   )r�   r8   )r�   r   ©rD   r   ©rD   rA   ©rD   r�   )rD   é	   ©r”   r   )rŽ   r‘   )rŽ   r’   )rŽ   r“   )rŽ   r•   )r�   r•   ))é   r   )r–   r
   )r–   r8   )r–   r   )r–   r   )r–   r   )r–   rA   )r–   r�   )r–   rD   )é   r   )r—   r
   c                    sR   z&t |  \‰‰ t‡ ‡fdd„tD ƒƒW S  tk
rL   tdd„ tD ƒƒ Y S X d S )Nc                    s(   g | ] }ˆ|  krˆ krn q|‘qS r   r   ©rr   Úcc©Zmax_ccZmin_ccr   r   rt   u  s
     
 ÿz(ccs_supported_by_ctk.<locals>.<listcomp>c                 S   s   g | ]}|t jkr|‘qS r   )r   ZCUDA_DEFAULT_PTX_CCr˜   r   r   r   rt   z  s    
ÿ)ÚCTK_SUPPORTEDÚtupleÚCOMPUTE_CAPABILITIESÚKeyError)Zctk_versionr   rš   r   Úccs_supported_by_ctkq  s
    rŸ   c                  C   s’   zddl m}  |  ¡ }W n   d}| Y S X ttƒ}||k r†d}|d › d|d › �}d|› d|d › d|d › d�}t |¡ |S t|ƒ}|S )	Nr   )Úruntimer   Ú.r
   zCUDA Toolkit z is unsupported by Numba - z! is the minimum required version.)Znumba.cuda.cudadrv.runtimer    rO   Úminr›   r„   r…   rŸ   )r    Zcudart_versionZ_supported_ccZ
min_cudartZctk_verZunsupported_verr   r   r   r<   ~  s    
 
r<   c                 C   sx   t ƒ j}|sd}t|ƒ‚t|ƒD ]N\}}|| kr8|  S || kr |dkr^d| |  }t|ƒ‚q ||d    S q |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   rI   r   Ú	enumerate)ZmyccrI   rU   Úir™   r   r   r   Úfind_closest_arch—  s    ÿ
r¦   c                 C   s"   t jrt j}nt| |fƒ}d| S )z1Matches with the closest architecture option
    zcompute_%d%d)r   ZFORCE_CUDA_CCr¦   )rL   rM   rh   r   r   r   Úget_arch_option¹  s    r§   z~Missing libdevice file.
Please ensure you have package cudatoolkit >= 11.0
Install package by:

    conda install cudatoolkit
c                   @   s    e Zd ZdZdd„ Zdd„ ZdS )Ú	LibDeviceNc                 C   s0   | j d kr$tƒ d krttƒ‚tƒ | _ | j | _d S rH   )Ú_cache_r   ÚRuntimeErrorÚMISSING_LIBDEVICE_FILE_MSGr   ÚbcrB   r   r   r   r@   Î  s
    

zLibDevice.__init__c                 C   s   | j S rH   )r¬   rB   r   r   r   rx   Ö  s    zLibDevice.get)rW   rX   rY   r©   r@   rx   r   r   r   r   r¨   Ë  s   r¨   zº
define internal {T} @___numba_atomic_{T}_cas_hack({T}* %ptr, {T} %cmp, {T} %val) alwaysinline {{
    %out = cmpxchg volatile {T}* %ptr, {T} %cmp, {T} %val monotonic
    ret {T} %out
}}
z’
    %cas_success = cmpxchg volatile {Ti}* %iptr, {Ti} %old, {Ti} %new monotonic monotonic
    %cas = extractvalue {{ {Ti}, i1 }} %cas_success, 0
zI
    %cas = cmpxchg volatile {Ti}* %iptr, {Ti} %old, {Ti} %new monotonic
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                 C   s$   t ƒ jrtj| d�S tj| d�S d S )N©ÚTi)r   rC   Ú
cas_nvvm70r€   Ú
cas_nvvm34r­   r   r   r   Úir_casK  s    r±   c                 C   s"   t | |||t|ƒd�}tjf |ŽS )N)ÚTr®   ÚOPÚFUNCÚCAS)Údictr±   Úir_numba_atomic_binary_templater€   )r²   r®   r³   r´   Úparamsr   r   r   Úir_numba_atomic_binaryR  s    r¹   c              	   C   s&   t | |||||t|ƒd�}tjf |ŽS )N)r²   r®   ÚNANr³   Ú
PTR_OR_VALr´   rµ   )r¶   r±   Úir_numba_atomic_minmax_templater€   )r²   r®   rº   r³   r»   r´   r¸   r   r   r   Úir_numba_atomic_minmaxW  s
     ÿr½   c                 C   s   t j| |t| ƒd�S ©N)r²   ÚTurµ   )Úir_numba_atomic_inc_templater€   r±   ©r²   r¿   r   r   r   Úir_numba_atomic_inc^  s    rÂ   c                 C   s   t j| |t| ƒd�S r¾   )Úir_numba_atomic_dec_templater€   r±   rÁ   r   r   r   Úir_numba_atomic_decb  s    rÄ   c                 C   sH   |   ¡ }t|ƒD ],\}}| d¡rd}| tƒ j¡||<  q>qd |¡S )z@
    Find the line containing the datalayout and replace it
    ztarget datalayoutztarget datalayout = "{0}"Ú
)Ú
splitlinesr¤   Ú
startswithr€   r   rG   r|   )ÚllvmirÚlinesr¥   ÚlnÚtmpr   r   r   Ú_replace_datalayoutf  s    
rÌ   c                 C   sz  dt ddddd�fdt dd	d
dd�fdt ddd
dd�fdtddd�fdtddd�fdtdd	ddddd�fdtddddddd�fdtdd	ddddd�fdtddddddd�fdtdd	ddddd�fdtddddddd�fd tdd	dd!ddd�fd"tdddd!ddd�fd#g}tƒ j�s>|d$tjd	d%�fd&tjdd%�fg7 }t| ƒ} |D ]\}}|  	||¡} �qBtƒ j�rnt
| ƒ} nt| ƒ} | S )'NzIdeclare double @"___numba_atomic_double_add"(double* %".1", double %".2")ÚdoubleZi64ZfaddÚadd)r²   r®   r³   r´   zEdeclare float @"___numba_atomic_float_sub"(float* %".1", float %".2")ÚfloatÚi32ZfsubÚsubzIdeclare double @"___numba_atomic_double_sub"(double* %".1", double %".2")z=declare i64 @"___numba_atomic_u64_inc"(i64* %".1", i64 %".2")Z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 oltZ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")ÚnanZultzLdeclare double @"___numba_atomic_double_nanmax"(double* %".1", double %".2")zHdeclare float @"___numba_atomic_float_nanmin"(float* %".1", float %".2")ZugtzLdeclare double @"___numba_atomic_double_nanmin"(double* %".1", double %".2"))ZimmargrŒ   zMdeclare i32 @"___numba_atomic_i32_cas_hack"(i32* %".1", i32 %".2", i32 %".3"))r²   zMdeclare i64 @"___numba_atomic_i64_cas_hack"(i64* %".1", i64 %".2", i64 %".3"))r¹   rÂ   rÄ   r½   r   rC   Úir_numba_cas_hackr€   rÌ   r{   Úllvm100_to_70_irÚllvm100_to_34_ir)rÈ   ZreplacementsÚdeclÚfnr   r   r   Úllvm_replaces  s¢    ÿÿÿ
ÿ
ÿ
 ÿÿ
 ÿÿ
 ÿÿ
 ÿÿ
 ÿÿ
 ÿÿ
 ÿÿ
 ÿÿÝ&

ÿ
ÿý

rÙ   c                 K   sx   t | tƒr| g} | dd¡r0| dddddœ¡ tƒ }tƒ }| D ]}t|ƒ}| | d¡¡ q@| 	| 
¡ ¡ |jf |ŽS )NZfastmathFT)ri   rl   rk   rj   rp   )Ú
isinstanceÚstrrw   Úupdater^   r¨   rÙ   re   rq   rf   rx   r‹   )rÈ   r‡   ÚcuZ	libdeviceÚmodr   r   r   Úllvm_to_ptx²  s     
ürß   z	\!\d+\s*=zmetadata\s*\![{'\"0-9]z\!\d+z,\!{i32 \d, \!\"Debug Info Version\", i32 \d}ú z\s+z"^attributes #\d+ = \{ ([\w\s]+)\ }ZalwaysinlineZcoldZ
inlinehintZminsizeZnoduplicateÚnoinlineZnoreturnZnounwindZoptnoneZoptiszeZreadnoneÚreadonlyz"\bgetelementptr\s(?:inbounds )?\(?z=\s*\bload\s(?:\bvolatile\s)?z(call\s[^@]+\))(\s@)z\s*!range\s+!\d+z
[,{}()[\]]z\bnonnull\bz"\b(local_unnamed_addr|writeonly)\bz\((.*)\)zspFlags: (.*),ZisDefinitionZisOptimized)ZDISPFlagDefinitionZDISPFlagOptimizedc                 C   sn   g }|   ¡ D ]V}| d¡rXt |¡}| d¡ ¡ }d dd„ |D ƒ¡}| | d¡|¡}| |¡ qd |¡S )z,
    Convert LLVM 10.0 IR for LLVM 7.0.
    úattributes #r
   rà   c                 s   s   | ]}|d kr|V  qdS )Z
willreturnNr   ©rr   Úar   r   r   Ú	<genexpr>ö  s      z#llvm100_to_70_ir.<locals>.<genexpr>rÅ   )	rÆ   rÇ   Úre_attributes_defÚmatchÚgroupÚsplitr|   r{   rv   )r	   ÚbufÚlineÚmÚattrsr   r   r   rÕ   ì  s    

rÕ   c                 C   sN  dd„ }g }|   ¡ D �],}| d¡r0| dd¡}| ¡  d¡rRd|krR| dd¡}t |¡rÌdt |¡krÌ| d	d
¡}| dd¡}| d¡}|d|d … ||d d…  }}dd„ }d 	|t
 ||¡f¡}| d¡rØqt |¡dk	röt dd„ |¡}| d¡�r@t |¡}| d¡ ¡ }	d 	dd„ |	D ƒ¡}	| | d¡|	¡}d|k�r�t |¡}|dk�rltd|f ƒ‚| ¡ }
|d|
… |||
d… ƒ }d|k�rÎt |¡}|�rÎ| ¡ }
|d|
… |||
d… ƒ }d|k�rt d|¡}t d|¡ d¡}d|k�rt t|¡}d |k�r,d|k�r,t t|¡}t d|¡}| |¡ qd! 	|¡S )"z,
    Convert LLVM 10.0 IR for LLVM 3.4.
    c                 S   s†   d}d}t  | |¡}|d kr,td| f ƒ‚qv| ¡ }| d¡}|dkrR|dkrtqvq|dkrd|d7 }q|dkr|d8 }q| |d …  ¡ S )Nr   zfailed parsing leading type: %sú,z{[(r
   z)]})Úre_type_tokÚsearchrª   Úendré   Úlstrip)ÚsZ	par_levelÚposrí   Útokr   r   r   Úparse_out_leading_type  s     


z0llvm100_to_34_ir.<locals>.parse_out_leading_typez!numba.llvm.dbg.cuz!llvm.dbg.cuz%tail call void asm sideeffect "// dbgz
!numba.dbgz!dbgNz!{zmetadata !{z!"zmetadata !"ú=r
   c                 S   s   d|   d¡ S )Nz	metadata r   )ré   ©rí   r   r   r   Úfix_metadata_ref/  s    z*llvm100_to_34_ir.<locals>.fix_metadata_refrà   zsource_filename =c                 S   s   dS )NrŒ   r   rù   r   r   r   Ú<lambda>6  ó    z"llvm100_to_34_ir.<locals>.<lambda>rã   c                 s   s   | ]}|t kr|V  qd S rH   )Úsupported_attributesrä   r   r   r   ræ   <  s      z#llvm100_to_34_ir.<locals>.<genexpr>zgetelementptr z failed parsing getelementptr: %szload zcall z\1*\2rŒ   rï   z@llvm.memsetZdeclarerÅ   )rÆ   rÇ   r{   ró   Úre_metadata_defrè   Úre_metadata_correct_usagerñ   Úfindr|   Úre_metadata_refrÑ   Úre_unsupported_keywordsrç   ré   rê   Úre_getelementptrrª   rò   Úre_loadÚre_callÚre_rangeÚrstripÚre_parenthesized_listÚ_replace_llvm_memset_usageÚ _replace_llvm_memset_declarationÚre_annotationsrv   )r	   r÷   rë   rì   ZassigposÚlhsÚrhsrú   rí   rî   rõ   r   r   r   rÖ   þ  sp    
ÿ

"
ÿ








þ

þrÖ   c                 C   sb   t |  d¡ d¡ƒ}t d|d ¡}|s2tdƒ‚n
| d¡}| dd |¡¡ d |¡}d	 |¡S )
zNReplace `llvm.memset` usage for llvm7+.

    Used as functor for `re.sub.
    r
   rï   zalign (\d+)r   z+No alignment attribute found on memset destr£   zi32 {}ro   ú({}))	Úlistré   rê   Úrerñ   Ú
ValueErrorÚinsertr€   r|   )rí   r¸   Z
align_attrÚalignÚoutr   r   r   r	  i  s    


r	  c                 C   s4   t |  d¡ d¡ƒ}| dd¡ d |¡}d |¡S )zTReplace `llvm.memset` declaration for llvm7+.

    Used as functor for `re.sub.
    r
   rï   r£   rÐ   ro   r  )r  ré   rê   r  r|   r€   )rí   r¸   r  r   r   r   r
  y  s    
r
  c                 C   sZ   | j }t |d¡}t t d¡d¡}| | ||f¡}t |d¡}| |¡ | j	 
d¡ d S )NZkernelé    r
   znvvm.annotationsrá   )Úmoduler	   ZMetaDataStringÚConstantÚIntTypeÚadd_metadatar   Zget_or_insert_named_metadatarÎ   Ú
attributesÚdiscard)ZlfuncrÞ   ZmdstrZmdvalueÚmdZnmdr   r   r   Úset_cuda_kernel„  s    
r  c                    s<   t  d¡‰ ‡ fdd„tƒ  ¡ D ƒ}|  |¡}|  d|¡ dS )zAdd NVVM IR version to moduler  c                    s   g | ]}ˆ |ƒ‘qS r   r   )rr   r‰   ©rÐ   r   r   rt   —  s     z"add_ir_version.<locals>.<listcomp>znvvmir.versionN)r	   r  r   r9   r  Zadd_named_metadata)rÞ   r?   Zmd_verr   r  r   Úadd_ir_version“  s    

r  )drZ   Úloggingr  rR   r„   Úctypesr   r   r   r   r   r   r   Ú	threadingZllvmliter	   rT   r   r   r   Zlibsr   r   r   Z
numba.corer   r   Ú	getLoggerrW   ÚloggerZADDRSPACE_GENERICZADDRSPACE_GLOBALZADDRSPACE_SHAREDZADDRSPACE_CONSTANTZADDRSPACE_LOCALr\   r[   rê   rP   r¤   r¥   rˆ   r0   ÚmodulesrE   rF   r   ÚLockr&   r(   r   r^   r�   r›   rŸ   r<   r¦   r§   r«   r¨   rÔ   r¯   r°   r·   rÀ   rÃ   r¼   r±   r¹   r½   rÂ   rÄ   rÌ   rÙ   rß   r‹   rþ   rÿ   r  Zdebuginfo_patternr{   Zre_metadata_debuginforç   rý   r  r  r  r  rð   r  r  r  Z
re_spflagsZ	spflagmaprÕ   rÖ   r	  r
  r  r  r   r   r   r   Ú<module>   sÂ   $
  	õ"
 ?



      þ








þk