+
    –°³j>Á  ã                   óp
  € ^ RI t ^ RIt^ RIHt ^ RIt^ RIHt ^ RIHt ^ RI	H
t
HtHtHt ^ RIHt ^ RIHt ^ RIHtHt ^ R	IHt ^ R
IHt ^ RIHt ^ RIHt ^RIHtH t H!t" ^RI#H$t$H%t% ^RI&H't' ^RI(H)t) ^RI*H+t+H,t, ^RI-H.t.H/t/H0t0H1t1 ^RI2H3t3 ^RI!H4t4H5t5H6t6H7t7H8t8H9t9 ^RI:H;t;H<t<H=t=H>t>H?t? ^RI@HAtAHBtBHCtCHDtDHEtEHFtFHGtGHHtHHItIHJtJHKtKHLtLHMtM ^RINHOtOHPtPHQtQHRtRHStSHTtTHUtU  ^ RIVtV]! ]VP®                  4      tXRtY]P¶                  ! ]\4      t]]P&                  P¼                  t^]P&                  P¾                  t_]?! R]S]PÀ                  PÂ                  e   ]XR8¼  d	   ]Q! R4      M]Q! R4      RRR7      tb]?! R ]T]Q! R!4      R"7      tc]?! R#]T]Q! R$4      R"7      td]?! R%]T]Q! R&4      R"7      te]?! R']T]Q! R(4      R"7      tf]?! R)]T]Q! R*4      R"7      tg] PÐ                  R+ 4       ti]<! ]PÔ                  R,]^PÔ                  PÖ                  R-7      tl]<! ]PÔ                  R.R/]^PÔ                  PÚ                  R07      tn]<! ]PÞ                  R1]^PÞ                  PÖ                  R-7      tp]<! ]Pâ                  R2]^Pâ                  PÖ                  R-7      tr]<! ]Pæ                  R3R]^Pæ                  Pè                  R47      tu]<! ]Pì                  R5]^Pì                  PÖ                  R-7      twR6 txR7RR8^R9^/R: ltyR; R< ltz]<! ]yR4      t{R= R> lt|R? t} ! R@ RA],4      t~]~! 4       t ! RB RC],4      t€RD t�RE t‚]€! RFRG]�4      tƒ]€! RHRI]‚4      t„]8! ]^PÔ                  RRJ7      RjRKR/RL ll4       t…]8! ]^Pâ                  RRJ7      RKR/RM l4       t†]8! ]^PÞ                  RRJ7      R8^R9^RKR/RN l4       t‡]8! ]^Pæ                  RRJ7      RORRKR/RP l4       tˆ]EP                  ]EP                  3]EP                  ]EP                  3]EP                  ]EP                  3]EP                  ]EP                  3]EP                  ]EP                  3.t�]EP                  ]EP                  .tŽ]EP                  ]EP                  .t�RQ RR lt�RS RT lt‘RU RV lt’RW RX lt“RkRY RZ llt”R[ R\ lt•R] R^ lt–]4! ]^EP.                  Pè                  RR_7      t˜]8! ]^EP.                  Pè                  RRJ7      RlR` Ra ll4       t™]8! ]^Pì                  Pè                  RRJ7      RlRb l4       tš] PÐ                  Rc Rd l4       t›Re tœRmRf Rg llt�Rh tžRi tŸR#   ]Z d    ]! R4      tXRtY EL>i ; i)né    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Úlowering)ÚCUTLASS2xGemmTemplateÚCUTLASS3xGemmTemplate)ÚCKTileGemmTemplate)ÚCKGemmTemplate)ÚSubgraphChoiceCallerÚSubgraphTemplate)ÚBufferÚChoiceCallerÚ	is_tritonÚLayout)ÚMMKernelInputs)Úfallback_handlerÚ	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Ú_use_small_mm_pointwiseÚ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)rC   rD   rE   Ú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                 ó   € \        V 4      # ©N)r(   )Úfns   &ÚZ/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/_inductor/kernel/mm.pyÚlazy_register_extern_choicerV   ˆ   s   € ä˜bÓ!Ð!ó    z
at::mm_out)Úop_overloadzat::mm_dtype_outÚmm_dtype)rC   rX   zat::addmm_outzat::_int_mm_outzat::_sparse_semi_structured_mm)Úhas_out_variantrX   zat::_scaled_mm_outc                 ód   € V P                  4       \        P                  \        P                  39   # rS   )Ú	get_dtypeÚtorchÚint8Úuint8)Úmats   &rU   Ú_is_int8_matra   ©   s    € Ø�=‰=‹?œuŸz™z¬5¯;©;Ð7Ñ7Ð7rW   ÚoutÚalphaÚbetac          	     óú   € V P                  ^ 4      ^ 8X  d   V P                  ^ 4      ^ 8w  g   V P                  ^ 4      ^8X  d"   \        P                  ! V ^ ,          WW4VR7      # \        P                  ! WW#WER7      # )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.
)rb   rc   rd   )ÚstrideÚsizer]   Úaddmm)ÚinpÚmat1Úmat2rb   rc   rd   s   &&&$$$rU   Ú
bias_addmmrl   ­   s^   € ð 	�
‰
�1‹˜Ô˜sŸx™x¨›{¨aÔ/°C·H±H¸Q³KÀ1Ô4DÜ�{Š{˜3˜q�6 4°3È$ÔOÐOÜ�;Š;�s $°uÔHÐHrW   c                ó   € V ^8„  d   QhRR/# ©é   ÚreturnN© )Úformats   "rU   Ú__annotate__rs   ¸   s   € ÷ ñ ¨dñ rW   c                 ób  a a€ R  R lpR R lpR R lp\         P                  ! V! S P                  4       4      ;'       g    V! S P                  4       4      V 3R l4       \         P                  ! V! SP                  4       4      ;'       g    V! SP                  4       4      V3R l4       R# )	c                ó$   € V ^8„  d   QhR\         /# ©ro   rp   ©Úbool)rr   s   "rU   rs   Ú.check_supported_striding.<locals>.__annotate__¹   ó   € ÷ Fñ F¤ñ FrW   c                 ód   € \         P                  P                  P                  V ^,          ^4      # ©é   ©r   ÚgraphÚsizevarsÚstatically_known_equals©rf   s   &rU   Úis_row_majorÚ.check_supported_striding.<locals>.is_row_major¹   ó#   € Ü�w‰w×Ñ×7Ñ7¸¸q½	À1ÓEÐErW   c                ó$   € V ^8„  d   QhR\         /# rv   rw   )rr   s   "rU   rs   ry   ¼   rz   rW   c                 ód   € \         P                  P                  P                  V ^ ,          ^4      # ©r   r~   r‚   s   &rU   Úis_col_majorÚ.check_supported_striding.<locals>.is_col_major¼   r…   rW   c                ó$   € V ^8„  d   QhR\         /# rv   rw   )rr   s   "rU   rs   ry   ¿   s   € ÷ 
ñ 
œdñ 
rW   c                 óè   € \        \        P                  P                  P	                  V ^ ,          ^ 4      ;'       g2    \        P                  P                  P	                  V ^,          ^ 4      4      # rˆ   )rx   r   r   r€   r�   )rg   s   &rU   Úhas_zero_dimÚ.check_supported_striding.<locals>.has_zero_dim¿   sW   € ÜÜ�G‰G×Ñ×4Ñ4°T¸!µW¸aÓ@÷ Dð DÜ�w‰w×Ñ×7Ñ7¸¸Q½ÀÓCó
ð 	
rW   c                  ó*   <€ R S P                  4        2# )z$mat_a must be row_major, got stride ©Ú
get_stride)Úmat_as   €rU   Ú<lambda>Ú*check_supported_striding.<locals>.<lambda>È   ó   ø€ Ð6°u×7GÑ7GÓ7IÐ6JÑKrW   c                  ó*   <€ R S P                  4        2# )z$mat_b must be col_major, got stride r�   )Úmat_bs   €rU   r“   r”   Î   r•   rW   N)r]   Ú_checkr‘   Úget_size)r’   r—   rƒ   r‰   r�   s   ff   rU   Úcheck_supported_stridingrš   ¸   s€   ù€ õFõFõ
ô 
‡L‚LÙ�U×%Ñ%Ó'Ó(×JÐJ©L¸¿¹Ó9IÓ,JÜKôô 
‡L‚LÙ�U×%Ñ%Ó'Ó(×JÐJ©L¸¿¹Ó9IÓ,JÜKörW   c                ó   € V ^8„  d   QhRR/# rn   rq   )rr   s   "rU   rs   rs   Õ   s   € ÷ ñ °Dñ rW   c                 ó†  € \         P                  ! V P                  4       VP                  4       8H  ;'       d"    V P                  4       VP                  4       8H  R  4       \         P                  ! V P                  4       VP                  4       8H  ;'       d"    V P                  4       VP                  4       8H  R 4       R# )c                  ó   € R # ©zinput dtypes must be the samerq   rq   rW   rU   r“   Ú-_check_addmm_input_metadata.<locals>.<lambda>Ø   s   € Ñ/rW   c                  ó   € R # )z%all inputs must be on the same devicerq   rq   rW   rU   r“   rŸ   Ü   s   € Ñ7rW   N)r]   r˜   r\   Ú
get_device)ri   rj   rk   s   &&&rU   Ú_check_addmm_input_metadatar¢   Õ   s‚   € Ü	‡L‚LØ�‰‹˜4Ÿ>™>Ó+Ñ+×SÐS°·±³À4Ç>Á>ÓCSÑ0SÙ/ôô 
‡L‚LØ�‰Ó˜DŸO™OÓ-Ñ-×WÐW°#·.±.Ó2BÀdÇoÁoÓFWÑ2WÙ7örW   c                 óª  € V P                   ^ ,          pVP                   ^,          pV P                   ^,          pWR,          pTp\        P                  ! V P                  W7V4      R4      pVP                  WvV4      p	\        P                  ! W‰\        P
                  R7      p
\        P                  ! V
^ 4      pVP                  V P                  4      # )r   ©Ú	out_dtype)r}   r   ro   )	Ú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   &&&         rU   Ú
decomposeKrº   à   s—   € Ø	�‰��
€AØ	�‰��
€AØ	�‰��
€Aà�m€GØ€AÜ—’˜qŸy™y¨¨wÓ7¸ÓC€JØ—‘˜1 qÓ)€JÜ�YŠY�z¼¿¹ÔG€FÜ—)’)˜F AÓ&€KØ�>‰>˜!Ÿ'™'Ó"Ð"rW   c                   óJ   a a€ ] tR t^ît oV 3R ltV3R lV 3R lltRtVtV ;t# )ÚDecomposeKSugraphTemplatec                ó(   <€ \         SV `  R R7       R# )Údecompose_k©rC   N)ÚsuperÚ__init__)ÚselfÚ	__class__s   &€rU   rÁ   Ú"DecomposeKSugraphTemplate.__init__ï   s   ø€ Ü‰ÑØð 	ö 	
rW   c                óB   <€ V ^8„  d   QhRS[ S[,          RS[RS[RS[/# )ro   Úinput_nodesÚlayoutÚk_splitrp   )Úlistr   r   Úintr   )rr   Ú__classdict__s   "€rU   rs   Ú&DecomposeKSugraphTemplate.__annotate__ô   s7   ø€ ÷ ñ á™&•\ðñ ðñ ð	ñ
 
ñrW   c           
     ó  <€ ^ RI Hp ^RIHp RV R2pRV: 2pV! 4       ;_uu_ 4        V! 4       p\	        \
        P                  ! \        VR7      V4      p	\        S
V `%  VVVV	VR7      uuRRR4       #   + '       g   i     R# ; i)	r   ©Úenable_python_dispatcher©Úselect_decomp_tableÚdecompose_k_mm_Ú_splitzk_split=)r°   ©rC   rÆ   rÇ   Úmake_fx_graphÚdescriptionN)
Útorch._dispatch.pythonrÏ   ÚdecompositionrÑ   r   Ú	functoolsÚpartialrº   rÀ   Úgenerate)rÂ   rÆ   rÇ   rÈ   rÏ   rÑ   rC   rÖ   ÚdecompositionsrT   rÃ   s   &&&&      €rU   rÛ   Ú"DecomposeKSugraphTemplate.generateô   s…   ø€ õ 	Då7à   	¨Ð0ˆØ!˜™
�mˆá%×'Õ'Ù0Ó2ˆNÜÜ×!Ò!¤*°wÔ?ØóˆBô
 ‘7Ñ#ØØ'ØØ Ø'ð $ó ÷ (×'×'Ó'ús   ©AA5Á5B	rq   ©	Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__rÁ   rÛ   Ú__static_attributes__Ú__classdictcell__Ú__classcell__©rÃ   rË   s   @@rU   r¼   r¼   î   s   ù‡ € õ
÷
÷ ó rW   r¼   c                   óV   a a€ ] tR tRt oV3R lV 3R lltV3R lV 3R lltRtVtV ;t# )ÚContiguousTemplatei  c                ó,   <€ V ^8„  d   QhRS[ RS[ RS[/# )ro   rC   rÖ   rT   )Ústrr   )rr   rË   s   "€rU   rs   ÚContiguousTemplate.__annotate__  s"   ø€ ÷ 
ñ 
™Sð 
©sð 
¹ñ 
rW   c                óL   <€ Wn         W n        W0n        \        SV `  VR 7       R# )r¿   N)rC   rÖ   rT   rÀ   rÁ   )rÂ   rC   rÖ   rT   rÃ   s   &&&&€rU   rÁ   ÚContiguousTemplate.__init__  s(   ø€ ØŒ	Ø&ÔØŒÜ‰ÑØð 	ö 	
rW   c                ó<   <€ V ^8„  d   QhRS[ S[,          RS[RS[/# )ro   rÆ   rÇ   rp   )rÉ   r   r   r   )rr   rË   s   "€rU   rs   rë     s-   ø€ ÷ ñ á™&•\ðñ ðñ 
ñ	rW   c           
     ó  <€ ^ RI Hp ^RIHp V! 4       ;_uu_ 4        V! 4       p\	        V P
                  V4      p\        SV `  V P                  VVVV P                  R7      uuRRR4       #   + '       g   i     R# ; i)r   rÎ   rÐ   rÔ   N)
r×   rÏ   rØ   rÑ   r   rT   rÀ   rÛ   rC   rÖ   )rÂ   rÆ   rÇ   rÏ   rÑ   rÜ   rT   rÃ   s   &&&    €rU   rÛ   ÚContiguousTemplate.generate  sm   ø€ õ
 	Då7á%×'Õ'Ù0Ó2ˆNÜØ—‘ØóˆBô
 ‘7Ñ#Ø—Y‘YØ'ØØ Ø ×,Ñ,ð $ó ÷ (×'×'Ó'ús   �AA-Á-A>	)rÖ   rT   rC   rÞ   ræ   s   @@rU   rè   rè     s   ù‡ € ÷
ó 
÷÷ ó rW   rè   c                 óJ   € \         P                  ! WP                  4       4      # rS   )r]   r@   Ú
contiguous)r®   r¯   s   &&rU   Úcontiguous_mmró   6  s   € Ü�8Š8�A—|‘|“~Ó&Ð&rW   c                 óL   € \         P                  ! WVP                  4       4      # rS   )r]   rh   rò   )ri   r®   r¯   s   &&&rU   Úcontiguous_addmmrõ   :  s   € Ü�;Š;�s˜qŸ|™|›~Ó.Ð.rW   ró   zcontiguous mmrõ   zcontiguous addmm)Útype_promotion_kindrÇ   c               ó|  a € VeÍ   S P                  4       p\        P                  ! VP                  4       V8H  R 4       \        P                  ! S P                  4       P                  R9   R 4       \        P                  ! W$8H  ;'       g?    V\        P
                  8H  ;'       d$    V\        P                  \        P                  39   R 4       \        S V4      '       dÿ   \        \        P                  ,          ! S R 4      o \        \        P                  ,          ! V^ 4      p\        S V./ RRRR7      w  rV\        P                  P                  '       dV   S P                   \        P                  \        P                  39   d'   V 3R lpV Uu. uF  p\#        V4      ! V4      NK  	  pp\#        \$        P&                  4      ! V!  p	\)        R	4      ! V	^4      p
V
# \+        S WVR
7      w  r¼rÓo pVfŽ   \-        W½WÃ4      '       d|   \.        R,          R;;,          ^,          uu&   \0        P                  ! S R 4      o \0        P                  ! V^ 4      p\0        P2                  ! \0        P4                  ! S V4      ^R7      # \7        V4      w  rïRp\9        S V.VR7      p\.        R,          RV RV RV 2;;,          ^,          uu&   \:        P=                  RVVVS P                  4       VP                  4       V4       . p\7        V4      w  rï\>        p/ pVe   \@        pRV/p. p/ p\C        4       '       d)   VPE                  V4       V'       d   VVVPF                  &   VEf#   V'       Ed   \I        VRR7      '       Ed   \K        W¼V4      '       d   VPE                  \L        4       \        PN                  R8H  pV'       g   \K        W¼V^R7      '       g›   VPE                  \P        4       \S        S WRR7      '       d   VPE                  \T        4       M[\W        S WRR7      '       dH   \        PX                  PZ                  f   VPE                  \\        4       MVPE                  \^        4       VPE                  \`        4       VPc                  \d        Pf                  Pi                  VVRVR7      4       VfR   V'       dJ   \k        W;WÍ4      '       d8   \m        R4      '       d'   \n        Pp                  ! VVVPs                  4       4       VfA   V'       d9   \u        W;WÍ4      '       d'   \v        Px                  ! VVVPs                  4       4       VfA   V'       d9   \{        W;WÍ4      '       d'   \|        P~                  ! VVVPs                  4       4       Vf-   V'       d%   \�        W;WÍS V4      '       d   ^RIAHBp V! VVV4       Vf:   \‡        VS V4      '       d'   \ˆ        P~                  ! VVVPs                  4       4       S V.pVEfA   V'       Ed8   \I        V4      '       Ed&   \        PŠ                  PŒ                  P�                  V4      '       d÷   \‘        S 4      '       dæ   . p\C        4       '       d   VPE                  R4       \“        V4      pVPc                  \d        Pf                  Pi                  V\P        .R4      4       \•        S VVVVVVV\—        4       R^
VR7      p\        PŠ                  PŒ                  P™                  V4      '       g5   Ve,   \“        V4      ^ 8”  d   V Uu. uF  pVV9   g   K  VNK  	  ppMVRV pVfP   \        Pš                   F;  pVPE                  \�        V4      PŸ                  VPs                  4       V4      4       K=  	  RpVf7   \        PŠ                  PŒ                  P                   '       d   \£        S V4      p\¤        P¦                  ! VVVPs                  4       V4      ;p'       d   V# \©        VVVPs                  4       VVR7      w  p p!V # u upi u upi )!zW
Lowering for autotuning aten.mm with different backends (Aten, Triton, CUTLASS, etc.)
Nc                  ó   € R # rž   rq   rq   rW   rU   r“   Útuned_mm.<locals>.<lambda>O  s   € Ñ3rW   c                  ó   € R # )z+out_dtype is only supported for CUDA or XPUrq   rq   rW   rU   r“   rù   S  s   € ÑArW   c                  ó   € R # )zFout_dtype must be the same as input dtype or fp32 for fp16/bf16 inputsrq   rq   rW   rU   r“   rù   [  s   € Ñ\rW   TF)ÚargsÚkwargsÚ	broadcaströ   Úconvert_input_to_boolc                 óJ   <€ \         P                  ! V SP                  R R7      # )F)Úuse_compute_types)r   Úto_dtyper­   )Úxrj   s   &€rU   Ú	_to_dtypeÚtuned_mm.<locals>._to_dtype|  s   ø€ Ü—|’| A t§z¡zÀUÔKÐKrW   Údot©rÇ   r¥   ÚinductorÚdecompose_mm_pointwise)Úaxisr@   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Ú
exhaustive)Úthreshold_multiple©Úoutput_layoutÚ
add_guards©Úkwarg_overrides)Úadd_nv_universal_gemm_choicesÚ	extern_mmzmm-ah)Útop_kÚalways_included)Úbest_config_future)ÚcudaÚxpuéÿÿÿÿ)Ur\   r]   r˜   r¡   Útyperª   Úfloat16Úbfloat16r?   r"   ÚatenÚ	unsqueezer&   Úinductor_configÚtritonÚcodegen_upcast_to_fp32r­   r#   r   r  r$   r<   r:   r   ÚLÚsum_Úmulr9   r    ÚlogÚinfoÚaten_mmÚaten_mm_dtyper.   ÚappendÚuidr7   r3   Údecompose_k_subgraph_templateÚmax_autotune_gemm_search_spaceÚmm_templater5   Ú.blackwell_ws_persistent_device_tma_mm_templater8   ÚversionÚhipÚpersistent_tma_mm_templateÚpersistent_mm_templateÚmm_contiguous_subgraph_templateÚextendr   ÚchoicesÚget_template_configsr2   r,   r   Úadd_cutlass_gemm_choicesÚnodesr/   r   Úadd_ck_gemm_choicesr0   r   Úadd_choicesr4   Úcodegen.nv_universal_gemmr  r1   r
   Ú	_inductorr   Úrun_autoheuristicr   ÚlenÚmm_autoheuristicr	   Úcollect_autoheuristicÚexternal_matmulrV   Ú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_nonzerorC   Úkernel_inputsr9  Úaten_handlerÚaten_extra_kwargsÚtemplates_to_user  Úis_exhaustiver  rÆ   r  Ú num_choices_before_extra_configsÚ
ah_choicesÚchoicer  ÚboxÚnoder  s"   f&&$                              rU   Útuned_mmrX  F  s  ø€ ð
 ÒØ—n‘nÓ&ˆÜ�ŠØ�N‰NÓ Ñ+Ù3ô	
ô 	�ŠØ�O‰OÓ×"Ñ" oÑ5ÙAô	
ô 	�ŠØÑ$÷ ð àœUŸ]™]Ñ*÷ Cð CØ¤E§M¡M´5·>±>Ð#BÑBá\ô	
ô. ˜˜t×$Ò$ÜœŸ™Ö(¨¨rÓ2ˆÜœŸ™Ö(¨¨qÓ1ˆÜ%Ø˜�ØØØ $Ø"'ô
‰ˆô ×!Ñ!×8×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¨Yô#Ñ€Aˆ!�T˜4ð ÒÔ4°Q¸1×EÒEÜ�ÕÐ5×6¸!Õ;Ó6Ü�{Š{˜4 Ó$ˆÜ�{Š{˜4 Ó#ˆÜ�vŠv”a—e’e˜D $Ó'¨aÔ0Ð0ä1°&Ó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×,Ñ,Ñ-ð 	ÓßˆJÜ ¸4×@Ô@ä! !¨×*Ò*Ø×#Ñ#Ô$AÔBô (×FÑFÈ,ÑVˆßÔ 6°q¸QÐST× UÓ UØ×#Ñ#¤KÔ0ä0Ø�d¸T÷ó ð !×'Ñ'Ô(VÕWÜ(Ø�d¸T÷ó ô —=‘=×$Ñ$Ò,Ø$×+Ñ+Ô,FÕGà$×+Ñ+Ô,BÔCà×ÑÔ ?Ô@à‡N�NÜ	�	‰	×&Ñ&ØØØØ+ð	 	'ó 	
ôð 	ÒßÜ  ¨A×1Ò1Ü ×%Ò%ä×6Ò6Ø�V˜]×0Ñ0Ó2ô	
ð ÒŸZÔ,@ÀÈA×,QÒ,QÜ×*Ò*¨7°F¸M×<OÑ<OÓ<QÔRØÒŸZÔ,EÀfÐQR×,VÒ,VÜ×&Ò& w°¸×8KÑ8KÓ8MÔNð 	ÒßÜ*¨6°a¸DÀ$×GÒGåMá% g¨v°}ÔEàÒÔ2°6¸4À×FÒFÜ×#Ò#ØØØ×ÑÓ!ô	
ð ˜�,€KàÓßˆJÜ ×'Ó'Ü�O‰O×"Ñ"×4Ñ4°T×:Ò:Ü�d�OŠOàˆÜ ×"Ò"Ø×"Ñ" ;Ô/Ü+.¨w«<Ð(Ø�‰Ü�I‰I×*Ñ*ð Ü�Øóô		
ô &ØØØØØØØØÜ‹OØØØ+ô
ˆ
ô �‰×%Ñ%×;Ñ;¸D×AÒAàÒ%¬#¨j«/¸AÔ*=ñ
 18ÓP± f¸6ÀZÑ;OŸ6˜6±�ÐP�à!Ð"CÐ#CÐD�àÒÜ ×0Ô0ˆAØ�N‰NÜ+¨AÓ.×3Ñ3°M×4GÑ4GÓ4IÈ6ÓRöñ 1ð
 ÐØÒœUŸ_™_×3Ñ3×N×NÐNô -¨T°4Ó8Ðä"×8Ò8Øˆg�}×*Ñ*Ó,¨fóð €sö ð ˆ
ä'ØØØ×ÑÓØØ-ô�G€Dˆ!ð €KùòE @ùòL Qs   Æ`4Ü5	`9Ý`9c          	      ój  € \        WV\        P                  R 7      w  r4rRrRp\        R,          RV RV RV 2;;,          ^,          uu&   \        P                  RVVVV P                  4       VP                  4       V4       \        V4      w  rxT;'       d    T;'       d    \        W#WE4      p	. p
\        W.\        P                  R7      p. p\        4       '       d   VP                  \        4       V'       d*   \        VRRR	7      '       d   VP                  \        4       V
P                  \         P"                  P%                  W¼V4      4       V	'       d:   \'        V4      '       d)   \(        P*                  ! W¢VP-                  4       RRR
7       \/        WjVP-                  4       V4      w  rÞV# )r  Úint_mmr  zaten._int_mm_r  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   r2   r    r.   r-  Úaten__int_mmr7   r1  r8  r   r9  r:  r,   r   r;  r<  r'   )rj   rk   rÇ   r±   r²   r³   rC   rL  rM  Úuse_cutlassr9  rN  rQ  rW  r  s   &&$            rU   Útuned_int_mmrb  D  s  € ô #*Ø˜6¬U¯[©[ô#Ñ€Aˆ!�Tð €Däˆ^Õ˜}¨Q¨C¨q°°°1°Q°CÐ8×9¸QÕ>Ó9Ü‡H�HØ^Ø	Ø	Ø	Ø�‰ÓØ�‰ÓØôô  2°&Ó9Ñ€LØ×WÐW :×WÐWÔ2FÀvÐRSÓ2W€KØ"$€Gô # D <¼5¿;¹;ÔG€Mð CEÐÜ×ÒØ×Ñ¤Ô-çÔ)Ø˜T°e÷ó ð 	×Ñ¤Ô,ð ‡N�NÜ	�	‰	×&Ñ& }ÈÓMô÷ Ô*¨4×0Ò0Ü×6Ò6Ø˜]×0Ñ0Ó2¸TÐPTõ	
ô (¨°}×7JÑ7JÓ7LÈfÓU�G€DØ€KrW   c          	     óD  € V^ 8X  dû   VP                  4       P                  R8X  dÜ   \        WV4       V^ 8X  dV   \        WVR7      w    rer\        \
        P                  ,          ! VP                  ^ VP                  VP                  R7      # Ve%   \        \
        P                  ,          ! WVR7      pM!\        \
        P                  ,          ! W4      pV^8w  d"   \        \
        P                  ,          ! W74      pV# \        W4      '       d¤   V^ 8X  d   \        WV4       ^ pM!\        \
        P                  ,          ! W@4      pV^ 8X  d   ^ p	MA\        \
        P                  ,          ! V\        \
        P                  ,          ! W4      4      p	\        \
        P                  ,          ! W‰4      # \        WWR7      w  r«rÅrp\        V 4      p \        V4      w  rïRp\!        WÑV.\#        W4R7      R7      p\!        WV.\#        W4R7      R7      p. p\$        R,          RV
 R	V R	V 2;;,          ^,          uu&   \&        P)                  R
V
VVVP+                  4       VP+                  4       V4       V'       d-   \,        P.                  '       gm   \,        P0                  '       gW   VP3                  \4        P6                  P9                  V\:        .V4      4       \=        VVVP?                  4       V4      w  ppV# . p\A        4       '       dÇ   \:        .pV PC                  4       ^ ,          ^ 8X  dt   \E        V PG                  4       4      ^8X  dV   \,        PH                  PJ                  '       d6   \4        PL                  PN                  '       g   VPQ                  \R        4       VP3                  \4        P6                  P9                  VVV4      4       V'       dã   \U        VRR7      '       dÐ   VPQ                  \V        4       \Y        WVRR7      '       d   VPQ                  \Z        4       M[\]        WVRR7      '       dH   \^        P`                  Pb                  f   VPQ                  \d        4       MVPQ                  \f        4       VP3                  \4        P6                  P9                  V\h        .V4      4       VP3                  \4        P6                  P9                  VVV4      4       V'       dT   \k        WZW¼4      '       dB   \m        V4      '       d1   \n        Pp                  ! VVVP?                  . ROR7      VV. ROR7       V'       d'   \s        WZW¼W4      '       d   ^RI:H;p V! VVVWVR7       V'       dC   \y        WZW¼4      '       d1   \z        P|                  ! VVVP?                  . ROR7      VV. ROR7       \        WQV4      '       d+   \€        P‚                  ! VVVP?                  4       VVRR7       \=        VVVP?                  4       V4      w  ppV# )zZ
Lowering for autotuning aten.addmm with different backends (Aten, Triton, CUTLASS, etc.)
r  )rÇ   )r­   Údevicerh   )rc   rd   )Úscalarsr  zaten.addmm_r  zRTuned aten.addmm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sFr  Tr  )Úreorder)rc   rd   Úinput_reorder)Úadd_nv_universal_addmm_choices)rc   rd   Úhas_bias)r}   ro   r   )ro   r   r}   )Br¡   r  r¢   r<   r"   r!  Úfullrg   r­   rd  r@   r(  r?   Úaddr*   r9   r    Údictr   r)  r*  r\   r#  Úmax_autotuneÚmax_autotune_gemmr8  r   r9  r:  Ú
aten_addmmr'   r<  r.   r‘   rB  r™   r$  Úautotune_cublasLtr   Úcpp_wrapperr-  Úaten_bias_addmmr7   r1  r5   r2  r8   r]   r3  r4  r5  r6  Ú"addmm_contiguous_subgraph_templater2   r,   r   r;  r4   r?  rh  r/   r   r=  r1   r
   r>  )ri   rj   rk   rc   rd   rÇ   r  r¸   Úarg1Úarg2r±   r²   r³   Úinp_expandedrL  rM  rC   rN  Úkernel_inputs_atenr9  rW  rQ  Úaten_templatesrh  s   &&&$$$                  rU   Útuned_addmmry  v  sí  € ð
 ˆq„y�T—_‘_Ó&×+Ñ+¨vÔ5Ü# C¨tÔ4Ø�AŒ:Ü*1°$ÀVÔ*LÑ'ˆAˆq�!˜TÜœTŸY™YÖ'Ø—‘ØØ—l‘lØ—}‘}ô	ð ð ÒÜœtŸw™wÖ'¨¸6ÔB‰FäœtŸw™wÖ'¨Ó3ˆFØ�AŒ:ÜœtŸx™xÖ(¨Ó7ˆFØˆä˜×$Ò$Ø�1Œ9Ü'¨°4Ô8Ø‰DäœTŸX™XÖ& tÓ1ˆDà�AŒ:Ø‰DäœTŸX™XÖ& u¬i¼¿¹Ö.@ÀÓ.LÓMˆDäœŸ™Ö" 4Ó.Ð.ô 18¸ÀCÔ0WÑ-€Aˆ!�T Ü
˜Ó
€CÜ1°&Ó9Ñ€LØ€Dô #Ø	˜TÐ"¬D°uÔ,Hô€Mô (Ø	�DÐ¤4¨eÔ#?ôÐð #%€Gô ˆ^Õ˜{¨1¨#¨Q¨q¨c°°1°#Ð6×7¸1Õ<Ó7Ü‡H�HØ\Ø	Ø	Ø	Ø�‰ÓØ�‰ÓØô÷ Ü×)×)Ð)¬_×-N×-NÐ-Nà�‰Ü�I‰I×*Ñ*Ø"Ü�Øóô	
ô ,Ø�'˜=×.Ñ.Ó0°&ó
‰ˆˆað ˆàBDÐä×ÒÜEOÀLˆà�N‰NÓ˜QÕ 1Ô$Ü�C—L‘L“NÓ# qÔ(Ü×&Ñ&×8×8Ð8Ü—G‘G×'×'Ð'à×!Ñ!¤/Ô2ð 	�‰Ü�I‰I×*Ñ*Ð+=¸~ÈtÓTô	
÷ Ô)¨&ÀU×KÓKØ×Ñ¤Ô,ä,Ø f¸÷
ó 
ð ×#Ñ#Ô$RÕSÜ$ T¸vÐRV×WÓWÜ�}‰}× Ñ Ò(Ø ×'Ñ'Ô(BÕCà ×'Ñ'Ô(>Ô?ð 	�‰Ü�I‰I×*Ñ*Ø"Ô%GÐ$HÈ$óô	
ð ‡N�NÜ	�	‰	×&Ñ& }Ð6FÈÓMô÷
 	Ü  ¨A×1Ò1Ü ×%Ò%ä×6Ò6ØØð ×Ñª	ÐÓ2ØØÚ#õ		
÷ Ô4°VÀÀd×QÒQÝNñ 	'Ø�V˜]¨CÀ4õ	
÷ Ô*¨6°a×;Ò;Ü×*Ò*ØØð ×Ñª	ÐÓ2ØØÚ#õ		
ô ˜V¨4×0Ò0Ü×#Ò#ØØØ×ÑÓ!ØØØõ	
ô (¨¨g°}×7JÑ7JÓ7LÈfÓU�G€Dˆ!Ø€KrW   r¥   c                óú  € ^ RI Hp V! WV4      w  rpV P                  4       w  rgVP                  4       w  r‰VP                  4       w  r«\        P                  P
                  P                  Wh4      p\        P                  P
                  P                  ^V,          V
4      pVf<   ^ RIHp T! VP                  4       V'       d   TMVP                  4       WË.V^.4      pMVe   \        R4      h\        4       '       d   \        P                  WV3WCR7      .M. pWË,          ^ 8w  d@   \        WLW½4      '       d.   \!        R4      '       d   \"        P$                  ! WôWV.RRR7       \'        RWðW3V4      w  pp	V# )r   )r*   )ÚFixedLayoutz,out_dtype is ignored if layout is specified.r¤   Úsparse_semi_structured_mmTr\  )Ú torch._inductor.select_algorithmr*   r™   r   r   r€   Úcheck_equals_and_simplifyÚtorch._inductor.irr{  r¡   r\   ÚAssertionErrorr.   Úaten__sparse_semi_structured_mmrF  r2   r,   r   r;  r'   )rj   Ú	mat1_metark   r¥   rÇ   r*   Úm1Úk1Úm2r  Úk2r²   r±   r³   r{  r9  rW  s   &&&$$            rU   Útuned_sparse_semi_structured_mmr‡  !  sf  € õ @ñ +¨4¸DÓAÑ€D�TØ�]‰]‹_�F€BØ×ÑÓ �E€BØ�M‰M‹O�E€BÜ	�‰×Ñ×2Ñ2°2Ó:€AÜ	�‰×Ñ×2Ñ2°1°rµ6¸2Ó>€AØ‚~Ý2áØ�O‰OÓß"‰I¨¯©Ó(8ØˆFØ�ˆFó	
‰ð Ò Ü Ð!OÓPÐPô !×"Ò"ô	 ,×0Ñ0Ø $Ð'¨ð 1ó ñ	
ð ð ð 	
��Œ
Ü  ¨A×1Ò1ÜÐ ;×<Ò<ä×6Ò6Ø˜d¨)Ð4¸tÐRVõ	
ô (Ø# W°YÐ.EÀvó�G€Dˆ!ð €KrW   c                ó0   € V ^8„  d   QhR\         R\        /# )ro   Úszrp   ©r   rx   )rr   s   "rU   rs   rs   a  s   € ÷ ñ œsð ¤tñ rW   c                 ó–   € \        V 4      ^ 8H  ;'       g5    \        ;QJ d    R V  4       F  '       d   K   R# 	  R# ! R V  4       4      # )r   c              3   óv   "  € T F/  p\         P                  P                  P                  V^4      x € K1  	  R# 5i)r}   Nr~   )Ú.0Úds   & rU   Ú	<genexpr>Ú)_is_tensorwise_scaling.<locals>.<genexpr>b  s-   é € ð !Ù@B¸1Œ�‰×Ñ×0Ñ0°°A×6Ð6Ãùs   ‚79FT)rB  Úall)r‰  s   &rU   Ú_is_tensorwise_scalingr’  a  sQ   € Ü�‹G�q‰L÷ ð ŸS›Sñ !Ù@Bó!ŸSœSð šSð ˜Sñ !Ù@Bó!ó ð rW   c                ó<   € V ^8„  d   QhR\         R\        R\        /# )ro   r‰  Ú	transposerp   rŠ  )rr   s   "rU   rs   rs   g  s&   € ÷ @ñ @œCð @¬Dð @´Tñ @rW   c                 óz   € V'       d   ^ MRp\         P                  P                  P                  W,          ^4      # )r   r  r~   )r‰  r”  Úidxs   && rU   Ú_is_rowwise_scalingr—  g  s*   € ß‰!˜b€CÜ�7‰7×Ñ×3Ñ3°BµG¸QÓ?Ð?rW   c          
      óT   € V ^8„  d   QhR\         R\         R\        R\        R\        /# )ro   r‰  Ú	tensor_szÚ	tile_sizer”  rp   )r   rÊ   rx   )rr   s   "rU   rs   rs   l  s4   € ÷ 	ñ 	Üð	Üð	Ü(+ð	Ü8<ð	ä	ñ	rW   c                 ó.  € V'       d   ^M^ pV'       d   ^ M^p\         P                  P                  P                  W,          W,          4      ;'       dA    \         P                  P                  P                  W,          \	        W,          V4      4      # r|   ©r   r   r€   r�   r-   )r‰  r™  rš  r”  ÚlhsÚrhss   &&&&  rU   Ú_is_blockwise1xTILESIZE_scalingrŸ  l  sn   € ÷ ‰!˜a€Cß‰!˜a€CÜ�7‰7×Ñ×3Ñ3Ø
��•ó÷ ð ä
�'‰'×
Ñ
×
2Ñ
2Ø
�”˜�¨Ó3óðrW   c                ó<   € V ^8„  d   QhR\         R\         R\        /# )ro   r‰  r™  rp   rŠ  )rr   s   "rU   rs   rs   x  s&   € ÷ Vñ V¤Sð V´Sð V¼Tñ VrW   c                 ó  € \         P                  P                  P                  V ^ ,          \	        V^ ,          ^€4      4      ;'       dC    \         P                  P                  P                  V ^,          \	        V^,          ^€4      4      # rˆ   rœ  )r‰  r™  s   &&rU   Ú_is_blockwise128x128_scalingr¢  x  sj   € Ü�7‰7×Ñ×3Ñ3Ø
ˆ1�Œw�y •| SÓ)ó÷ Vð Vä
�'‰'×
Ñ
×
2Ñ
2°2°aµ5¼'À)ÈAÅ,ÐPSÓ:TÓ
UðVrW   c          
      óh   € V ^8„  d   QhR\         R\        P                  R\        R\        R\        /# )ro   ÚtÚ
scale_sizeÚscaling_typer”  rp   )r   r]   ÚTensorr   rx   )rr   s   "rU   rs   rs   ~  sE   € ÷ Mñ MÜ
ðMä—‘ðMô ðMô ð	Mô
 
ñMrW   c                 ób  € T;\         P                  8X  d    \        V4      # ;\         P                  8X  d    \	        W4      # ;\         P
                  8X  d    \        WP                  4       ^€V4      # \         P                  8X  d   \        WP                  4       4      #  \        RV 24      h)é€   úUnsupported scaling type )r   Ú
TensorWiser’  ÚRowWiser—  ÚBlockWise1x128rŸ  r™   ÚBlockWise128x128r¢  r€  )r¤  r¥  r¦  r”  s   &&&&rU   Úis_desired_scalingr¯  ~  sŠ   € ð Ø#Œ[×#Ö#Ý)¨*Ó5Ð5Ø Œ[× Ö Ý& zÓ=Ð=Ø'Œ[×'Ö'Ý2ØŸJ™J›L¨#¨yóð ô ×)Ö)Ü/°
¿J¹J»LÓIÐIØÜ Ð#<¸\¸NÐ!KÓLÐLrW   c                ó$   € V ^8„  d   QhR\         /# rv   ©rÊ   )rr   s   "rU   rs   rs   “  s   € ÷ 	ñ 	¤3ñ 	rW   c                 ó€   € T ;\         P                  8X  d    ^€# \         P                  8X  d   ^€#  \        RV  R24      h)r©  rª  z in get_tile_size)r   r®  r­  r€  )Úscale_options   &rU   Úget_tile_sizer´  “  s<   € Ø
Ø)Œ[×)Ö)ÚÜ×'Ö'ÙØÜ Ø+¨L¨>Ð9JÐKóð rW   c                óž   € V ^8„  d   QhR\         R\         R\        P                  R\        P                  R\        \        \        3,          /# )ro   r’   r—   Úscale_a_sizeÚscale_b_sizerp   )r   r]   r§  Útupler   )rr   s   "rU   rs   rs   Ÿ  sL   € ÷ ñ Üðäðô —,‘,ðô —,‘,ð	ô
 Œ;œÐ#Õ$ñrW   c                 óœ   € \          F2  w  rE\        WV4      '       g   K  \        WVR R7      '       g   K/  WE3u # 	  \        RV RV 24      h)T)r”  z1Inductor Triton does not support scale_a.shape = z, scale_b.shape = )Úscaling_pairsr¯  r€  )r’   r—   r¶  r·  Úscale_option_aÚscale_option_bs   &&&&  rU   Úget_scaling_optionsr½  Ÿ  s^   € ÷ +8Ñ&ˆÜØ ÷
ô 
ä  °nÐPT×UÕUØ!Ð1Ò1ñ	 +8ô Ø
;¸L¸>ÐI[Ð\hÐ[iÐjóð rW   )Úadd_to_fallback_setc                óä   € V ^8„  d   QhR\         \        ,          R\         \        ,          R\         \        ,          R\         \        ,          R\         \        ,          R\         \        ,          /# )ro   Úscale_aÚrecipe_aÚ	swizzle_aÚscale_bÚrecipe_bÚ	swizzle_b)rÉ   r   rÊ   )rr   s   "rU   rs   rs   º  sh   € ÷ Yñ Yô ”#�YðYô ”3�ið	Yô
 ”C�yðYô ”#�YðYô ”3�iðYô ”C�yñYrW   c                ó`
  € R R lp\        V4      ^8H  ;'       d    \        V4      ^8H  pV! V4      ;'       d	    V! V4      p\        ;QJ d    R V 4       F  '       g   K   RM	  RM! R V 4       4      '       gq   \        ;QJ d    R V 4       F  '       g   K   RM	  RM! R V 4       4      '       g7   V'       d/   V'       d'   V^ ,          P                  \        P                  8w  d   V
f   . MT
p\        V VVVVVVVVV	VV4      # \        WWÉR7      w  ppprÀp\        R,          R	V R
V R
V 2;;,          ^,          uu&   \        P                  RVVVV P                  4       VP                  4       V4       Rp\        W4       \        V^ ,          V^ ,          4      w  ppV'       g   WVV.pM\        V4      pWVVV.p\        V^ ^V	R7      p. p. p/ p\        4       '       d3   VP                  \         4       \#        W›R7      V\         P$                  &   \'        V4      w  ppV'       Edã   V'       EdÚ   V^ ,          P                  \        P                  8X  Ed³   V'       Edª   \)        VRRR7      '       Ed•   \#        VR7      p\+        V^ ,          4      \+        V^ ,          4      p!p \-        WVRR7      '       dÎ   V'       gÆ   V P.                  VR&   V!P.                  VR&   \1        V V!\2        4      '       d*   VP                  \4        4       VV\4        P$                  &   Mg\1        V V!\6        4      '       dF   \9        V 4      VR&   \9        V!4      VR&   VP                  \:        4       VV\:        P$                  &   M\=        R4      h\?        WVRR7      '       d1   V'       g)   VP                  \@        4       VV\@        P$                  &   \1        V V!\2        4      '       d)   VP                  \B        4       VV\B        P$                  &   VPE                  \F        PH                  PK                  VVVVR7      4       V'       d(   \M        VVVVW4      '       d   ^RI'H(p" V"! VVVVR7       V'       dN   \S        VVVV4      '       d:   \U        V4      '       d)   \V        PX                  ! VVVP[                  4       VR7       V'       d;   \]        VVVV4      '       d'   \^        P`                  ! VVVP[                  4       4       \c        VVVP[                  4       V4      w  p#pV## )a6  
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.
c                óF   € V ^8„  d   QhR\         \        ,          R\        /# )ro   Úreciperp   )rÉ   rÊ   rx   )rr   s   "rU   rs   Ú(tuned_scaled_mm_v2.<locals>.__annotate__Ý  s"   € ÷ Eñ E¤t¬C¥yð E´Tñ ErW   c                 óÊ   a€ \        \        P                  \        P                  .4      o\        ;QJ d    V3R  lV  4       F  '       d   K   R# 	  R# ! V3R  lV  4       4      # )c              3   ó@   <"  € T F  p\        V4      S9  x € K  	  R # 5irS   r   )r�  ÚrÚ
disalloweds   & €rU   r�  ÚEtuned_scaled_mm_v2.<locals>.check_supported_recipe.<locals>.<genexpr>ß  s   øé € ÐD¹V¸”;˜q“>¨Ö3»Vùs   ƒFT)r   r   ÚBlockWise1x16ÚBlockWise1x32r‘  )rÈ  rÍ  s   &@rU   Úcheck_supported_recipeÚ2tuned_scaled_mm_v2.<locals>.check_supported_recipeÝ  sG   ø€ Ü¤×!:Ñ!:¼K×<UÑ<UÐ VÓWˆ
ß‹sÔD¹VÓD�sŒsÐDŠsÐDˆsÔD¹VÓDÓDÐDrW   c              3   ó*   "  € T F	  q^ 8g  x € K  	  R# 5i©r   Nrq   ©r�  Úss   & rU   r�  Ú%tuned_scaled_mm_v2.<locals>.<genexpr>æ  s   é € Ð&™I�q�ŽF›Iùó   ‚TFc              3   ó*   "  € T F	  q^ 8g  x € K  	  R# 5irÔ  rq   rÕ  s   & rU   r�  r×  ç  s   é € Ð)™y˜!�AŽv›yùrØ  r  r  zaten._scaled_mm_v2.default_r  zbTuned aten._scaled_mm_v2.default: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sÚ	scaled_mm©Úmat1_idxÚmat2_idxr¥   ©r¥   Úuse_fast_accum©Úenable_float8r  ©ÚUSE_FAST_ACCUMr  ÚSCALE_RECIPE_AÚSCALE_RECIPE_BÚTILE_SIZE_AÚTILE_SIZE_Br  ©Ú$add_nv_universal_scaled_gemm_choices©rN  ©rß  úpInductor Triton does not support scaling options that are present in both epilogue scaling and main loop scaling)2rB  Úanyr­   r]   rª   Úscaled_mm_v2_fallbackr<   r   r)  r*  r\   rš   r*   r    r.   r-  Úaten__fp8_mmrl  r.  r9   r7   r   r8   Úvaluer6   Úepilogue_scaling_typesÚ.scaled_mm_device_tma_epilogue_scaling_templateÚmain_loop_scaling_typesr´  Ú/scaled_mm_device_tma_main_loop_scaling_templater€  r5   r2  r1  r8  r   r9  r:  r4   r?  ré  r2   r,   r   r;  r<  r/   r   r=  r'   )$r’   r—   rÀ  rÁ  rÂ  rÃ  rÄ  rÅ  Úbiasr¥   Úcontraction_dimrß  rÇ   rÑ  Úis_single_level_scaleÚsupported_recipeÚfallback_contraction_dimr±   r²   r³   rC   Úscale_a_realÚscale_b_realrÆ   Ú	bias_realrN  r9  rQ  r  r  rM  Ú
overridersr»  r¼  ré  rW  s$   &&&&&&&&&&&&&                       rU   Útuned_scaled_mm_v2rþ  ¹  sÉ  € õHEô   ›L¨AÑ-×CÐC´#°g³,À!Ñ2CÐÙ-¨hÓ7÷ ð Ñ<RØó=Ð÷ 	‹Ñ&™IÓ&��ŠÑ&™IÓ&×&Ò&ß‹3Ñ)™yÓ)�3�3Š3Ñ)™yÓ)×)Ò)ßß$Ø�1�:×ÑœuŸ}™}Ô,ð *9Ò)@¡2ÀoÐ Ü$ØØØØØØØØØØØ$Øó
ð 	
ô %,Ø˜Vô%Ñ!€A€qˆ!ˆV˜Eô ˆ^ÕÐ:¸1¸#¸Q¸q¸cÀÀ1À#ÐF×GÈ1ÕLÓGÜ‡H�HØlØ	Ø	Ø	Ø�‰ÓØ�‰ÓØôð €DÜ˜UÔ*ä!/°¸µ
¸GÀA½JÓ!GÑ€L�,÷ Ø \°<Ð@‰ä" 4Ó(ˆ	Ø \°<ÀÐKˆô #Ø˜a¨!°yô€Mð #%€Gð CEÐØ€Oä×ÒØ×Ñ¤Ô-Ü,0Øô-
ˆœ×(Ñ(Ñ)ô ' vÓ.�M€A€z÷ 	ÐßÐØ�A�J×Ñ¤§¡Õ-ßˆJÜ °dÈu×UÔUä¨Ô8ˆ
ô
 ˜ �Ó$Ü˜ �Ó$ð 'ˆô $ EÀÐSW×XÓXßà+9×+?Ñ+?ˆJÐ'Ñ(Ø+9×+?Ñ+?ˆJÐ'Ñ(ä*Ø Ô0F÷ò ð !×'Ñ'Ô(VÔWàð  Ô N× RÑ RÒSô -Ø Ô0G÷ò ô -:¸.Ó,I�
˜=Ñ)Ü,9¸.Ó,I�
˜=Ñ)à ×'Ñ'Ô(WÔXàð  Ô O× SÑ SÒTô %ðGóð ô .Ø¨F¸t÷ó ÷ à×#Ñ#Ô$RÔSàð ÔJ×NÑNÑOô 'Ø˜NÔ,B÷
ò 
ð ×#Ñ#¤KÔ0Ø/9ˆOœKŸO™OÑ,ð ‡N�NÜ	�	‰	×&Ñ&ØØØØ+ð	 	'ó 	
ô÷ Ô4°V¸QÀÀ1Àe×SÒSÝTá,ØØØØ'õ		
÷ 	Ü  ¨¨A¨q×1Ò1Ü ×%Ò%ä×6Ò6ØØØ×ÑÓ!Ø)õ		
÷ Ô*¨6°1°a¸×;Ò;Ü×*Ò*¨7°F¸M×<OÑ<OÓ<QÔRä'¨¨g°}×7JÑ7JÓ7LÈfÓU�G€Dˆ!Ø€KrW   c	           	     óP  € \        WW†R7      w  ršr¸r\        R,          RV	 RV
 RV 2;;,          ^,          uu&   \        P                  RV	V
VV P	                  4       VP	                  4       V4       Rp\        W4       \        W#4      w  rÞV'       g   WWÞ.pM\        V4      pWWÞV.p\        V^ ^VR7      p. p. p/ p\        4       '       d3   VP                  \        4       \        WgR7      V\        P                  &   \        V4      w  ppVP                  \        P                   8X  Ed·   V'       Ed®   \#        VR	R
R7      '       Ed™   \        VR7      pVP$                  VP$                  pp\'        WVV4      w  pp\)        WVR	R7      '       dÎ   V'       gÆ   VP*                  VR&   VP*                  VR&   \-        VV\.        4      '       d*   VP                  \0        4       VV\0        P                  &   Mg\-        VV\2        4      '       dF   \5        V4      VR&   \5        V4      VR&   VP                  \6        4       VV\6        P                  &   M\9        R4      h\;        WVR	R7      '       d1   V'       g)   VP                  \<        4       VV\<        P                  &   \-        VV\.        4      '       d)   VP                  \>        4       VV\>        P                  &   VPA                  \B        PD                  PG                  VVVVR7      4       V'       d&   \I        W‰W«W4      '       d   ^RI%H&p V! VVVVR7       VP                  \        P                   8w  d   \O        VVWø4      w  ppV# V'       dL   \Q        W‰W«4      '       d:   \S        V4      '       d)   \T        PV                  ! VVVPY                  4       VR7       V'       d9   \[        W‰W«4      '       d'   \\        P^                  ! VVVPY                  4       4       \O        VVVPY                  4       V4      w  ppV# )a	  
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  z_Tuned aten._scaled_mm.default: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%srÚ  rÛ  rÞ  TFrà  râ  r  rä  rå  ræ  rç  r  rè  rê  rë  rì  )0r<   r   r)  r*  r\   rš   r*   r    r.   r-  rï  rl  r.  r9   r­   r]   rª   r7   r¦   r½  r8   rð  r6   rñ  rò  ró  r´  rô  r€  r5   r2  r1  r8  r   r9  r:  r4   r?  ré  r'   r2   r,   r   r;  r<  r/   r   r=  )r’   r—   rÀ  rÃ  rõ  Úscale_resultr¥   rß  rÇ   r±   r²   r³   rC   rú  rû  rÆ   rü  rN  r9  rQ  r  r  rM  rý  r¶  r·  r»  r¼  ré  rW  s   &&&&&&&&&                     rU   Útuned_scaled_mmr  –  sá  € ô8 %,Ø˜Vô%Ñ!€Aˆ!�Uô ˆ^ÕÐ7¸°s¸!¸A¸3¸aÀ¸sÐC×DÈÕIÓDÜ‡H�HØiØ	Ø	Ø	Ø�‰ÓØ�‰ÓØôð €DÜ˜UÔ*ä!/°Ó!AÑ€L÷ Ø \Ð@‰ä" 4Ó(ˆ	Ø \ÀÐKˆô #Ø˜a¨!°yô€Mð #%€Gð CEÐØ€Oä×ÒØ×Ñ¤Ô-Ü,0Øô-
ˆœ×(Ñ(Ñ)ô ' vÓ.�M€A€zð 	�‰œŸ™Õ&ßˆJÜ °dÈu×UÔUä¨Ô8ˆ
à%1×%7Ñ%7¸×9KÑ9K�lˆä)<Ø˜,¨ó*
Ñ&ˆ˜ô $ EÀÐSW×XÓXßà+9×+?Ñ+?ˆJÐ'Ñ(Ø+9×+?Ñ+?ˆJÐ'Ñ(ä*Ø Ô0F÷ò ð !×'Ñ'Ô(VÔWàð  Ô N× RÑ RÒSô -Ø Ô0G÷ò ô -:¸.Ó,I�
˜=Ñ)Ü,9¸.Ó,I�
˜=Ñ)à ×'Ñ'Ô(WÔXàð  Ô O× SÑ SÒTô %ðGóð ô .Ø¨F¸t÷ó ÷ à×#Ñ#Ô$RÔSàð ÔJ×NÑNÑOô 'Ø˜NÔ,B÷
ò 
ð ×#Ñ#¤KÔ0Ø/9ˆOœKŸO™OÑ,ð ‡N�NÜ	�	‰	×&Ñ&ØØØØ+ð	 	'ó 	
ô÷ Ô4°VÀÀe×SÒSÝTá,ØØØØ'õ		
ð ‡}�}œŸ™Ô%Ü+¨D°'¸;ÓO‰ˆˆaØˆ÷ 	Ü  ¨A×1Ò1Ü ×%Ò%ä×6Ò6ØØØ×ÑÓ!Ø)õ		
÷ Ô*¨6°a×;Ò;Ü×*Ò*¨7°F¸M×<OÑ<OÓ<QÔRä'¨¨g°}×7JÑ7JÓ7LÈfÓU�G€Dˆ!Ø€KrW   c                ó>   € V ^8„  d   QhR\         R,          R\        /# )ro   ÚindexNrp   )rÊ   rx   )rr   s   "rU   rs   rs   M  s   € ÷ ñ ¤ t¥ð ´ñ rW   c                 ór   € \         P                  P                  T ;'       g    ^ 4      pVP                  ^8*  # rˆ   )r]   r  Úget_device_propertiesÚmajor)r  Úpropss   & rU   Ú_is_sm7x_or_older_gpur  L  s+   € ä�J‰J×,Ñ,¨U¯Z¨Z°aÓ8€EØ�;‰;˜!ÑÐrW   c                 ój   € \         ;QJ d    R  V  4       F  '       d   K   R# 	  R# ! R  V  4       4      # )c              3   óB   "  € T F  p\        V\        4      x € K  	  R # 5irS   )Ú
isinstancerÊ   )r�  Údims   & rU   r�  Údims_are_int.<locals>.<genexpr>S  s   é € Ð4©t¨Œz˜#œs×#Ð#«tùs   ‚FT)r‘  )Údimss   &rU   Údims_are_intr  R  s+   € ß‹3Ñ4©tÓ4�3Œ3Ð4Š3Ð4ˆ3Ñ4©tÓ4Ó4Ð4rW   c                ó2   € V ^8„  d   QhR\         R,          /# )ro   r  Nr±  )rr   s   "rU   rs   rs   V  s   € ÷ :-ñ :-ô ��:ñ:-rW   c           
      ó   a€ \        WW#V4      w  r#p\        W#V.4      '       g   R # \        W4      w  rÍV3R lpR pV! W$W0WV4      p\        VVVVSVV	R7      pV
e   VP	                  W«R7      # VP                  4       # )Nc                 óX  <€ \        4       pVP                  R V 4       VP                  RV4       VP                  RV4       VP                  RVP                  P                  RR7       VP                  RVP                  P                  RR7       \	        VRV4       \	        VRV4       VP                  R	VP                  P                  4       RR7       VP                  R
VP                  P                  4       RR7       SR8X  d    \        WsP                  P                  4       V# )r±   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ÚcontextrC   s	   &&&&&&& €rU   Úget_contextÚ%mm_autoheuristic.<locals>.get_contexti  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¯K©K×,=Ñ,=Ô>ØˆrW   c                  ó   € R # rS   rq   rq   rW   rU   ÚfallbackÚ"mm_autoheuristic.<locals>.fallback|  s   € ÙrW   )r   r9  rÆ   r  rC   Úaugment_contextÚprecondition)r  )Úget_size_hintsr  Úget_size_hints_stridesr   Úget_top_k_choices_callerÚget_choice_caller)rj   rk   r±   r²   r³   r9  rC   rÆ   r   r#  r  r  r  r  r  r   r  Úautoheuristics   &&&&&&f&&&&&      rU   rC  rC  V  s¥   ø€ ô ˜T¨¨qÓ1�G€Aˆ!Ü˜˜q˜	×"Ò"ÙÜ5°dÓAÑ€Kõò&ñ ˜! ¨¸KÓH€GÜ0ØØØØØØØ!ô€Mð Òà×5Ñ5Øð 6ó 
ð 	
ð ×*Ñ*Ó,Ð,rW   c                 óž  € \        V\        4      '       d   \        V\        4      '       g:   \        P                  P                  P                  V P                  4       4      w  r$\        V\        4      '       d   \        V\        4      '       g:   \        P                  P                  P                  VP                  4       4      w  rCW#V3# rS   )r  rÊ   r   r   r€   Úoptimization_hintsr™   )rj   rk   r±   r²   r³   s   &&&&&rU   r$  r$  “  s�   € Ü�aœ×Ò¤Z°´3×%7Ò%7Ü—‘×!Ñ!×4Ñ4°T·]±]³_ÓE‰ˆä�aœ×Ò¤Z°´3×%7Ò%7Ü—‘×!Ñ!×4Ñ4°T·]±]³_ÓE‰ˆØ�ˆ7€NrW   c                 ó:  € V P                   P                  pVP                   P                  pW#.p. pV FS  p\        V\        4      '       g*   \        P
                  P                  P                  V4      pVP                  V4       KU  	  V^ ,          V^,          3# rˆ   )	rÇ   rf   r  rÊ   r   r   r€   r*  r-  )rj   rk   r  r  ÚstridesÚstrides_hintsrf   s   &&     rU   r%  r%  œ  sƒ   € Ø—+‘+×$Ñ$€KØ—+‘+×$Ñ$€KØÐ(€GØ€MÛˆÜ˜&¤#×&Ò&Ü—W‘W×%Ñ%×8Ñ8¸Ó@ˆFØ×Ñ˜VÖ$ñ ð ˜Õ˜]¨1Õ-Ð-Ð-rW   rS   )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   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   rN  r    r!   r"   r#   r$   r%   r&   Úselect_algorithmr'   r(   r)   r*   r+   Úutilsr,   r-   r.   r/   r0   r1   r2   r3   r4   r5   r6   r7   r8   Ú	mm_commonr9   r:   r;   r<   r=   r>   r?   r$  Ú__version__Útriton_versionÚ
has_tritonÚImportErrorÚ	getLoggerrß   r)  r!  Úprimsr3  r4  r1  r5  r6  rò  rô  r2  ÚcacherV   r@   rb   r+  Ú	dtype_outr,  rh   ro  Ú_int_mmr`  Ú_sparse_semi_structured_mmÚdefaultr�  Ú
_scaled_mmrï  ra   rl   rš   rr  r¢   rº   r¼   r/  rè   ró   rõ   r7  rs  rX  rb  ry  r‡  r«  r¬  r­  r®  rº  rñ  ró  r’  r—  rŸ  r¢  r¯  r´  r½  Ú_scaled_mm_v2rî  rþ  r  r  r  rC  r$  r%  rq   rW   rU   Ú<module>rP     sk  ðã Û Ý ã Ý (Ý T÷ó õ FÝ Fß .Ý 6Ý +Ý ,Ý /ç MÑ Mß XÝ MÝ Dß Eß 8Ó 8Ý *÷÷ ÷õ ÷÷ ÷ õ ÷÷ ñ ðÛá! &×"4Ñ"4Ó5€NØ€Jð
 ×Ò˜Ó!€Ø‡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ðI tð I°1ð I¸1ô Iõñ4 % Z°Ó6€õò#ô Ð 0ô  ñF !:Ó ;Ð ôÐ)ô òD'ò/ñ #5Ø�_ mó#Ð ñ &8ØÐ*Ð,<ó&Ð "ñ
 �4—7‘7°Ô5ñz°4ô zó 6ðzñz �4—<‘<°TÔ:ð. tó .ó ;ð.ñb �4—:‘:°4Ô8ðg¨!ð g°!ð g¸Dó gó 9ðgñT �4×2Ñ2ÈÔMð/Ø(,ð/Ø59ó/ó Nð/ðf ×Ò˜[×3Ò3Ð4Ø×Ò˜+×-Ò-Ð.Ø×Ò ×!=Ò!=Ð>Ø×Ò ×!;Ò!;Ð<Ø×!Ò! ;×#=Ò#=Ð>ð€ð &×0Ò0°+×2EÒ2EÐFÐ Ø&×5Ò5°{×7SÒ7SÐTÐ õõ@õ
	õV÷Mõ*	õñ* )Ø×Ò×Ñ°EôÐ ñ
 �4×%Ò%×-Ñ-À4ÔHöYó IðYñx �4—?‘?×*Ñ*ÀÔEóró Fðrðj ‡�ôó ðò
5÷:-òzô	.øða* ô Ù! 'Ó*€NØƒJðús   ÃT  Ô T5Ô4T5