ó
    “~i-  ã                   óf   • S SK r S SKJr  S SKJr  S SKJr  SSKJr  SSK	J
r
  SS	K	Jr  S
 rS rS rg)é    N)Úknobs)Ú_core)Ú_checké   )Ú_unwrap_if_constexpr)ÚDotOperandLayouté   )ÚAMDWMMALayoutc                 ó€  ^ • [        US LS 5        UR                  R                  n[        [        U[        5      =(       a    UR
                  T :H  U 4S j5        UR                  R                  n[        [        U[        5      =(       a:    [        UR                  [        5      =(       a    UR                  R
                  T :H  S 5        UR                  R                  n[        [        U[        5      =(       a:    [        UR                  [        5      =(       a    UR                  R
                  T :H  S 5        g )Nc                  ó   • g)Nzacc is required© r   ó    Úh/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/experimental/gluon/language/amd/_ops.pyÚ<lambda>Ú_verify_wmma.<locals>.<lambda>   s   € Ð$5r   c                  ó   >• ST  3$ )Nz@Expected layout to be an instance of AMDWMMALayout with version r   )Úversions   €r   r   r      s   ø€ ÐRÐSZÐR[Ñ\r   c                  ó   • g)NzOExpected a's layout to be a DotOperandLayout with parent matching AMDWMMALayoutr   r   r   r   r   r      ó   € Ðar   c                  ó   • g)NzOExpected b's layout to be a DotOperandLayout with parent matching AMDWMMALayoutr   r   r   r   r   r      r   r   )r   ÚtypeÚlayoutÚ
isinstancer
   r   r   Úparent)r   ÚaÚbÚaccr   Úa_layoutÚb_layouts   `      r   Ú_verify_wmmar       sá   ø€ Ü
ˆ3�dˆ?Ñ5Ô6à�X‰X�_‰_€FÜ
Ü�6œ=Ó)×G¨f¯n©nÀÑ.GÜ\ô^ð �v‰v�}‰}€HÜ
Ü�8Ô-Ó.÷ 	/´:¸h¿o¹oÌ}Ó3]÷ 	/Ø�O‰O×#Ñ# wÑ.Ùaôcð
 �v‰v�}‰}€HÜ
Ü�8Ô-Ó.÷ 	/´:¸h¿o¹oÌ}Ó3]÷ 	/Ø�O‰O×#Ñ# wÑ.Ùaõcr   c           	      óØ   • [        XX#5        UR                  XU[        R                  R                  SUR
                  S9R                  n[        R                  " XSR                  5      $ )zAShared implementation for AMD WMMA operations for Gluon builtins N)Úinput_precisionÚmax_num_imprecise_accÚ	out_dtype)
r    Údotr   ÚlanguageÚfp32_defaultÚdtypeÚhandleÚttglÚtensorr   )r   r   r   r   Úsemanticr)   s         r   Ú_wmmar-   !   sT   € ä�˜QÔ$à�\‰\˜! ´U·^±^×5PÑ5PÐhlØ$'§I¡Ið ð /ß/5©vð ä�;Š;�vŸx™xÓ(Ð(r   c	                 óê   ^ ^^^^• S mUU UUU4S jn	U	" SX5      nU	" SXE5      nTR                  T XTXEUSSS[        R                  S9n
[        R                  " U
R                  UR
                  5      $ )zEShared implementation for AMD WMMA scaled and MFMA scaled operation. c                 óø   • UR                   R                   Vs/ s H  o3PM     nnUnUR                  S:X  a  SOSnU S:X  a  US   U-  nUS-  US'   U$ US   U-  nUS-  US'   US   US   sUS'   US'   U$ s  snf )NÚe2m1r   r	   r   éÿÿÿÿé    éþÿÿÿ)r   ÚshapeÚvalue)Úop_idxÚoperandÚformatÚsÚoperand_shapeÚscale_shapeÚunpack_factorÚks           r   Ú_get_scale_shapeÚ%_mma_scaled.<locals>._get_scale_shape-   s§   € Ø$+§L¡L×$6Ò$6Ó7Ò$6˜qšÑ$6ˆÐ7Ø#ˆØ#Ÿ\™\¨VÓ3™¸ˆØ�Q‹;Ø˜B‘ -Ñ/ˆAØ 2™gˆK˜‰Oð
 Ðð ˜B‘ -Ñ/ˆAØ 2™gˆK˜‰OØ/:¸2©ÀÈBÁÐ,ˆK˜‰O˜[¨™_ØÐùò 8s   ™A7c                 óª  >• U S:X  a  TOT	nT" XU5      n[        U[        R                  5      (       aQ  UR                  R                  S:w  a7  [
        R                  " U5      UR                  R                  :X  d   S5       eU$ T
" UR                  R                  U5      n[        U5      nUc  SOUnTR                  XF[        R                  U5      $ )Nr   r	   zIncompatible scale shapeé   )r   r*   r+   Únumelr5   ÚmathÚprodr   r   r   ÚfullÚuint8)r6   Úscaler8   r7   r;   Úscale_layoutÚscale_valuer>   r   r   Úscale_fnr,   s          €€€€€r   Ú#_create_and_broadcast_default_scaleÚ8_mma_scaled.<locals>._create_and_broadcast_default_scale:   s¯   ø€ Ø “{‘!¨ˆá& v¸Ó?ˆÜ�eœTŸ[™[×)Ñ)¨e¯k©k×.?Ñ.?À1Ó.Dô —9’9˜[Ó)¨U¯[©[×->Ñ->Ó>ÐZÐ@ZÓZÐ>ØˆLá §¡× 3Ñ 3°[ÓAˆÜ*¨5Ó1ˆØ)Ñ1‘d°{ˆØ�}‰}˜[´t·z±zÀ<ÓPÐPr   r   r	   FT)Ú	fast_mathÚ
lhs_k_packÚ
rhs_k_packr$   )Ú
dot_scaledr*   Úfloat32r+   r)   r   )r   Úa_scaleÚa_formatr   Úb_scaleÚb_formatr   rJ   r,   rK   Úoutputr>   s   `  `   ``  @r   Ú_mma_scaledrW   *   sv   ü€ ò÷Qñ Qñ 2°!°WÓG€GÙ1°!°WÓG€GØ× Ñ   G°q¸'ÈSÐ\aÐnrØ,0¼D¿L¹Lð !ð J€Fä�;Š;�v—}‘} c§h¡hÓ/Ð/r   )rC   Útritonr   Ú"triton.experimental.gluon.languager   r*   Ú,triton.experimental.gluon.language._semanticr   r   Ú_layoutsr   r
   r    r-   rW   r   r   r   Ú<module>r\      s*   ðÛ å Ý <Ý ?å (Ý 'Ý #òcò*)ó#0r   