
    ^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    RSZR[\ 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    
3d?56XX__F
6=)Gfnn.G\^ vv}}H
8-. 	/:hoo}3] 	/OO##w.ac
 vv}}H
8-. 	/:hoo}3] 	/OO##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^^5P5Phl$'II  //5v ;;vsxx((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    $+LL$6$67q77##\\V3Q;B-/A2gKO
  B-/A2gK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    {!&vw?eT[[)ekk.?.?1.D 99[)U[[->->>Z@ZZ>L 3 3[A*51)1d{}}[+tzz<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!WhGG1!WhGG  GXq'8S\anr,0DLL ! JF;;v}}chh//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   