Ë
    ÄÓ^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„«       Z„e�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Z‰e�j
                  e�j                  gZŠe�j                  e�j                  gZ‹dPed;eŒfdQ„Z�dPedReŒd;eŒfdS„ZŽdPedTedUe�dReŒd;eŒf
dV„Z�dPedTed;eŒfdW„Z‘	 dsdXedYe�j$                  dZedReŒd;eŒf
d[„Z“d;e�fd\„Z”d]ed^ed_e�j$                  d`e�j$                  d;e•eef   f
da„Z– e6e[�j.                  jâ                  d¬H«      	 	 	 	 	 dtdbe˜e   dce˜e�   dde˜e�   dee˜e   dfe˜e�   dge˜e�   fdh„«       Z™ e6e[jæ                  jâ                  d¬H«      	 	 	 	 	 dtdi„«       Zše jÊ                  dje�dz  d;eŒfdk„«       Z›dl„ Zœ	 	 dudme�dz  fdn„Z�do„ Zždp„ ZŸy# 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    € Ø�=‰=‹?œuŸz™z¬5¯;©;Ð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‹˜Ò˜sŸx™x¨›{¨aÒ/°C·H±H¸Q³KÀ1Ò4DÜ�{‰{˜3˜q™6 4¨°3¸eÈ$ÔOÐOÜ�;‰;�s˜D $¨C°uÀ4Ô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·   ó#   € Ü�w‰w×Ñ×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   € ÜÜ�G‰G×Ñ×4Ñ4°T¸!±W¸aÓ@ò DÜ�w‰w×Ñ×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>Æ   ó   ø€ Ð6°u×7GÑ7GÓ7IÐ6JÐ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ó 
ô 
‡L�LÙ�U×%Ñ%Ó'Ó(ÒJ©L¸¿¹Ó9IÓ,JÛKôô 
‡L�LÙ�U×%Ñ%Ó'Ó(ÒJ©L¸¿¹Ó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à�8‰m€GØ€AÜ—‘˜qŸy™y¨¨A¨wÓ7¸ÓC€JØ—‘˜1˜g qÓ)€JÜ�Y‰Y�z :¼¿¹ÔG€FÜ—)‘)˜F AÓ&€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Ó2ˆNÜÜ×!Ñ!¤*°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Ó2ˆNÜØ—‘ØóˆBô
 ‘7Ñ#Ø—Y‘YØ'ØØ Ø ×,Ñ,ð $ó ÷	÷ 	ò 	ú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   € Ü�8‰8�A�q—|‘|“~Ó&Ð&rV   c                 óL   — t        j                  | ||j                  «       «      S rR   )r\   rh   rÉ   )ri   r“   r”   s      rT   Úcontiguous_addmmrÌ   -  s   € Ü�;‰;�s˜A˜qŸ|™|›~Ó.Ð.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        ‰ |«      rèt        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>                  <   |€ø|rötA        |d¬«      rétC        |||«      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‡                  |«      rÚt‰        ‰ «      rÏg }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   ø€ Ü—|‘| A t§z¡zÀUÔ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�  ø€ ð
 ÐØ—n‘nÓ&ˆÜ�‰Ø�N‰NÓ Ñ+Ù3ô	
ô 	�‰Ø�O‰OÓ×"Ñ" oÐ5ÙAô	
ô 	�‰Ø˜Ñ$ò àœUŸ]™]Ñ*ò CØ¤E§M¡M´5·>±>Ð#BÐBá\ô	
ô. ˜˜tÔ$ÜœŸ™Ñ(¨¨rÓ2ˆÜœŸ™Ñ(¨¨qÓ1ˆÜ%Ø˜�ØØØ $Ø"'ô
‰ˆˆfô ×!Ñ!×8Ò8¸T¿Z¹ZÜ�M‰MÜ�N‰NðL
ñ >
ô
Lð ;?Ö?°QÐ-”N 9Ó-¨aÕ0Ð?ˆDÐ?Ø/œ¤s§w¡wÓ/°Ð6ˆØ-œ uÓ-¨m¸QÓ?ˆàÐô #*Øˆd˜6¨Yô#Ñ€A€qˆ!ˆV�T˜4ô  2°&Ó9Ñ€L�*Ø€Dô # D¨$ <¸9ÔE€Mô ˆ^Ñ˜x¨ s¨!¨A¨3¨a°¨sÐ3Ó4¸Ñ9Ó4Ü‡H�HØYØ	Ø	Ø	Ø�‰ÓØ�‰ÓØôð #%€GÜ1°&Ó9Ñ€L�*ä'.€LØ(*ÐØÐÜ$ˆØ(¨)Ð4ÐàBDÐØ13€OÜÔØ×Ñ Ô-ÙØ0AˆO˜L×,Ñ,Ñ-ð 	ÐÙÜ ¸4Õ@ä! ! Q¨Ô*Ø×#Ñ#Ô$AÔBô (×FÑFÈ,ÑVˆÙÔ 6°q¸!¸QÐSTÕ UØ×#Ñ#¤KÔ0ä0Ø�d¨&¸Tõð !×'Ñ'Ô(VÕWÜ(Ø�d¨&¸Tõô —=‘=×$Ñ$Ð,Ø$×+Ñ+Ô,FÕGà$×+Ñ+Ô,BÔCà×ÑÔ ?Ô@à‡N�NÜ	�	‰	×&Ñ&ØØØØ+ð	 	'ó 	
ôð 	ÐÙÜ  ¨¨A¨qÔ1Ü Ô%ä×6Ñ6Ø�V˜]×0Ñ0Ó2ô	
ð Ð™ZÔ,@ÀÈÈAÈqÔ,QÜ×*Ñ*¨7°F¸M×<OÑ<OÓ<QÔRØÐ™ZÔ,EÀfÈaÐQRÐTUÔ,VÜ×&Ñ& w°¸×8KÑ8KÓ8MÔNð 	ÐÙÜ*¨6°1°a¸¸DÀ$ÔGåMá% g¨v°}ÔEàÐÔ2°6¸4ÀÔFÜ×#Ñ#ØØØ×ÑÓ!ô	
ð ˜�,€KàÑÚÜ Õ'Ü�O‰O×"Ñ"×4Ñ4°TÔ:Ü�dŒOàˆÜ Ô"Ø×"Ñ" ;Ô/Ü+.¨w«<Ð(Ø�‰Ü�I‰I×*Ñ*ð Ü�Øóô		
ô &ØØØØØØØØÜ‹OØØØ+ô
ˆ
ô �‰×%Ñ%×;Ñ;¸DÔAàÐ%¬#¨j«/¸AÒ*=ð
 18ÖP f¸6ÀZÒ;Oš6ÐP�ÑPà!Ð"CÐ#CÐD�àÐÜ ×0Ñ0ò 	ˆAØ�N‰NÜ+¨AÓ.×3Ñ3°M×4GÑ4GÓ4IÈ6ÓRõð	ð
 ÐØÐœUŸ_™_×3Ñ3×NÒNô -¨T°4Ó8Ðä"×8Ñ8Øˆg�}×*Ñ*Ó,¨fóð €sð ð ˆ
ä'ØØØ×ÑÓØØ-ô�G€Dˆ!ð €Kùòw @ùò~ 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ˆ  € ô #*Øˆd˜6¬U¯[©[ô#Ñ€A€qˆ!ˆV�T˜4ð €Däˆ^Ñ˜}¨Q¨C¨q°°°1°Q°CÐ8Ó9¸QÑ>Ó9Ü‡H�HØ^Ø	Ø	Ø	Ø�‰ÓØ�‰ÓØôô  2°&Ó9Ñ€L�*ØÒW :ÒWÔ2FÀvÈqÐRSÐUVÓ2W€KØ"$€Gô # D¨$ <¼5¿;¹;ÔG€Mð CEÐÜÔØ×Ñ¤Ô-áÔ)Ø˜T°eõð 	×Ñ¤Ô,ð ‡N�NÜ	�	‰	×&Ñ& }Ð6FÈÓMôñ Ô*¨4Ô0Ü×6Ñ6Ø�V˜]×0Ñ0Ó2¸TÐPTõ	
ô (¨¨g°}×7JÑ7JÓ7LÈfÓU�G€Dˆ!Ø€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        «       r²t,        g}| j5                  «       d   dk(  ret7        | j9                  «       «      dk(  rIt        j:                  j<                  r/t&        j>                  j@                  s|jC                  tD        «       |j%                  t&        j(                  j+                  |||«      «       |rÐtG        |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Ô$Ø�1Š9Ø‰DäœTŸX™XÑ& t¨SÓ1ˆDà�AŠ:Ø‰DäœTŸX™XÑ& u¬i¼¿¹Ñ.@ÀÀtÓ.LÓMˆDäœŸ™Ñ" 4¨Ó.Ð.ô 18¸¸dÀCÐPVÔ0WÑ-€A€qˆ!ˆV�T˜4 Ü
˜Ó
€CÜ1°&Ó9Ñ€L�*Ø€Dô #Ø	�t˜TÐ"¬D°uÀ4Ô,Hô€Mô (Ø	ˆd�DÐ¤4¨e¸$Ô#?ôÐð #%€Gô ˆ^Ñ˜{¨1¨#¨Q¨q¨c°°1°#Ð6Ó7¸1Ñ<Ó7Ü‡H�HØ\Ø	Ø	Ø	Ø�‰ÓØ�‰ÓØôñ Ü×)Ò)¬_×-NÒ-Nà�‰Ü�I‰I×*Ñ*Ø"Ü�Øóô	
ô ,Ø�'˜=×.Ñ.Ó0°&ó
‰ˆˆað ˆàBDÐäÔÜEOÀLˆà�N‰NÓ˜QÑ 1Ò$Ü�C—L‘L“NÓ# qÒ(Ü×&Ñ&×8Ò8Ü—G‘G×'Ò'à×!Ñ!¤/Ô2ð 	�‰Ü�I‰I×*Ñ*Ð+=¸~ÈtÓTô	
ñ Ô)¨&ÀUÕKØ×Ñ¤Ô,ä,Ø�$ f¸õ
ð ×#Ñ#Ô$RÕSÜ$ T¨4¸vÐRVÕWÜ�}‰}× Ñ Ð(Ø ×'Ñ'Ô(BÕCà ×'Ñ'Ô(>Ô?ð 	�‰Ü�I‰I×*Ñ*Ø"Ô%GÐ$HÈ$óô	
ð ‡N�NÜ	�	‰	×&Ñ& }Ð6FÈÓMôñ
 	Ü  ¨¨A¨qÔ1Ü Ô%ä×6Ñ6ØØð ×Ñª	ÐÓ2ØØÚ#õ		
ñ Ô*¨6°1°a¸Ô;Ü×*Ñ*ØØð ×Ñª	ÐÓ2ØØÚ#õ		
ô ˜V T¨4Ô0Ü×#Ñ#ØØØ×ÑÓ!ØØØõ	
ô (¨¨g°}×7JÑ7JÓ7LÈfÓU�G€Dˆ!Ø€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Ø�]‰]‹_�F€BˆØ×ÑÓ �E€BˆØ�M‰M‹O�E€BˆÜ	�‰×Ñ×2Ñ2°2°rÓ:€AÜ	�‰×Ñ×2Ñ2°1°r±6¸2Ó>€AØ€~Ý2áØ�O‰OÓÙ"‰I¨¯©Ó(8Ø�ˆFØ�ˆFó	
‰ð Ð ÐPÐ"PÓPÐ ô !Ô"ô	 ,×0Ñ0Ø�y $Ð'¨¸9ð 1ó ñ	
ð ð ð 	
ˆA‰�Š
Ü  ¨¨A¨qÔ1ÜÐ ;Ô<ä×6Ñ6Ø�V˜d D¨)Ð4¸tÐRVõ	
ô (Ø# W¨t°YÀÐ.EÀvó�G€Dˆ!ð €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+   € Ü�‹G�q‰Lò œ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,   € Ù‰!˜b€CÜ�7‰7×Ñ×3Ñ3°B°s±G¸QÓ?Ð?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   € ñ ‰!˜a€CÙ‰!˜a€CÜ�7‰7×Ñ×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   € Ü�7‰7×Ñ×3Ñ3Ø
ˆ1‰Œw�y ‘| SÓ)óò Vä
�'‰'×
Ñ
×
2Ñ
2°2°a±5¼'À)È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Ø Œ[× Õ Ü& z°9Ó=Ð=Ø'Œ[×'Õ'Ü2Ø˜AŸJ™J›L¨#¨yóð ô ×)Ô)Ü/°
¸A¿J¹J»LÓ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Ñ&ˆ˜ÜØ�< õ
ä  ¨°nÐPTÖ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   è ø€ Ò)˜!ˆq�A�vÑ)ùó   ‚c              3   ó&   K  — | ]	  }|d k(  –— Œ y­wr�  rÑ   rŽ  s     rT   r^  z%tuned_scaled_mm_v2.<locals>.<genexpr>²  s   è ø€ Ò1L¸Q°!°qµ&Ñ1Lùr�  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ÑDùs   ƒ)r   r   ÚBlockWise1x16ÚBlockWise1x32r_  )r“  r—  s    @rT   Úcheck_supported_recipez2tuned_scaled_mm_v2.<locals>.check_supported_recipe·  s/   ø€ Ü¤×!:Ñ!:¼K×<UÑ<UÐ VÓWˆ
ÜÓD¸VÔ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_BúpInductor 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 %,Øˆu˜V¨yô%Ñ!€A€qˆ!ˆV�U˜Eô ˆ^ÑÐ:¸1¸#¸Q¸q¸cÀÀ1À#ÐFÓGÈ1ÑLÓGÜ‡H�HØlØ	Ø	Ø	Ø�‰ÓØ�‰ÓØôð €DÜ˜U EÔ*äˆw‹<˜1Ò¤ W£°Ò!2ð Ø?óÐ2ô   ›L¨AÑ-ÒC´#°g³,À!Ñ2CÐô
 Ñ)˜yÔ)Ô)¬cÑ1LÀ)Ô1LÔ.Lð ð	#Ü#'¨	£?Ð"3°<ÄÀYÃÐ?PÐPQð	SóÐLð
E¤t¬C¡yð E´Tó Eñ .¨hÓ7ò Ñ<RØó=Ðô
 "0°¸±
¸GÀA¹JÓ!GÑ€L�,ñ Ø˜e \°<Ð@‰ä" 4Ó(ˆ	Ø˜e \°<ÀÐKˆô #Ø˜a¨!°yô€Mð #%€Gð CEÐØ€OäÔØ×Ñ¤Ô-Ü,0Ø°ô-
ˆœ×(Ñ(Ñ)ô ' vÓ.�M€A€zò 	ÚØ�A‰J×Ñ¤§¡Ó-ÚÜ °dÈuÖUä¨Ô8ˆ
ô
 ˜ ™Ó$Ü˜ ™Ó$ð 'ˆô $ E¨5ÀÐSWÕXÙà+9×+?Ñ+?ˆJÐ'Ñ(Ø+9×+?Ñ+?ˆJÐ'Ñ(ä*Ø Ô0Fôð !×'Ñ'Ô(VÔWàð  Ô N× RÑ RÒSô -Ø Ô0Gôô -:¸.Ó,I�
˜=Ñ)Ü,9¸.Ó,I�
˜=Ñ)à ×'Ñ'Ô(WÔXàð  Ô O× SÑ SÒTô %ðGóð ô .Ø�u¨F¸tõñ à×#Ñ#Ô$RÔSàð ÔJ×NÑNÑOô 'Ø˜NÔ,Bô
ð ×#Ñ#¤KÔ0Ø/9ˆOœKŸO™OÑ,ð ‡N�NÜ	�	‰	×&Ñ&ØØØØ+ð	 	'ó 	
ôñ Ô4°V¸QÀÀ1ÀeÈUÔSÝTá,ØØØØ'õ		
ð 	�‰
×ÑœEŸM™MÒ)Ù Ù%ä+¨D°'¸;ÈÓO‰ˆˆaØˆñ 	Ü  ¨¨A¨qÔ1Ü Ô%ä×6Ñ6ØØØ×ÑÓ!Ø)õ		
ñ Ô*¨6°1°a¸Ô;Ü×*Ñ*¨7°F¸M×<OÑ<OÓ<QÔRä'¨¨g°}×7JÑ7JÓ7LÈfÓU�G€Dˆ!Ø€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ƒ|�r€t#        |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 %,Øˆu˜V¨yô%Ñ!€A€qˆ!ˆV�U˜Eô ˆ^ÑÐ7¸°s¸!¸A¸3¸aÀ¸sÐCÓDÈÑIÓDÜ‡H�HØiØ	Ø	Ø	Ø�‰ÓØ�‰ÓØôð €DÜ˜U EÔ*ä!/°¸Ó!AÑ€L�,ñ Ø˜e \°<Ð@‰ä" 4Ó(ˆ	Ø˜e \°<ÀÐKˆô #Ø˜a¨!°yô€Mð #%€Gð CEÐØ€OäÔØ×Ñ¤Ô-Ü,0Ø°ô-
ˆœ×(Ñ(Ñ)ô ' vÓ.�M€A€zð 	�‰œŸ™Ó&ÚÜ °dÈuÖUä¨Ô8ˆ
à%1×%7Ñ%7¸×9KÑ9K�lˆä)<Ø�5˜,¨ó*
Ñ&ˆ˜ô $ E¨5ÀÐSWÕXÙà+9×+?Ñ+?ˆJÐ'Ñ(Ø+9×+?Ñ+?ˆJÐ'Ñ(ä*Ø Ô0Fôð !×'Ñ'Ô(VÔWàð  Ô N× RÑ RÒSô -Ø Ô0Gôô -:¸.Ó,I�
˜=Ñ)Ü,9¸.Ó,I�
˜=Ñ)à ×'Ñ'Ô(WÔXàð  Ô O× SÑ SÒTô %ðGóð ô .Ø�u¨F¸tõñ à×#Ñ#Ô$RÔSàð ÔJ×NÑNÑOô 'Ø˜NÔ,Bô
ð ×#Ñ#¤KÔ0Ø/9ˆOœKŸO™OÑ,ð ‡N�NÜ	�	‰	×&Ñ&ØØØØ+ð	 	'ó 	
ôñ Ô4°V¸QÀÀ1ÀeÈUÔSÝTá,ØØØØ'õ		
ð ‡}�}œŸ™Ò%Ü+¨D°'¸;ÈÓO‰ˆˆaØˆñ 	Ü  ¨¨A¨qÔ1Ü Ô%ä×6Ñ6ØØØ×ÑÓ!Ø)õ		
ñ Ô*¨6°1°a¸Ô;Ü×*Ñ*¨7°F¸M×<OÑ<OÓ<QÔRä'¨¨g°}×7JÑ7JÓ7LÈfÓU�G€Dˆ!Ø€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)   € ä�J‰J×,Ñ,¨UªZ°aÓ8€EØ�;‰;˜!ÑÐ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×#Ñ4ùs   ‚)r_  )Údimss    rT   Údims_are_intrË    s   € ÜÑ4¨tÔ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Ü˜G V¨[Ô9Ü˜G V¨[Ô9Ø×ÑØ˜TŸ[™[×6Ñ6Ó8Èð 	ô 	
ð 	×ÑØ˜TŸ[™[×6Ñ6Ó8Èð 	ô 	
ð �4Š<Ü" 7¨D¯K©K×,=Ñ,=Ô>Øˆ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¶   ø€ ô ˜T 4¨¨A¨qÓ1�G€A€qˆ!Ü˜˜A˜q˜	Ô"ØÜ5°d¸DÓAÑ€K�ôò&ñ ˜!˜Q  4¨¨{¸KÓH€GÜ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Ñ4°T·]±]³_ÓE‰ˆˆAä�aœÔ¤Z°´3Ô%7Ü—‘×!Ñ!×4Ñ4°T·]±]³_ÓE‰ˆˆAØˆa�ˆ7€NrV   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Øò %ˆÜ˜&¤#Ô&Ü—W‘W×%Ñ%×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Ó5€NØ€Jð
 €g×Ñ˜Ó!€Ø‡y�y‡~�~€Ø�	‰	�‰€ñ
 Ø	Ø	à�‰×ÑÐ! n¸Ò&?ñ   Ô,ñ
 
Ð.Ó	/Ø'+Ø"ô€ñ ,Ø	Ø	ÙÐ :Ó;ôÐ ñ (Ø	Ø	ÙÐ 6Ó7ôÐ ñ 2@Ø	0Ø	ÙÐ ;Ó<ô2Ð .ñ 3AØ	1Ø	ÙÐ <Ó=ô3Ð /ñ 2@Ø	-Ø	ÙÐ NÓOô2Ð .ð ‡�ñ"ó ð"ñ ˜UŸX™X |ÀÇÁÇÁÔ
M€Ù"Ø	‡H�HØØ	Ø—‘×!Ñ!ô	€ñ  Ø	‡K�K�¨d¯j©j¯n©nô€
ñ "Ø	‡M�MÐ$°$·,±,×2BÑ2Bô€ñ #5Ø	×$Ñ$Ø$ØØ×/Ñ/×7Ñ7ô	#Ð ñ "Ø	×ÑÐ*¸¿¹×8KÑ8Kô€ò
8ð (,°1¸1ô Ióñ4 % Z°Ó6€ò#ô Ð 0ô  ñF !:Ó ;Ð ôÐ)ô òD'ò/ñ #5Ø�_ mó#Ð ñ &8ØÐ*Ð,<ó&Ð "ñ
 �4—7‘7°Ô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°+×2EÒ2EÐFÐ Ø&×5Ò5°{×7SÒ7SÐ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ð �3‰ið	Mð
 �C‰yðMð �#‰YðMð �3‰iðMð �C‰yòMó IðMñ` �4—?‘?×*Ñ*ÀÔEð 
ØØØØòró Fðrðj ‡�ð  t¡ð °ò ó ðò
5ð Øñ:-ð �‰:ó:-òzó	.øðY( ò Ù! 'Ó*€NØƒJðús   ÃU ÕU+Õ*U+