ó
    Eñic”  ã                   óî	  • S SK r S SKrS SKJrJrJr  S SKrS SKJr  S SK	J
r
  S SKJrJrJrJr  S SKJr  S SKJr  S SKJrJr  S S	KJr  S S
KJr  S SKJr  SSKJrJ r   SSK!J"r"J#r#  SSK$J%r%  SSK&J'r'  SSK(J)r)J*r*  SSK+J,r,J-r-J.r.J/r/  SSK0J1r1  SSK2J3r3J4r4J5r5J6r6J7r7  SSK8J9r9J:r:J;r;J<r<J=r=  SSK>J?r?J@r@JArAJBrBJCrCJDrDJErEJFrFJGrGJHrHJIrIJJrJJKrK  SSKLJMrMJNrNJOrOJPrPJQrQJRrR   S SKSrS\" \SR¨                  5      rUSrV\R°                  " \Y5      rZ\R*                  R¶                  r[\R*                  R¸                  r\\=" S\P\Rº                  R¼                  b  \US:¼  a  \N" S5      O\N" S5      SSS 9r_\=" S!\Q\N" S"5      S#9r`\=" S$\Q\N" S%5      S#9ra\=" S&\Q\N" S'5      S#9rb\=" S(\Q\N" S)5      S#9rc\ RÈ                  S* 5       re\:" \RÌ                  S+\[RÌ                  RÎ                  S,9rh\:" \RÌ                  S-S.\[RÌ                  RÒ                  S/9rj\:" \RÖ                  S0\[RÖ                  RÎ                  S,9rl\:" \RÚ                  S1\[RÚ                  RÎ                  S,9rn\:" \RÞ                  S2S\[RÞ                  Rà                  S39rq\:" \Rä                  S4\[Rä                  RÎ                  S,9rsS5 rtSSSS6.S7 jruSgS9 jrv\:" \uS5      rwS: rx " S; S<\*5      ry\y" 5       rz " S= S>\*5      r{S? r|S@ r}\{" SASB\|5      r~\{" SCSD\}5      r\6" \[RÌ                  SSE9ShSSF.SG jj5       r€\6" \[RÚ                  SSE9SSF.SH j5       r�\6" \[RÖ                  SSE9SSSSI.SJ j5       r‚\6" \[RÞ                  SSE9SSSK.SL j5       rƒ\GR                  \GR                  4\GR
                  \GR
                  4\GR                  \GR                  4\GR                  \GR                  4\GR                  \GR                  4/rˆ\GR                  \GR
                  /r‰\GR                  \GR                  /rŠSM\S8\‹4SN jrŒSM\SO\‹S8\‹4SP jr�SM\SQ\SR\ŽSO\‹S8\‹4
SS jr�SM\SQ\S8\‹4ST jr� SiSU\SV\GR"                  SW\SO\‹S8\‹4
SX jjr’S8\Ž4SY jr“SZ\S[\S\\GR"                  S]\GR"                  S8\”\\4   4
S^ jr•\6" \[Rä                  Rà                  SSE9     SjS_ j5       r–\ RÈ                  S`\\Ž   S8\‹4Sa j5       r—Sb r˜  SkSc\\Ž   4Sd jjr™Se ršSf r›g! \W a    \" S5      rUSrV GNf = f)lé    N)ÚAnyÚOptionalÚUnion)Úcounters)ÚAutoHeuristicSelectAlgorithm)Ú	AHContextÚcontext_add_stridesÚcontext_add_using_tf32Úmm_operations)ÚCppGemmTemplate)Úgen_best_config)ÚopsÚV)Úmake_fx)ÚScalingType)ÚTorchVersioné   )Ú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   Ú%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                 ó   • [        U 5      $ ©N)r'   )Úfns    ÚV/home/mande/repo/quber/.venv/lib/python3.13/site-packages/torch/_inductor/kernel/mm.pyÚlazy_register_extern_choicerS   }   s   € ä˜bÓ!Ð!ó    z
at::mm_out)Úop_overloadzat::mm_dtype_outÚmm_dtype)rB   rU   zat::addmm_outzat::_int_mm_outzat::_sparse_semi_structured_mm)Úhas_out_variantrU   zat::_scaled_mm_outc                 ód   • U R                  5       [        R                  [        R                  4;   $ rP   )Ú	get_dtypeÚtorchÚint8Úuint8)Úmats    rR   Ú_is_int8_matr^   ž   s    € Ø�=‰=‹?œuŸz™z¬5¯;©;Ð7Ñ7Ð7rT   ©ÚoutÚalphaÚbetac          	      óà   • U R                  S5      S:X  a  U R                  S5      S:w  d  U R                  S5      S:X  a  [        R                  " U S   XX4US9$ [        R                  " XX#XES9$ )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   r_   )ÚstrideÚsizerZ   Úaddmm)ÚinpÚmat1Úmat2r`   ra   rb   s         rR   Ú
bias_addmmrj   ¢   s^   € ð 	�
‰
�1‹˜Ó˜sŸx™x¨›{¨aÓ/°C·H±H¸Q³KÀ1Ó4DÜ�{Š{˜3˜q™6 4°3È$ÑOÐOÜ�;Š;�s $°uÑHÐHrT   Úreturnc                 óz  ^ ^• S[         4S jnS[         4S jnS[         4S jn[        R                  " U" T R                  5       5      =(       d    U" T R	                  5       5      U 4S j5        [        R                  " U" TR                  5       5      =(       d    U" TR	                  5       5      U4S j5        g )Nrk   c                 ó\   • [         R                  R                  R                  U S   S5      $ )Nr8   ©r   ÚgraphÚsizevarsÚstatically_known_equals©rd   s    rR   Úis_row_majorÚ.check_supported_striding.<locals>.is_row_major®   ó#   € Ü�w‰w×Ñ×7Ñ7¸¸q¹	À1ÓEÐErT   c                 ó\   • [         R                  R                  R                  U S   S5      $ ©Nr   r8   rn   rr   s    rR   Úis_col_majorÚ.check_supported_striding.<locals>.is_col_major±   ru   rT   c                 óÖ   • [        [        R                  R                  R	                  U S   S5      =(       d-    [        R                  R                  R	                  U S   S5      5      $ rw   )Úboolr   ro   rp   rq   )re   s    rR   Úhas_zero_dimÚ.check_supported_striding.<locals>.has_zero_dim´   sQ   € ÜÜ�G‰G×Ñ×4Ñ4°T¸!±W¸aÓ@÷ DÜ�w‰w×Ñ×7Ñ7¸¸Q¹ÀÓCó
ð 	
rT   c                  ó*   >• ST R                  5        3$ )Nz$mat_a must be row_major, got stride ©Ú
get_stride)Úmat_as   €rR   Ú<lambda>Ú*check_supported_striding.<locals>.<lambda>½   ó   ø€ Ð6°u×7GÑ7GÓ7IÐ6JÑKrT   c                  ó*   >• ST R                  5        3$ )Nz$mat_b must be col_major, got stride r   )Úmat_bs   €rR   r‚   rƒ   Ã   r„   rT   )r{   rZ   Ú_checkr€   Úget_size)r�   r†   rs   rx   r|   s   ``   rR   Ú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õrT   c                 ó„  • U R                   S   nUR                   S   nU R                   S   nXR-  nUn[        R                  " U R                  X7U5      S5      nUR                  XvU5      n	[        R                  " X‰[        R
                  S9n
[        R                  " U
S5      nUR                  U R                  5      $ )Nr   r8   )r8   r   r   ©Ú	out_dtype)	ÚshaperZ   ÚpermuteÚreshapeÚbmmÚfloat32ÚsumÚtoÚdtype)ÚaÚbÚk_splitsÚmÚnÚkÚk_partsÚBÚ
a_reshapedÚ
b_reshapedÚresultÚreduced_bufs               rR   Ú
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Ø�>‰>˜!Ÿ'™'Ó"Ð"rT   c                   óN   ^ • \ rS rSrU 4S jrS\\   S\S\S\	4U 4S jjr
SrU =r$ )	ÚDecomposeKSugraphTemplateéØ   c                 ó    >• [         TU ]  SS9  g )NÚdecompose_k©rB   )ÚsuperÚ__init__)ÚselfÚ	__class__s    €rR   r©   Ú"DecomposeKSugraphTemplate.__init__Ù   s   ø€ Ü‰ÑØð 	ò 	
rT   Úinput_nodesÚlayoutÚk_splitrk   c           	      óî   >• SSK Jn  SSKJn  SU S3nSU< 3nU" 5          U" 5       n[	        [
        R                  " [        US9U5      n	[        T
U ]%  UUUU	US	9sS S S 5        $ ! , (       d  f       g = f)
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¹   ÚdecompositionsrQ   r«   s             €rR   r¾   Ú"DecomposeKSugraphTemplate.generateÞ   s   ø€ õ 	Då7à   	¨Ð0ˆØ!˜™
�mˆá%Õ'Ù0Ó2ˆNÜÜ×!Ò!¤*°wÑ?ØóˆBô
 ‘7Ñ#ØØ'ØØ Ø'ð $ð ÷ (×'×'ús   ¡;A&Á&
A4© )Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__r©   Úlistr   r   Úintr   r¾   Ú__static_attributes__Ú__classcell__©r«   s   @rR   r£   r£   Ø   s<   ø† õ
ð
à˜&‘\ðð ðð ð	ð
 
÷õ rT   r£   c                   óZ   ^ • \ rS rSrS\S\S\4U 4S jjrS\\   S\	S\
4U 4S	 jjrS
rU =r$ )ÚContiguousTemplateéþ   rB   r¹   rQ   c                 óD   >• Xl         X l        X0l        [        TU ]  US9  g )Nr§   )rB   r¹   rQ   r¨   r©   )rª   rB   r¹   rQ   r«   s       €rR   r©   ÚContiguousTemplate.__init__ÿ   s(   ø€ ØŒ	Ø&ÔØŒÜ‰ÑØð 	ò 	
rT   r­   r®   rk   c           	      óä   >• SSK Jn  SSKJn  U" 5          U" 5       n[	        U R
                  U5      n[        TU ]  U R                  UUUU R                  S9sS S S 5        $ ! , (       d  f       g = f)Nr   r±   r   r³   r·   )
rº   r²   r»   r´   r   rQ   r¨   r¾   rB   r¹   )rª   r­   r®   r²   r´   r¿   rQ   r«   s          €rR   r¾   ÚContiguousTemplate.generate  sg   ø€ õ
 	Då7á%Õ'Ù0Ó2ˆNÜØ—‘ØóˆBô
 ‘7Ñ#Ø—Y‘YØ'ØØ Ø ×,Ñ,ð $ð ÷ (×'×'ús   •AA!Á!
A/)r¹   rQ   rB   )rÂ   rÃ   rÄ   rÅ   Ústrr   r©   rÆ   r   r   r   r¾   rÈ   rÉ   rÊ   s   @rR   rÌ   rÌ   þ   sG   ø† ð
˜Sð 
¨sð 
¸÷ 
ðà˜&‘\ðð ðð 
÷	õ rT   rÌ   c                 óJ   • [         R                  " XR                  5       5      $ rP   )rZ   r?   Ú
contiguous)r•   r–   s     rR   Úcontiguous_mmrÕ      s   € Ü�8Š8�A—|‘|“~Ó&Ð&rT   c                 óL   • [         R                  " XUR                  5       5      $ rP   )rZ   rf   rÔ   )rg   r•   r–   s      rR   Úcontiguous_addmmr×   $  s   € Ü�;Š;�s˜qŸ|™|›~Ó.Ð.rT   rÕ   zcontiguous mmr×   zcontiguous addmm)Útype_promotion_kind©r®   c                ó¸  ^ • UbÊ  T R                  5       n[        R                  " UR                  5       U:H  S 5        [        R                  " T R                  5       R                  S;   S 5        [        R                  " X$:H  =(       d=    U[        R
                  :H  =(       a#    U[        R                  [        R                  4;   S 5        [        T U5      (       añ  [        [        R                     " T S5      m [        [        R                     " US5      n[        T U/0 SSS	S
9u  pV[        R                  R                  (       aU  T R                   [        R                  [        R                  4;   a'  U 4S jnU Vs/ s H  n[#        U5      " U5      PM     nn[#        [$        R&                  5      " U6 n	[)        S5      " U	S5      n
U
$ [+        T XUS9u  p¼pÓm n[-        U5      u  pïSn[/        T U/US9n[0        S   SU SU SU 3==   S-  ss'   [2        R5                  SUUUT R                  5       UR                  5       U5        / n[-        U5      u  pï[6        n0 nUb
  [8        nSU0n/ n0 n[;        5       (       a'  UR=                  U5        U(       a  UUUR>                  '   UGc"  U(       Ga  [A        USS9(       Ga
  [C        X¼U5      (       a  UR=                  [D        5        [        RF                  S:H  nU(       d  [C        X¼USS9(       d¥  UR=                  [H        5        [K        T XS9(       a  UR=                  [L        5        O$[O        T XS9(       a  UR=                  [P        5        [        RR                  " 5       (       a-  [        R                  RT                  (       a  SSK+J,n  U" U5      nUR=                  [Z        5        UR]                  [^        R`                  Rc                  UUSUS95        UcN  U(       aG  [e        X;XÍ5      (       a6  [g        S5      (       a&  [h        Rj                  " UUURm                  5       5        Uc>  U(       a7  [o        X;XÍ5      (       a&  [p        Rr                  " UUURm                  5       5        Uc>  U(       a7  [u        X;XÍ5      (       a&  [v        Rx                  " UUURm                  5       5        Uc*  U(       a#  [{        X;XÍT U5      (       a  SSK>J?n  U" UUU5        Uc8  [�        UT U5      (       a&  [‚        Rx                  " UUURm                  5       5        T U/nUGc5  U(       Ga-  [A        U5      (       Ga  [        R„                  R†                  R‰                  U5      (       aî  [‹        T 5      (       aÞ  / n[;        5       (       a  UR=                  S5        [�        U5      nUR]                  [^        R`                  Rc                  U[H        /S5      5        [�        T UUUUUUU[‘        5       SS US!9n[        R„                  R†                  R“                  U5      (       d2  Ub*  [�        U5      S:”  a  U Vs/ s H  nUU;   d  M  UPM     nnOUSU nUcO  [        R”                   H;  nUR=                  [—        U5      R™                  URm                  5       U5      5        M=     SnUc5  [        R„                  R†                  Rš                  (       a  [�        T U5      n[ž        R                   " UUURm                  5       U5      =n (       a  U $ [£        UUURm                  5       UUS"9$ s  snf s  snf )#zW
Lowering for autotuning aten.mm with different backends (Aten, Triton, CUTLASS, etc.)
Nc                  ó   • g)Nzinput dtypes must be the samerÁ   rÁ   rT   rR   r‚   Útuned_mm.<locals>.<lambda>9  s   € Ð3rT   )ÚcudaÚxpuc                  ó   • g)Nz+out_dtype is only supported for CUDA or XPUrÁ   rÁ   rT   rR   r‚   rÜ   =  s   € ÐArT   c                  ó   • g)NzFout_dtype must be the same as input dtype or fp32 for fp16/bf16 inputsrÁ   rÁ   rT   rR   r‚   rÜ   E  s   € Ð\rT   éÿÿÿÿr   TF)ÚargsÚkwargsÚ	broadcastrØ   Úconvert_input_to_boolc                 óD   >• [         R                  " U TR                  SS9$ )NF)Úuse_compute_types)r   Úto_dtyper”   )Úxrh   s    €rR   Ú	_to_dtypeÚtuned_mm.<locals>._to_dtypef  s   ø€ Ü—|’| A t§z¡zÀUÑKÐKrT   Ú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)Ú
append_tlx©Úkwarg_overrides)Úadd_nv_universal_gemm_choicesÚ	extern_mmzmm-ahé
   )Útop_kÚalways_included)Úbest_config_future)RrY   rZ   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   Úpersistent_tma_mm_templateÚ	is_fbcodeÚenable_tlx_templatesÚ-torch._inductor.fb.tlx_templates.mm_templatesrö   Ú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_matmulrS   ÚbindÚremote_gemm_autotune_cacher   r   Úmaybe_autotune_remoter&   )!rh   ri   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­   rý   Ú num_choices_before_extra_configsÚ
ah_choicesÚchoicerþ   Úboxs!   `                                rR   Útuned_mmr6  0  sx  ø€ ð
 ÑØ—n‘nÓ&ˆÜ�ŠØ�N‰NÓ Ñ+Ù3ô	
ô 	�ŠØ�O‰OÓ×"Ñ" oÑ5ÙAô	
ô 	�ŠØÑ$÷ àœUŸ]™]Ñ*÷ CØ¤E§M¡M´5·>±>Ð#BÑBá\ô	
ô. ˜˜t×$Ñ$ÜœŸ™Ò(¨¨rÓ2ˆÜœŸ™Ò(¨¨qÓ1ˆÜ%Ø˜�ØØØ $Ø"'ñ
‰ˆô ×!Ñ!×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ô  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ÐàHJÐØ13€OÜ×ÑØ×Ñ Ô-ÞØ0AˆO˜L×,Ñ,Ñ-ð 	ÒßÜ ¸4×@Ð@ä! !¨×*Ñ*Ø×#Ñ#Ô$AÔBô (×FÑFÈ,ÑVˆÞÔ 6°q¸QÐST× UØ×#Ñ#¤KÔ0ä0°°t×RØ ×'Ñ'Ô(VÕWÜ(¨¨t×JØ ×'Ñ'Ô(BÔCô  ×)Ò)×+Ñ+Ü#×*Ñ*×?×?åTá#-Ð.>Ó#?Ð à×ÑÔ ?Ô@à‡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àÒßÜ ×'Ò'Ü�O‰O×"Ñ"×4Ñ4°T×:Ñ:Ü�d�O‰OàˆÜ ×"Ñ"Ø×"Ñ" ;Ô/Ü+.¨w«<Ð(Ø�‰Ü�I‰I×*Ñ*ð Ü�Øóô		
ô &ØØØØØØØØÜ‹OØØØ+ñ
ˆ
ô �‰×%Ñ%×;Ñ;¸D×AÑAàÑ%¬#¨j«/¸AÓ*=ñ
 18ÓP² f¸6ÀZÑ;OŸ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ô -¨T°4Ó8Ðä"×8Ò8Øˆg�}×*Ñ*Ó,¨fóð €sõ ð ˆ
ä$ØØØ×ÑÓØØ-ñð ùòm @ùòB Qs   Æ]Ù
]Ù-]c          	      ó$  • [        XU[        R                  S9u  p4pRpSn[        S   SU SU SU 3==   S-  ss'   [        R                  SUUUU R                  5       UR                  5       U5        [        U5      u  pxU=(       a    U=(       a    [        X#XE5      n	/ n
[        X/[        R                  S9n/ n[        5       (       a  UR                  [        5        U(       a%  [        US	S
S9(       a  UR                  [        5        U
R                  [         R"                  R%                  X¼U5      5        U	(       a5  ['        U5      (       a%  [(        R*                  " X¢UR-                  5       S	S	S9  [/        XjUR-                  5       U5      $ )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;   rZ   Úint32r   r  r	  rY   r9   r1   r    r-   r  Úaten__int_mmr6   r  r  r   r  r  r+   r   r  r  r&   )rh   ri   r®   r˜   r™   rš   rB   r+  r,  Úuse_cutlassr  r-  r0  s                rR   Útuned_int_mmr@  (  sn  € ô #*Ø˜6¬U¯[©[ñ#Ñ€Aˆ!�Tð €Däˆ^Ñ˜}¨Q¨C¨q°°°1°Q°CÐ8Ó9¸QÑ>Ó9Ü‡H�HØ^Ø	Ø	Ø	Ø�‰ÓØ�‰ÓØôô  2°&Ó9Ñ€LØ×W :×WÔ2FÀvÐRSÓ2W€KØ"$€Gô # D <¼5¿;¹;ÑG€Mð IKÐÜ×ÑØ×Ñ¤Ô-æÔ)Ø˜T°e÷ð 	×Ñ¤Ô,ð ‡N�NÜ	�	‰	×&Ñ& }ÈÓMôö Ô*¨4×0Ñ0Ü×6Ò6Ø˜]×0Ñ0Ó2¸TÐPTò	
ô % T°M×4GÑ4GÓ4IÈ6ÓRÐRrT   )ra   rb   r®   c          	      óø  • [        X5      (       a…  US:X  a  SnO[        [        R                     " X@5      nUS:X  a  SnO9[        [        R                     " U[        [        R                     " X5      5      n[        [        R
                     " Xg5      $ [        XXS9u  p‰p¥pn[        U5      u  pÍSn[        X±U/[        X4S9S9n/ n[        S   SU SU	 SU
 3==   S	-  ss'   [        R                  S
UU	U
UR                  5       UR                  5       U5        U(       a*  [        R                  (       dy  [        R                   (       dd  [        XU/[        X4S9S9nUR#                  [$        R&                  R)                  U[*        /U5      5        [-        UUUR/                  5       U5      $ / n[1        5       (       a  UR#                  [2        [*        /5        U(       a‚  [5        USS9(       as  UR7                  [8        5        [;        XUS9(       a  UR7                  [<        5        O$[?        XUS9(       a  UR7                  [@        5        UR7                  [B        5        UR#                  [$        R&                  R)                  UUU5      5        U(       aH  [E        XXXš5      (       a7  [G        U5      (       a'  [H        RJ                  " UUUR/                  / SQS9UUS9  U(       a;  [M        XXXš5      (       a*  [N        RP                  " UUUR/                  / SQS9UU/ SQS9  [S        XQU5      (       a'  [T        RV                  " UUUR/                  5       UUSS9  [-        UUUR/                  5       U5      $ )zZ
Lowering for autotuning aten.addmm with different backends (Aten, Triton, CUTLASS, etc.)
r   rÙ   rf   )ra   rb   )Úscalarsrî   zaten.addmm_rï   r8   zRTuned aten.addmm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sFrð   rô   )r8   r   r   )Úreorder)r   r   r8   )ra   rb   Úinput_reorderT)ra   rb   Úhas_bias),r>   r!   r  Úmulr?   Úaddr;   r9   r    Údictr   r  r	  rY   r  Úmax_autotuneÚmax_autotune_gemmr  r   r  r  Ú
aten_addmmr&   r  r-   Úaten_bias_addmmr6   r  r  r4   r  r7   r  Ú"addmm_contiguous_subgraph_templater1   r+   r   r  r.   r   r  r0   r   r  )rg   rh   ri   ra   rb   r®   Úarg1Úarg2r˜   r™   rš   Úinp_expandedr+  r,  rB   r-  r  r0  s                     rR   Útuned_addmmrQ  Y  s  € ô
 ˜×$Ñ$Ø�1‹9Ø‰DäœTŸX™XÒ& tÓ1ˆDà�A‹:Ø‰DäœTŸX™XÒ& u¬i¼¿¹Ò.@ÀÓ.LÓMˆDäœŸ™Ò" 4Ó.Ð.ô 18¸ÀCÑ0WÑ-€Aˆ!�T Ü1°&Ó9Ñ€LØ€Dô #Ø	˜TÐ"¬D°uÑ,Hñ€Mð #%€Gô ˆ^Ñ˜{¨1¨#¨Q¨q¨c°°1°#Ð6Ó7¸1Ñ<Ó7Ü‡H�HØ\Ø	Ø	Ø	Ø�‰ÓØ�‰ÓØôö Ü×)×)¬_×-N×-Nô
 'Ø˜Ð¤t°%Ñ'Cñ
ˆð 	�‰Ü�I‰I×*Ñ*ØÜ�Øóô	
ô )¨¨w¸×8KÑ8KÓ8MÈvÓVÐVð IKÐÜ×ÑØ×Ñ¤´*Ð =Ô>æÔ)¨&ÀU×KØ×Ñ¤Ô,ä,¨TÀv×NØ×#Ñ#Ô$RÕSÜ$ T¸v×FØ×#Ñ#Ô$>Ô?à×ÑÔ BÔCð ‡N�NÜ	�	‰	×&Ñ& }Ð6FÈÓMôö
 	Ü  ¨A×1Ñ1Ü ×%Ñ%ä×6Ò6ØØð ×Ñª	ÐÐ2ØØò	
ö Ô*¨6°a×;Ñ;Ü×*Ò*ØØð ×Ñª	ÐÐ2ØØÚ#ò		
ô ˜V¨4×0Ñ0Ü×#Ò#ØØØ×ÑÓ!ØØØò	
ô % T¨7°M×4GÑ4GÓ4IÈ6ÓRÐRrT   )rŒ   r®   c                ó¾  • SSK Jn  U" XU5      u  pnU R                  5       u  pgUR                  5       u  p‰UR                  5       u  p«[        R                  R
                  R                  Xh5      n[        R                  R
                  R                  SU-  U
5      nUc:  SSKJn  U" UR                  5       U(       a  UOUR                  5       XË/US/5      nO
Ub   S5       e[        5       (       a  [        R                  XU4XCS9/O/ nXË-  S:w  a:  [        XLX½5      (       a)  [        S5      (       a  [         R"                  " XôXU/S	S	S
9  [%        SXðX4U5      $ )Nr   )r)   r   )ÚFixedLayoutr8   z,out_dtype is ignored if layout is specified.r‹   Úsparse_semi_structured_mmTr:  )Ú torch._inductor.select_algorithmr)   rˆ   r   ro   rp   Úcheck_equals_and_simplifyÚtorch._inductor.irrS  rÿ   rY   r-   Úaten__sparse_semi_structured_mmr%  r1   r+   r   r  r&   )rh   Ú	mat1_metari   rŒ   r®   r)   Úm1Úk1Úm2rï   Úk2r™   r˜   rš   rS  r  s                   rR   Útuned_sparse_semi_structured_mmr^  Ï  s^  € õ @ñ +¨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ó	
‰ð Ñ ÐPÐ"PÓPÐ ô !×"Ñ"ô	 ,×0Ñ0Ø $Ð'¨ð 1ð ñ	
ð ð ð 	
‰�‹
Ü  ¨A×1Ñ1ÜÐ ;×<Ñ<ä×6Ò6Ø˜d¨)Ð4¸tÐRVò	
ô %Ø# W°YÐ.EÀvóð rT   Úszc                 óP   • [        U 5      S:H  =(       d    [        S U  5       5      $ )Nr   c              3   óv   #   • U  H/  n[         R                  R                  R                  US 5      v •  M1     g7f)r8   Nrn   )Ú.0Úds     rR   Ú	<genexpr>Ú)_is_tensorwise_scaling.<locals>.<genexpr>  s-   é € ð !Ú@B¸1Œ�‰×Ñ×0Ñ0°°A×6Ð6Âùs   ‚79)r!  Úall)r_  s    rR   Ú_is_tensorwise_scalingrg    s+   € Ü�‹G�q‰L÷ œSñ !Ù@Bó!ó ð rT   Ú	transposec                 óp   • U(       a  SOSn[         R                  R                  R                  X   S5      $ )Nr   rá   r8   rn   )r_  rh  Úidxs      rR   Ú_is_rowwise_scalingrk    s*   € Þ‰!˜b€CÜ�7‰7×Ñ×3Ñ3°B±G¸QÓ?Ð?rT   Ú	tensor_szÚ	tile_sizec                 ó  • U(       a  SOSnU(       a  SOSn[         R                  R                  R                  X   X   5      =(       a8    [         R                  R                  R                  X   [	        X   U5      5      $ )Nr8   r   ©r   ro   rp   rq   r,   )r_  rl  rm  rh  ÚlhsÚrhss         rR   Ú_is_blockwise1xTILESIZE_scalingrr    si   € ö ‰!˜a€CÞ‰!˜a€CÜ�7‰7×Ñ×3Ñ3Ø
‰�‘ó÷ ä
�'‰'×
Ñ
×
2Ñ
2Ø
‰”˜™¨Ó3óðrT   c                 óø   • [         R                  R                  R                  U S   [	        US   S5      5      =(       a:    [         R                  R                  R                  U S   [	        US   S5      5      $ )Nr   é€   r8   ro  )r_  rl  s     rR   Ú_is_blockwise128x128_scalingru  $  sd   € Ü�7‰7×Ñ×3Ñ3Ø
ˆ1‰Œw�y ‘| SÓ)ó÷ Vä
�'‰'×
Ñ
×
2Ñ
2°2°a±5¼'À)ÈAÁ,ÐPSÓ:TÓ
UðVrT   ÚtÚ
scale_sizeÚscaling_typec                 óZ  • U=[         R                  :X  a    [        U5      $ =[         R                  :X  a    [	        X5      $ =[         R
                  :X  a    [        XR                  5       SU5      $ [         R                  :X  a  [        XR                  5       5      $  [        SU 35      e)Nrt  úUnsupported scaling type )r   Ú
TensorWiserg  ÚRowWiserk  ÚBlockWise1x128rr  rˆ   ÚBlockWise128x128ru  ÚAssertionError)rv  rw  rx  rh  s       rR   Úis_desired_scalingr€  *  sŠ   € ð Ø#Œ[×#Ö#Ü)¨*Ó5Ð5Ø Œ[× Ö Ü& zÓ=Ð=Ø'Œ[×'Ö'Ü2ØŸJ™J›L¨#¨yóð ô ×)Õ)Ü/°
¿J¹J»LÓIÐIØÜ Ð#<¸\¸NÐ!KÓLÐLrT   c                 óx   • U =[         R                  :X  a    g[         R                  :X  a  g [        SU  S35      e)Nrt  rz  z in get_tile_size)r   r~  r}  r  )Úscale_options    rR   Úget_tile_sizerƒ  ?  s<   € Ø
Ø)Œ[×)Ö)ØÜ×'Õ'ØØÜ Ø+¨L¨>Ð9JÐKóð rT   r�   r†   Úscale_a_sizeÚscale_b_sizec                 ó’   • [          H-  u  pE[        XU5      (       d  M  [        XUSS9(       d  M*  XE4s  $    [        SU SU 35      e)NT)rh  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         rR   Úget_scaling_optionsrŠ  K  s^   € ÷ +8Ñ&ˆÜØ ÷
ó 
ä  °nÐPT×UÑUØ!Ð1Ò1ñ	 +8ô Ø
;¸L¸>ÐI[Ð\hÐ[iÐjóð rT   c	           	      ó¼  • [        XX†S9u  pšp¸p[        S   SU	 SU
 SU 3==   S-  ss'   [        R                  SU	U
UU R	                  5       UR	                  5       U5        Sn[        X5        [        X#5      u  pÞU(       d  XXÞ/nO[        U5      nXXÞU/n[        USSUS	9n/ n/ n0 n[        5       (       a/  UR                  [        5        [        XgS
9U[        R                  '   [        U5      u  nnUR                  [        R                   :X  GaŸ  U(       Ga—  [#        USSS9(       Ga†  [        US9nUR$                  UR$                  nn['        XUU5      u  nn[)        XUS9(       aÊ  U(       dÃ  UR*                  US'   UR*                  US'   [-        UU[.        5      (       a)  UR                  [0        5        UU[0        R                  '   Of[-        UU[2        5      (       aE  [5        U5      US'   [5        U5      US'   UR                  [6        5        UU[6        R                  '   O[9        S5      e[;        XUS9(       a/  U(       d(  UR                  [<        5        UU[<        R                  '   [-        UU[.        5      (       a(  UR                  [>        5        UU[>        R                  '   URA                  [B        RD                  RG                  UUUUS95        U(       a!  [I        X‰X«X5      (       a  SSK%J&n  U" UUUUS9  UR                  [        R                   :w  a  [O        UUXø5      $ U(       aF  [Q        X‰X«5      (       a5  [S        U5      (       a%  [T        RV                  " UUURY                  5       US9  U(       a7  [[        X‰X«5      (       a&  [\        R^                  " UUURY                  5       5        [O        UUURY                  5       U5      $ )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ï   r8   z_Tuned aten._scaled_mm.default: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sÚ	scaled_mmr   )Úmat1_idxÚmat2_idxrŒ   )rŒ   Úuse_fast_accumTF)Úenable_float8rñ   )ÚUSE_FAST_ACCUMrô   ÚSCALE_RECIPE_AÚSCALE_RECIPE_BÚTILE_SIZE_AÚTILE_SIZE_BzpInductor 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�  )0r;   r   r  r	  rY   r‰   r)   r    r-   r  Úaten__fp8_mmrH  r  r9   r”   rZ   r‘   r6   r�   rŠ  r7   Úvaluer5   Úepilogue_scaling_typesÚ.scaled_mm_device_tma_epilogue_scaling_templateÚmain_loop_scaling_typesrƒ  Ú/scaled_mm_device_tma_main_loop_scaling_templater  r4   r  r  r  r   r  r  r3   r  r–  r&   r1   r+   r   r  r  r.   r   r  )r�   r†   Úscale_aÚscale_bÚbiasÚscale_resultrŒ   r�  r®   r˜   r™   rš   rB   Úscale_a_realÚscale_b_realr­   Ú	bias_realr-  r  r0  rø   rï   r,  Ú
overridersr„  r…  rˆ  r‰  r–  s                                rR   Ú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ð IKÐØ€Oä×ÑØ×Ñ¤Ô-Ü,0Øñ-
ˆœ×(Ñ(Ñ)ô ' vÓ.�M€A€zð 	�‰œŸ™Ô&ßÜ °dÈu×UÐUä¨Ñ8ˆ
à%1×%7Ñ%7¸×9KÑ9K�lˆä)<Ø˜,¨ó*
Ñ&ˆ˜ô # 5¸v×FÎtØ+9×+?Ñ+?ˆJÐ'Ñ(Ø+9×+?Ñ+?ˆJÐ'Ñ(ä*Ø Ô0F÷ñ ð !×'Ñ'Ô(VÔWàð  Ô N× RÑ RÒSô -Ø Ô0G÷ñ ô -:¸.Ó,I�
˜=Ñ)Ü,9¸.Ó,I�
˜=Ñ)à ×'Ñ'Ô(WÔXàð  Ô O× SÑ SÒTô %ðGóð ô .¨eÈ&×QÞà×#Ñ#Ô$RÔSàð ÔJ×NÑNÑOô 'Ø˜NÔ,B÷
ñ 
ð ×#Ñ#¤KÔ0Ø/9ˆOœKŸO™OÑ,ð ‡N�NÜ	�	‰	×&Ñ&ØØØØ+ð	 	'ð 	
ôö Ô4°VÀÀe×SÑSÝTá,ØØØØ'ò		
ð ‡}�}œŸ™Ó%Ü(¨¨w¸ÓLÐLö 	Ü  ¨A×1Ñ1Ü ×%Ñ%ä×6Ò6ØØØ×ÑÓ!Ø)ò		
ö Ô*¨6°a×;Ñ;Ü×*Ò*¨7°F¸M×<OÑ<OÓ<QÔRä$ T¨7°M×4GÑ4GÓ4IÈ6ÓRÐRrT   Úindexc                 óp   • [         R                  R                  U =(       d    S5      nUR                  S:*  $ )Nr   é   )rZ   rÝ   Úget_device_propertiesÚmajor)r¦  Úpropss     rR   Ú_is_sm7x_or_older_gpur¬    s)   € ä�J‰J×,Ñ,¨U¯Z°aÓ8€EØ�;‰;˜!ÑÐrT   c                 ó&   • [        S U  5       5      $ )Nc              3   óB   #   • U  H  n[        U[        5      v •  M     g 7frP   )Ú
isinstancerÇ   )rb  Údims     rR   rd  Údims_are_int.<locals>.<genexpr>  s   é € Ð4ªt¨Œz˜#œs×#Ð#ªtùs   ‚)rf  )Údimss    rR   Údims_are_intr³    s   € ÜÑ4©tÓ4Ó4Ð4rT   rü   c           
      óî   ^• [        XX#U5      u  p#n[        X#U/5      (       d  g [        X5      u  pÍU4S jnS nU" X$X0XU5      n[        UUUUTUU	S9nU
b  UR	                  X«S9$ UR                  5       $ )Nc                 ó>  >• [        5       nUR                  SU 5        UR                  SU5        UR                  SU5        UR                  SUR                  R                  SS9  UR                  SUR                  R                  SS9  [	        USU5        [	        US	U5        UR                  S
UR                  R                  5       SS9  UR                  SUR                  R                  5       SS9  TS:X  a  [        XsR                  R                  5        U$ )Nr˜   rš   r™   Ú
mat1_dtypeT)Úis_categoricalÚ
mat2_dtyperh   ri   Úmat1_iscontigÚmat2_iscontigr?   )r   Úadd_featurer®   r”   r	   Úis_contiguousr
   )	r˜   rš   r™   rh   ri   Úmat1_strideÚmat2_strideÚcontextrB   s	           €rR   Úget_contextÚ%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¯K©K×,=Ñ,=Ô>ØˆrT   c                  ó   • g rP   rÁ   rÁ   rT   rR   ÚfallbackÚ"mm_autoheuristic.<locals>.fallback;  s   € ØrT   )rÃ  r  r­   r¿  rB   Úaugment_contextÚprecondition)rý   )Úget_size_hintsr³  Úget_size_hints_stridesr   Úget_top_k_choices_callerÚget_choice_caller)rh   ri   r˜   r™   rš   r  rB   r­   r   rÆ  rü   rý   r½  r¾  rÀ  rÃ  r¿  Úautoheuristics         `           rR   r"  r"    s¥   ø€ ô ˜T¨¨qÓ1�G€Aˆ!Ü˜˜q˜	×"Ñ"ØÜ5°dÓAÑ€Kõò&ñ ˜! ¨¸KÓH€GÜ0ØØØØØØØ!ñ€Mð Ñà×5Ñ5Øð 6ð 
ð 	
ð ×*Ñ*Ó,Ð,rT   c                 ó–  • [        U[        5      (       a  [        U[        5      (       d9  [        R                  R                  R                  U R                  5       5      u  p$[        U[        5      (       a  [        U[        5      (       d9  [        R                  R                  R                  UR                  5       5      u  pCX#U4$ rP   )r¯  rÇ   r   ro   rp   Úoptimization_hintsrˆ   )rh   ri   r˜   r™   rš   s        rR   rÇ  rÇ  R  s�   € Ü�aœ×Ñ¤Z°´3×%7Ñ%7Ü—‘×!Ñ!×4Ñ4°T·]±]³_ÓE‰ˆä�aœ×Ñ¤Z°´3×%7Ñ%7Ü—‘×!Ñ!×4Ñ4°T·]±]³_ÓE‰ˆØ�ˆ7€NrT   c                 ó(  • U R                   R                  nUR                   R                  nX#/n/ nU HR  n[        U[        5      (       d)  [        R
                  R                  R                  U5      nUR                  U5        MT     US   US   4$ rw   )	r®   rd   r¯  rÇ   r   ro   rp   rÍ  r  )rh   ri   r½  r¾  ÚstridesÚstrides_hintsrd   s          rR   rÈ  rÈ  [  sƒ   € Ø—+‘+×$Ñ$€KØ—+‘+×$Ñ$€KØÐ(€GØ€MÛˆÜ˜&¤#×&Ñ&Ü—W‘W×%Ñ%×8Ñ8¸Ó@ˆFØ×Ñ˜VÖ$ñ ð ˜Ñ˜]¨1Ñ-Ð-Ð-rT   )rk   NrP   )F)NNNFN)NN)œr¼   ÚloggingÚtypingr   r   r   rZ   Ú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   Ú 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  ÚprimsÚversionÚhipr  r  rš  rœ  r  ÚcacherS   r?   r`   r
  Ú	dtype_outr  rf   rK  Ú_int_mmr>  Ú_sparse_semi_structured_mmÚdefaultrX  Ú
_scaled_mmr—  r^   rj   r‰   rL  r¡   r£   r  rÌ   rÕ   r×   r  rM  r6  r@  rQ  r^  r{  r|  r}  r~  r‡  r™  r›  r{   rg  rk  rÇ   rr  ru  ÚTensorr€  rƒ  ÚtuplerŠ  r¥  r¬  r³  r"  rÇ  rÈ  rÁ   rT   rR   Ú<module>rö     sò  ðã Û ß 'Ñ 'ã Ý (Ý T÷ó õ FÝ Fß .Ý 6Ý +Ý ,ç >ß XÝ MÝ Dß Eß 8Ó 8Ý *÷õ ÷õ ÷÷ ÷ õ ÷÷ ðÛá! &×"4Ñ"4Ó5€NØ€Jð
 ×Ò˜Ó!€Ø‡y�y‡~�~€Ø�	‰	�‰€ñ
 Ø	Ø	à�‰×ÑÑ! n¸Ó&?ñ   Ô,ñ
 
Ð.Ó	/Ø'+Ø"ñ€ñ ,Ø	Ø	ÙÐ :Ó;ñÐ ñ 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ðt°4õ tó 6ðtñn �4—<‘<°TÑ:Ø'+ô -Só ;ð-Sñ` �4—:‘:°4Ñ8Ø*+°!¸Dô rSó 9ðrSñj �4×2Ñ2ÈÑMà(,°Tô-ó Nð-ðb ×Ò˜[×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—?‘?×*Ñ*ÀÑEð 
ØØØØókSó FðkSð\ ‡�ð ¨#¡ð °4ó ó ðò
5ð  Øñ:-ð �C‰=õ:-òzó	.øðe  ó Ù! 'Ó*€NØƒJðús   ÃS  Ó S4Ó3S4