9jٗ < ddlZddlZddlmZddlZddlmZddlmZddl m Z m Z m Z m Z ddlmZddlmZddlmZmZdd lmZdd lmZdd lmZd d lmZmZd dlm Z m!Z!d dl"m#Z#d dl$m%Z%d dl&m'Z'm(Z(d dl)m*Z*m+Z+m,Z,m-Z-d dl.m/Z/d dl0m1Z1m2Z2m3Z3m4Z4m5Z5d dl6m7Z7m8Z8m9Z9m:Z:m;Z;d dlZ>m?Z?m@Z@mAZAmBZBmCZCmDZDmEZEmFZFmGZGmHZHmIZIddlJmKZKmLZLmMZMmNZNmOZOmPZP ddlQZQeeQjZSdZTejeWZXej&jZYej&jZZe;deNejjeSdk\reLdneLddd Z]e;d!eOeLd"#Z^e;d$eOeLd%#Z_e;d&eOeLd'#Z`e;d(eOeLd)#Zae;d*eOeLd+#Zbejd,Zde8ejd-eYjj.Zge8ejd/d0eYjj1Zie8ejd2eYjj.Zke8ejd3eYjj.Zme8ejd4deYjj5Zpe8ejd6eYjj.Zrd7Zsdddd8d9Ztdid;Zue8etdZvd<ZwGd=d>e(ZxexZyGd?d@e(ZzdAZ{dBZ|ezdCdDe{Z}ezdEdFe|Z~e4eYjdGdjddHdIZe4eYjdGddHdJZe4eYjdGddddKdLZe4eYjdGdddMdNZejejfejejfej ej fej ej fej ej fgZejejgZej ej gZdOed:efdPZdOedQed:efdRZdOedSedTedQed:ef dUZdOedSed:efdVZ dkdWedXej dYedQed:ef dZZd:efd[Zd\ed]ed^ej d_ej d:eeeff d`Ze4eYjjdG dldaZejdbedzd:efdcZddZ dmdeedzfdfZdgZdhZy#eU$redZSdZTY0wxYw)nN)Any)counters)AutoHeuristicSelectAlgorithm) AHContextcontext_add_stridescontext_add_using_tf32 mm_operations)CppGemmTemplate)gen_best_config)opsV)make_fx) ScalingType) TorchVersion)configdistributed_autotune)CUTLASS2xGemmTemplateCUTLASS3xGemmTemplate)CKTileGemmTemplate)CKGemmTemplate)SubgraphChoiceCallerSubgraphTemplate)Buffer ChoiceCaller is_tritonLayout)MMKernelInputs) loweringsmake_pointwisemake_reductionregister_loweringtransform_args)autotune_select_algorithmExternKernelChoiceKernelTemplaterealize_inputsTritonTemplate) _use_cutlass_for_opceildivuse_aten_gemm_kernelsuse_ck_gemm_templateuse_ck_tile_gemm_templateuse_cpp_gemm_templateuse_cutlass_templateuse_decompose_k_choiceuse_nv_universal_gemm_template!use_triton_blackwell_tma_templateuse_triton_scaling_templateuse_triton_templateuse_triton_tma_template)_is_static_problemload_kernel_templatemm_argsmm_gridpersistent_mm_griduse_native_matmulTz0.0.0Fmmz3.3.0 triton_mmtriton_mm_rocm)namegridsource"cache_codegen_enabled_for_templateprologue_loads_all_inputsmm_persistent_tmatriton_persistent_tma_mm)r@rArB mm_persistenttriton_persistent_mm%scaled_mm_device_tma_epilogue_scalingtriton_epilogue_scaled_mm&scaled_mm_device_tma_main_loop_scalingtriton_main_loop_scaled_mm"blackwell_ws_persistent_device_tma,triton_blackwell_ws_persistent_device_tma_mmct|SN)r%)fns Y/media/conek/DATA/Code/OCR/venv/lib/python3.12/site-packages/torch/_inductor/kernel/mm.pylazy_register_extern_choicerSs b !!z at::mm_out) op_overloadzat::mm_dtype_outmm_dtype)r@rUz at::addmm_outzat::_int_mm_outzat::_sparse_semi_structured_mm)has_out_variantrUzat::_scaled_mm_outcb|jtjtjfvSrP) get_dtypetorchint8uint8)mats rR _is_int8_matr^s ==?uzz5;;7 77rToutalphabetac|jddk(r|jddk7s|jddk(rtj|d|||||Stj||||||S)z Giving torch.addmm a 1D tensor calls a different (faster) cublasLt kernel under the hood. There are a few shapes where this is slower, but they are rare. rr6r_)stridesizerZaddmm)inpmat1mat2r`rarbs rR bias_addmmrjsh 1 sxx{a/CHHQK14D{{3q643e$OO ;;sD$Cu4 HHrTreturncXdtfd}dtfd}dtfd}tj|jxs|j fdtj|jxs|j fdy)Nrkc\tjjj|ddS)Nr6r graphsizevarsstatically_known_equalsrds rR is_row_majorz.check_supported_striding..is_row_major#ww77q 1EErTc\tjjj|ddSNrr6rnrrs rR is_col_majorz.check_supported_striding..is_col_majorrtrTcttjjj |ddxs-tjjj |ddSrv)boolr rorprq)res rR has_zero_dimz.check_supported_striding..has_zero_dimsQ GG   4 4T!Wa @ Dww77QC  rTc*djS)Nz$mat_a must be row_major, got stride  get_stride)mat_asrRz*check_supported_striding..6u7G7G7I6JKrTc*djS)Nz$mat_b must be col_major, got stride r|)mat_bsrRrz*check_supported_striding..rrT)ryrZ_checkr}get_size)r~rrsrwrzs`` rRcheck_supported_stridingrsFFFF d  LLU%%'(JL9I,JK  LLU%%'(JL9I,JKrTc|jd}|jd}|jd}||z}|}tj|j|||d}|j|||} tj|| tj } tj | d} | j|jS)Nrr6)r6rr out_dtype) shaperZpermutereshapebmmfloat32sumtodtype) abk_splitsmnkk_partsB a_reshaped b_reshapedresult reduced_bufs rR decomposeKrs  A  A  A8mGAqyyAw7CJ1gq)J YYz: GF))FA&K >>!'' ""rTc@eZdZfdZdeedededeffd Z xZ S)DecomposeKSugraphTemplatec&t|dy)N decompose_kr@)super__init__)self __class__s rRrz"DecomposeKSugraphTemplate.__init__s   rT input_nodeslayoutk_splitrkcddlm}ddlm}d|d}d|}|5|}t t j t||} t |%|||| | cdddS#1swYyxYw) Nrenable_python_dispatcherrselect_decomp_tabledecompose_k_mm__splitzk_split=)rr@rr make_fx_graph description) torch._dispatch.pythonr decompositionrr functoolspartialrrgenerate) rrrrrrr@rdecompositionsrQrs rRrz"DecomposeKSugraphTemplate.generates D7  0! m % ' 02N!!*w?B 7#' ' $   s AA**A3) __name__ __module__ __qualname__rlistrrintrr __classcell__rs@rRrrs< &\   rTrcJeZdZdededeffd Zdeedede ffd Z xZ S) ContiguousTemplater@rrQcP||_||_||_t||y)Nr)r@rrQrr)rr@rrQrs rRrzContiguousTemplate.__init__s. &   rTrrrkcddlm}ddlm}|5|}t |j |}t ||j||||jcdddS#1swYyxYw)Nrrrrr) rrrrrrQrrr@r)rrrrrrrQrs rRrzContiguousTemplate.generatesp D7 % ' 02NB 7#YY'  ,, $   s AA$$A-) rrrstrrrrrrrrrrs@rRrrsG S s  &\  rTrcJtj||jSrP)rZr= contiguous)rrs rR contiguous_mmr(s 88Aq||~ &&rTcLtj|||jSrP)rZrfr)rgrrs rRcontiguous_addmmr,s ;;sAq||~ ..rTrz contiguous mmrzcontiguous addmm)type_promotion_kindrcp |j}tj|j|k(dtjjjdvdtj||k(xs7|tj k(xr"|tj tjfvdt|rttjdttj|d}t|giddd \}}tjjrOj tj tjfvr#fd }|Dcgc]}t#||}}t#t$j&|} t)d | d } | St+|||\} } } }}t-|\}}d}t/|g|}t0dd| d| d| xxd z cc<t2j5d| | | j|j|g}t-|\}}t6}i}| t8}d|i}g}i}t;r"|j=||r|||j><||rtA|drtC| | | r|j=tDtjFdk(}|stC| | | ds|j=tHtK||dr|j=tLnTtO||drEtjPjR|j=tTn|j=tV|j=tX|j[t\j^ja||d||@|r>tc|| | | r0tedr%tgjh|||jk|5|r3tm|| | | r%tojp|||jk|5|r3ts|| | | r%tujv|||jk|"|r ty|| | | |rddl=m>}|||||2t||r%tjv|||jk|g}||rtA|rtjjj|rtrg}t;r|j=dt|}|j[t\j^ja|tHgdt|| | | |||tdd| }tjjj|s*|#t|dkDr|Dcgc] }||vs| }}n|d|}|MtjD]:} |j=t| j|jk|<d}|0tjjjr t|}tj|||jk|x}r|St|||jk||!\} }!| Scc}wcc}w)"z_ Lowering for autotuning aten.mm with different backends (Aten, Triton, CUTLASS, etc.) Ncy)Nzinput dtypes must be the samerrTrRrztuned_mm..ArT)cudaxpucy)Nz+out_dtype is only supported for CUDA or XPUrrrTrRrztuned_mm..ErrTcy)NzFout_dtype must be the same as input dtype or fp32 for fp16/bf16 inputsrrrTrRrztuned_mm..MrrTrTF)argskwargs broadcastrconvert_input_to_boolcHtj|jdS)NF)use_compute_types)r to_dtyper)xrhs rR _to_dtypeztuned_mm.._to_dtypens||AtzzUKKrTdotr6rrr=r aten_mm_infozaten.mm__zOTuned aten.mm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%srcheck_max_autotune exhaustiver)threshold_multiple output_layout add_guardskwarg_overrides)add_nv_universal_gemm_choices extern_mmzmm-ah )top_kalways_included)best_config_future)QrYrZr get_devicetyperfloat16bfloat16r<raten unsqueezer#inductor_configtritoncodegen_upcast_to_fp32rr r rr!r9r7rrloginfoaten_mm aten_mm_dtyper+appenduidr4r0decompose_k_subgraph_templatemax_autotune_gemm_search_space mm_templater2.blackwell_ws_persistent_device_tma_mm_templater5versionhippersistent_tma_mm_templatepersistent_mm_templatemm_contiguous_subgraph_templateextendr choicesget_template_configsr/r)radd_cutlass_gemm_choicesnodesr,radd_ck_gemm_choicesr-r add_choicesr1codegen.nv_universal_gemmrr.r _inductorrrun_autoheuristicrlenmm_autoheuristicr collect_autoheuristicexternal_matmulrSbindremote_gemm_autotune_cacher rmaybe_autotune_remoter$)"rhrirr input_dtyperrrr mul_pointwise dot_reductionrrr static_shape is_nonzeror@ kernel_inputsr  aten_handleraten_extra_kwargstemplates_to_user is_exhaustiverrr num_choices_before_extra_configs ah_choiceschoicerboxnoders"` rRtuned_mmr*8s nn&  NN  + 3   OO  " "o 5 A    $ U]]*CEMM5>>#BB \  .t$(r2(q1% $"'  f  ! ! 8 8TZZ MM NNL >  L;??Q-N9-a0?D?/sww/6 -u-mQ? #* d6Y#Aq!VT4 2&9L* D#D$<9EM ^xs!A3as3494HHY     #%G1&9L*'.L(*$ ()4BD13O - 0AOL,, -   4 @ !!Q *  # #$A B (FF,V  6q!QST U  # #K 00d&T!''(VW(d&T==$$,$++,FG$++,BC ?@ NN &&   + '    Aq 1  %66 V]002 Z,@Aq,Q**7FM9HH^      2&9L*W:W2FvqRSUV2WK"$G#D$<5;;GMCE -)Te  , NN &&}6FM*4066 V]002TPT (g}7J7J7LfUGD! KrT)rarbrc tt||r|dk(rd}nttj||}|dk(rd}n8ttj|ttj||}ttj ||St ||||\}} } }}}} t|\} } d}t| ||gt||}t|||gt||}g}tdd|d| d| xxd z cc<tjd || | |j|j|| r tjsftj sV|j#t$j&j)|t*g|t-|||j/|\}}|Sg}t1rt*g}|j3ddk(ret5|j7d k(rItj8j:r/t$j<j>s|jAtB|j#t$j&j)|||| rtE|d r|jAtFtI|||dr|jAtJnTtM|||drEtNjPjR|jAtTn|jAtV|jAtX|j#t$j&j)|||| rHt[||| | r:t]|r/t_j`|||j/gd||gd| r=tc||| | r/tejf|||j/gd||gdti|||r)tkjl|||j/||dt-|||j/|\}}|S)zb Lowering for autotuning aten.addmm with different backends (Aten, Triton, CUTLASS, etc.) rrrf)rarb)scalarsrz aten.addmm_rr6zRTuned aten.addmm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%srFrTr)r6rr)reorder)rrr6)rarb input_reorder)rarbhas_bias)7r<rrmulr=addr9r7rdictrrrrYr max_autotunemax_autotune_gemmr r r r  aten_addmmr$rr+r}rrrautotune_cublasLtro cpp_wrapperraten_bias_addmmr4rr2rr5rZrrrr"addmm_contiguous_subgraph_templater/r)rr r,rrr.r r)rgrhrirarbrarg1arg2rrr inp_expandedrrr@r kernel_inputs_atenr r)rr#aten_templatess rR tuned_addmmrIas t$ 19DTXX&tS1D A:DTXX&ui.@t.LMD"4..18dCPV0W-Aq!VT41&9L* D# tT"Du4,HM( dD4e$#?#%G ^{1#Qqc1#671<7HH\        ) )_-N-N II * *"   , '=..0& a BDEOL NN Q 1 $CLLN#q(&&88GG''  ! !/ 2  II * *+=~t T )&UK , , $f   # #$R S $T4vRV W}}  ( ''(BC ''(>? BC NN &&}6FM  Aq 1  %66       2# *61a;**       2# VT40##      !  (g}7J7J7LfUGD! KrT)rrcddlm}||||\}}}|j\}}|j\}} |j\} } tjj j ||} tjj j d|z| } |6ddlm}||j|r|n|j| | g| dg}n |Jdtrtj|||f||gng}| | zdk7r6t|| | | r(tdrt!j"|||||gd d t%d||||f|\}} |S) Nr)r'r) FixedLayoutr6z,out_dtype is ignored if layout is specified.rsparse_semi_structured_mmTr.) torch._inductor.select_algorithmr'rr rorpcheck_equals_and_simplifytorch._inductor.irrKrrYr+aten__sparse_semi_structured_mmrr/r)rr r$)rh mat1_metarirrr'm1k1m2rk2rrrrKr r)s rRtuned_sparse_semi_structured_mmrVs}@+4DAD)T ]]_FB    EB MMOEB 222r:A 221r62>A ~2 OO "I(8 F F   P"PP  ! " , 0 0y$'9 1    A Aq 1  ; <66 VdD)4tRV (#WtY.EvGD! KrTszcFt|dk(xstd|DS)Nrc3pK|].}tjjj|d0yw)r6Nrn).0ds rR z)_is_tensorwise_scaling..+s,!;<00A6!s46)rall)rWs rR_is_tensorwise_scalingr^*s+ GqL S!@B!rT transposech|rdnd}tjjj||dS)Nrrr6rn)rWr_idxs rR_is_rowwise_scalingrb0s,!bC 77   3 3BsGQ ??rT tensor_sz tile_sizec|rdnd}|rdnd}tjjj||||xr:tjjj||t |||S)Nr6rr rorprqr*)rWrcrdr_lhsrhss rR_is_blockwise1xTILESIZE_scalingri5sq!aC!aC 77   3 3 33   ''   2 2 333 rTctjjj|dt |ddxr:tjjj|dt |ddS)Nrr6rf)rWrcs rR_is_blockwise128x128_scalingrlAsd 77   3 3 1wy|S)  V ''   2 22a5')A,PS:T UVrTt scale_size scaling_typecX|xtjk(r t|Sxtjk(r t ||Sxtj k(rt ||jd|Stjk(rt||jS td|)NrkUnsupported scaling type ) r TensorWiser^RowWiserbBlockWise1x128rirBlockWise128x128rlAssertionError)rmrnror_s rRis_desired_scalingrwGs  #[ # #)*5 5 [ &z9= = '[ ' '2AJJL#y  ) )/ AJJLI I  #<\N!KL LrTct|xtjk(rytjk(ry td|d)Nrkrqz in get_tile_size)rrurtrv) scale_options rR get_tile_sizerz\s<  )[ ) )  ' '  +L>9JK rTr~r scale_a_size scale_b_sizectD](\}}t|||st|||ds$||fcStd|d|)NT)r_z1Inductor Triton does not support scale_a.shape = z, scale_b.shape = ) scaling_pairsrwrv)r~rr{r|scale_option_ascale_option_bs rRget_scaling_optionsrhsc +82&  <   nPTU!>1 1 2  ;L>I[\h[ij rTc t||||\} } } }}}tdd| d| d| xxdz cc<tjd| | | |j |j |d} t ||t ||\} }|s||| |g}nt |}||| ||g}t|dd| }g}g}i}tr3|jtt|| |tj<t|\}}|jtj k(r|rt#|d d rqt|}| j$|j$}}t'||||\}}t)|||d r|s|j*|d<|j*|d<t-||t.r)|jt0||t0j<nat-||t2rEt5||d<t5||d<|jt6||t6j<n t9dt;|||d r*|s(|jt<||t<j<t-||t.r(|jt>||t>j<|jAtBjDjG||| ||r"tI|| | | ||rddl%m&}||||||jtj k7rtO| |||\}}|S|r@tQ|| | | r2tS| r'tUjV|||jY||r3t[|| | | r%t]j^|||jYtO| ||jY|\}}|S)a9 Performs an optimized matrix multiplication where scaling factors are applied to the inputs and/or output. Args: mat1 (Tensor): First input matrix mat2 (Tensor): Second input matrix scale1 (Tensor): Scale factor applied to mat1 (supports broadcasting) scale2 (Tensor): Scale factor applied to mat2 (supports broadcasting) bias (Tensor, optional): Optional bias tensor to add to the result layout: Layout hint for optimization Returns: Tensor: The result of the scaled matrix multiplication rrzaten._scaled_mm.default_rr6z_Tuned aten._scaled_mm.default: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%s scaled_mmr)mat1_idxmat2_idxr)ruse_fast_accumTF) enable_float8r)USE_FAST_ACCUMrSCALE_RECIPE_ASCALE_RECIPE_B TILE_SIZE_A TILE_SIZE_BzpInductor Triton does not support scaling options that are present in both epilogue scaling and main loop scalingrr)$add_nv_universal_scaled_gemm_choices)r )r)0r9rrrrYrr'rr+r aten__fp8_mmr<rr7rrZrr4rrr5valuer3epilogue_scaling_types.scaled_mm_device_tma_epilogue_scaling_templatemain_loop_scaling_typesrz/scaled_mm_device_tma_main_loop_scaling_templatervr2rrr r r r r1rrr$r/r)rr rr,rr)r~rscale_ascale_bbias scale_resultrrrrrrr@ scale_a_real scale_b_realr bias_realr r r#rrr overridersr{r|rrrr)s rRtuned_scaled_mmrys8%, uVy%!Aq!VUE ^7s!A3asCDIDHHi      DUE*!/!AL, e\<@ "4( e\<K #a!yM#%GCEO -,0-  (()'v.MAz  &  du U8 %1%7%79K9Kl )< 5, * & $E5SW X+9+?+?J' (+9+?+?J' (*0F!''(VW N R RS-0G-:.,I =),9.,I =) ''(WX O S ST%G .uFt   # #$R S JNN O ' N,B   # #K 0/9OKOO , NN &&   + ' 4VQ1eUST,   '  }} %+D';Oa   Aq 1  %66      !)  *61a;**7FM.6s4z#s#4s)r])dimss rR dims_are_intr5s 4t4 44rTrc t|||||\}}}t|||gsyt||\} } fd}d}||||||| | }t|||||| }| |j | | S|j S)NcVt}|jd||jd||jd||jd|jjd|jd|jjdt |d|t |d ||jd |jj d|jd |jj dd k(r t ||jj|S) Nrrr mat1_dtypeT)is_categorical mat2_dtyperhri mat1_iscontig mat2_iscontigr=)r add_featurerrr is_contiguousr) rrrrhri mat1_stride mat2_stridecontextr@s rR get_contextz%mm_autoheuristic..get_contextLs+C#C#C#L$++*;*;DQL$++*;*;DQGV[9GV[9 T[[668    T[[668   4< "7DKK,=,= >rTcyrPrrrTrRfallbackz"mm_autoheuristic..fallback_srT)rr rrr@augment_context precondition)r)get_size_hintsrget_size_hints_stridesrget_top_k_choices_callerget_choice_caller)rhrirrrr r@rr rrrrrrrr autoheuristics ` rRrr9sT4Aq1GAq! Aq "5dDAK&!Q4{KHG0 !M 55 ?6    * * ,,rTctt|trt|ts:tjjj |j \}}t|trt|ts:tjjj |j \}}|||fSrP)rrr rorpoptimization_hintsr)rhrirrrs rRrrvs{ a Z3%7!!44T]]_EA a Z3%7!!44T]]_EA a7NrTc|jj}|jj}||g}g}|D]L}t|ts)tj j j|}|j|N|d|dfSrv) rrdrrr rorprr)rhrirrstrides strides_hintsrds rRrrs++$$K++$$KK(GM%&#&WW%%88@FV$%  ]1- --rT)rkNrP)F)NNNFN)NN)rloggingtypingrrZtorch._dynamo.utilsr+torch._inductor.autoheuristic.autoheuristicr1torch._inductor.autoheuristic.autoheuristic_utilsrrrr )torch._inductor.codegen.cpp_gemm_templater *torch._inductor.remote_gemm_autotune_cacher torch._inductor.virtualizedr r "torch.fx.experimental.proxy_tensorrtorch.nn.functionalrtorch.torch_versionrrrrcodegen.cutlass.gemm_templaterr,codegen.rocm.ck_tile_universal_gemm_templater'codegen.rocm.ck_universal_gemm_templatercodegen.subgraphrrirrrrrr rloweringrr r!r"r#select_algorithmr$r%r&r'r(utilsr)r*r+r,r-r.r/r0r1r2r3r4r5 mm_commonr7r8r9r:r;r<r __version__triton_version has_triton ImportError getLoggerrrrprimsrrrrrrrrcacherSr=r`r dtype_outrrfr?_int_mmr2_sparse_semi_structured_mmdefaultrP _scaled_mmrr^rjrrBrrrrrrr rCr*r4rIrVrrrsrtrur~rrryr^rbrrirlTensorrwrztuplerrrrrrrrrTrRrs (T FF.6+,>XMDE88*!&"4"45NJ g! yy~~    !n&? , . /'+"  ,    : ;(    6 72@ 0   ; <2.3A 1   < =3/ 2@ -   N O2. "" UXX| M" HH !!    KKdjjnn " MM$$,,2B2B #5 $$$//77 #" *8K8K 8(,11I4%Z6 #  0 F!: ;)D'/#5_m#&8*,<&" 4775s4s6sl4<!;!;<!!;#=#=>  &00+2E2EF&55{7S7STst @C@D@T@    (+ 8<   VSVSVTV M M MM M  M* 3   ,,,,   ; #$ "4??**E  rFrj t 5:- ::-z .m!!'*NJsTTT