o Tdz @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 /Z0Gddde1Z2Gddde1Z3dZ4dddddddddddd Z5ddZ6ddZ7d d!Z8d"d#Z9d$Z:Gd%d&d&e1Z;d'Zd*Z?d+Z@d,ZAd-ZBd.d/ZCd0d1ZDd2d3ZEd4d5ZFd6d7ZGd8d9ZHd:d;ZIdZLeKd?ZMeKd@ZNdAZOeKeOPdBdCZQeKdDZRhdEZSeKdFZTeKdGZUeKdHZVeKdIZWeKdJZXeKdKZYeKdLZZeKdMZ[eKdNZ\dOdPdQZ]dRdSZ^dTdUZ_dVdWZ`dXdYZadZd[Zbd\d]ZcdS)^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 rrj/home/ncw/WWW/www-new/content/articles/pi-bbp/venv/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 Zed d ZddZddZdddZdS)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|jfdkS)N)r )r<r=rBrrr is_nvvm70szNVVM.is_nvvm70cCs|j|jfdkr tStS)N)r )r<r=_datalayout_original_datalayout_i128rFrrr data_layoutszNVVM.data_layoutcC|jSN)rArFrrr 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$rrOrP)rBmajorIRminorIRmajorDbgminorDbgrSrrrr;s  zNVVM.get_ir_versionFcCs4|rt|t|}|rt|tddS|dS)Nr )r RESULT_CODE_NAMESprintsysexit)rBerrormsgr\excrrrrOszNVVM.check_error)F)__name__ __module__ __qualname____doc__ nvvm_resultrr nvvm_programrrr-r(r*rDpropertyrGrKrNrTr;rOrrrrrKsF        6    rc@sDeZdZddZddZddZddZd d Zd d Zd dZ dS)CompilationUnitcCs4t|_t|_|jt|j}|j|ddS)NzFailed to create CU)rr+re_handlerrrO)rBrSrrrrDszCompilationUnit.__init__cCs*t}|t|j}|j|ddddS)NzFailed to destroy CUT)r\)rrrrhrO)rBr+rSrrr__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+rrhlenrOrBbufferrSrrr add_module zCompilationUnit.add_modulecCrj)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. Nrk)r+rrhrlrOrmrrrlazy_add_modulerpzCompilationUnit.lazy_add_modulec Ks~g}d|vr|d|d|dr|d|dd}|D]}||vr?tt||}|d|dd|fq#|rSd tt| }t d |t t |d d |D}|j|jt ||}||d |j|jt ||}||dt} |j|jt| }||dt| j} |j|j| }||d||_|jrtj|jtd| ddS)ajPerform 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=%darchz-arch=%s)ftz prec_sqrtprec_divfmaz-%s=%d_-, zunsupported option {0}cSsg|] }t|dqS)utf8)rencode).0xrrr +sz+CompilationUnit.compile..zFailed to verify zFailed to compile z&Failed to get size of compiled result.zFailed to get compiled result.)categoryN)appendpopgetintboolreplacejoinmapreprkeysr formatrrlr+r%rh _try_errorrrr rrrPr!get_loglogwarningswarnr ) rBoptionsopts other_optionskvoptstrc_optsrSreslenptxbufrrrcompiles>"        zCompilationUnit.compilecCs|j|d||fdS)Nz%s %s)r+rOr)rBrSr^rrrrFszCompilationUnit._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"rhrrOrPrr#decode)rBrrSlogbufrrrrIs   zCompilationUnit.get_logN) r`rarbrDrirorqrrrrrrrrgs  V rg)rr)rrr)rr:)rr)rEr)rEr )rEr:)rr)rr:)rrrHrrHrErHr)rH rr)rr)rr)rr)rr)rr) ) r)rr )rr:)rr)rr)rr)rrE)rr)rrH) r)rr csLzt|\tfddtDWSty%tddtDYSw)Ncs(g|]}|krkrnn|qSrrr}ccmax_ccmin_ccrrrus z(ccs_supported_by_ctk..cSsg|] }|tjkr|qSr)rCUDA_DEFAULT_PTX_CCrrrrrzs ) 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.runtimerrTminrrrr)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 )rrNr enumerate)myccrNr^irrrrfind_closest_archs  rcCs(tjr tj}d|St||f}d|S)z1Matches with the closest architecture option z compute_%d%d)r FORCE_CUDA_CCr)rQrRrsrrrget_arch_options  rz~Missing libdevice file. Please ensure you have package cudatoolkit >= 11.0 Install package by: conda install cudatoolkit c@s eZdZdZddZddZdS) LibDeviceNcCs0|jdurtdurttt|_|j|_dSrM)_cache_r RuntimeErrorMISSING_LIBDEVICE_FILE_MSGrbcrFrrrrDs   zLibDevice.__init__cCrLrM)rrFrrrrsz LibDevice.get)r`rarbrrDrrrrrrs rz 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 }} cCs tjr tj|dStj|dS)NTi)rrG cas_nvvm70r cas_nvvm34rrrrir_casKs  rcCs&t||||t|d}tjdi|S)N)TrOPFUNCCASr)dictrir_numba_atomic_binary_templater)rrrrparamsrrrir_numba_atomic_binaryRsrc Cs*t||||||t|d}tjdi|S)N)rrNANr PTR_OR_VALrrr)rrir_numba_atomic_minmax_templater)rrrrrrrrrrir_numba_atomic_minmaxWs rcCtj||t|dSN)rTur)ir_numba_atomic_inc_templaterrrrrrrir_numba_atomic_inc^rcCrr)ir_numba_atomic_dec_templaterrrrrrir_numba_atomic_decbrrcCsH|}t|D]\}}|drd}|tj||<nqd|S)z@ Find the line containing the datalayout and replace it ztarget datalayoutztarget datalayout = "{0}" ) splitlinesr startswithrrrKr)llvmirlinesrlntmprrr_replace_datalayoutfs  rcCsvdtdddddfdtdd d d dfd tddd d dfd tdddfdtdddfdtdd dddddfdtdddddddfdtdd dddddfdtdddddddfdtdd dddddfdtdddddddfd tdd dd!dddfd"tdddd!dddfd#g}tjs|d$tjd d%fd&tjdd%fg7}t|}|D] \}}| ||}qtjrt |}|St |}|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"))immargrzMdeclare i32 @"___numba_atomic_i32_cas_hack"(i32* %".1", i32 %".2", i32 %".3"))rzMdeclare i64 @"___numba_atomic_i64_cas_hack"(i64* %".1", i64 %".2", i64 %".3")) rrrrrrGir_numba_cas_hackrrrllvm100_to_70_irllvm100_to_34_ir)r replacementsdeclfnrrr llvm_replacess          &   rcKs|t|tr|g}|ddr|dddddt}t}|D]}t|}||dq | | |j di|S)NfastmathFT)rtrwrvrur{r) isinstancestrrupdatergrrror|rqrr)rrcu libdevicemodrrr llvm_to_ptxs  rz \!\d+\s*=zmetadata\s*\![{'\"0-9]z\!\d+z,\!{i32 \d, \!\"Debug Info Version\", i32 \d} z\s+z"^attributes #\d+ = \{ ([\w\s]+)\ }> coldminsizeoptiszeoptnonenoinlinenoreturnnounwindreadnonereadonly 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\((.*)\)zspFlags: (.*), isDefinition isOptimized)DISPFlagDefinitionDISPFlagOptimizedcCsng}|D]+}|dr,t|}|d}ddd|D}||d|}||qd|S)z, Convert LLVM 10.0 IR for LLVM 7.0. attributes #r r css|] }|dkr|VqdS) willreturnNrr}arrr z#llvm100_to_70_ir..r) rrre_attributes_defmatchgroupsplitrrr)r buflinemattrsrrrrs     rc Cs@dd}g}|D]}|dr|dd}|dr)d|vr)|dd}t|rfdt|urf|d d }|d d }|d }|d|d||dd}}dd}d |t ||f}|drlq t |dur{t dd|}|drt |}|d} d dd| D} ||d| }d|vrt|}|durtd|f|} |d| ||| d}d|vrt|}|r|} |d| ||| d}d|vrt d|}t d|d}d|vrt t|}d |vrd|vrt t|}t d|}||q d! |S)"z, Convert LLVM 10.0 IR for LLVM 3.4. cSsd}d} t||}|durtd|f|}|d}|dkr)|dkr(nn|dvr2|d7}n|dvr:|d8}q||dS)NrTzfailed parsing leading type: %s,z{[(r z)]}) re_type_toksearchrendr"lstrip)s par_levelposr&tokrrrparse_out_leading_types$   z0llvm100_to_34_ir..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 cSsd|dS)Nz metadata r)r"r&rrrfix_metadata_ref/sz*llvm100_to_34_ir..fix_metadata_refr zsource_filename =cSsdS)Nrrr3rrr6sz"llvm100_to_34_ir..rcss|] }|tvr|VqdSrM)supported_attributesrrrrr<rz#llvm100_to_34_ir..zgetelementptr z failed parsing getelementptr: %szload zcall z\1*\2rr(z @llvm.memsetdeclarer)rrrr,re_metadata_defr!re_metadata_correct_usager*findrre_metadata_refrre_unsupported_keywordsr r"r#re_getelementptrrr+re_loadre_callre_rangerstripre_parenthesized_list_replace_llvm_memset_usage _replace_llvm_memset_declarationre_annotationsr) r r1r$r%assigposlhsrhsr4r&r'r/rrrrsp       "            rcCs`t|dd}td|d}|std|d}|dd|d|}d |S) zNReplace `llvm.memset` usage for llvm7+. Used as functor for `re.sub. r r(z align (\d+)rz+No alignment attribute found on memset destrzi32 {}rz({})) listr"r#rer* ValueErrorinsertrr)r&r align_attralignoutrrrrCis   rCcCs4t|dd}|ddd|}d|S)zTReplace `llvm.memset` declaration for llvm7+. Used as functor for `re.sub. r r(rrrzrI)rJr"r#rMrr)r&rrPrrrrDys   rDcCsZ|j}t|d}ttdd}||||f}t|d}|||j ddS)Nkernel r znvvm.annotationsr) moduler MetaDataStringConstantIntType add_metadatarget_or_insert_named_metadatar attributesdiscard)lfuncrmdstrmdvaluemdnmdrrrset_cuda_kernels   r`cs<tdfddtD}||}|d|dS)zAdd NVVM IR version to modulerRcsg|]}|qSrr)r}rrrrrsz"add_ir_version..znvvmir.versionN)r rVrr;rWadd_named_metadata)rrCmd_verrraradd_ir_versions  rd)drcloggingrKr[rctypesrrrrrrr threadingllvmliter r]r r r libsrrr numba.corerr getLoggerr`loggerADDRSPACE_GENERICADDRSPACE_GLOBALADDRSPACE_SHAREDADDRSPACE_CONSTANTADDRSPACE_LOCALrerdr#rYrrrr2modulesrIrJrLockr'r)rrgrrrr@rrrrrrrrrrrrrrrrrrrrr8r9r;debuginfo_patternrre_metadata_debuginfor r6r=r>r?r@r)rEr<rB re_spflags spflagmaprrrCrDr`rdrrrrs$        "   ?             k