o
    8ήcv                     @   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 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' qcdd Z+e, Z-G dd de.Z/G dd de.Z0dZ1da2dd Z3dd Z4dd Z5dZ6G dd de.Z7dZ8dZ9dZ:d Z;d!Z<d"Z=d#Z>d$d% Z?d&d' Z@d(d) ZAd*d+ ZBd,d- ZCd.d/ ZDd0d1 ZEd2d3 ZFeGd4ZHeGd5ZIeGd6ZJd7ZKeGeKLd8d9ZMeGd:ZNh d;ZOeGd<ZPeGd=ZQeGd>ZReGd?ZSeGd@ZTeGdAZUeGdBZVeGdCZWeGdDZXdEdFdGZYdHdI ZZdJdK Z[dLdM Z\dNdO Z]dPdQ Z^dRdS Z_dTdU Z`dS )Vz(
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)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
c                   C   s"   zt   W dS  ty   Y dS w )z(
    Return if libNVVM is available
    FT)NVVMr    r   r   >/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/cudadrv/nvvm.pyis_available3   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dZ
dZdd Z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nvvmIRVersionNc                 C   s   t a | jd u rQt|  | _}ztd|_W n ty. } zd | _d}t|| d }~ww |j	 D ]%\}}t
|j|}|d |_|dd  |_t||| q4W d    | jS W d    | jS 1 sfw   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funcr   r   r   r(   y   s0   



zNVVM.__new__c                 C   s4   |   }|d | _|d | _|d | _|d | _d S )Nr   r
      r   )get_ir_version_majorIR_minorIR	_majorDbg	_minorDbg)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   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   rD   rE   )r>   majorIRminorIRmajorDbgminorDbgrH   r   r   r   r9      s   zNVVM.get_ir_versionFc                 C   s4   |rt |t| }|rt| td d S |d S )Nr
   )r   RESULT_CODE_NAMESprintsysexit)r>   errormsgrQ   excr   r   r   rD      s   zNVVM.check_error)F)__name__
__module____qualname____doc__nvvm_resultr   r   nvvm_programr   r   r+   r&   r(   r@   propertyrC   rI   r9   rD   r   r   r   r   r   B   s8    




2

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)   rZ   _handler   r   rD   )r>   rH   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)rQ   )r   r   r   r]   rD   )r>   r)   rH   r   r   r   __del__   s   zCompilationUnit.__del__c                 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)r)   r   r]   lenrD   r>   bufferrH   r   r   r   
add_module      zCompilationUnit.add_modulec                 C   r_   )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.
        Nr`   )r)   r   r]   ra   rD   rb   r   r   r   lazy_add_module   re   zCompilationUnit.lazy_add_modulec                 K   sx  g }d|v r| dr|d | ddr|d d|v r)|d| d  |dr8|d	| d  d
}|D ]}||v rXtt| |}|d|dd|f  q<|rldtt|	 }t
d|tt| dd |D  }| j| jt||}| |d t }	| j| jt|	}| |d t|	j  }
| j| j|
}| |d |  | _|
dd S )aF  Perform Compilation

        The valid compiler options are

         *   - -g (enable generation of debugging information)
         *   - -opt=
         *     - 0 (disable optimizations)
         *     - 3 (default, enable optimizations)
         *   - -arch=
         *     - compute_20 (default)
         *     - compute_30
         *     - compute_35
         *   - -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)
         *
         debugz-glineinfoFz-generate-line-info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 compile
z&Failed to get size of compiled result.zFailed to get compiled result.N)popappendgetintboolreplacejoinmapreprkeysr   formatr   ra   r)   r   r]   
_try_errorr   r   r   r   rE   r    get_loglog)r>   optionsoptsother_optionskvoptstrc_optsrH   reslenptxbufr   r   r   compile   s@   #




zCompilationUnit.compilec                 C   s   | j |d||  f  d S )Nz%s
%s)r)   rD   r   )r>   rH   rS   r   r   r   r   .  s   zCompilationUnit._try_errorc                 C   sl   t  }| j| jt|}| j|d |jdkr4t|j  }| j| j|}| j|d |j	dS dS )Nz#Failed to get compilation log size.r
   zFailed to get compilation log.rr    )
r   r)   r!   r]   r   rD   rE   r   r"   decode)r>   r   rH   logbufr   r   r   r   1  s   
zCompilationUnit.get_logN)
rU   rV   rW   r@   r^   rd   rf   r   r   r   r   r   r   r   r\      s    

Wr\   ze-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:64c                  C   s   t rt S zddlm}  |  }W n   d}Y |d  d|d  }d| dd }|dkr2d	a t S |d
kr:da t S |dkrBda t S |dkrJda t S |dkrRda t S d	a t| t S )Nr   )runtime)r   r   .r
   zCUDA Toolkit z is unsupported by Numba - z%10.2 is the minimum required version.r   )
   r8   )r   r   r      r   r   r   r8   r   r   rA   r   rA   r
   rA   r8   r   r   r   r8   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   rA   )r   r   )r   r   r   r   r   r   r   r   r   r   r   r   r   )r   r   )_supported_ccnumba.cuda.cudadrv.runtimer   rI   warningswarn)r   cudart_versionctk_verunsupported_verr   r   r   get_supported_ccsI  s:   
	
r   c                 C   st   t  }|sd}t|t|D ]&\}}|| kr|  S || kr5|dkr-d| |  }t|||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   r   	enumerate)myccsupported_ccrS   iccr   r   r   find_closest_arch{  s    r   c                 C   s(   t jr
t j}d| S t| |f}d| S )z1Matches with the closest architecture option
    zcompute_%d%d)r   FORCE_CUDA_CCr   )rF   rG   rj   r   r   r   get_arch_option  s
   r   z~Missing libdevice file.
Please ensure you have package cudatoolkit >= 10.2
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 u rt d u rttt | _ | j | _d S N)_cache_r   RuntimeErrorMISSING_LIBDEVICE_FILE_MSGr   bcrB   r   r   r   r@     s
   

zLibDevice.__init__c                 C   s   | j S r   )r   rB   r   r   r   ry     s   zLibDevice.get)rU   rV   rW   r   r@   ry   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 )NTi)r   rC   
cas_nvvm70r   
cas_nvvm34r   r   r   r   ir_cas/  s   r   c                 C   s&   t | |||t|d}tjdi |S )N)Tr   OPFUNCCASr   )dictr   ir_numba_atomic_binary_templater   )r   r   r   r   paramsr   r   r   ir_numba_atomic_binary6  s   r   c              	   C   s*   t | |||||t|d}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   r   r   r   ir_numba_atomic_minmax;  s   r   c                 C      t j| |t| dS N)r   Tur   )ir_numba_atomic_inc_templater   r   r   r   r   r   r   ir_numba_atomic_incB     r   c                 C   r   r   )ir_numba_atomic_dec_templater   r   r   r   r   r   ir_numba_atomic_decF  r   r   c                 C   sD   |   }t|D ]\}}|drd}|t||<  nqd|S )z@
    Find the line containing the datalayout and replace it
    ztarget datalayoutztarget datalayout = "{0}"
)
splitlinesr   
startswithr   data_layoutr}   )llvmirlinesr   lntmpr   r   r   _replace_datalayoutJ  s   

r   c                 C   s  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 ]
\}}| 	||} q| 	d'd(} t jrt
| } | S 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 ogtminzIdeclare 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   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")zllvm.numba_nvvm.atomiczllvm.nvvm.atomic)r   r   r   r   r   rC   ir_numba_cas_hackr   r   r|   llvm100_to_70_irllvm100_to_34_ir)r   replacementsdeclfnr   r   r   llvm_replaceW  s   









&

r   c                 K   s|   t | tr| g} |ddr|ddddd t }t }| D ]}t|}||d q |	|
  |jdi |S )NfastmathFT)rk   rn   rm   rl   rr   r   )
isinstancestrrw   updater\   r   r   rd   rs   rf   ry   r   )r   r   cu	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]+)\ }>   coldminsizeoptiszeoptnonenoinlinenoreturnnounwindreadnonereadonly
inlinehintnoduplicatealwaysinlinez"\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: (.*),isDefinitionisOptimized)DISPFlagDefinitionDISPFlagOptimizedc                 C   sn   g }|   D ]+}|dr,t|}|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 )
willreturnNr   rt   ar   r   r   	<genexpr>      z#llvm100_to_70_ir.<locals>.<genexpr>r   )	r   r   re_attributes_defmatchgroupsplitr}   r|   rx   )r	   buflinemattrsr   r   r   r     s   


r   c                 C   s@  dd }g }|   D ]}|dr|dd}| dr)d|v r)|dd}t|rfdt|u rf|d	d
}|dd}|d}|d|d  ||d d }}dd }d	|t
||f}|drlq
t|dur{tdd |}|drt|}|d }	d	dd |	D }	||d|	}d|v rt|}|du rtd|f | }
|d|
 |||
d  }d|v rt|}|r| }
|d|
 |||
d  }d|v rtd|}td|d}d|v rtt|}d |v rd|v 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 u rtd| f | }|d}|dkr)|dkr(nn|dv r2|d7 }n|dv r:|d8 }q| |d   S )Nr   Tzfailed parsing leading type: %s,z{[(r
   z)]})re_type_toksearchr   endr  lstrip)s	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>   s    z"llvm100_to_34_ir.<locals>.<lambda>r  c                 s   s    | ]	}|t v r|V  qd S r   )supported_attributesr  r   r   r   r  &  r  z#llvm100_to_34_ir.<locals>.<genexpr>zgetelementptr z failed parsing getelementptr: %szload zcall z\1*\2r   r  z@llvm.memset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_annotationsrx   )r	   r&  r  r  assigposlhsrhsr)  r  r  r$  r   r   r   r     sp   


"








r   c                 C   s`   t | dd}td|d }|std|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 {}rq   ({}))	listr  r  rer  
ValueErrorinsertr   r}   )r  r   
align_attralignoutr   r   r   r8  S  s   


r8  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   rq   r>  )r?  r  r  rB  r}   r   )r  r   rE  r   r   r   r9  c  s   

r9  c                 C   sN   | j }t|d}ttdd}|| ||f}t|d}|| d S )Nkernel    r
   znvvm.annotations)	moduler	   MetaDataStringConstantIntTypeadd_metadatar   get_or_insert_named_metadatar   )lfuncr   mdstrmdvaluemdnmdr   r   r   set_cuda_kerneln  s   rS  c                 C   sf   t d}t jr|d|d|d|dg}n|d|d|d|dg}| |}| d| dS )	zAdd NVVM IR version to modulerG  r
   rA   r   r   r8   znvvmir.versionN)r	   rK  r   rC   rL  add_named_metadata)r   r   r?   md_verr   r   r   add_ir_versiony  s   

rV  c                 C   s
   t | _ d S r   )r   )rH  r   r   r   fix_data_layout  s   
rW  )arX   loggingr@  rP   r   ctypesr   r   r   r   r   r   r   	threadingllvmliter	   rR   r   r   libsr   r   r   
numba.corer   r   	getLoggerrU   loggerADDRSPACE_GENERICADDRSPACE_GLOBALADDRSPACE_SHAREDADDRSPACE_CONSTANTADDRSPACE_LOCALrZ   rY   r  rN   r   r   r   r0   modulesr   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.  r0  debuginfo_patternr|   re_metadata_debuginfor  r+  r2  r3  r4  r5  r  r:  r1  r7  
re_spflags	spflagmapr   r   r8  r9  rS  rV  rW  r   r   r   r   <module>   s    $
u 2"
 E












k