o Y۷i+^ @sdZddlZddlZddlZddlZddlmZmZmZm Z m Z m Z m Z ddl Z ddlmZddlmZmZmZddlmZmZmZddlmZmZeeZdZdZd Z d Z!d Z"eZ#eZ$d %Z&e'e&D] \Z(Z)e*ej+ee)e(qed Z,dZ-ddZ.e /Z0Gddde1Z2Gddde1Z3dZ4ddddddddddddd Z5ddZ6ddZ7dd Z8d!d"Z9d#Z:Gd$d%d%e1Z;d&Zd)Z?d*Z@d+d,ZAd-d.ZBd/d0ZCd1d2ZDd3d4ZEd5d6ZFd7d8ZGeHd9ZId:d;ZJdd?ZLdS)@z( 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:64cCs"ztWdStyYdSw)z( Return if libNVVM is available FT)NVVMr rrM/home/ubuntu/vllm_env/lib/python3.10/site-packages/numba/cuda/cudadrv/nvvm.py is_available<s  rc@seZdZdZeeeeefeeefeeefeeee efeeee efeeeeefeeee feeefeeee feeefeeeeeeeeefeeeeefd Z dZ ddZ ddZ edd Zed d Zd d ZddZdddZdS)rzProcess-wide singleton. ) nvvmVersionnvvmCreateProgramnvvmDestroyProgramnvvmAddModuleToProgramnvvmLazyAddModuleToProgramnvvmCompileProgramnvvmGetCompiledResultSizenvvmGetCompiledResultnvvmGetProgramLogSizenvvmGetProgramLog nvvmIRVersionnvvmVerifyProgramNc Csta|jdurQt||_}ztd|_Wnty.}z d|_d}t||d}~ww|j D]%\}}t |j|}|d|_ |dd|_ t |||q4Wd|jSWd|jS1sfwY|jS)Nnvvmz;libNVVM cannot be found. Do `conda install cudatoolkit`: %srr ) _nvvm_lock_NVVM__INSTANCEobject__new__rdriverOSErrorr _PROTOTYPESitemsgetattrrestypeargtypessetattr)clsinsteerrmsgnameprotofuncrrrr*s0      z NVVM.__new__cCs<|}|d|_|d|_|d|_|d|_t|_dS)Nrr r)get_ir_version_majorIR_minorIR _majorDbg _minorDbgget_supported_ccs_supported_ccs)self ir_versionsrrr__init__s      z NVVM.__init__cCs|j|jfdkr tStS)N)r )r<r=_datalayout_original_datalayout_i128rBrrr data_layoutszNVVM.data_layoutcC|jSN)rArHrrr supported_ccsszNVVM.supported_ccscCs8t}t}|t|t|}||d|j|jfS)NzFailed to get version.)rrr check_errorvalue)rBmajorminorerrrrr get_versions   zNVVM.get_versioncCsXt}t}t}t}|t|t|t|t|}||d|j|j|j|jfS)NzFailed to get IR version.)rr$rrMrN)rBmajorIRminorIRmajorDbgminorDbgrQrrrr;s  zNVVM.get_ir_versionFcCs4|rt|t|}|rt|tddS|dS)Nr )r RESULT_CODE_NAMESprintsysexit)rBerrormsgrZexcrrrrMszNVVM.check_error)F)__name__ __module__ __qualname____doc__ nvvm_resultrr nvvm_programrrr-r(r*rDpropertyrIrLrRr;rMrrrrrKsB        6   rc@sDeZdZddZddZddZddZd d Zd d Zd dZ dS)CompilationUnitcCs4t|_t|_|jt|j}|j|ddS)NzFailed to create CU)rr+rc_handlerrrM)rBrQrrrrDszCompilationUnit.__init__cCs*t}|t|j}|j|ddddS)NzFailed to destroy CUT)rZ)rrrrfrM)rBr+rQrrr__del__szCompilationUnit.__del__cC*|j|j|t|d}|j|ddS)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)r+rrflenrMrBbufferrQrrr add_module zCompilationUnit.add_modulecCrh)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. Nri)r+rrfrjrMrkrrrlazy_add_modulernzCompilationUnit.lazy_add_modulec sddfdd|D}tt|dd|D}|j|jt||}||d|j|jt||}||dt}|j |jt |}||dt |j }|j |j|}||d ||_|jrutj|jtd |d d S) aPerform 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 "-=". - 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 cSs@|dd}|durd|St|trt|}d|d|S)N_-=)replace isinstanceboolint)kvrrrstringify_options   z1CompilationUnit.compile..stringify_optioncsg|] \}}||qSrr).0rwrxryrr sz+CompilationUnit.compile..cSsg|] }t|dqS)utf8)rencode)rzxrrrr| szFailed to verify zFailed to compile z&Failed to get size of compiled result.zFailed to get compiled result.)categoryN)r.rrjr+r%rf _try_errorrrr rrrNr!get_loglogwarningswarnr )rBoptionsc_optsrQreslen output_bufferrr{rcompiles&        zCompilationUnit.compilecCs|j|d||fdS)Nz%s %s)r+rMr)rBrQr\rrrr%szCompilationUnit._try_errorcCslt}|j|jt|}|j|d|jdkr4t|j}|j|j|}|j|d|j dSdS)Nz#Failed to get compilation log size.r zFailed to get compilation log.r}) rr+r"rfrrMrNrr#decode)rBrrQlogbufrrrr(s   zCompilationUnit.get_logN) r^r_r`rDrgrmrorrrrrrrres  < re)rr)rrr)rr:)rr)r)rr )rr:)rr)rr:)rr)rErrErrEr)rE rr)rr)rr)rr)rr) ) r:)rr)rr)rr)rr)rr)rrE) r)rr )rr:)rr)rrcsLzt|\tfddtDWSty%tddtDYSw)Ncs(g|]}|krkrnn|qSrrrzccmax_ccmin_ccrrr|Usz(ccs_supported_by_ctk..cSsg|] }|tjkr|qSr)rCUDA_DEFAULT_PTX_CCrrrrr|Zs ) CTK_SUPPORTEDtupleCOMPUTE_CAPABILITIESKeyError) ctk_versionrrrccs_supported_by_ctkQs   rcCsz ddlm}|}Wn d}|YStt}||krBd}|dd|d}d|d|dd|dd}t||St|}|S) Nr)runtimer.r z CUDA Toolkit z is unsupported by Numba - z! is the minimum required version.)numba.cuda.cudadrv.runtimerrRminrrrr)rcudart_version _supported_cc min_cudartctk_verunsupported_verrrrr@^s&   r@cCsvtj}|s d}t|t|D]&\}}||kr|S||kr6|dkr.d||}t|||dSq|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 )rrLr enumerate)myccrLr\irrrrfind_closest_archws  rcCs(tjr tj}d|St||f}d|S)z1Matches with the closest architecture option z compute_%d%d)r FORCE_CUDA_CCr)rOrParchrrrget_arch_options  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" c@s eZdZdZddZddZdS) LibDeviceNcCs0|jdurtdurttt|_|j|_dSrK)_cache_r RuntimeErrorMISSING_LIBDEVICE_FILE_MSGrbcrHrrrrDs   zLibDevice.__init__cCrJrK)rrHrrrgetsz LibDevice.get)r^r_r`rrDrrrrrrs rz %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 }} cCs tj|dS)NTi)cas_nvvmformatrrrrir_cas#s rcCs&t||||t|d}tjdi|S)N)TrOPFUNCCASr)dictrir_numba_atomic_binary_templater)rrrrparamsrrrir_numba_atomic_binary'src Cs*t||||||t|d}tjdi|S)N)rrNANr PTR_OR_VALrrr)rrir_numba_atomic_minmax_templater)rrrrrrrrrrir_numba_atomic_minmax,s rcCtj||t|dSN)rTur)ir_numba_atomic_inc_templaterrrrrrrir_numba_atomic_inc3rcCrr)ir_numba_atomic_dec_templaterrrrrrir_numba_atomic_dec7rrcCs.dtdddddfdtdd d d dfd tddd d dfd tdddfdtdddfdtdd dddddfdtdddddddfdtdd dddddfdtdddddddfdtdd dddddfdtdddddddfd tdd dd!dddfd"tdddd!dddfd#g}|D] \}}|||}qt|}|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)rrrrrsllvm140_to_70_ir)llvmir replacementsdeclfnrrr llvm_replace;sz           &rcKs|t|tr|g}|ddr|dddddt}t}|D]}t|}||dq | | |j di|S)NfastmathFT)ftzfmaprec_div prec_sqrtr}r) rtstrpopupdatererrrmr~rorr)roptscu libdevicemodrrr compile_irjs  rz"^attributes #\d+ = \{ ([\w\s]+)\ }cCsng}|D]+}|dr,t|}|d}ddd|D}||d|}||qd|S)z, Convert LLVM 14.0 IR for LLVM 7.0. z attributes #r  css|] }|dkr|VqdS) willreturnNr)rzarrr sz#llvm140_to_70_ir.. ) splitlines startswithre_attributes_defmatchgroupsplitjoinrsappend)r buflinemattrsrrrrs     rc Cs|j}t|d}ttdd}||||f}t|d}||td }t |d}| |}t ||d} d| _ d| _t||g| _|jd 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 znvvm.annotationsrEz llvm.used appendingz llvm.metadatanoinlineN)moduler MetaDataStringConstantIntType add_metadatarget_or_insert_named_metadatar as_pointer ArrayTypebitcastGlobalVariablelinkagesection initializer attributesdiscard) functionr mdstrmdvaluemdnmdptrtyusedtyfnptr llvm_usedrrrset_cuda_kernels     r$cs<tdfddtD}||}|d|dS)zAdd NVVM IR version to moduler csg|]}|qSrr)rzrxrrrr|sz"add_ir_version..znvvmir.versionN)r rrr;radd_named_metadata)rrCmd_verrr%radd_ir_versions  r()MraloggingrerYrctypesrrrrrrr threadingllvmliter r[r r r libsrrr numba.corerr getLoggerr^loggerADDRSPACE_GENERICADDRSPACE_GLOBALADDRSPACE_SHAREDADDRSPACE_CONSTANTADDRSPACE_LOCALrcrbrrWrrrwr2modulesrFrGrLockr'r)rrerrrr@rrrrrrrrrrrrrrrrrrrr$r(rrrrs$     ~n  "   /  '