+
    "Ý³j//  ã                   óº  € ^ RI t ^ RIt^ RIt^ RIHt ^ RIHt ]! ] P                  ! . RO^R7       U UUu. uFS  w  rpWV^^3WV^^3WV^^3WV^^3WV^^3WV^^3WV^^3WV^^3WV^^3WV^^3WV^^3WV^^3WV^^3WV^^3WV^^3.NKU  	  uppp . 4      t
]P                  P                  P                  R8X  dG   ] P                  ! . RO^R7       UUUUUu. uF  w  r4pR F  pR F  pW4WVV3NK
  	  K  	  K  	  upppppt
]
 U UUUU	u. uF#  w  rr(p	]P                  RV RVRVR^/VV	R	7      NK%  	  up	pppp t
]P                  R
 R l4       t]P                  ]P"                  3R R ll4       tR tR t]P*                  P-                  RR4      t]P1                  R4       ]P1                  R4       ]P*                  P3                  ]RR4      R 4       t]P*                  P3                  ]RR4      R 4       t]P*                  P3                  ]RR4      R 4       t]P*                  P3                  ]RR4      R 4       tR# u uppp i u upppppi u up	pppp i ) é    N)Úget_best_config_fn)ÚrepeatÚ
EXHAUSTIVEÚBLOCK_MÚBLOCK_NÚBLOCK_KÚGROUP_M)Ú
num_stagesÚ	num_warpsc                ó˜   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  /# )é   r   r   r   r	   ©ÚtlÚ	constexpr)Úformats   "Ú\/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torchao/kernel/intmm_triton.pyÚ__annotate__r   N   sJ   € ÷ _4ñ _4ô& �\‰\ð'_4ô( �\‰\ð)_4ô* �\‰\ð+_4ô, �\‰\ñ-_4ó    c           	     óÐ  € \         P                  ! ^ R7      p\         P                  ! W<4      p\         P                  ! WM4      pVV,          pVV,          pVV,          p\        VV,
          V4      pVVV,          ,           pVV,          V,          p\         P                  ! V W53Wg3VV,          ^ 3WÎ3RR7      p\         P                  ! VWT3W‰3^ VV,          3Wí3RR7      p\         P
                  ! WÍ3\         P                  R7      p\        ^ W^4       Fƒ  p\         P                  ! VRR7      p\         P                  ! VRR7      pV\         P                  ! VV4      ,          p\         P                  ! V^ V34      p\         P                  ! VV^ 34      pK…  	  Tp\         P                  ! VW43W«3VV,          VV,          3WÍ3RR7      p\         P                  ! VVRR7       R# )ziKernel for computing the matmul C = A x B.
A has shape (M, K), B has shape (K, N) and C has shape (M, N)
)Úaxis)ÚbaseÚshapeÚstridesÚoffsetsÚblock_shapeÚorder©Údtype)Úboundary_checkN)é   r   )r   r    )r   Ú
program_idÚcdivÚminÚmake_block_ptrÚzerosÚint32ÚrangeÚloadÚdotÚadvanceÚstore) Úa_ptrÚb_ptrÚc_ptrÚMÚNÚKÚ	stride_amÚ	stride_akÚ	stride_bkÚ	stride_bnÚ	stride_cmÚ	stride_cnr   r   r   r	   ÚpidÚ	num_pid_mÚ	num_pid_nÚnum_pid_in_groupÚgroup_idÚfirst_pid_mÚpid_mÚpid_nÚa_block_ptrÚb_block_ptrÚaccumulatorÚkÚaÚbÚcÚc_block_ptrs    &&&&&&&&&&&&&&&&                r   Ú!matmul_kernel_with_block_pointersrH   M   sÄ  € ô@ �-Š-˜QÔ
€CÜ—’˜Ó#€IÜ—’˜Ó#€IØ Õ*ÐØÐ&Õ&€HØ˜WÕ$€KÜ�)˜kÕ)¨7Ó3€GØ˜3 �=Õ)€EØÐ#Õ#¨Õ/€Eô ×#Ò#ØØˆfØÐ&Ø˜• !Ð$ØÐ&Øô€Kô ×#Ò#ØØˆfØÐ&Ø�E˜G•OÐ$ØÐ&Øô€Kô —(’(˜GÐ-´R·X±XÔ>€KÜ�1�aÖ!ˆô �GŠG�K°Ô7ˆÜ�GŠG�K°Ô7ˆà”r—v’v˜a “|Õ#ˆô —j’j ¨q°'¨lÓ;ˆÜ—j’j ¨w¸¨lÓ;Šñ "ð 	€Aô
 ×#Ò#ØØˆfØÐ&Ø˜• %¨'¥/Ð2ØÐ&Øô€Kô ‡H‚Hˆ[˜!¨F×3r   c                óØ   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  /# )r   r   r   r   r	   ÚEVEN_KÚACC_TYPEr   )r   s   "r   r   r   ±   sp   € ÷ KNñ KNô, �\‰\ð-KNô. �\‰\ð/KNô0 �\‰\ð1KNô2 �\‰\ð3KNô4 �L‰Lð5KNô6 �l‰lñ7KNr   c                 ó>  € \         P                  ! ^ 4      pWO,           ^,
          V,          pVV,           ^,
          V,          pVV,          pVV,          p\        VVV,          ,
          V4      pVV,          VV,          ,           pVV,          V,          pVV,          \         P                  ! ^ V4      ,           pVV,          \         P                  ! ^ V4      ,           p\         P                  ! \         P
                  ! VV,          V4      V4      p\         P                  ! \         P
                  ! VV,          V4      V4      p \         P                  ! ^ V4      p!V VR,          V,          V!R,          V,          ,           ,           p"VV!R,          V	,          V R,          V
,          ,           ,           p#\         P                  ! VV3VR7      p$\        V^ V) 4       F¼  p%V'       d.   \         P                  ! V"4      p&\         P                  ! V#4      p'MF\         P                  ! V"V!R,          V%8  RR7      p&\         P                  ! V#V!R,          V%8  RR7      p'V$\         P                  ! V&V'4      ,          p$V"VV,          ,          p"V#VV	,          ,          p#K¾  	  VV,          \         P                  ! ^ V4      ,           pVV,          \         P                  ! ^ V4      ,           pVR,          p(VR,          p)V(V8  V)V8  ,          p*V)VV(,          ,           p+\         P                  ! V\         P                  ! V(V*P                  4      ,           V*RR7      p,\         P                  ! V\         P                  ! V+V*P                  4      ,           V$V,,          V*4       R# )	r   Nr   g        )ÚmaskÚotherÚ
evict_last)Úeviction_policy)ºNNNN)NrQ   )r   r!   r#   ÚarangeÚmax_contiguousÚmultiple_ofr%   r'   r(   r)   Úbroadcast_tor   r+   )-r,   r-   r.   Ús1_ptrr/   r0   r1   r2   r3   r4   r5   r6   r7   Ú
stride_s1mÚ
stride_s1nr   r   r   r	   rJ   rK   r8   Úgrid_mÚgrid_nÚwidthr<   Ú
group_sizer>   r?   ÚrmÚrnÚramÚrbnÚrkÚAÚBÚaccrC   rD   rE   Úidx_mÚidx_nrM   ÚxindexÚtmp0s-   &&&&&&&&&&&&&&&&&&&&&                        r   Ú(scaled_matmul_kernel_with_block_pointersri   °   s±  € ô> �-Š-˜Ó
€CØ�k˜A�o 'Õ)€FØ�'�k˜A�o 'Õ)€Fð �fÕ€EØ�e�|€HÜ�V˜h¨Õ0Õ0°'Ó:€JØ�wÕ #¨
Õ"2Õ3€EØ�5�[˜jÕ)€Eà	��œ2Ÿ9š9 Q¨Ó0Õ	0€BØ	��œ2Ÿ9š9 Q¨Ó0Õ	0€BÜ
×
Ò
œBŸNšN¨2°­6°7Ó;¸WÓ
E€CÜ
×
Ò
œBŸNšN¨2°­6°7Ó;¸WÓ
E€CÜ	�Š�1�gÓ	€BØ��W• 	Õ)¨B¨w­K¸)Õ,CÕCÕD€AØ��G•˜yÕ(¨3¨w­<¸)Õ+CÕCÕD€Aä
�(Š(�G˜WÐ%¨XÔ
6€CÜ�1�a˜'˜Ö"ˆßÜ—’˜“
ˆAÜ—’˜“
‰Aä—’˜  7¥¨a¡°sÔ;ˆAÜ—’˜  7¥¨a¡°sÔ;ˆAØŒr�vŠv�a˜‹|ÕˆØ	ˆW�yÕ Õ ˆØ	ˆW�yÕ Õ Šñ #ð 
��œ2Ÿ9š9 Q¨Ó0Õ	0€BØ	��œ2Ÿ9š9 Q¨Ó0Õ	0€BØˆw�K€EØˆw�K€EØ�A‰I˜% !™)Õ$€Dð �a˜%•iÕ €FÜ�7Š7Ø”"—/’/ %¨¯©Ó4Õ5ØØ$ô€Dô
 ‡H‚HˆU”b—o’o f¨d¯j©jÓ9Õ:¸CÀ$½JÈÖMr   c                 ó¢  aa€ V P                   w  opVP                   w  poVV3R  lp\        V,          ! V VVSSVV P                  ^ 4      V P                  ^4      VP                  ^ 4      VP                  ^4      VP                  ^ 4      VP                  ^4      3RVP                  RVP                  RVP
                  /VP                  B  V# )c                 óˆ   <€ \         P                  ! SV R ,          4      \         P                  ! SV R,          4      ,          3# ©r   r   ©Útritonr"   ©ÚMETAr/   r0   s   &€€r   Ú<lambda>Ú#int_matmul_kernel.<locals>.<lambda>  ó.   ø€ Ü�Š�A�t˜I•Ó'¬&¯+ª+°a¸¸i½Ó*IÕIñr   r   r
   Únum_ctas)r   rH   Ústrider   r
   rt   Úkwargs)rD   rE   rF   Úconfigr1   Úgridr/   r0   s   &&&&  @@r   Úint_matmul_kernelry   ÿ   sÊ   ù€ Ø�7‰7�D€A€qØ�7‰7�D€A€qõ€Dô & dÖ+Ø	Ø	Ø	Ø	Ø	Ø	Ø	�‰�‹Ø	�‰�‹Ø	�‰�‹Ø	�‰�‹Ø	�‰�‹Ø	�‰�‹ñð ×"Ñ"ðð ×$Ñ$ðð —‘ðð  �-‰-ò!ð$ €Hr   c                 óü  aa€ V P                   w  opVP                   w  poVV3R  lp\        V,          ! V VVVSSVV P                  ^ 4      V P                  ^4      VP                  ^ 4      VP                  ^4      VP                  ^ 4      VP                  ^4      VP                  ^ 4      VP                  ^4      3RVP                  RVP                  RVP
                  RV^,          ^ 8H  /VP                  B  V# )c                 óˆ   <€ \         P                  ! SV R ,          4      \         P                  ! SV R,          4      ,          3# rl   rm   ro   s   &€€r   rq   Ú*int_scaled_matmul_kernel.<locals>.<lambda>!  rs   r   r   r
   rt   rJ   )r   ri   ru   r   r
   rt   rv   )	rD   rE   Úscales1rF   rw   r1   rx   r/   r0   s	   &&&&&  @@r   Úint_scaled_matmul_kernelr~     s÷   ù€ Ø�7‰7�D€A€qØ�7‰7�D€A€qõ
€Dô -¨TÖ2Ø	Ø	Ø	ØØ	Ø	Ø	Ø	�‰�‹Ø	�‰�‹Ø	�‰�‹Ø	�‰�‹Ø	�‰�‹Ø	�‰�‹Ø�‰�qÓØ�‰�qÓñð  ×"Ñ"ð!ð" ×$Ñ$ð#ð$ —‘ð%ð& �A•˜‘
ð'ð( �-‰-ò)ð, €Hr   ÚtorchaoÚFRAGMENTz(int_matmul(Tensor a, Tensor b) -> Tensorz?int_scaled_matmul(Tensor a, Tensor b, Tensor scales1) -> TensorÚ
int_matmulÚMetac                 óž   € V P                   w  r#VP                   w  r4\        P                  ! W$3V P                  \        P                  R 7      # ©©Údevicer   )r   ÚtorchÚemptyr†   r&   )rD   rE   r/   r1   r0   s   &&   r   Úint_matmul_metar‰   B  s5   € à�7‰7�D€AØ�7‰7�D€AÜ�;Š;˜�v a§h¡h´e·k±kÔBÐBr   ÚCUDAc                 óx  € V P                   ^,          VP                   ^ ,          8X  g   Q R4       hV P                   w  r#VP                   w  r4\        P                  ! W$3V P                  \        P                  R7      p\        \        WV.\        4      pVf   \        P                  ! . 4      # \        WWV4      # ©r    zIncompatible dimensionsr…   )	r   r‡   rˆ   r†   r&   r   ry   Úint8_mm_kernel_configsÚtensor)rD   rE   r/   r1   r0   rF   Úbest_configs   &&     r   Úint_matmul_cudar�   I  s–   € ð �7‰7�1�:˜Ÿ™ �Ô#Ð>Ð%>Ó>Ð#ð �7‰7�D€AØ�7‰7�D€AÜ�Š�Q�F 1§8¡8´5·;±;Ô?€Aä$Ü˜A !˜9Ô&<ó€Kð Òä�|Š|˜BÓÐÜ˜Q 1Ó2Ð2r   Úint_scaled_matmulc                 ó–   € V P                   w  r4VP                   w  rE\        P                  ! W53V P                  VP                  R 7      # r„   )r   r‡   rˆ   r†   r   )rD   rE   r}   r/   r1   r0   s   &&&   r   Úint_scaled_matmul_metar“   ]  s5   € à�7‰7�D€AØ�7‰7�D€AÜ�;Š;˜�v a§h¡h°g·m±mÔDÐDr   c                 ó>  € V P                   ^,          VP                   ^ ,          8X  g   Q R4       hV P                   w  r4VP                   w  rE\        P                  ! W53V P                  VP                  R7      p\        \        WW&.\        4      p\        WW&V4      # rŒ   )r   r‡   rˆ   r†   r   r   r~   r�   )rD   rE   r}   r/   r1   r0   rF   r�   s   &&&     r   Úint_scaled_matmul_cudar•   d  s„   € ð �7‰7�1�:˜Ÿ™ �Ô#Ð>Ð%>Ó>Ð#ð �7‰7�D€AØ�7‰7�D€AÜ�Š�Q�F 1§8¡8°7·=±=ÔA€Aä$Ü  1¨Ð"4Ô6Ló€Kô $ A¨'°kÓBÐBr   )é    é@   é€   é   )é   r–   r—   r˜   r™   )r    r   é   é   é   é   é   é   )r   rœ   r    )Ú	itertoolsr‡   rn   Útriton.languageÚlanguager   Útorchao.kernel.autotunerr   ÚsumÚproductr�   Ú	_inductorrw   Úmax_autotune_gemm_search_spaceÚConfigÚjitrH   r&   ri   ry   r~   ÚlibraryÚLibraryÚlibÚdefineÚimplr‰   r�   r“   r•   )
ÚiÚjrC   r   r   r   r
   r   ÚsÚws
   0000000000r   Ú<module>r´      s  ðó ã Û Ý å 7ñ ð( #×*Ò*Ò+=ÀaÕHõ'ñ& I‰IˆQ�1ð! �1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOØ�1�a˜ˆOó	
ñ" Ió'ð* ó-Ð ð2 	‡?�?×Ñ×8Ñ8¸LÔHð *3×):Ò):Ú"¨1õ*
÷ñ*
Ñ%ˆG˜gó 3ˆJÛ"ˆIð 
˜7°	Ó:ñ
 #ñ 	;ñ 3ñ	 	;ñ*
õÐñ@ 2÷ñ 2‰ˆˆq�Qð ‡M�MØ	�A�y ! Y°°9¸aÐ@ØØð ö ñ
 2õÐ ð ‡�ô_4ó ð_4ðD ‡�ð8  ŸX™Xö7KNó ðKNò\ò6 ðF ‡m�m×Ñ˜I zÓ2€Ø ‡
�
Ð5Ô 6Ø ‡
�
ÐLÔ Mð ‡�×Ñ�C˜ vÓ.ñCó /ðCð ‡�×Ñ�C˜ vÓ.ñ3ó /ð3ð& ‡�×Ñ�CÐ,¨fÓ5ñEó 6ðEð ‡�×Ñ�CÐ,¨fÓ5ñCó 6òCùôk
ùö2ùö4s   µAIÃ#IÄ)I