
    ^jV                        d dl mZ d dlmZmZ d dlm	Z	m
Z
 g dZe
d        Ze
d        Ze
d        Ze
d        Ze G d	 d
             Ze
d        Ze
d        Ze	d        Ze	d        Ze
d        Ze
d        Ze
d        Ze
d        Zy)    )
_aggregate)_core	_standard)constexpr_functionjit)pack2unpack2packunpackfmaFloat2Tensorc                 X    t        j                  dd| |gt         j                  dd      S )Nz'
        add.f32x2 $0, $1, $2;
        =l,l,lT   dtypeis_purer
   ttglinline_asm_elementwiseint64abs     /var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/nvidia/blackwell/float2.py
_add_f32x2r      3    &&	 		
Ajj	 	    c                 X    t        j                  dd| |gt         j                  dd      S )Nz'
        sub.f32x2 $0, $1, $2;
        r   Tr   r   r   r   s     r   
_sub_f32x2r       r   r   c                 X    t        j                  dd| |gt         j                  dd      S )Nz'
        mul.f32x2 $0, $1, $2;
        r   Tr   r   r   r   s     r   
_mul_f32x2r"   +   r   r   c                 Z    t        j                  dd| ||gt         j                  dd      S )Nz.
        fma.rn.f32x2 $0, $1, $2, $3;
        z=l,l,l,lTr   r   r   r   r   cs      r   
_fma_f32x2r&   9   s5    &&	 		
Aq	jj	 	r   c                       e Zd ZU ej                  ed<   edej                  fd       Zed        Z	ed        Z
ed        Zedej                  fd       Zy)	r   valuec                     || _         y N)r(   )selfr(   s     r   __init__zFloat2Tensor.__init__K   s	    
r   c                     t        j                  t        |t              d       t        t	        | j
                  |j
                              S Nzrhs must be a Float2Tensor)r   static_assert
isinstancer   r   r(   r+   rhss     r   __add__zFloat2Tensor.__add__O   4    :c<8:VWJtzz399=>>r   c                     t        j                  t        |t              d       t        t	        | j
                  |j
                              S r.   )r   r/   r0   r   r    r(   r1   s     r   __sub__zFloat2Tensor.__sub__T   r4   r   c                     t        j                  t        |t              d       t        t	        | j
                  |j
                              S r.   )r   r/   r0   r   r"   r(   r1   s     r   __mul__zFloat2Tensor.__mul__Y   r4   r   axisc                 `    t        t        j                  | j                  |t                    S )N)r9   
combine_fn)r   r   reducer(   r   )r+   r9   s     r   sumzFloat2Tensor.sum^   s    DKK

*UVVr   N)__name__
__module____qualname__r   tensor__annotations__r   r,   r   r3   r6   r8   	constexprr=    r   r   r   r   G   s    ;;dkk   	? 	? 	? 	? 	? 	? 	W W 	Wr   r   c                 n    t        j                  dd| |gt         j                  dd      }t        |      S )Nz)
        mov.b64 $0, { $1, $2 };
        z=l,r,rTr   r   )r   r   r   r   )x0x1r(   s      r   r   r   c   s=    ''	 		Rjj	E r   c                     t        j                  dd| j                  gt         j                  t         j                  gdd      S )Nz)
        mov.b64 { $0, $1 }, $2;
        z=r,=r,lTr   r   )r   r   r(   float32)xs    r   r	   r	   r   s>    &&	 		
	||T\\*	 	r   c                 ~   | D cg c]  }| } }| |   dk\  s
J d|        | |xx   dz  cc<   | j                  |dz   d       t        t        t        |                   }|t        |      dz
     ||dz      c||dz   <   |t        |      dz
  <   t	        j
                  |       t	        j
                  |      fS c c}w )N   z'not enough elements to pack along axis r   )insertlistrangelenr   tupleshaper9   dpermutes       r   _get_split_shaperV      s    1QE;!MFtfMM	$KAK	LL15U$%G3:3w<!;K3LgVZ]^V^N_0GD1Hws7|a/0::edjj111 s   	B:c                    | D cg c]  }| } }| |xx   dz  cc<   t        t        t        |                   }|j                  |dz   t        |             t	        j
                  |       t	        j
                  |      fS c c}w )NrL   r   )rN   rO   rP   rM   r   rQ   rR   s       r   _get_join_shaperX      sn    1QE	$K1K5U$%GNN4!8S\*::edjj111	 s   	Bc                     t        | j                  |      }  | j                  |d    j                  |d    j	                         \  }}t        ||      S )Nr   r   )rV   rS   reshaperU   splitr   )rJ   r9   sprF   rG   s        r   r
   r
      sP    )!''48B&YQYY1&&1.446FBR=r   c                     | j                   j                  }t        ||      }t        |       \  }}  t	        j
                  ||      j                  |d    j                  |d    S )Nr   r   )r(   rS   rX   r	   r   joinrU   rZ   )rJ   r9   rS   r\   rF   rG   s         r   r   r      s\    GGMME(5BQZFB4$499R$$be,44be<<r   c                     t        j                  |j                  t         j                  k(  d       t	        j
                  | j                  |t         j                        }t        ||      S )Nzfill_value must be a float32)r   )r   r/   r   rI   stdlib	full_liker(   r   )rJ   
fill_valuefills      r   ra   ra      sK    z''4<<79WXAGGZt||DDtr   c                 j    t        t        | j                  |j                  |j                              S r*   )r   r&   r(   r$   s      r   r   r      s"    
177AGGQWW=>>r   N)triton.language.corer   	aggregate"triton.experimental.gluon.languager   r   r   r`   "triton.experimental.gluon._runtimer   r   __all__r   r    r"   r&   r   r   r	   rV   rX   r
   r   ra   r   rD   r   r   <module>rj      s!   8 Q F 
 
 
 
 
 
 
 
 W W W6   
 
 2 2 2 2   = =   ? ?r   