xKg+^ jdZddlZddlZddlZddlZddlmZmZmZm Z m Z m Z m Z ddl Z ddlmZddlmZmZmZddlmZmZmZddlmZmZej6eZdZdZd Z d Z!d Z"eZ#eZ$d jKZ&e'e&D]\Z(Z)e*ejVee)e(d Z,dZ-dZ.e j^Z0Gdde1Z2Gdde1Z3dZ4ddddddddddddd Z5dZ6dZ7dZ8dZ9dZ:Gdd e1Z;d!Zd$Z?d%Z@d&ZAd'ZBd(ZCd)ZDd*ZEd+ZFd,ZGejd-ZId.ZJd/ZKd0ZLy)1z( This is a direct translation of nvvm.h N)c_void_pc_intPOINTERc_char_pc_size_tbyrefc_char)ir) NvvmErrorNvvmSupportError NvvmWarning) get_libdeviceopen_libdevice open_cudalib)cgutilsconfiga 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:64c8 ty#t$rYywxYw)z( Return if libNVVM is available TF)NVVMr [/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/cudadrv/nvvm.py is_availabler<s%  s cVeZdZdZeeeeefeeefeeefeeee efeeee efeeeeefeeee feeefeeee feeefeeeeeeeeefeeeeefd Z dZ dZ dZ edZedZdZd Zd d Zy) rzProcess-wide singleton. ) nvvmVersionnvvmCreateProgramnvvmDestroyProgramnvvmAddModuleToProgramnvvmLazyAddModuleToProgramnvvmCompileProgramnvvmGetCompiledResultSizenvvmGetCompiledResultnvvmGetProgramLogSizenvvmGetProgramLog nvvmIRVersionnvvmVerifyProgramNct5|jtj|x|_} t d|_|jjD]<\}}t|j |}|d|_ |dd|_ t|||>ddd|jS#t $r}d|_d}t||zd}~wwxYw#1swY|jSxYw)Nnvvmz;libNVVM cannot be found. Do `conda install cudatoolkit`: %srr ) _nvvm_lock_NVVM__INSTANCEobject__new__rdriverOSErrorr _PROTOTYPESitemsgetattrrestypeargtypessetattr)clsinsteerrmsgnameprotofuncs rr/z NVVM.__new__s ~~%(.s(;;7".v"6DK$(#3#3#9#9#;KD%"4;;5D#(8DL$)!"IDMD$- $<$~~7%)CN2F*6A:66 7 $~~s/)CB.AC. C7CCCC*c|j}|d|_|d|_|d|_|d|_t |_y)Nrr r)get_ir_version_majorIR_minorIR _majorDbg _minorDbgget_supported_ccs_supported_ccs)self ir_versionss r__init__z NVVM.__init__sG))+ #A #A $Q$Q/1rcP|j|jfdkrtStS)N)r )rBrC_datalayout_original_datalayout_i128rHs r data_layoutzNVVM.data_layouts# MM4== )F 2' '# #rc|jSN)rGrOs r supported_ccszNVVM.supported_ccss"""rct}t}|jt|t|}|j|d|j|jfS)NzFailed to get version.)rrr check_errorvalue)rHmajorminorerrs r get_versionzNVVM.get_versionsNuU|U5\: 67{{EKK''rcBt}t}t}t}|jt|t|t|t|}|j|d|j|j|j|jfS)NzFailed to get IR version.)rr(rrUrV)rHmajorIRminorIRmajorDbgminorDbgrYs rrAzNVVM.get_ir_versionsx''77  ww!&x%/C 9:}}gmmX^^X^^KKrcx|r8t|t|}|r!t|tjdy|y)Nr )r RESULT_CODE_NAMESprintsysexit)rHerrormsgrdexcs rrUzNVVM.check_errors6 C!25!9:Cc   r)F)__name__ __module__ __qualname____doc__ nvvm_resultrr nvvm_programrrr2r-r/rJpropertyrPrSrZrArUrrrrrKs< $WU^WU^D*7<+@A +GL,AB x8#E x8'E ugh.?A wx'8&:#.|X!F#.|WX=N!O*<B&wu~wu~!%.'%.:*<%h/1c3KlJ*2$$ ##(Lrrc6eZdZdZdZdZdZdZdZdZ y) CompilationUnitct|_t|_|jj t |j}|jj |dy)NzFailed to create CU)rr0rm_handlerrrU)rHrYs rrJzCompilationUnit.__init__sFf #~ kk++E$,,,?@ %:;rct}|jt|j}|j |ddy)NzFailed to destroy CUT)rd)rr rrrrU)rHr0rYs r__del__zCompilationUnit.__del__s8''dll(;<3 6TBrc|jj|j|t|d}|jj |dy)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. NFailed to add module)r0r!rrlenrUrHbufferrYs r add_modulezCompilationUnit.add_modules? kk00v14VdD %;=". - 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 "-". For documentation on NVVM compilation options, see the CUDA Toolkit Documentation: https://docs.nvidia.com/cuda/libnvvm-api/index.html#_CPPv418nvvmCompileProgram11nvvmProgramiPPKc cz|jdd}|d|St|tr t|}d|d|S)N_-=)replace isinstanceboolint)kvs rstringify_optionz1CompilationUnit.compile..stringify_optionsG #s#Ay1#w!T"Fqc1#; rutf8zFailed to verify zFailed to compile z&Failed to get size of compiled result.zFailed to get compiled result.)categoryN)r3rrwencoder0r)rr _try_errorr#rr$rr rVr%get_loglogwarningswarnr) rHoptionsrrrxc_optsrYreslen output_buffers rcompilezCompilationUnit.compiles|( 7>mmoFoda#Aq)oFS\)6=-?6=.6ahhv6F-G6=-?@kk++DLL#g,O 12kk,,T\\3w<P 23kk33DLL%-P EF&,,.1 kk// mL =><<> 88 MM$(([ 9Q7G-?s G!Gcb|jj||d|jy)N )r0rUr)rHrYrfs rrzCompilationUnit._try_error%s! dlln%EFrct}|jj|jt |}|jj |d|j dkDrtt|j z}|jj|j|}|jj |d|j jdSy)Nz#Failed to get compilation log size.r zFailed to get compilation log.r) rr0r&rrrrUrVr r'decode)rHrrYlogbufs rrzCompilationUnit.get_log(skk// eFmL %JK <9>>BA1 0 B?Bc ddlm}|j}tt}||kr8d}|dd|d}d|d|dd|dd}t j ||St|}|S#d}|cYSxYw) Nr)runtimer.r z CUDA Toolkit z is unsupported by Numba - z! is the minimum required version.)numba.cuda.cudadrv.runtimerrZminrrrr)rcudart_version _supported_cc min_cudartctk_verunsupported_vers rrFrF^s6 ,,.]#J " #A&'q):(;<*7)3N(m_Ajm_=//  o&(8M % s A11A9ctj}|s d}t|t|D]5\}}||k(r|cS||kDs|dk(rd||zz}t|||dz cS|dS)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.rz?GPU compute capability %d.%d is not supported(requires >=%d.%d)r )rrSr enumerate)myccrSrfirs rfind_closest_archrwsF((M Qs##=)2 :I $YAv+.2Ri9&s++%QU++*  rcptjrtj}d|zSt||f}d|zS)z1Matches with the closest architecture option z compute_%d%d)r FORCE_CUDA_CCr)rWrXarchs rget_arch_optionrs@## D  !%0 D  raNMissing 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" ceZdZdZdZdZy) LibDeviceNc|j(tttt |_|j|_yrR)_cache_r RuntimeErrorMISSING_LIBDEVICE_FILE_MSGrbcrOs rrJzLibDevice.__init__s5 << &"#=>>)+DL,,rc|jSrR)rrOs rgetz LibDevice.gets wwr)rhrirjrrJrrrrrrsGrrz %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.tj|S)NTi)cas_nvvmformatrs rir_casr#s ??b? !!rc ^t||||t|}tjdi|S)N)TrOPFUNCCASr)dictrir_numba_atomic_binary_templater)rrrrparamss rir_numba_atomic_binaryr's, A"$F2J ?F * 1 1 ;F ;;rc bt||||||t|}tjdi|S)N)rrNANr PTR_OR_VALrrr)rrir_numba_atomic_minmax_templater)rrrrrrrs rir_numba_atomic_minmaxr,s5 A"#"-F + 1 1 ;F ;;rcDtj||t|SN)rTur)ir_numba_atomic_inc_templaterrrrs rir_numba_atomic_incr3 ' . .rvay . IIrcDtj||t|Sr)ir_numba_atomic_dec_templaterrrs rir_numba_atomic_decr7rrcdtddddfdtdd d d fd tddd d fd tddfdtddfdtdd ddddfdtddddddfdtdd ddddfdtddddddfdtdd ddddfdtddddddfd tdd dd!ddfd"tdddd!ddfd#g}|D]\}}|j ||}t |}|S)$NzIdeclare double @"___numba_atomic_double_add"(double* %".1", double %".2")doublei64faddadd)rrrrzEdeclare float @"___numba_atomic_float_sub"(float* %".1", float %".2")floati32fsubsubzIdeclare double @"___numba_atomic_double_sub"(double* %".1", double %".2")z=declare i64 @"___numba_atomic_u64_inc"(i64* %".1", i64 %".2")u64rz=declare i64 @"___numba_atomic_u64_dec"(i64* %".1", i64 %".2")zEdeclare float @"___numba_atomic_float_max"(float* %".1", float %".2")rznnan oltptrmax)rrrrrrzIdeclare double @"___numba_atomic_double_max"(double* %".1", double %".2")zEdeclare float @"___numba_atomic_float_min"(float* %".1", float %".2")znnan ogtrzIdeclare double @"___numba_atomic_double_min"(double* %".1", double %".2")zHdeclare float @"___numba_atomic_float_nanmax"(float* %".1", float %".2")nanultzLdeclare 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)rrrrrllvm140_to_70_ir)llvmir replacementsdeclfns r llvm_replacer;s T (ue L N P 'eU K M T (ue L N H u / 1 H u / 1 P 'e +0u > ? U (u"+0u > ? Q 'e +0u > ? U (u"+0u > ? T 'e5+-E ; < X (u%E+-E ; < T 'e5+-E ; < X (u%E+-E ; < G$LL!bb)!f %F Mrc ft|tr|g}|jddr|jdddddt }t }|D]-}t |}|j|jd/|j|j|jdi|S)NfastmathFT)ftzfmaprec_div prec_sqrtrr) rstrpopupdaterprrrzrr|rr)roptscu libdevicemods r compile_irrjs&# xx E"      B I3 cjj()y}}' 2::  rz"^attributes #\d+ = \{ ([\w\s]+)\ }crg}|jD]}|jdrmtj|}|j dj }dj d|D}|j|j d|}|j|dj |S)z, Convert LLVM 14.0 IR for LLVM 7.0. z attributes #r  c3,K|] }|dk7s |yw) willreturnNr).0as r z#llvm140_to_70_ir..sC1l1BQs r) splitlines startswithre_attributes_defmatchgroupsplitjoinrappend)r buflinemattrss rrrs C  ??> *!''-AGGAJ$$&EHHCCCE<< E2D 4  99S>rc\|j}tj|d}tjtjdd}|j |||f}t j|d}|j|tjdj}tj|d}|j|}tj||d} d| _ d| _tj||g| _|j j#d y ) 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 znvvm.annotationsrLz llvm.used appendingz llvm.metadatanoinlineN)moduler MetaDataStringConstantIntType add_metadatarget_or_insert_named_metadatar as_pointer ArrayTypebitcastGlobalVariablelinkagesection initializer attributesdiscard) functionr)mdstrmdvaluemdnmdptrtyusedtyfnptr llvm_useds rset_cuda_kernelrAs__F   fh /Ekk"**R.!,G   hw7 8B  . .v7I JCGGBK JJqM $ $ &E \\% #F   U #E!!&&+>I#I'IKK8I  +rctjd}tjDcgc] }|| }}|j |}|j d|ycc}w)zAdd NVVM IR version to moduler&znvvmir.versionN)r r,rrAr-add_named_metadata)rrrrImd_vers radd_ir_versionrEs\ **R.C#'6#8#8#:;#:a3q6#:K;   k *F+V4rVs ;;==&   8 $    EG ' (DAq CKK !1a( ) ;7 Y^^  {6{|kfk\             <2D!    #, ( (#@"< <JJ,^.BJJDE$$,N5r