
    J/Ph+^                     ~   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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
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:64ze-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        W/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/cudadrv/nvvm.pyis_availabler   <   s;     t    uu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"6DKK 7 7 7%)CN2F*6A:666	7 $(#3#9#9#;#; . .KD%"4;55D#(8DL$)!""IDM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"    M4=)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[    uU||U5\\::6777{EK''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    ''''7777  ww!&x%//C C9:::}gmX^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!25!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        
 $WWU^^WWU^^D *77<+@+@A  +GGL,A,AB
 x8#E x8'E uggh.?.?A wwx'8'8&: #.|X!F #.|WWX=N=N!O *<B &wwu~~wwu~~!'%..''%..: *<%gh//1c3 3Kl 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    ff#~~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    ''dl(;(;<<3 6T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v14V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4T\658[[$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1ww!T"" FFq;;1;;r   c                 .    g | ]\  }} ||          S r   r   ).0r   r   r   s      r   
<listcomp>z+CompilationUnit.compile.<locals>.<listcomp>  s+    FFFda##Aq))FFFr   c                 R    g | ]$}t          |                    d                     %S )utf8)r   encode)r   xs     r   r   z+CompilationUnit.compile.<locals>.<listcomp>
  s<     -? -? -?12 .6ahhv6F6F-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output_bufferr   s         @r   compilezCompilationUnit.compile   s   (		 		 		 GFFFgmmooFFFS\\) -? -?6=-? -? -? @ k++DL#g,,OO1222 k,,T\3w<<PP2333 k33DL%--PPEFFF&,.11k//mLL=>>> <<>>8 	:M$([9999QQQr   c                 j    | j                             ||d|                                            d S )N
)r0   rU   r   )rH   rY   rf   s      r   r   zCompilationUnit._try_error%  s3    dllnnn%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//eFmmLL%JKKK<!v|+..F+//fEEC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@   )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>U  sD     1 1 1R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>Z  s/     ; ; ;Rv999 999r   )CTK_SUPPORTEDtupleCOMPUTE_CAPABILITIESKeyError)ctk_versionr   r   s    @@r   ccs_supported_by_ctkr   Q  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8M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   w  s     FF(M $Qs###=)) , ,2::III$YYAvv+.2Ri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   aN  Missing libdevice file.
Please ensure you have a CUDA Toolkit 11.2 or higher.
For CUDA 12, ``cuda-nvcc`` and ``cuda-nvrtc`` are required:

    $ conda install -c conda-forge cuda-nvcc cuda-nvrtc "cuda-version>=12.0"

For CUDA 11, ``cudatoolkit`` is required:

    $ conda install -c conda-forge cudatoolkit "cuda-version>=11.2,<12.0"
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=    <&"#=>>>)++DL,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"$F2JJ???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   3  #    '..rvayy.IIIr   c                 V    t                               | |t          |                     S r   )ir_numba_atomic_dec_templater   r   r   s     r   ir_numba_atomic_decr   7  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  ;  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
+0u
> 
> 
>	? 
U	(u"+0u
> 
> 
>	? 
Q	'e
+0u
> 
> 
>	? 
U	(u"+0u
> 
> 
>	? 
T	'e5+-E
; 
; 
;	< 
X	(u%E+-E
; 
; 
;	< 
T	'e5+-E
; 
; 
;	< 
X	(u%E+-E
; 
; 
;	< 	G$LL ! * *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   
compile_irr  j  s    &# xx
E"" 	
 
 	 	 	 
		BI * *3
cjj(())))y}}'''2:r   z"^attributes #\d+ = \{ ([\w\s]+)\ }c                    g } | j                     D ]}|                    d          rt                              |          }|                    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C1l1B1BQ1B1B1B1B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GGAJJ$$&&EHHCCCCCCCE<<

E22D

499S>>r   c                 b   | j         }t          j        |d          }t          j        t          j        d          d          }|                    | ||f          }t          j        |d          }|                    |           t          j        d          	                                }t          j
        |d          }|                     |          }t          j        ||d          }	d|	_        d|	_        t          j        ||g          |	_        | j                            d	           d
S )al  
    Mark a function as a CUDA kernel. Kernels have the following requirements:

    - Metadata that marks them as a kernel.
    - Addition to the @llvm.used list, so that they will not be discarded.
    - The noinline attribute is not permitted, because this causes NVVM to emit
      a warning, which counts as failing IR verification.

    Presently it is assumed that there is one kernel per module, which holds
    for Numba-jitted functions. If this changes in future or this function is
    to be used externally, this function may need modification to add to the
    @llvm.used list rather than creating it.
    kernel    r   znvvm.annotationsrL   z	llvm.used	appendingzllvm.metadatanoinlineN)moduler
   MetaDataStringConstantIntTypeadd_metadatar   get_or_insert_named_metadatar   
as_pointer	ArrayTypebitcastGlobalVariablelinkagesectioninitializer
attributesdiscard)
functionr.  mdstrmdvaluemdnmdptrtyusedtyfnptr	llvm_useds
             r   set_cuda_kernelrF    s    _F fh//Ek"*R..!,,G			hw7	8	8B

.v7I
J
JCGGBKKK JqMM$$&&E\%##FU##E!&&+>>I#I'IK88I 
+++++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!    ;;;a33q66;;;r   znvvmir.versionN)r
   r1  r   rA   r2  add_named_metadata)r  rI   md_verr   s      @r   add_ir_versionrK    sk     *R..C;;;;466#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  rF  rK  r   r   r   <module>r\     s     				 



                             ; ; ; ; ; ; ; ; ; ; = = = = = = = = = = & & & & & & & & 
	8	$	$       
EGG  I'(( ) )DAqGCK!1a((((; 7 
	 	 	 Y^
{ { { { {6 { { {|k k k k kf k k k\    
< 
< 
<  2  D! ! !	        # ,  (  (# @" " "< < <
< < <J J JJ J J, , ,^  . BJDEE   $$, $, $,N5 5 5 5 5r   