&ThKXSSKJrJr SSKJrJrJrJr SSKJ r SSK J r SSK r SSK JrJrJrJr SSKJr SSKrSSKrSSKrSSKrSSKrSSKrSSKJr SSKrS \4S jr\ R<"5S \4S j5r \ R<"5S \!4Sj5r"\ R<"5S \!4Sj5r#\ R<"5S\!4Sj5r$S \!4Sjr%\ R<"5S \!4Sj5r&\ R<"S5S5r'S\!4Sjr(\ "SS9"SS55r)"SS\5r*g)) BaseBackend GPUTarget)irpassesllvmnvidia) PTXASError) dataclassN)AnyDictTupleOptional) ModuleType)Pathtargetc@S[[[[44SjnU$)NreturncURRnURRnX#:XdS5eUS:Xagg)Nz%lhs and rhs bitwidth must be the same)r )rrr)scalarprimitive_bitwidth)lhs_typerhs_type lhs_bitwidth rhs_bitwidths W/var/www/auris/envauris/lib/python3.13/site-packages/triton/backends/nvidia/compiler.pycheck_dot_compatibility-min_dot_size..check_dot_compatibilitys@99 99 +T-TT+ 1 )r int)rrs r min_dot_sizer#s! uS#s]7K  #"r!binarycU[R"S5- n[RR SUR 5S3S5[R R[R R[5SU5/nUHn[R RU5(dM)[R RU5(dMO[R"US/[RS9nUcMy[R "SUR#S 5[R$S 9nUcMX$R'S 54s $ [)S U35e) NEXETRITON__PATHbin --version)stderrz.*release (\d+\.\d+).*utf-8flagsz Cannot find ) sysconfigget_config_varosenvirongetupperpathjoindirname__file__existsisfile subprocess check_outputSTDOUTresearchdecode MULTILINEgroup RuntimeError)r$pathsr7resultversions r_path_to_binaryrI!s i&&u--F  06;  RWW__X.v> E  77>>$  BGGNN4$8$8,,dK-@IZIZ[F!))$=v}}W?U]_]i]ij&q!111  fX. //r!archc,US:aSOSn[U5$)Ndzptxas-blackwellptxas)rI)rJnames r get_ptxasrO3s $ D 4  r!c[RRS5nUbU$[R"[ U5SS/5R S5nU$)NTRITON_MOCK_PTX_VERSIONrr+r-)r3r4r5r=r>rOrB)rJmock_verrHs rget_ptxas_versionrS9sOzz~~78H%%yq'9;&GHOOPWXG Nr!rc[U[5(de[[UR S55upUS:XaUS:aSU-$SU-S- $US:XaSU-$US:XaS U-$[ S U-5e) zC Get the highest PTX version supported by the current CUDA driver. . Pr0 F ?z?Triton only support CUDA 10.0 or higher, but got CUDA version: ) isinstancestrmapr"splitrE) cuda_versionmajorminors rptx_get_versionrdBs lC ( (( (sL..s34LE { 19: :> ! {Ez {Ez X[gg hhr!cTURnUc[U5up4[U5nU$N) ptx_versionrOrd)optionsrJrg_ras rget_ptx_version_from_optionsrjUs.%%K#D/%l3 r!c>[X5n[SU5nSU3nU$)NVz+ptx)rjmin)rhrJrgllvm_ptx_versionfeaturess r get_featuresrp]s..w=K2{+&'(H Or!c[US5n[R"UR55R 5sSSS5 $!,(df  g=f)Nrb)openhashlibsha256read hexdigest)r7fs r file_hashryks5 dD Q~~affh'113   s 2A  A capabilityc$US:aSOSnSUU3$)NZar)sm_)rzsuffixs rsm_arch_from_capabilityrqs!"$S"F  VH %%r!T)frozencn\rSrSr%Sr\\S'Sr\\S'Sr\\S'Sr \\S 'Sr \\S 'Sr \\S 'Sr \\S 'S r \\\S'Sr\\S'S r\\S'Sr\\S'Sr\\S'Sr\\\S'Sr\\\S'Sr\\S'Sr\\\S'S r\\S'S r\\S'Sr\\S 'S!r\\S"'Sr \\S#'S r!\\S$'S%r"S&r#Sr$g )' CUDAOptionsw num_warpsr0num_ctas num_stagesrnum_buffers_warp_specnum_consumer_groupsreg_dec_producerreg_inc_consumerNmaxnreg)r0r0r0 cluster_dimsrgTenable_fp_fusionFlaunch_cooperative_grid)fp8e5fp8e4b15supported_fp8_dtypesrdeprecated_fp8_dtypestf32default_dot_input_precision)rtf32x3ieeeallowed_dot_input_precisionsmax_num_imprecise_acc_default extern_libsdebugcuda backend_namesanitize_overflowrJc[[5RS- nURc0O[ UR5nUR SS5(d&[ R"S[US- 55US'[RUS[UR555 URS:a URURS- -S:XdS5eg) Nlib libdeviceTRITON_LIBDEVICE_PATHzlibdevice.10.bcrrr0znum_warps must be a power of 2)rr:parentrdictr5r3getenvr^object __setattr__tupleitemsr)selfdefault_libdirrs r __post_init__CUDAOptions.__post_init__sh..6 ,,4b$t?O?O:P {D11')yy1H#n_pNpJq'rK $4k6G6G6I0JK~~!t~~!9K'LQR&R 0/ 0R&Rr!c d[UR5n[S[US555US'SR [UR 55VVs/sH up#USU3PM snn5n[ R"URS55R5$s snnf)Nc3@# UHupU[U54v M g7frf)ry).0kvs r #CUDAOptions.hash..s(hGgtq!Yq\):Ggsrri-r-) r__dict__rsortedr8rrtruencoderw)r hash_dictrNvalkeys rhashCUDAOptions.hashs' #((hviXeNfGg(h#h - hh @Q9RS9RID4&#9RST~~cjj12<<>>TsB, )%__name__ __module__ __qualname____firstlineno__rr"__annotations__rrrrrrrrrrrgrboolrrr r^rrrrrrrrrrJrr__static_attributes__rr!rrrwsIsHcJ!"3"  cc"GXc]!#L%#K!d!$)T)'<%*<(*5:*'--/I %*I*.!4.KE4L#"t"D#0?r!rc^\rSrSr\S\4Sj5rSrS\SS4U4SjjrS\ 4Sjr S r S r S\ \\44S jrS r\S 5r\S5rSrSrSrSr\R2"5S5rSrU=r$) CUDABackendrc URS:H$)Nr)backend)rs rsupports_targetCUDABackend.supports_targets~~''r!cSn[R"X!5nU(d[SU35e[UR S55$)Nz ^sm(\d+)$z(TRITON_OVERRIDE_ARCH must have the form r0)r@ fullmatch ValueErrorr"rD)rrJpatternmatchs r _parse_archCUDABackend._parse_archs> W+GyQR R5;;q>""r!rNc2>[TU]U5 SUlg)Ncubin)super__init__ binary_ext)rr __class__s rrCUDABackend.__init__s  !r!cvS[R"SSURR350nUR [ R R5Vs0sHo3U;dM XcMX1U_M sn5 [URUS55nSU;aG[[ R5nUS:aURS5 [[U55US'SU;a US:aS US'S U;a[R"S S 5S :HUS 'US:XaS OSUS'[ S0UD6$s snf)NrJTRITON_OVERRIDE_ARCHsmrYfp8e4nvrr|)rrTRITON_DEFAULT_FP_FUSION1i@rrr)r3rrrJupdater__dataclass_fields__keysr"rsetraddrr)roptsargsrrzrs r parse_optionsCUDABackend.parse_optionss2 "8Bt{{?O?O>P:QRS )I)I)N)N)Pu)PAY]T]ZaeahZQQZ)Puv))$v,78 ! -#&{'G'G#H R$((3+08L1M+ND' ( "$ .R0>,- T )')yy1KS'QUX'XD# $9Cr9Iq ,-"T""%vs D6)D60 D6cURURURURSURSURS4$)Nrr0)rrsharedr)rmetadatas r pack_metadataCUDABackend.pack_metadatasO       OO  ! !! $  ! !! $  ! !! $   r!cSSKJs Js Jn [ UR UR 55nUS:a URO UR[UR5S.nU$)NrrX)convert_custom_typesr#) triton.language.extra.cudalanguageextrarr"rrJconvert_custom_float8_sm80convert_custom_float8_sm70r#r)rrhrrz codegen_fnss rget_codegen_implementation&CUDABackend.get_codegen_implementationsV11))',,78 0:R/?D + +TEdEd  % r!cSSKJn SU0$)Nr)rztriton.language.extra.libdevice)rr)rrs rget_module_mapCUDABackend.get_module_maps819==r!c0[R"U5 grf)r load_dialects)rctxs rrCUDABackend.load_dialectssS!r!cx[R"UR5nUR5 [R R U5 [RRU5 [R RU5 [RRU5 [RRU5 [R RU5 [R RU5 [RRU5 URU5 U$rf)r pass_managercontext enable_debugrcommon add_inlinerttiradd_rewrite_tensor_pointeradd_canonicalizer add_combineadd_reorder_broadcastadd_cseadd_symbol_dceadd_loop_unrollrun)modroptpms r make_ttirCUDABackend.make_ttirs __S[[ )  !!"% ..r2 ''+ # ))"- b! $$R( ##B' s  r!cV[R"5nURb<URSUlURSUlURSUl[ R"UR5nUR5n[RRUSU3URSUR5 [RR!U5 US-S:a[RR#U5 [RR$R'XT5 [RR)U5 [RR+U5 [RR-U5 [RR)U5 [RR/XSS:5 [R0R3U5 US-S ;Ga[RR5U5 [R0R7U5 [R0R9U5 [RR;U5 [R0R7U5 [RR=U5 [RR?XRR@5 [RRCXRR@5 [RREXRR@5 [RRGXRRHUR@URJURL5 [RROXRRPU5 [RRSXRR@5 [RRUXRR@5 GONUS-S:Ga%[RR5U5 [R0R7U5 [R0R9U5 [RR;U5 [RR?XRR@5 [RRCXRR@5 [RREXRR@5 [RRGXRRHUR@URJURL5 [RROXRRPU5 [RR=U5 [RR$RWU5 [RR$RYU5 [RRUXRR@5 [R0R7U5 O[R0R9U5 [RR[U5 [RR/XSS:5 [RR]U5 [RR)U5 [RR_U5 [RRaU5 [R0R3U5 [R0RcU5 US-S :aR[RR$ReU5 [RR$RgU5 [R0R7U5 US-S :a)[RRiXRR@5 URkU5 URURUR 4US 'U$) Nrr0rzcuda:rr[rrX)r rr)6r ClusterInfor clusterDimX clusterDimY clusterDimZrrrrrr add_convert_to_ttgpuirrrttgpuir add_coalesceadd_f32_dot_tc ttnvgpuir add_plan_ctaadd_remove_layout_conversionsadd_optimize_thread_localityadd_accelerate_matmuladd_optimize_dot_operandsrradd_fuse_nested_loopsr add_licmadd_optimize_accumulator_init add_combine_tensor_select_and_ifadd_ws_task_partitionradd_taskid_propagateadd_ws_data_partitionadd_ws_code_partitionrrr add_pipelineradd_ping_pong_syncadd_ws_loweringadd_promote_lhs_to_tmemadd_keep_acc_in_tmem add_prefetchadd_coalesce_async_copyadd_reduce_data_duplicationadd_reorder_instructionsradd_fence_insertionadd_tma_loweringadd_ws_canonicalizationr)rrrrz cluster_infor dump_enableds r make_ttgirCUDABackend.make_ttgirs=))+    ''*'7'7':L $'*'7'7':L $'*'7'7':L $ __S[[ )(  **2zl/CS]]TVX[XdXde##B'  q NN ) )" - ,,R>44R833B7,,R044R80025EF b!  v % NN 0 0 4 MM + +B / MM " "2 & NN 8 8 < MM + +B / NN ; ;B ? NN 0 05L5L M NN / /4K4K L NN 0 05L5L M NN 0 05N5NPSPgPg141E1EsG[G[ ] NN ' 'NNL I NN - -b2I2I J NN * *2/F/F G 2  # NN 0 0 4 MM + +B / MM " "2 & NN 8 8 < NN 0 05L5L M NN / /4K4K L NN 0 05L5L M NN 0 05N5NPSPgPg141E1EsG[G[ ] NN ' 'NNL I NN ; ;B ? MM # # ; ;B ? MM # # 8 8 < NN * *2/F/F G MM + +B / MM " "2 &##B'0025EF..r244R82226//3 b! $$R(  q MM # # 7 7 ; MM # # 4 4R 8 ''+  q NN 2 227N7N O s $0$<$V>VXdXpXp#q  r!c [X0RR5nUn[R"UR 5nUR 5 [RRRU5 [RRU5 [RRU5 [RRU5 [RR!U5 [RRR#U5 [RR%U5 [RRR'XtU5 [R(R+U5 [R(R-U5 [RRR/U5 [RRR1U5 [R(R+U5 [R(R-U5 [R(R3U5 [4R6R9SS5S:Xa[R:R=U5 UR?U5 [@RB"5 [@R "5n[4R6R9SS5S:Xa [ES5e[@RF"Xh5n [IU5n [KX0RR5n Sn [@RL"XX5 [RN"U 5 URPb`U RS5HLn U RU5(aMU RW5(dM1U RYURP5 MN URZ(a7URZVVs/sHupUPM nnn[@R\"U U5 [@R^"U [@R`5 URcS5nUbUUS'URcS 5US 'URcS 5US 'URcS 5US'URcS5US'[eU 5nA AU$s snnf)NTRITON_DISABLE_LINE_INFO0TRITON_ENABLE_ASANrzYAddress Sanitizer Error: Address sanitizer is currently only supported on the AMD backendnvptx64-nvidia-cudazttg.total-num-warpsrz ttg.sharedrzttg.tensor_memory_size tmem_sizezttg.global_scratch_memory_sizeglobal_scratch_sizez#ttg.global_scratch_memory_alignmentglobal_scratch_align)3rjrrJrrrrrrr" add_lower_mmarr+add_allocate_warp_groupsconvert add_scf_to_cfadd_allocate_shared_memoryadd_allocate_tensor_memory"add_allocate_global_scratch_memory add_to_llvmirrr radd_nvgpu_to_llvmadd_warp_specialize_to_llvmrr3r4r5llvmir add_di_scoperr init_targetsrE to_modulerrpattach_datalayoutset_nvvm_reflect_ftzr get_functionsis_declarationis_external_linkageset_nvvm_maxnregrlink_extern_libsoptimize_module OPTIMIZE_O3 get_int_attrr^)rsrcrrhrzrgrrrllvm_modprocrotriplerrNr7rFtotal_num_warpsrets r make_llirCUDABackend.make_llir<s]27KK99"= ++BKH ''+ b! 11"5 ;;B? ''+ b! $$R( ::>>4c :c A MM & &r * s  ,,. ::>>. 4 ;km m>>#/&z2)9)9:& x@##H- ?? &++-''))a.C.C.E.E&&w7.   .5.A.AB.AltT.AEB  ! !(E 2 Xt'7'78**+@A  &$3H[ ! --l; # 0 01I J*-*:*:;[*\&'+.+;+;*:; ff*i },EsRTR^R^_ff';zl)CSPRP\P\]ff+R5 ::>>-s 3s : 4 5 #J r!c [URR5upV[R"SSSS9n[R"SSSS9nUR U5 UR 5 URS-n [RRSS 5S :XaS S /OS /n UR(a/OS /n [U5n [RRSS 5S :XaSS /O/n U/U QU QSPU QSU 3PURPSPU Pn[R"USSUS9 [RR!UR5(a [R""UR5 [RR!UR5(a [R""UR5 ['U S5nUR)5nSSS5 [RR!U 5(a[R""U 5 SSS5 SSS5 W$![R$an['UR5nUR)5nSSS5 O!,(df  O=f[RR!UR5(a [R""UR5 UR*S:XaSnO3UR*S[,R.-:XaSnOSUR*3n[1USWSSR3U5S35eSnAff=f!,(df  GNn=f!,(df  GN>=f!,(df  W$=f)NFwz.ptx)deletemoderrz.logz.orArBrz -lineinfoz-suppress-debug-infoz --fmad=falseDISABLE_PTXAS_OPTz --opt-levelz-vz --gpu-name=z-oT)check close_fdsr,z!Internal Triton PTX codegen errorz`ptxas` raised SIGSEGVz`ptxas` failed with error code z `ptxas` stderr: z Repro command:   rr)rOrrJtempfileNamedTemporaryFilewriteflushrNr3r4r5rrr=rr7r;removeCalledProcessErrorrsrv returncodesignalSIGSEGVr r8)rr`rrrzrMrifsrcflogfbin line_infofmadrJ opt_level ptxas_cmdelog_filelogerrorrxrs r make_cubinCUDABackend.make_cubinsT[[--.  ( (C OSW  ' 'u3v NRV JJsO JJL99t#DACPjPSBUX[B\&<=bman --2N3CD*:6D02 ?RTW0X\_0_,egIqqTq4q)q{SWRXEYq[_[d[dqfjqlpqI LydS77>>$)),,IIdii(77>>$)),,IIdii($dD!Q"ww~~d## $K OPN +00 L$))_"--/C%__77>>$)),,IIdii(<<3&?E\\S6>>%994E=all^LE E7+558E:33688I3F2Gr"KLL L""!E O NP ON s{M2CM B3H; M M'AM )M2;M M$I> 5 M> J B>MM  M  M M  M/ *M22 Nc^^^TRTR5mUU4SjUS'UUU4SjUS'UUU4SjUS'UU4SjUS'UU4S jUS 'g) Nc(>TRXT5$rf)rr`rrhrs r(CUDABackend.add_stages..st~~cW/Ur!r c*>TRXTT5$rf)r>r`rrzrhrs rrrswXb0cr!ttgirc*>TRXTT5$rf)rfrs rrrst~~cWV`/ar!llircR>TRXTTRR5$rf)rprrJrs rrrsdmmC7TXT_T_TdTd.er!ptxcR>TRXTTRR5$rf)rrrJrs rrrswX\XcXcXhXh0ir!r)rrJ)rstagesrhrzs` `@r add_stagesCUDABackend.add_stagessD%%gll3 Uvcwaveu iwr!cv[URR5nUSURR3$)Nr)rSrrJ)rrHs rrCUDABackend.hashs2#DKK$4$45!DKK,,-..r!)r)rrrr staticmethodrrrrr rrrr r^rrrrr>rfrprr functools lru_cacherr __classcell__)rs@rrrs( ((#"y"T"#S#, >S*_ 5>"  FFP?B,)Vj//r!r)+triton.backends.compilerrrtriton._C.libtritonrrrrtriton.runtime.errorsr dataclassesr rtypingr r r rtypesrrtr@r~rr3r=pathlibrr1r#rr^rIr"rOrSrdrjrpryrrrrr!rrs_;88,!--   # # 0C00" !C!!  C iSii$     T44 &&  $'?'?'?Tg/+g/r!