Ë
    ŒÓ^j-  ã                   óf   — d dl Z d dlmZ d dlmZ d dlmZ ddlmZ ddl	m
Z
 dd	l	mZ d
„ Zd„ Zd„ Zy)é    N)Úknobs)Ú_core)Ú_checké   )Ú_unwrap_if_constexpr)ÚDotOperandLayouté   )ÚAMDWMMALayoutc                 óL  ‡ — t        |d ud„ «       |j                  j                  }t        t        |t        «      xr |j
                  ‰ k(  ˆ fd„«       |j                  j                  }t        t        |t        «      xr5 t        |j                  t        «      xr |j                  j
                  ‰ k(  d„ «       |j                  j                  }t        t        |t        «      xr5 t        |j                  t        «      xr |j                  j
                  ‰ k(  d„ «       y )Nc                   ó   — y)Nzacc is required© r   ó    úv/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/amd/_ops.pyú<lambda>z_verify_wmma.<locals>.<lambda>   ó   � r   c                  ó   •— d‰ › �S )Nz@Expected layout to be an instance of AMDWMMALayout with version r   )Úversions   €r   r   z_verify_wmma.<locals>.<lambda>   s   ø€ ÐRÐSZÐR[Ð\€ r   c                   ó   — y)NzOExpected a's layout to be a DotOperandLayout with parent matching AMDWMMALayoutr   r   r   r   r   z_verify_wmma.<locals>.<lambda>   r   r   c                   ó   — y)NzOExpected b's layout to be a DotOperandLayout with parent matching AMDWMMALayoutr   r   r   r   r   z_verify_wmma.<locals>.<lambda>   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                 óä   — t        | |||«       |j                  |||t        j                  j                  d|j
                  ¬«      j                  }t        j                  ||j                  «      S )zB Shared 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,   !   s\   € ä�˜!˜Q Ô$à�\‰\˜!˜Q ´U·^±^×5PÑ5PÐhlØ$'§I¡Ið ó /ß/5©vð ä�;‰;�v˜sŸx™xÓ(Ð(r   c	                 óô   ‡ ‡‡‡‡— d„ Šˆˆ ˆˆˆfd„}	 |	d||«      } |	d||«      }‰j                  ‰ ||‰|||dddt        j                  ¬«      }
t        j                  |
j                  |j
                  «      S )zF Shared implementation for AMD WMMA scaled and MFMA scaled operation. c                 óð   — |j                   j                  D �cg c]  }|‘Œ }}|}|j                  dk(  rdnd}| dk(  r|d   |z  }|dz  |d<   |S |d   |z  }|dz  |d<   |d   |d   c|d<   |d<   |S c c}w )NÚe2m1r   r	   r   éÿÿÿÿé    éþÿÿÿ)r   ÚshapeÚvalue)Úop_idxÚoperandÚformatÚsÚoperand_shapeÚscale_shapeÚunpack_factorÚks           r   Ú_get_scale_shapez%_mma_scaled.<locals>._get_scale_shape-   s¤   € Ø$+§L¡L×$6Ñ$6Ö7˜qšÐ7ˆÐ7Ø#ˆØ#Ÿ\™\¨VÒ3™¸ˆØ�QŠ;Ø˜B‘ -Ñ/ˆAØ 2™gˆK˜‰Oð
 Ðð ˜B‘ -Ñ/ˆAØ 2™gˆK˜‰OØ/:¸2©ÀÈBÁÐ,ˆK˜‰O˜[¨™_ØÐùò 8s   ™	A3c                 óš  •— | dk(  r‰n‰	} ‰| ||«      }t        |t        j                  «      rN|j                  j                  dk7  r5t        j                  |«      |j                  j                  k(  sJ d«       ‚|S  ‰
|j                  j                  |«      }t        |«      }|€dn|}‰j                  ||t        j                  |«      S )Nr   r	   zIncompatible scale shapeé   )r   r)   r*   Únumelr4   ÚmathÚprodr   r   r   ÚfullÚuint8)r5   Úscaler7   r6   r:   Úscale_layoutÚscale_valuer=   r   r   Úscale_fnr+   s          €€€€€r   Ú#_create_and_broadcast_default_scalez8_mma_scaled.<locals>._create_and_broadcast_default_scale:   s°   ø€ Ø ’{‘!¨ˆá& v¨w¸Ó?ˆÜ�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   rH   r+   rI   Úoutputr=   s   `  `   ``  @r   Ú_mma_scaledrT   *   s~   ü€ ò÷Qð Qñ 2°!°W¸hÓG€GÙ1°!°W¸hÓG€GØ× Ñ   G¨X°q¸'À8ÈSÐ\aÐnrØ,0¼D¿L¹Lð !ó J€Fä�;‰;�v—}‘} c§h¡hÓ/Ð/r   )rA   Útritonr   Ú"triton.experimental.gluon.languager   r)   Ú,triton.experimental.gluon.language._semanticr   r   Ú_layoutsr   r
   r   r,   rT   r   r   r   ú<module>rY      s*   ðÛ å Ý <Ý ?å (Ý 'Ý #òcò*)ó#0r   