
    ^jd                     
   d dl Z d dlZd dlmZ d dlZd dlmZ d dlmZ d dl	m
Z
mZmZmZ d dlmZ d dlmZ d dlmZmZ d d	lmZ d d
lmZ d dlmZ d dlmZ ddlmZm Z  ddl!m"Z"m#Z# ddl$m%Z% ddl&m'Z' ddl(m)Z)m*Z* ddl+m,Z,m-Z-m.Z.m/Z/ ddl0m1Z1 ddl2m3Z3m4Z4m5Z5m6Z6m7Z7 ddl8m9Z9m:Z:m;Z;m<Z<m=Z= ddl>m?Z?m@Z@mAZAmBZBmCZCmDZDmEZEmFZFmGZGmHZHmIZImJZJmKZK ddlLmMZMmNZNmOZOmPZPmQZQmRZR 	 d dlSZS eeSj                        ZUdZV ej                  eY      ZZej&                  j                  Z[ej&                  j                  Z\ e=dePej                  j                  eUdk\  r eNd      n eNd       dd!      Z_ e=d"eQ eNd#      $      Z` e=d%eQ eNd&      $      Za e=d'eQ eNd(      $      Zb e=d)eQ eNd*      $      Zc e=d+eQ eNd,      $      Zde j                  d-        Zf e:ej                  d.e[j                  j                  /      Zi e:ej                  d0d1e[j                  j                  2      Zk e:ej                  d3e[j                  j                  /      Zm e:ej                  d4e[j                  j                  /      Zo e:ej                  d5de[j                  j                  6      Zr e:ej                  d7e[j                  j                  /      Ztd8 Zudddd9d:Zvdqd<Zw e:evd      Zxd= Zy G d> d?e*      Zz ez       Z{ G d@ dAe*      Z|dB Z}dC Z~ e|dDdEe}      Z e|dFdGe~      Z e6e[j                  dH      drddIdJ       Z e6e[j                  dH      ddIdK       Z e6e[j                  dH      ddddLdM       Z e6e[j                  dH      dddNdO       Zej
                  ej
                  fej                  ej                  fej                  ej                  fej                  ej                  fej                  ej                  fgZej
                  ej                  gZej                  ej                  gZdPed;efdQZdPedRed;efdSZdPedTedUedRed;ef
dVZdPedTed;efdWZ	 dsdXedYej$                  dZedRed;ef
d[Zd;efd\Zd]ed^ed_ej$                  d`ej$                  d;eeef   f
daZ e6e[j.                  j                  dH      	 	 	 	 	 dtdbee   dcee   ddee   deee   dfee   dgee   fdh       Z e6e[j                  j                  dH      	 	 	 	 	 dtdi       Ze j                  djedz  d;efdk       Zdl Z	 	 dudmedz  fdnZdo Zdp Zy# eW$ r  ed      ZUdZVY zw xY w)v    N)Any)counters)AutoHeuristicSelectAlgorithm)	AHContextcontext_add_stridescontext_add_using_tf32mm_operations)CppGemmTemplate)gen_best_config)opsV)make_fxScalingType)TorchVersion)
OrderedSet   )configdistributed_autotune)CUTLASS2xGemmTemplateCUTLASS3xGemmTemplate)CKTileGemmTemplate)CKGemmTemplate)SubgraphChoiceCallerSubgraphTemplate)BufferChoiceCaller	is_tritonLayout)MMKernelInputs)	loweringsmake_pointwisemake_reductionregister_loweringtransform_args)autotune_select_algorithmExternKernelChoiceKernelTemplaterealize_inputsTritonTemplate)_use_cutlass_for_opceildivuse_aten_gemm_kernelsuse_ck_gemm_templateuse_ck_tile_gemm_templateuse_cpp_gemm_templateuse_cutlass_templateuse_decompose_k_choiceuse_nv_universal_gemm_template!use_triton_blackwell_tma_templateuse_triton_scaling_templateuse_triton_templateuse_triton_tma_template   )_is_static_problemload_kernel_templatemm_argsmm_gridpersistent_mm_griduse_native_matmulTz0.0.0Fmmz3.3.0	triton_mmtriton_mm_rocm)namegridsource"cache_codegen_enabled_for_templateprologue_loads_all_inputsmm_persistent_tmatriton_persistent_tma_mm)rB   rC   rD   mm_persistenttriton_persistent_mm%scaled_mm_device_tma_epilogue_scalingtriton_epilogue_scaled_mm&scaled_mm_device_tma_main_loop_scalingtriton_main_loop_scaled_mm"blackwell_ws_persistent_device_tma,triton_blackwell_ws_persistent_device_tma_mmc                     t        |       S N)r'   )fns    d/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/kernel/mm.pylazy_register_extern_choicerU      s    b!!    z
at::mm_out)op_overloadzat::mm_dtype_outmm_dtype)rB   rW   zat::addmm_outzat::_int_mm_outzat::_sparse_semi_structured_mm)has_out_variantrW   zat::_scaled_mm_outc                 b    | j                         t        j                  t        j                  fv S rR   )	get_dtypetorchint8uint8)mats    rT   _is_int8_matr`      s     ==?uzz5;;777rV   outalphabetac                    | j                  d      dk(  r| j                  d      dk7  s| j                  d      dk(  rt        j                  | d   |||||      S t        j                  | |||||      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.
    r   r8   ra   )stridesizer\   addmm)inpmat1mat2rb   rc   rd   s         rT   
bias_addmmrl      sh     	

1sxx{a/CHHQK14D{{3q643e$OO;;sD$Cu4HHrV   returnc                 X    dt         fd}dt         fd}dt         fd}t        j                   | j                               xs  | j	                                fd       t        j                   |j                               xs  |j	                               fd       y )Nrm   c                 \    t         j                  j                  j                  | d   d      S )Nr8   r   graphsizevarsstatically_known_equalsrf   s    rT   is_row_majorz.check_supported_striding.<locals>.is_row_major   #    ww77q	1EErV   c                 \    t         j                  j                  j                  | d   d      S Nr   r8   rp   rt   s    rT   is_col_majorz.check_supported_striding.<locals>.is_col_major   rv   rV   c                     t        t        j                  j                  j	                  | d   d      xs- t        j                  j                  j	                  | d   d            S rx   )boolr   rq   rr   rs   )rg   s    rT   has_zero_dimz.check_supported_striding.<locals>.has_zero_dim   sQ    GG44T!Wa@ Dww77QC
 	
rV   c                  *    d j                          S )Nz$mat_a must be row_major, got stride 
get_stride)mat_as   rT   <lambda>z*check_supported_striding.<locals>.<lambda>       6u7G7G7I6JK rV   c                  *    d j                          S )Nz$mat_b must be col_major, got stride r~   )mat_bs   rT   r   z*check_supported_striding.<locals>.<lambda>   r   rV   )r{   r\   _checkr   get_size)r   r   ru   ry   r|   s   ``   rT   check_supported_stridingr      s    F FF F
d 
 
LLU%%'(JL9I,JK 
LLU%%'(JL9I,JKrV   c                    | j                   d   }|j                   d   }| j                   d   }||z  }|}t        j                  | j                  |||      d      }|j                  |||      }	t        j                  ||	t        j
                        }
t        j                  |
d      }|j                  | j                        S )Nr   r8   )r8   r   r   	out_dtype)	shaper\   permutereshapebmmfloat32sumtodtype)abk_splitsmnkk_partsB
a_reshaped
b_reshapedresultreduced_bufs               rT   
decomposeKr      s    	
A	
A	
A8mGAqyyAw7CJ1gq)JYYz:GF))FA&K>>!''""rV   c                   @     e Zd Z fdZdee   dededef fdZ	 xZ
S )DecomposeKSugraphTemplatec                 &    t         |   d       y )Ndecompose_krB   )super__init__)self	__class__s    rT   r   z"DecomposeKSugraphTemplate.__init__   s     	 	
rV   input_nodeslayoutk_splitrm   c                     ddl m} ddlm} d| d}d|} |       5   |       }t	        t        j                  t        |      |      }	t        
| %  ||||	|	      cd d d        S # 1 sw Y   y xY w)
Nr   enable_python_dispatcherr   select_decomp_tabledecompose_k_mm__splitzk_split=)r   rB   r   r   make_fx_graphdescription)
torch._dispatch.pythonr   decompositionr   r   	functoolspartialr   r   generate)r   r   r   r   r   r   rB   r   decompositionsrS   r   s             rT   r   z"DecomposeKSugraphTemplate.generate   s     	D7 	0!
m%' 	02N!!*w?B
 7#' ' $ 	 	 	s   A A**A3)__name__
__module____qualname__r   listr   r   intr   r   __classcell__r   s   @rT   r   r      s<    

&\  	
 
 rV   r   c                   J     e Zd Zdededef fdZdee   dede	f fdZ
 xZS )	ContiguousTemplaterB   r   rS   c                 P    || _         || _        || _        t        |   |       y )Nr   )rB   r   rS   r   r   )r   rB   r   rS   r   s       rT   r   zContiguousTemplate.__init__  s.    	& 	 	
rV   r   r   rm   c                     ddl m} ddlm}  |       5   |       }t	        | j
                  |      }t        |   | j                  |||| j                        cd d d        S # 1 sw Y   y xY w)Nr   r   r   r   r   )
r   r   r   r   r   rS   r   r   rB   r   )r   r   r   r   r   r   rS   r   s          rT   r   zContiguousTemplate.generate  sp    
 	D7%' 	02NB
 7#YY'  ,, $ 	 	 	s   AA$$A-)r   r   r   strr   r   r   r   r   r   r   r   r   s   @rT   r   r     sG    
S 
s 
 
&\  
	 rV   r   c                 J    t        j                  | |j                               S rR   )r\   r?   
contiguous)r   r   s     rT   contiguous_mmr   )  s    88Aq||~&&rV   c                 L    t        j                  | ||j                               S rR   )r\   rh   r   )ri   r   r   s      rT   contiguous_addmmr   -  s    ;;sAq||~..rV   r   zcontiguous mmr   zcontiguous addmm)type_promotion_kindr   c                p    | j                         }t        j                  |j                         |k(  d        t        j                   j                         j                  dv d        t        j                  ||k(  xs7 |t        j
                  k(  xr" |t        j                  t        j                  fv d        t         |      rt        t        j                      d       t        t        j                     |d      }t         |gi ddd	
      \  }}t        j                  j                  rO j                   t        j                  t        j                  fv r# fd}|D cg c]  } t#        |      |       }} t#        t$        j&                        | }	 t)        d      |	d      }
|
S t+         |||      \  }}}} }t-        |      \  }}d}t/         |g|      }t0        d   d| d| d| xx   dz  cc<   t2        j5                  d||| j                         |j                         |       g }t-        |      \  }}t6        }i }|
t8        }d|i}g }i }t;               r"|j=                  |       |r|||j>                  <   ||rtA        |d      rtC        |||      r|j=                  tD               t        jF                  dk(  }|stC        |||d      s|j=                  tH               tK         ||d      r|j=                  tL               nTtO         ||d      rEt        jP                  jR                  |j=                  tT               n|j=                  tV               |j=                  tX               |j[                  t\        j^                  ja                  ||d|             |@|r>tc        ||||      r0te        d      r%tg        jh                  |||jk                                |5|r3tm        ||||      r%to        jp                  |||jk                                |5|r3ts        ||||      r%tu        jv                  |||jk                                |"|r ty        |||| |      rddl=m>}  ||||       |2t        | |      r%t        jv                  |||jk                                 |g}||rtA        |      rt        j                  j                  j                  |      rt               rg }t;               r|j=                  d       t        |      }|j[                  t\        j^                  ja                  |tH        gd             t         |||||||t               dd|       }t        j                  j                  j                  |      s*|#t        |      dkD  r|D cg c]	  }||v s| }}n|d| }|Mt        j                  D ]:  }|j=                  t        |      j                  |jk                         |             < d}|0t        j                  j                  j                  rt         |      }t        j                  |||jk                         |      x}r|S t        |||jk                         ||!      \  } }!| S c c}w c c}w )"z_
    Lowering for autotuning aten.mm with different backends (Aten, Triton, CUTLASS, etc.)
    Nc                       y)Nzinput dtypes must be the same r   rV   rT   r   ztuned_mm.<locals>.<lambda>B      rV   )cudaxpuc                       y)Nz+out_dtype is only supported for CUDA or XPUr   r   rV   rT   r   ztuned_mm.<locals>.<lambda>F  r   rV   c                       y)NzFout_dtype must be the same as input dtype or fp32 for fp16/bf16 inputsr   r   rV   rT   r   ztuned_mm.<locals>.<lambda>N  r   rV   r   TF)argskwargs	broadcastr   convert_input_to_boolc                 H    t        j                  | j                  d      S )NF)use_compute_types)r   to_dtyper   )xrj   s    rT   	_to_dtypeztuned_mm.<locals>._to_dtypeo  s    ||AtzzUKKrV   dotr8   r   r   r?   r   aten_mm_infozaten.mm__zOTuned aten.mm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sr   check_max_autotune
exhaustiver   )threshold_multipleoutput_layout
add_guardskwarg_overrides)add_nv_universal_gemm_choices	extern_mmzmm-ah
   )top_kalways_included)best_config_future)Qr[   r\   r   
get_devicetyper   float16bfloat16r>   r!   aten	unsqueezer%   inductor_configtritoncodegen_upcast_to_fp32r   r"   r   r   r#   r;   r9   r    r   loginfoaten_mmaten_mm_dtyper-   appenduidr6   r2   decompose_k_subgraph_templatemax_autotune_gemm_search_spacemm_templater4   .blackwell_ws_persistent_device_tma_mm_templater7   versionhippersistent_tma_mm_templatepersistent_mm_templatemm_contiguous_subgraph_templateextendr   choicesget_template_configsr1   r+   r   add_cutlass_gemm_choicesnodesr.   r   add_ck_gemm_choicesr/   r   add_choicesr3   codegen.nv_universal_gemmr   r0   r
   	_inductorr   run_autoheuristicr   lenmm_autoheuristicr	   collect_autoheuristicexternal_matmulrU   bindremote_gemm_autotune_cacher   r   maybe_autotune_remoter&   )"rj   rk   r   r   input_dtyper   r   r   r   mul_pointwisedot_reductionr   r   r   static_shape
is_nonzerorB   kernel_inputsr  aten_handleraten_extra_kwargstemplates_to_user   is_exhaustiver   r   r    num_choices_before_extra_configs
ah_choiceschoicer   boxnoder   s"   `                                 rT   tuned_mmr,  9  s   
 nn&NN+3	
 	OO""o5A	
 	$ U]]* CEMM5>>#BB\	
. t$(r2(q1% $"'
f !!88TZZMMNNL
 >

L ;??Q-N9-a0?D?/sww/6-u-mQ? #*d6Y#Aq!VT4  2&9L*D #D$<9EM ^xs!A3as3494HHY			 #%G1&9L*'.L(*$()4BD13O-0AOL,,- 	4@!!Q*##$AB (FF,V 6q!QST U##K00d&T !''(VW(d&T ==$$,$++,FG$++,BC ?@NN			&&+	 	' 	
 	 Aq1%66V]002	
 Z,@Aq,Q**7FM<O<O<QRZ,EfaQRTU,V&&w8K8K8MN 	*61aD$GM%gv}E264F##!	
 ,K'OO""44T:dO """;/+.w<(II** 		
 &O+

 %%;;DA%#j/A*=
 18Pf6Z;O6PP!"C#CD 00 	ANN+A.33M4G4G4I6R	
 U__33NN -T48"88g}**,f s  
'-GD! Kw @~ Qs   *Z.	Z3Z3c          	         t        | ||t        j                        \  }}}}} }d}t        d   d| d| d| xx   dz  cc<   t        j                  d|||| j                         |j                         |       t        |      \  }}|xr |xr t        ||||      }	g }
t        | |gt        j                        }g }t               r|j                  t               |r#t        |d	d
      r|j                  t               |
j                  t         j"                  j%                  |||             |	r3t'        |      r(t)        j*                  |
||j-                         d	d	       t/        ||
|j-                         |      \  }}|S )Nr   int_mmr   zaten._int_mm_r   r8   zTTuned aten._int_mm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sr   TF)enable_int32r   fuseablenon_fuseable)r;   r\   int32r   r   r   r[   r9   r1   r    r-   r  aten__int_mmr6   r  r  r   r  r  r+   r   r  r  r&   )rj   rk   r   r   r   r   rB   r   r!  use_cutlassr  r"  r%  r+  r   s                  rT   tuned_int_mmr6  0  s    #*d6U[[#Aq!VT4 D^}QCq1QC89Q>9HH^			  2&9L*W:W2FvqRSUV2WK"$G #D$<5;;GM CE-)Te 	, NN			&&}6FM *4066V]002TPT	
 (g}7J7J7LfUGD!KrV   )rc   rd   r   c          	         t        ||      r|dk(  rd}nt        t        j                     ||       }|dk(  rd}n8t        t        j                     |t        t        j                     ||            }t        t        j
                     ||      S t        ||| |      \  }}	}
}}}}t        |       } t        |      \  }}d}t        |||gt        ||            }t        | ||gt        ||            }g }t        d   d| d|	 d|
 xx   d	z  cc<   t        j                  d
||	|
|j                         |j                         |       |r t        j                   sft        j"                  sV|j%                  t&        j(                  j+                  |t,        g|             t/        |||j1                         |      \  }}|S g }t3               rt,        g}| j5                         d   dk(  ret7        | j9                               dk(  rIt        j:                  j<                  r/t&        j>                  j@                  s|jC                  tD               |j%                  t&        j(                  j+                  |||             |rtG        |d      r|jC                  tH               tK        |||d      r|jC                  tL               nTtO        |||d      rEtP        jR                  jT                  |jC                  tV               n|jC                  tX               |j%                  t&        j(                  j+                  |tZ        g|             |j%                  t&        j(                  j+                  |||             |rHt]        |||	|
      r:t_        |      r/ta        jb                  |||j1                  g d      ||g d       |r=te        |||	|
      r/tg        jh                  |||j1                  g d      ||g d       tk        |||      r)tm        jn                  |||j1                         ||d       t/        |||j1                         |      \  }}|S )zb
    Lowering for autotuning aten.addmm with different backends (Aten, Triton, CUTLASS, etc.)
    r   r   rh   )rc   rd   )scalarsr   zaten.addmm_r   r8   zRTuned aten.addmm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sr   Fr   Tr   )r8   r   r   )reorder)r   r   r8   )rc   rd   input_reorder)rc   rd   has_bias)8r>   r!   r   mulr?   addr;   r)   r9   r    dictr   r   r   r[   r   max_autotunemax_autotune_gemmr  r   r  r  
aten_addmmr&   r  r-   r   r  r   r   autotune_cublasLtrq   cpp_wrapperr  aten_bias_addmmr6   r  r4   r  r7   r\   r  r  r	  r
  "addmm_contiguous_subgraph_templater1   r+   r   r  r.   r   r  r0   r
   r  )ri   rj   rk   rc   rd   r   arg1arg2r   r   r   inp_expandedr   r!  rB   r"  kernel_inputs_atenr  r+  r   r%  aten_templatess                         rT   tuned_addmmrK  b  s   
 t$19DTXX&tS1DA:DTXX&ui.@t.LMD"4.. 18dCPV0W-Aq!VT4

C1&9L*D #	tT"Du4,HM (	dD4e$#? #%G ^{1#Qqc1#671<7HH\			 ))_-N-NII**"	
 ,'=..0&
a BDEOLNNQ1$CLLN#q(&&88GG''!!/2 	II**+=~tT	
 )&UK,,$f
 ##$RS$T4vRVW}}  ( ''(BC ''(>? 	II**"%G$H$	
 NN			&&}6FM
 	 Aq1%66 	2#		
 *61a;** 	2#		
 VT40##!	
 (g}7J7J7LfUGD!KrV   )r   r   c                   ddl m}  || ||      \  } }}| j                         \  }}|j                         \  }}	|j                         \  }
}t        j                  j
                  j                  ||      }t        j                  j
                  j                  d|z  |
      }|6ddlm}  ||j                         |r|n|j                         ||g|dg      }n	|J d       t               rt        j                  | ||f||      gng }||z  dk7  r6t        ||||      r(t        d      rt!        j"                  ||| ||gd	d	
       t%        d|| ||f|      \  }}	|S )Nr   )r)   r   )FixedLayoutr8   z,out_dtype is ignored if layout is specified.r   sparse_semi_structured_mmTr0  ) torch._inductor.select_algorithmr)   r   r   rq   rr   check_equals_and_simplifytorch._inductor.irrM  r   r[   r-   aten__sparse_semi_structured_mmr  r1   r+   r   r  r&   )rj   	mat1_metark   r   r   r)   m1k1m2r   k2r   r   r   rM  r  r+  s                    rT   tuned_sparse_semi_structured_mmrX    s}    @ +4DAD)T]]_FB EBMMOEB	222r:A	221r62>A~2OO"I(8FF	
  P"PP  !"	 ,00y$'9 1 	
   	
A
 Aq1 ;<66VdD)4tRV	
 (#WtY.EvGD! KrV   szc                 F    t        |       dk(  xs t        d | D              S )Nr   c              3   p   K   | ].  }t         j                  j                  j                  |d        0 yw)r8   Nrp   ).0ds     rT   	<genexpr>z)_is_tensorwise_scaling.<locals>.<genexpr>1  s,      !;<00A6!s   46)r  all)rY  s    rT   _is_tensorwise_scalingr`  0  s+    GqL S !@B!  rV   	transposec                 h    |rdnd}t         j                  j                  j                  | |   d      S )Nr   r   r8   rp   )rY  ra  idxs      rT   _is_rowwise_scalingrd  6  s,    !bC7733BsGQ??rV   	tensor_sz	tile_sizec                     |rdnd}|rdnd}t         j                  j                  j                  | |   ||         xr: t         j                  j                  j                  | |   t	        ||   |            S )Nr8   r   r   rq   rr   rs   r,   )rY  re  rf  ra  lhsrhss         rT   _is_blockwise1xTILESIZE_scalingrk  ;  sq     !aC!aC7733
33 
''


2
2
333rV   c                     t         j                  j                  j                  | d   t	        |d   d            xr: t         j                  j                  j                  | d   t	        |d   d            S )Nr      r8   rh  )rY  re  s     rT   _is_blockwise128x128_scalingrn  G  sd    7733
1wy|S) V
''


2
22a5')A,PS:T
UVrV   t
scale_sizescaling_typec                 X   |xt         j                  k(  r t        |      S xt         j                  k(  r t	        ||      S xt         j
                  k(  r t        || j                         d|      S t         j                  k(  rt        || j                               S 	 t        d|       )Nrm  Unsupported scaling type )r   
TensorWiser`  RowWiserd  BlockWise1x128rk  r   BlockWise128x128rn  AssertionError)ro  rp  rq  ra  s       rT   is_desired_scalingry  M  s     #[##)*55 [  &z9=='[''2AJJL#y  ))/
AJJLII #<\N!KLLrV   c                 t    | xt         j                  k(  r yt         j                  k(  ry	 t        d|  d      )Nrm  rs  z in get_tile_size)r   rw  rv  rx  )scale_options    rT   get_tile_sizer|  b  s<    
)[))'' +L>9JK rV   r   r   scale_a_sizescale_b_sizec                     t         D ](  \  }}t        | ||      st        |||d      s$||fc S  t        d| d|       )NT)ra  z1Inductor Triton does not support scale_a.shape = z, scale_b.shape = )scaling_pairsry  rx  )r   r   r}  r~  scale_option_ascale_option_bs         rT   get_scaling_optionsr  n  sc     +8 2&<
 nPTU!>11	2 
;L>I[\h[ij rV   scale_arecipe_a	swizzle_ascale_brecipe_b	swizzle_bc           	         t        | |||	      \  }}}}} }t        d   d| d| d| xx   dz  cc<   t        j                  d|||| j	                         |j	                         |       d}t        | |       t        |      dk\  rt        |      dk\  sJ d       t        |      dk(  xr t        |      dk(  }t        d	 |D              rt        d
 |D              s J dt        |       dt        |       d       dt        t           dt        fd} ||      xr  ||      }t        |d   |d         \  }}|s| |||g}nt        |      }| ||||g}t        |dd|	      }g }g }i }t               r3|j                  t               t!        |	|      |t        j"                  <   t%        |      \  }}|r|r|d   j&                  t(        j*                  k(  rv|rst-        |dd      rdt!        |      }t/        |d         t/        |d         } }t1        | ||d      r|s|j2                  |d<   | j2                  |d<   t5        || t6              r)|j                  t8               ||t8        j"                  <   nat5        || t:              rEt=        |      |d<   t=        |       |d<   |j                  t>               ||t>        j"                  <   ntA        d      tC        | ||d      r*|s(|j                  tD               ||tD        j"                  <   t5        || t6              r(|j                  tF               ||tF        j"                  <   |jI                  tJ        jL                  jO                  ||||             |r"tQ        ||||| |      rdd l)m*}!  |!||||!       |d   j&                  t(        j*                  k7  s|r|stW        ||||      \  }"}|"S |r@tY        ||||      r2t[        |      r't]        j^                  |||ja                         |"       |r3tc        ||||      r%te        jf                  |||ja                                tW        |||ja                         |      \  }"}|"S )#aN  
    Performs an optimized matrix multiplication where scaling factors are
    applied to the inputs per the supplied recipes, and optionally swizzled.

    This is the _scaled_mm_v2 API, which takes scale recipes (ScalingType) and
    swizzle patterns alongside the scale tensors, and supports multi-level
    scaling via lists.
    r   r   zaten._scaled_mm_v2.default_r   r8   zbTuned aten._scaled_mm_v2.default: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%s	scaled_mmz5scale_a and scale_b must each have at least one entryc              3   &   K   | ]	  }|d k(    ywr   Nr   r\  ss     rT   r^  z%tuned_scaled_mm_v2.<locals>.<genexpr>  s     )!qAv)   c              3   &   K   | ]	  }|d k(    ywr  r   r  s     rT   r^  z%tuned_scaled_mm_v2.<locals>.<genexpr>  s     1LQ!q&1Lr  zYInductor _scaled_mm_v2 lowering does not yet support non-trivial swizzles (got swizzle_a=z, swizzle_b=)reciperm   c                 ~    t        t        j                  t        j                  g      t	        fd| D              S )Nc              3   8   K   | ]  }t        |      v  y wrR   r   )r\  r
disalloweds     rT   r^  zEtuned_scaled_mm_v2.<locals>.check_supported_recipe.<locals>.<genexpr>  s     D;q>3Ds   )r   r   BlockWise1x16BlockWise1x32r_  )r  r  s    @rT   check_supported_recipez2tuned_scaled_mm_v2.<locals>.check_supported_recipe  s/    !:!:K<U<U VW
DVDDDrV   r   mat1_idxmat2_idxr   r   use_fast_accumTFenable_float8r   USE_FAST_ACCUMr   SCALE_RECIPE_ASCALE_RECIPE_BTILE_SIZE_ATILE_SIZE_BpInductor Triton does not support scaling options that are present in both epilogue scaling and main loop scalingr   r   $add_nv_universal_scaled_gemm_choicesr"  r  )4r;   r   r   r   r[   r   r  r_  r   r   r{   r)   r    r-   r  aten__fp8_mmr>  r  r9   r   r\   r   r6   r   r7   valuer5   epilogue_scaling_types.scaled_mm_device_tma_epilogue_scaling_templatemain_loop_scaling_typesr|  /scaled_mm_device_tma_main_loop_scaling_templaterx  r4   r  r  r  r   r  r  r3   r  r  r&   r1   r+   r   r  r  r.   r   r  )#r   r   r  r  r  r  r  r  biasr   contraction_dimr  r   r   r   r   rB   is_single_level_scaler  supported_recipescale_a_realscale_b_realr   	bias_realr"  r  r%  r   r   r!  
overridersr  r  r  r+  s#                                      rT   tuned_scaled_mm_v2r    s   2 %,uVy%!Aq!VUE ^:1#Qqc1#FG1LGHHl			 DUE*w<1W!2 ?2  LA-C#g,!2C
 )y))c1L)1L.L 	##'	?"3<Y?PPQ	SL
EtCy ET E .h7 <R=
 "0
GAJ!GL, e\<@"4(	e\<K #a!yM #%G CEO-,0-
(() 'v.MAz 	AJ-duU8

 $$ ' $E5SWX+9+?+?J'(+9+?+?J'(*0F !''(VW   N R RS -0G -:.,I
=),9.,I
=) ''(WX   O S ST %G  .uFt ##$RS JNNO 'N,B
 ##K0/9OKOO, NN			&&+	 	' 	
 4VQ1eUST,'		
 	
EMM) %+D';Oa 	 Aq1%66!)		
 *61a;**7FM<O<O<QR'g}7J7J7LfUGD!KrV   c	           	         t        | |||      \  }	}
}}} }t        d   d|	 d|
 d| xx   dz  cc<   t        j                  d|	|
|| j	                         |j	                         |       d}t        | |       t        ||      \  }}|s| |||g}nt        |      }| ||||g}t        |dd|	      }g }g }i }t               r3|j                  t               t        ||
      |t        j                  <   t        |      \  }}|j                  t        j                   k(  r|rt#        |dd      rqt        |      }|j$                  |j$                  }}t'        | |||      \  }}t)        | ||d      r|s|j*                  |d<   |j*                  |d<   t-        ||t.              r)|j                  t0               ||t0        j                  <   nat-        ||t2              rEt5        |      |d<   t5        |      |d<   |j                  t6               ||t6        j                  <   nt9        d      t;        | ||d      r*|s(|j                  t<               ||t<        j                  <   t-        ||t.              r(|j                  t>               ||t>        j                  <   |jA                  tB        jD                  jG                  ||||             |r"tI        ||	|
|| |      rddl%m&}  |||||       |j                  t        j                   k7  rtO        ||||      \  }}|S |r@tQ        ||	|
|      r2tS        |      r'tU        jV                  |||jY                         |       |r3t[        ||	|
|      r%t]        j^                  |||jY                                tO        |||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
    r   r   zaten._scaled_mm.default_r   r8   z_Tuned aten._scaled_mm.default: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sr  r   r  r  TFr  r  r   r  r  r  r  r  r   r   r  r  r  )0r;   r   r   r   r[   r   r)   r    r-   r  r  r>  r  r9   r   r\   r   r6   r   r  r7   r  r5   r  r  r  r|  r  rx  r4   r  r  r  r   r  r  r3   r  r  r&   r1   r+   r   r  r  r.   r   r  )r   r   r  r  r  scale_resultr   r  r   r   r   r   rB   r  r  r   r  r"  r  r%  r   r   r!  r  r}  r~  r  r  r  r+  s                                 rT   tuned_scaled_mmr  P  s   8 %,uVy%!Aq!VUE ^7s!A3asCDIDHHi			 DUE*!/!AL, e\<@"4(	e\<K #a!yM #%G CEO-,0-
(() 'v.MAz 	&duU8
%1%7%79K9Kl)<5,*
& $E5SWX+9+?+?J'(+9+?+?J'(*0F !''(VW   N R RS -0G -:.,I
=),9.,I
=) ''(WX   O S ST %G  .uFt ##$RS JNNO 'N,B
 ##K0/9OKOO, NN			&&+	 	' 	
 4VQ1eUST,'		
 }}%+D';Oa 	 Aq1%66!)		
 *61a;**7FM<O<O<QR'g}7J7J7LfUGD!KrV   indexc                 f    t         j                  j                  | xs d      }|j                  dk  S )Nr      )r\   r   get_device_propertiesmajor)r  propss     rT   _is_sm7x_or_older_gpur    s)    JJ,,UZa8E;;!rV   c                 &    t        d | D              S )Nc              3   <   K   | ]  }t        |t                y wrR   )
isinstancer   )r\  dims     rT   r^  zdims_are_int.<locals>.<genexpr>  s     4z#s#4s   )r_  )dimss    rT   dims_are_intr    s    4t444rV   r   c           	          t        | ||||      \  }}}t        |||g      sy t        | |      \  }}fd}d } ||||| |||      }t        ||||||	      }|
|j	                  |
|      S |j                         S )Nc                 V   t               }|j                  d|        |j                  d|       |j                  d|       |j                  d|j                  j                  d       |j                  d|j                  j                  d       t	        |d|       t	        |d	|       |j                  d
|j                  j                         d       |j                  d|j                  j                         d       dk(  r t        ||j                  j                         |S )Nr   r   r   
mat1_dtypeT)is_categorical
mat2_dtyperj   rk   mat1_iscontigmat2_iscontigr?   )r   add_featurer   r   r   is_contiguousr   )	r   r   r   rj   rk   mat1_stridemat2_stridecontextrB   s	           rT   get_contextz%mm_autoheuristic.<locals>.get_context#  s   +C#C#C#L$++*;*;DQL$++*;*;DQGV[9GV[9T[[668 	 	
 	T[[668 	 	
 4<"7DKK,=,=>rV   c                       y rR   r   r   rV   rT   fallbackz"mm_autoheuristic.<locals>.fallback6  s    rV   )r  r  r   r  rB   augment_contextprecondition)r   )get_size_hintsr  get_size_hints_stridesr   get_top_k_choices_callerget_choice_caller)rj   rk   r   r   r   r  rB   r   r   r  r   r   r  r  r  r  r  autoheuristics         `           rT   r  r    s     T4Aq1GAq!Aq	"5dDAK& !Q4{KHG0!M 55? 6 
 	
 **,,rV   c                 t   t        |t              rt        |t              s:t        j                  j                  j                  | j                               \  }}t        |t              rt        |t              s:t        j                  j                  j                  |j                               \  }}|||fS rR   )r  r   r   rq   rr   optimization_hintsr   )rj   rk   r   r   r   s        rT   r  r  M  s{    aZ3%7!!44T]]_EAaZ3%7!!44T]]_EAa7NrV   c                    | j                   j                  }|j                   j                  }||g}g }|D ]L  }t        |t              s)t        j
                  j                  j                  |      }|j                  |       N |d   |d   fS rx   )	r   rf   r  r   r   rq   rr   r  r  )rj   rk   r  r  stridesstrides_hintsrf   s          rT   r  r  V  s    ++$$K++$$KK(GM %&#&WW%%88@FV$% ]1---rV   )rm   NrR   )F)NNNFN)NN)r   loggingtypingr   r\   torch._dynamo.utilsr   +torch._inductor.autoheuristic.autoheuristicr   1torch._inductor.autoheuristic.autoheuristic_utilsr   r   r   r	   )torch._inductor.codegen.cpp_gemm_templater
   *torch._inductor.remote_gemm_autotune_cacher   torch._inductor.virtualizedr   r   "torch.fx.experimental.proxy_tensorr   torch.nn.functionalr   torch.torch_versionr   torch.utils._ordered_setr    r   r   r   codegen.cutlass.gemm_templater   r   ,codegen.rocm.ck_tile_universal_gemm_templater   'codegen.rocm.ck_universal_gemm_templater   codegen.subgraphr   r   irr   r   r   r   r"  r    loweringr!   r"   r#   r$   r%   select_algorithmr&   r'   r(   r)   r*   utilsr+   r,   r-   r.   r/   r0   r1   r2   r3   r4   r5   r6   r7   	mm_commonr9   r:   r;   r<   r=   r>   r   __version__triton_version
has_tritonImportError	getLoggerr   r   r   primsr  r  r  r	  r
  r  r  r  cacherU   r?   rb   r   	dtype_outr   rh   rA  _int_mmr4  _sparse_semi_structured_mmdefaultrR  
_scaled_mmr  r`   rl   r   rD  r   r   r  r   r   r   r  rE  r,  r6  rK  rX  rt  ru  rv  rw  r  r  r  r{   r`  rd  r   rk  rn  Tensorry  r|  tupler  _scaled_mm_v2r   r  r  r  r  r  r  r  r   rV   rT   <module>r     s       ( T  F F . 6 + , / > X M D E 8 8 *      !&"4"45NJ
 g!yy~~		
 		!n&?  ,
 
.	/'+" ,		 :;  (		 67  2@	0	 ;<2 . 3A	1	 <=3 / 2@	-	 NO2 . " " UXX|
M"	HH	!!	  	KKdjjnn
 "	MM$$,,2B2B #5	$$$//77	#  "	*8K8K
8 (,11 I4 %Z6#  0  F !: ; ) D'/ #5_m#  &8*,<& "
 4775s4 s 6sl 4<<T:'+ . ;.b 4::48*+!D K 9K\ 422M(,T. N.d [334+--.!=!=>!;!;<!!;#=#=> &00+2E2EF &55{7S7ST s t @C @D @T @
			(+	8<			VS VS VT V 	M
MM M 	M
 
M*	3 	 ,, ,,	
 ;#$" 4%%--4H 
M #YM 3i	M
 CyM #YM 3iM CyM IM` 4??**E 
r Frj t   
5 :- ::-z	.Y(  !'*NJs   U U+*U+