sh+^ nSrSSKrSSKrSSKrSSKrSSKJrJrJrJ r J r J r J r SSK r SSKJr SSKJrJrJr SSKJrJrJr SSKJrJr \R6"\5rSrSrS r S r!S r"\r#\r$S RK5r&\'"\&5Hur(r)\*"\RV\\)\(5 M S r,Sr-Sr.\ R^"5r0"SS\15r2"SS\15r3Sr4SSSSSSSSSSSSS. r5Sr6Sr7Sr8Sr9Sr:"SS \15r;S!rS$r?S%r@S&rAS'rBS(rCS)rDS*rES+rFS,rG\R"S-5rIS.rJS/rKS0rLg)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:64c:[5 g![a gf=f)z Return if libNVVM is available TF)NVVMr k/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/cudadrv/nvvm.py is_availabler<s%  s c`\rSrSrSr\\"\5\"\54\\"\54\\"\54\\\ \ \ 4\\\ \ \ 4\\\\"\ 54\\\"\ 54\\\ 4\\\"\ 54\\\ 4\\"\5\"\5\"\5\"\54\\\\"\ 54S. r Sr Sr Sr\S5r\S5rS rS rS S jrS rg)rKzProcess-wide singleton. ) nvvmVersionnvvmCreateProgramnvvmDestroyProgramnvvmAddModuleToProgramnvvmLazyAddModuleToProgramnvvmCompileProgramnvvmGetCompiledResultSizenvvmGetCompiledResultnvvmGetProgramLogSizenvvmGetProgramLog nvvmIRVersionnvvmVerifyProgramNc[ URc[RU5=Uln[ S5UlURR5H;upE[UR U5nUSUl USSUl [XU5 M= SSS5 UR$![ anSUlSn[X2-5eSnAff=f!,(df  UR$=f)Nnvvmz;libNVVM cannot be found. Do `conda install cudatoolkit`: %srr ) _nvvm_lock_NVVM__INSTANCEobject__new__rdriverOSErrorr _PROTOTYPESitemsgetattrrestypeargtypessetattr)clsinsteerrmsgnameprotofuncs rr0 NVVM.__new__s ~~%(.s(;;7".v"6DK$(#3#3#9#9#;KD"4;;5D#(8DL$)!"IDMD- $<$~~7%)CN2F*6:66 7 Z$~~s/*CB/AC/ C9CCC C0cUR5nUSUlUSUlUSUlUSUl[ 5Ulg)Nrr r)get_ir_version_majorIR_minorIR _majorDbg _minorDbgget_supported_ccs_supported_ccs)self ir_versionss r__init__ NVVM.__init__sG))+ #A #A $Q$Q/1rcRURUR4S:a[$[$)N)r )rDrE_datalayout_original_datalayout_i128rJs r data_layoutNVVM.data_layouts# MM4== )F 2' '# #rcUR$N)rIrRs r supported_ccsNVVM.supported_ccss"""rc[5n[5nUR[U5[U55nURUS5 URUR4$)NzFailed to get version.)rrr check_errorvalue)rJmajorminorerrs r get_versionNVVM.get_versionsNuU|U5\: 67{{EKK''rcB[5n[5n[5n[5nUR[U5[U5[U5[U55nURUS5 URURURUR4$)NzFailed to get IR version.)rr)rrZr[)rJmajorIRminorIRmajorDbgminorDbgr^s rrCNVVM.get_ir_versionsx''77  ww!&x%/C 9:}}gmmX^^X^^KKrcU(a>[U[U5nU(a"[U5 [R"S5 gUeg)Nr )r RESULT_CODE_NAMESprintsysexit)rJerrormsgrkexcs rrZNVVM.check_errors6 C!25!9:Cc   r)rFrDrGrErI)F)__name__ __module__ __qualname____firstlineno____doc__ nvvm_resultrr nvvm_programrrr3r.r0rLpropertyrSrWr_rCrZ__static_attributes__rrrrrKs< $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$$ ##(Lrrc>\rSrSrSrSrSrSrSrSr Sr S r g ) CompilationUnitc[5Ul[5UlURR [ UR55nURR US5 g)NzFailed to create CU)rr1rv_handler rrZ)rJr^s rrLCompilationUnit.__init__sFf #~ kk++E$,,,?@ %:;rc[5nUR[UR55nUR USSS9 g)NzFailed to destroy CUT)rk)rr!rr}rZ)rJr1r^s r__del__CompilationUnit.__del__s8''dll(;<3 6TBrcURRURU[U5S5nURR US5 g)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)r1r"r}lenrZrJbufferr^s r add_moduleCompilationUnit.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 cURSS5nUcSU3$[U[5(a [U5nSUSU3$)N_-=)replace isinstanceboolint)kvs rstringify_option1CompilationUnit.compile..stringify_optionsJ #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)r4rrencoder1r*r} _try_errorr$rr%rr r[r&get_loglogwarningswarnr) rJoptionsrrrxc_optsr^reslen output_buffers rcompileCompilationUnit.compilesx( 7>mmoFoda#A)oFS\)6=-?6=.6ahhv6F-G6=-?@kk++DLL#g,O 12kk,,T\\3w<P 23kk33DLL%-P EF&,,.1 kk// mL =><<> 88 MM$(([ 9Q7G-?s G$GcdURRX<SUR5<35 g)N )r1rZr)rJr^rms rrCompilationUnit._try_error%s! dlln%EFrc[5nURRUR[ U55nURR US5 UR S:au[UR -"5nURRURU5nURR US5 UR RS5$g)Nz#Failed to get compilation log size.r zFailed to get compilation log.r) rr1r'r}rrZr[r r(decode)rJrr^logbufs rrCompilationUnit.get_log(skk// eFmL %JK < 7A> = B  B cSSKJn UR5n[[5nX:a9SnUSSUS3nSUSUSSUSS3n[ R "U5 U$[U5nU$! SnUs$=f) Nr)runtimer.r z CUDA Toolkit z is unsupported by Numba - z! is the minimum required version.)numba.cuda.cudadrv.runtimerr_minrrrr)rcudart_version _supported_cc min_cudartctk_verunsupported_vers rrHrH^s6 ,,.]#J" #A&'q):(;<*7)3N(m_Ajm_=//  o&(8M % s A22A:c[5RnU(d Sn[U5e[U5H4up4X@:XaUs $X@:dMUS:XaSX--n[U5eXS- s $ US$)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 )rrWr enumerate)myccrWrmirs rfind_closest_archrwsF((M Qs##=) :I YAv+.2i9&s++%U++*  rcx[R(a[RnSU-$[X45nSU-$)z1Matches with the closest architecture option z compute_%d%d)r FORCE_CUDA_CCr)r\r]archs 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" c$\rSrSrSrSrSrSrg) LibDeviceiNcURc)[5c[[5e[ 5UlURUlgrV)_cache_r RuntimeErrorMISSING_LIBDEVICE_FILE_MSGrbcrRs rrLLibDevice.__init__s5 << &"#=>>)+DL,,rcUR$rV)rrRs rget LibDevice.gets wwr)rr)rprqrrrsrrLrrxrrrrrsGrrz %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([RUS9$)NTi)cas_nvvmformatrs rir_casr#s ??b? !!rc V[XX#[U5S9n[R"S0UD6$)N)TrOPFUNCCASr)dictrir_numba_atomic_binary_templater)rrrrparamss rir_numba_atomic_binaryr's( AF2J ?F * 1 1 ;F ;;rc Z[XX#UU[U5S9n[R"S0UD6$)N)rrNANr PTR_OR_VALrrr)rrir_numba_atomic_minmax_templater)rrrrrrrs rir_numba_atomic_minmaxr,s1 A#-F + 1 1 ;F ;;rc<[RX[U5S9$N)rTur)ir_numba_atomic_inc_templaterrrrs rir_numba_atomic_incr3 ' . .vay . IIrc<[RX[U5S9$r)ir_numba_atomic_dec_templaterrrs rir_numba_atomic_decr7rrcS[SSSSS94S[SS S S S94S [SSS S S94S [SSS94S[SSS94S[SS SSSSS94S[SSSSSSS94S[SS SSSSS94S[SSSSSSS94S[SS SSSSS94S[SSSSSSS94S [SS SS!SSS94S"[SSSS!SSS94S#/nUHup#UR X#5nM [ U5nU$)$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!)!f %F Mrc ~[U[5(aU/nURSS5(aURSSSSS.5 [ 5n[ 5nUH.n[ U5nURURS55 M0 URUR55 UR"S0UD6$)NfastmathFT)ftzfmaprec_div prec_sqrtrr) rstrpopupdaterzrrrrrrr)roptscu libdevicemods r compile_irr%js&# xx E""      B I3 cjj()y}}' ::  rz"^attributes #\d+ = \{ ([\w\s]+)\ }c/nUR"5HnURS5(am[RU5nUR S5R 5nSR SU55nURUR S5U5nURU5 M SR U5$)z$ Convert LLVM 14.0 IR for LLVM 7.0. z attributes #r  c34# UHoS:wdM Uv M g7f) willreturnNr).0as r #llvm140_to_70_ir..sC1l1BQQs  r) splitlines startswithre_attributes_defmatchgroupsplitjoinrappend)r buflinemattrss rrrs C  ??> * *!''-AGGAJ$$&EHHCCCE<< E2D 4  99S>rcfURn[R"US5n[R"[R"S5S5nUR XU45n[ R"US5nURU5 [R"S5R5n[R"US5nURU5n[R"XS5n SU l SU l[R"Xx/5U lUR R#S 5 g ) aD 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.annotationsrOz llvm.used appendingz llvm.metadatanoinlineN)moduler MetaDataStringConstantIntType add_metadatarget_or_insert_named_metadatar as_pointer ArrayTypebitcastGlobalVariablelinkagesection initializer attributesdiscard) functionr?mdstrmdvaluemdnmdptrtyusedtyfnptr llvm_useds rset_cuda_kernelrWs__F   fh /Ekk"**R.!,G   hw7 8B  . .v7I JCGGBK JJqM $ $ &E \\% #F   U #E!!&+>I#I'IKK8I  +rc[R"S5n[5R5Vs/sH o!"U5PM nnUR U5nUR SU5 gs snf)zAdd NVVM IR version to moduler<znvvmir.versionN)r rBrrCrCadd_named_metadata)r$rrrKmd_vers radd_ir_versionr[s\ **R.C#'6#8#8#:;#:a3q6#:K;   k *F+V4rls ;;==&   8 $    EG ' (DAq CKK !1a( ) ;7 ^^  {6{|kfk\             <2D!    #, ( (#@"< <JJ,^.JJDE$$,N5r