6^(hKNddlmZmZddlmZmZmZmZddlm Z ddl m Z ddl Z ddl mZmZmZmZddlmZddlZddlZddlZddlZddlZddlZddlmZddlZd efd Ze j<d efd Z e j<d e!fdZ"e j<d e!fdZ#e j<de!fdZ$d e!fdZ%e j<d e!fdZ&e j<ddZ'de!fdZ(e dGddZ)GddeZ*y)) BaseBackend GPUTarget)irpassesllvmnvidia) PTXASError) dataclassN)AnyDictTupleOptional) ModuleType)Pathtargetc>dttttffd}|S)Nreturnc|jj}|jj}||k(sJd|dk(ryy)Nz%lhs and rhs bitwidth must be the same)r )rrr)scalarprimitive_bitwidth)lhs_typerhs_type lhs_bitwidth rhs_bitwidths ]/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/triton/backends/nvidia/compiler.pycheck_dot_compatibilityz-min_dot_size..check_dot_compatibilitysB99 99 |+T-TT+ 1 )r int)rrs r min_dot_sizer"s! uS#s]7K  #"r binaryc|tjdz }tjj d|j ddtj jtj jtd|g}|D]}tj j|s#tj j|sCtj|dgtj}|mtj d|j#d tj$ }|||j'd fcSt)d |) NEXETRITON__PATHbin --version)stderrz.*release (\d+\.\d+).*utf-8flagsz Cannot find ) sysconfigget_config_varosenvirongetupperpathjoindirname__file__existsisfile subprocess check_outputSTDOUTresearchdecode MULTILINEgroup RuntimeError)r#pathsr6resultversions r_path_to_binaryrH!s i&&u--F  06;  RWW__X.v> E 2 77>>$ BGGNN4$8,,dK-@IZIZ[F!))$=v}}W?U]_]i]ij&q!111 2 fX. //r archc*|dk\rdnd}t|S)Ndzptxas-blackwellptxas)rH)rInames r get_ptxasrN3s $ D 4  r ctjjd}||Stjt |ddgj d}|S)NTRITON_MOCK_PTX_VERSIONrr*r,)r2r3r4r<r=rNrA)rImock_verrGs rget_ptxas_versionrR9sOzz~~78H%%yq'9;&GHOOPWXG Nr rct|tsJtt|j d\}}|dk(r|dkrd|zSd|zdz S|dk(rd|zS|dk(rd |zSt d |z) zK Get the highest PTX version supported by the current CUDA driver. . Pr/ F ?z?Triton only support CUDA 10.0 or higher, but got CUDA version: ) isinstancestrmapr!splitrD) cuda_versionmajorminors rptx_get_versionrcBs lC (( (sL..s34LE5 { 19: :> ! {Ez {Ez X[gg hhr cT|j}|t|\}}t|}|SN) ptx_versionrNrc)optionsrIrf_r`s rget_ptx_version_from_optionsriUs0%%K#D/<%l3 r c@t||}td|}d|}|S)NVz+ptx)rimin)rgrIrfllvm_ptx_versionfeaturess r get_featuresro]s0.w=K2{+&'(H Or ct|d5}tj|jj cdddS#1swYyxYw)Nrb)openhashlibsha256read hexdigest)r6fs r file_hashrxks> dD 4Q~~affh'113444s 1AA capabilityc"|dk\rdnd}d||S)NZar(sm_)rysuffixs rsm_arch_from_capabilityrqs!"$S"F  VH %%r T)frozencfeZdZUdZeed<dZeed<dZeed<dZeed<dZ eed <dZ eed <dZ eed <d Z e eed <dZeed<d Zeed<dZeed<dZeed<dZeeed<dZeeed<dZeed<dZeeed<d Zeed<d Zeed<dZeed<d Zeed!<dZeed"<d Z eed#<d$Z!d%Z"y )& CUDAOptions num_warpsr/num_ctas num_stagesrnum_buffers_warp_specnum_consumer_groupsreg_dec_producerreg_inc_consumerNmaxnreg)r/r/r/ cluster_dimsrfTenable_fp_fusionFlaunch_cooperative_grid)fp8e5fp8e4b15supported_fp8_dtypesr~deprecated_fp8_dtypestf32default_dot_input_precision)rtf32x3ieeeallowed_dot_input_precisionsmax_num_imprecise_acc_default extern_libsdebugcuda backend_namesanitize_overflowrIcttjdz }|jint |j}|j dds%t jdt|dz |d<tj|dt|j|jdkDr|j|jdz zdk(sJdy) Nlib libdeviceTRITON_LIBDEVICE_PATHzlibdevice.10.bcrrr/znum_warps must be a power of 2)rr9parentrdictr4r2getenvr]object __setattr__tupleitemsr)selfdefault_libdirrs r __post_init__zCUDAOptions.__post_init__sh..6 ,,4b$t?O?O:P {D1')yy1H#n_pNpJq'rK $4k6G6G6I0JK~~!t~~!9K'LQR&R 0/ 0R&Rr c ^t|j}tdt|dD|d<dj t|j Dcgc] \}}|d|c}}}t j|jdjScc}}w)Nc3<K|]\}}|t|fywre)rx).0kvs r z#CUDAOptions.hash..s(htq!!Yq\):(hsrrh-r,) r__dict__rsortedr7rrsrtencoderv)r hash_dictrMvalkeys rhashzCUDAOptions.hashs' #((hviXeNfGg(h#h - hh @Q9RSID#4&#ST~~cjj12<<>>TsB) )#__name__ __module__ __qualname__rr!__annotations__rrrrrrrrrrrfrboolrrr r]rrrrrrrrrrIrrr~r rrrwsIsHcJ!"3"  cc"GXc]!#L%#K!d!$)T)'<%*<(*5:*'--/I %*I*.!4.KE4L#"t"D#0?r rceZdZedefdZdZdeddffd ZdefdZ dZ d Z de e effd Zd Zed Zed ZdZdZdZdZej0dZxZS) CUDABackendrc |jdk(S)Nr)backend)rs rsupports_targetzCUDABackend.supports_targets~~''r cd}tj||}|std|t|j dS)Nz ^sm(\d+)$z(TRITON_OVERRIDE_ARCH must have the form r/)r? fullmatch ValueErrorr!rC)rrIpatternmatchs r _parse_archzCUDABackend._parse_archs@ Wd+GyQR R5;;q>""r rNc2t||d|_y)Ncubin)super__init__ binary_ext)rr __class__s rrzCUDABackend.__init__s  !r cVdtjdd|jji}|j t j jDcic]}||vs|||||c}t|j|d}d|vrFtt j}|dk\r|jdtt||d<d|vr |dk\rd |d<d |vrtjd d d k(|d <|dk(rd nd|d<t di|Scc}w)NrITRITON_OVERRIDE_ARCHsmrYfp8e4nvrr{)rrTRITON_DEFAULT_FP_FUSION1i@rrr~)r2rrrIupdater__dataclass_fields__keysr!rsetraddrr)roptsargsrryrs r parse_optionszCUDABackend.parse_optionss3 "8Bt{{?O?O>P:QRS )I)I)N)N)PuATUY]T]aefgahatQQZuv))$v,78 ! -#&{'G'G#H R$((3+08L1M+ND' ( "$ .R0>,- T )')yy1KS'QUX'XD# $9Cr9Iq ,-"T""%vs D&%D&+D&c|j|j|j|jd|jd|jdfS)Nrr/)rrsharedr)rmetadatas r pack_metadatazCUDABackend.pack_metadatasO       OO  ! !! $  ! !! $  ! !! $   r cddlmcmcm}t |j |j }|dk\r |jn |jt|jd}|S)NrrW)convert_custom_typesr") triton.language.extra.cudalanguageextrarr!rrIconvert_custom_float8_sm80convert_custom_float8_sm70r"r)rrgrry codegen_fnss rget_codegen_implementationz&CUDABackend.get_codegen_implementationsV11))',,78 0:R/?D + +TEdEd  % r cddlm}d|iS)Nr)rztriton.language.extra.libdevice)rr)rrs rget_module_mapzCUDABackend.get_module_maps819==r c.tj|yre)r load_dialects)rctxs rrzCUDABackend.load_dialectssS!r cvtj|j}|jtj j |tjj|tj j|tjj|tjj|tj j|tj j|tjj||j||Sre)r pass_managercontext enable_debugrcommon add_inlinerttiradd_rewrite_tensor_pointeradd_canonicalizer add_combineadd_reorder_broadcastadd_cseadd_symbol_dceadd_loop_unrollrun)modroptpms r make_ttirzCUDABackend.make_ttirs __S[[ )  !!"% ..r2 ''+ # ))"- b! $$R( ##B' s  r cftj}|j<|jd|_|jd|_|jd|_t j|j}|j}tjj|d||jd|jtjj!||dzdk\rtjj#|tjj$j'||tjj)|tjj+|tjj-|tjj)|tjj/||dk\tj0j3||dzd vrtjj5|tj0j7|tj0j9|tjj;|tj0j7|tjj=|tjj?||j@tjjC||j@tjjE||j@tjjG||jH|j@|jJ|jLtjjO||jP|tjjS||j@tjjU||j@nS|dzdk\r+tjj5|tj0j7|tj0j9|tjj;|tjj?||j@tjjC||j@tjjE||j@tjjG||jH|j@|jJ|jLtjjO||jP|tjj=|tjj$jW|tjj$jY|tjjU||j@tj0j7|ntj0j9|tjj[|tjj/||dk\tjj]|tjj)|tjj_|tjja|tj0j3|tj0jc||dzd k\rRtjj$je|tjj$jg|tj0j7||dzd k\r*tjji||j@|jk||j|j|j f|d <|S) Nrr/rzcuda:rrZrrW)r rr)6r ClusterInfor clusterDimX clusterDimY clusterDimZrrrrrradd_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_loopsradd_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)rrrry cluster_infor dump_enableds r make_ttgirzCUDABackend.make_ttgirs_))+    ''*'7'7':L $'*'7'7':L $'*'7'7':L $ __S[[ )(  **2zl/CS]]TVX[XdXde##B'  q NN ) )" - ,,R>44R833B7,,R044R800Z25EF b!  v % NN 0 0 4 MM + +B / MM " "2 & NN 8 8 < MM + +B / NN ; ;B ? NN 0 0S5L5L M NN / /C4K4K L NN 0 0S5L5L M NN 0 0S5N5NPSPgPg141E1EsG[G[ ] NN ' 'CNNL I NN - -b#2I2I J NN * *2s/F/F G 2  # NN 0 0 4 MM + +B / MM " "2 & NN 8 8 < NN 0 0S5L5L M NN / /C4K4K L NN 0 0S5L5L M NN 0 0S5N5NPSPgPg141E1EsG[G[ ] NN ' 'CNNL I NN ; ;B ? MM # # ; ;B ? MM # # 8 8 < NN * *2s/F/F G MM + +B / MM " "2 &##B'00Z25EF..r244R82226//3 b! $$R(  q MM # # 7 7 ; MM # # 4 4R 8 ''+  q NN 2 22s7N7N O s $0$<$V>VXdXpXp#q  r c^ t||jj}|}tj|j }|j tjjj|tjj|tjj|tjj|tjj!|tjjj#|tjj%|tjjj'|||tj(j+|tj(j-|tjjj/|tjjj1|tj(j+|tj(j-|tj(j3|t4j6j9dddk(rtj:j=||j?|tAjBtAj }t4j6j9dddk(r tEdtAjF||} tI|} tK||jj} d} tAjL| | | | tjN| |jPR| jSD]?} | jUr| jWs%| jY|jPA|jZr4|jZDcgc]\}}| }}}tAj\| |tAj^| t@j`|jcd}|||d<|jcd |d <|jcd |d <|jcd |d<|jcd|d<te| }~ ~|Scc}}w)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)3rirrIrrrrrrr add_lower_mmarradd_allocate_warp_groupsconvert add_scf_to_cfadd_allocate_shared_memoryadd_allocate_tensor_memory"add_allocate_global_scratch_memory add_to_llvmirrrradd_nvgpu_to_llvmadd_warp_specialize_to_llvmrr2r3r4llvmir add_di_scoperr init_targetsrD to_modulerroattach_datalayoutset_nvvm_reflect_ftzr get_functionsis_declarationis_external_linkageset_nvvm_maxnregrlink_extern_libsoptimize_module OPTIMIZE_O3 get_int_attrr])rsrcrrgryrfrrrllvm_modprocrntriplerrMr6rEtotal_num_warpsrets r make_llirzCUDABackend.make_llir<sg27DKK99"= ++B KH ''+ b! 11"5 ;;B? ''+ b! $$R( ::>>4c :c A MM & &r * s  ,,. ::>>. 4 ;km m>>#w/&z2)9)9:& xx@##H- ?? &++- 8'')a.C.C.E&&w7 8   .5.A.ABltTTBEB  ! !(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 t|jj\}}tjddd5}tjddd5}|j ||j |jdz} tjjdd d k(rd d gnd g} |jrgnd g} t|} tjjdd d k(rdd gng} |g| | d| d| |jd| } tj|dd|tjj!|jrtj"|jtjj!|jrtj"|jt'| d5}|j)}dddtjj!| rtj"| ddddddS#tj$$r}t'|j5}|j)}dddn #1swYnxYwtjj!|jrtj"|j|j*dk(rd}n2|j*dt,j.zk(rd}nd|j*}t1|dddj3|dd}~wwxYw#1swYUxYw#1swY&xYw#1swYSxYw)NFwz.ptx)deletemoderrz.logz.or/r0rz -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:   rq)rNrrItempfileNamedTemporaryFilewriteflushrMr2r3r4rrr<rr6r:removeCalledProcessErrorrrru returncodesignalSIGSEGVr r7)rrNrrryrLrhfsrcflogfbin line_infofmadrI opt_level ptxas_cmdelog_filelogerrorrwrs r make_cubinzCUDABackend.make_cubinsT[[--.q  ( (C O& SW  ' 'u3v N& RV JJsO JJL99t#DACPjPSBUX[B\&<=bman --2N3CD*:6D02 ?RTW0X\_0_,egIqqTq4q)q{SWRXEYq[_[d[dqfjqlpqI LydS77>>$)),IIdii(77>>$)),IIdii($dD! !Q !ww~~d# $M& & N +00 L$))_*"--/C***77>>$)),IIdii(<<3&?E\\S6>>%994E=all^LE E7+558E:33688I3F2Gr"KLL L" ! !G& & & N szMCL;B)H)> L; L.z(CUDABackend.add_stages..st~~c8W/Ur rc,j||Sre)r-rNrryrgrs rrz(CUDABackend.add_stages..sXwXb0cr ttgirc,j||Sre)rTrs rrz(CUDABackend.add_stages..st~~c8WV`/ar llircTj||jjSre)r]rrIrs rrz(CUDABackend.add_stages..s#dmmC7TXT_T_TdTd.er ptxcTj||jjSre)r~rrIrs rrz(CUDABackend.add_stages..s#XwX\XcXcXhXh0ir r)rrI)rstagesrgrys` `@r add_stageszCUDABackend.add_stagessD%%gll3 Uvcwaveu iwr cvt|jj}|d|jjS)Nr)rRrrI)rrGs rrzCUDABackend.hashs2#DKK$4$45!DKK,,-..r )rrr staticmethodrrrrr rrrr r]rrrrr-rTr]r~r functools lru_cacher __classcell__)rs@rrrs( ((#"y"T"#S#, >S*_ 5>"  FFP?B,)VjY//r r)+triton.backends.compilerrrtriton._C.libtritonrrrrtriton.runtime.errorsr dataclassesr rtypingr r r rtypesrrsr?rjrqr2r<pathlibrr0r"rr]rHr!rNrRrcrirorxrrrr~r rrsk;88,!--   # #0C00"!C!! CiSii$   T44 &&  $'?'?'?Tg/+g/r