o
    پi+                     @   sv  d dl Z d dlmZ 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 dlmZmZ d dlmZ d d	lmZ g d
ZerPg dng dZdededefddZdedefddZ				dde jde jde jde jde jdede jdede j dB d e jdB d!ede!e je jf fd"d#Z"						dde jde jde jde jde jdede jded!ede j dB de!e je jf fd$d%Z#ej$d&ej%d'ej%d(ej%d)ej%d*ej%d+ej%d,ej%d-ej%d.ej%fd/d0Z&ej$d&ej%d(ej%d)ej%d*ej%d+ej%d,ej%d-ej%d.ej%fd1d2Z'	3	4				dd5e jde jd6e jd7e jd8e(d9ed:e jd;e j)d<e j)d=efd>d?Z*	3				@dd5e jde jd6e jd7e jd8e(d:e jdB dAedBed9efdCdDZ+G dEdF dFej,Z-ej.dGdH dID dJdKgdLej$dMgdNdOej%dPej%d'ej%dJej%dQej%dRej%dKej%fdSdTZ/ej.dUdH dVD g dWdLej$dMgdNdOej%dPej%d'ej%dJej%dQej%dKej%fdXdYZ0dddddZe j1fde jde jd[e jdB de jdB dedB de j dB d\ed]e j)de!e je jf fd^d_Z2ej.d`dH daD g dbdLej$dMgdNdOej%dPej%dcej%d'ej%dQej%ddej%deej%dfej%dKej%dgej%fdhdiZ3			dde jde jde jdje jde jdB d[e jdB de j dB de!e je jf fdkdlZ4ej.dmdH dID d'dKgdLej$dMgdNdOej%dPej%dcej%d'ej%dQej%ddej%dKej%fdndoZ5		Zdde jde jde jdje jdpe jdqe jdede j dB d\efdrdsZ6	dde jde jde jde jde jdede jdte jde j dB fdudvZ7					dde jde jde jde jde jdede jdte jd!ede j dB fdwdxZ8ej.dydH eD dOd&gdLej$dej%dzej%d&ej%d'ej%d(ej%d.ej%fd{d|Z9		}	~dde jdje jdede jdB dedzede jfddZ:dS )    N)chunk_gated_delta_rule_fwd_h)chunk_local_cumsum)+fused_recurrent_gated_delta_rule_fwd_kernel)prepare_chunk_indices)
l2norm_fwd)explog)
solve_tril)is_amd)    @      )            )r   r   r   r   abreturnc                 C   s   | |   S )zCeiling division. )r   r   r   r   W/home/ubuntu/.local/lib/python3.10/site-packages/sglang/srt/layers/attention/fla/kda.pycdiv   s   r   nc                 C   s   | dk rdS d| d   > S )zThe next power of 2 (inclusive)   )
bit_length)r   r   r   r   next_power_of_2    s   r   TFqkvgbetascaleinitial_stateinplace_final_state
cu_seqlensnum_accepted_tokensuse_qk_l2norm_in_kernelc                 C   s  g |j |j d R \}}}}}|j d }|d u r|nt|d }t|tt|d}}t||t||}}|dks@J dd}d}t|}|rN|}n| j|||||jd}|	d}|	d}|||| f}t
| d$i d	| d
|d|d|d|d|d|d|d|d|d|d|d|d|d|d|d|d|d|d ud|d ud|j|jkd|
d|d ud d!d"|d#| ||fS )%Nr   r   r   zNK > 1 is not supported yet   dtyper   r   r   r   r   r    oh0htr$   r!   TBHHVKVBKBVUSE_INITIAL_STATESTORE_FINAL_STATEIS_BETA_HEADWISEUSE_QK_L2NORM_IN_KERNEL	IS_VARLENIS_KDAT	num_warps
num_stagesr   )shapelenr   minr   torch
empty_like	new_emptyr*   strider   ndim)r   r   r   r   r    r!   r"   r#   r$   r%   r&   r/   r.   r0   r2   r3   r1   Nr4   r5   NKNVr=   r<   r+   final_statestride_init_state_tokenstride_final_state_tokengridr   r   r   fused_recurrent_kda_fwd'   s    



		 !"%rM   c
                 K   s   |	d ur| j d dkrtd| j d  d|d u r!|j d d }t|  | | | | ||||	d |d\}}||fS )Nr   r   z/The batch size is expected to be 1 rather than zQ when using `cu_seqlens`.Please flatten variable-length inputs before processing.r'         )r   r   r   r   r    r!   r"   r#   r$   r%   r&   )r>   
ValueErrorrM   
contiguous)r   r   r   r   r    r!   r"   r#   r&   r$   kwargsr+   rI   r   r   r   fused_recurrent_kdaw   s(   
rR   DBTBD
ACTIVATIONIS_RMS_NORMSTORE_RESIDUAL_OUTHAS_RESIDUAL
HAS_WEIGHTHAS_BIASc           (      C   sl  t d}t d|}||k }t | |
|f|df|| df||fd}t j|ddt j}|rQt ||
|f|df|| df||fd}|t j|ddt j7 }|rtt ||
|f|df|| df||fd}t j|||jj	dd |st j
|dd| }t ||
fd|| f|fd}t j|||jj	dd t |d d d f ||d d d f  d	}t j
|| dd| }nt |d d d f |d	}t j
|| dd| }dt ||	  }t ||
fd|| f|fd} t j| || jj	dd |r	t j|| |d
t j}!|rt j|| |d
t j}"|s.||d d d f  |d d d f  n	||d d d f  }#|rE|#|!d d d f  n|#}$|rT|$|"d d d f  }$t ||
|f|df|| df||fd}%t j|%ddt j}&|dks}|dkr|$|& t |& }$n|dkr|$t |& }$t ||
|f|df|| df||fd}'t j|'|$|'jj	dd d S )Nr   r   r   r   r   r   boundary_checkaxis)r   r           maskswishsilusigmoid)tl
program_idarangemake_block_ptrloadtofloat32storer*   
element_tysumwheresqrtrh   )(xr   ywr   residualresidual_outmeanrstdepsr.   rS   rT   rU   rV   rW   rX   rY   rZ   r[   i_to_dm_dp_xb_xp_res	p_res_outb_meanp_meanb_xbarb_varb_rstdp_rstdb_wb_bb_x_hatb_yp_gb_gp_yr   r   r   layer_norm_gated_fwd_kernel   sZ   
(  *((
(r   c                 C   s2  t d}| ||
 7 } |||
 7 }|||
 7 }|r|||
 7 }|r'|||
 7 }t d|}||
k }t j| | |ddt j}|rQ|t j|| |ddt j7 }|r]t j|| ||d |st j|dd|
 }t || | t ||| d}t j|| dd|
 }nt ||d}t j|| dd|
 }dt 	||	  }t || | |rt j|| |dt j}|rt j|| |dt j}|s|| | n|| }|r|| n|}|r|| }t j|| |ddt j}|dks|dkr|| t 
| }n|d	kr|t 
| }t j|| ||d d S )
Nr   rc   re   otherrd   r`   r   rf   rg   rh   )ri   rj   rk   rm   rn   ro   rp   rr   rs   rt   rh   )ru   r   rv   rw   r   rx   ry   rz   r{   r|   rS   rU   rV   rW   rX   rY   rZ   r[   r}   r~   r   r   r   r   r   r   r   r   r   r   r   r   r   r   layer_norm_gated_fwd_kernel1   sL   
 
r   rf   h㈵>ru   weightbias
activationr|   rx   	out_dtyperesidual_dtypeis_rms_normc
                 C   sn  |d ur|j }| j\}
}|d ur|j|
|fksJ |d ur%|j|fks%J |d ur1|j|fks1J |d u r7| ntj| |d}|d usK|d urV|| j krVtj|
|| j|d}nd }|	setj|
ftj| jdnd }tj|
ftj| jd}d|   }t|t	|}||krt
d|dkrd}tt|
|f di d| d	|d
|d|d|d|d|d|d|d|d|
d|d|d|d|d|	d|d ud|d ud|d ud|d udd nIt|
f di d| d	|d
|d|d|d|d|d|d|d|d|d|d|d|	d|d ud|d ud|d ud|d udd ||||d ur4|fS | fS )Nr)   devicer*   )r*   r   i   z4This layer norm doesn't support feature dim >= 64KB.i   r   ru   r   rv   rw   r   rx   ry   rz   r{   r|   r.   rS   rU   rT   rV   rW   rX   rY   rZ   r[   r<   r   r   )r*   r>   rA   rB   emptyr   floatelement_sizer@   r   RuntimeErrorr   r   r   )ru   r   r   r   r   r|   rx   r   r   r   r.   rS   rv   ry   rz   r{   MAX_FUSED_SIZErU   rT   r   r   r   layer_norm_gated_fwd2  s   
	
	
r   ư>prenormresidual_in_fp32c	                 C   s   | j }	|  d| j d } | d|j d }|d ur/|j |	ks$J | d|j d }|d ur6|jn|r;tjnd }
t| |||||||
dd	\}}}}||	}|sW|S |||	fS )Nr'   T)	ru   r   r   r   r   r|   rx   r   r   )r>   rP   reshaper*   rA   r   r   )ru   r   r   r   r   rx   r   r   r|   
x_shape_ogr   rv   _ry   r   r   r   rms_norm_gated  s.   
r   c                       s   e Zd Z					ddedededed	ejdB d
ej	dB ddf fddZ
			ddejdejdejdB dededejfddZ  ZS )FusedRMSNormGatedTr   rf   Nhidden_sizeelementwise_affiner|   r   r   r*   r   c                    s   ||d}t    || _|| _|| _|| _| jdvr#td| j |r3tt	j
|fi || _n| dd  | dd  d S )Nr   )rf   rg   rh   zUnsupported activation: r   r   )super__init__r   r   r|   r   rO   nn	ParameterrA   r   r   register_parameter)selfr   r   r|   r   r   r*   factory_kwargs	__class__r   r   r     s   
	

zFusedRMSNormGated.__init__Fru   r   rx   r   r   c                 C   s"   t ||| j| j| j|| j||d	S )N)rx   r|   r   r   )r   r   r   r   r|   )r   ru   r   rx   r   r   r   r   r   forward  s   zFusedRMSNormGated.forward)Tr   rf   NN)NFF)__name__
__module____qualname__intboolr   strrA   r   r*   r   Tensorr   __classcell__r   r   r   r   r     sJ    r   c              	   C   s4   g | ]}d D ]}dD ]}t jd|i||dq
qqS )r   r   r   r   r   r(   r   r4   r<   r=   tritonConfig).0r4   r<   r=   r   r   r   
<listcomp>      
r   )r   r   BCr:   )configskeyr.   )do_not_specializer0   r2   r4   NCc           0   
   C   s  t dt dt d}}}||
 ||
 }}|| || }}|r`t ||d  t jt ||d  d t j}}t || t jt || d t j}}|| }	n||	 ||	 |	 }}|| ||  |	krwd S ||kr}d S | ||
 | | 7 } |||
 | | 7 }|||
 | | 7 }|||
 | | 7 }|||
 | | 7 }t |||
  | |	f|
f|| ||  f|fd}t j|dd}t j||gt jd}t j||gt jd}tt 	||D ]}t | |	|f|
| df|| ||  || f||fd} t ||	|f|
| df|| ||  || f||fd}!t ||	|f|
| df|| ||  || f||fd}"t |||	fd|
| f|| || ||  f||fd}#t |||	fd|
| f|| || ||  f||fd}$|| t 
d| }%|%|k }&t j||| ||  |
 |  |% |&dd	}'t j|"dd}(t j|!ddt|(|'d d d f   })t j|$dd}*t j|#dd}#|#t|'d d d f |*  }+|t |)|+7 }t j| dd},|,t|(|'d d d f   | }-|t |-|+7 }q||d d d f 9 }t ||	|f|
| df|| ||  || f||fd}.t j|.||jjdd t ||	|f|
| df|| ||  || f||fd}/t j|/||jjdd d S )
Nr   r   r   rb   r^   r)   r\   r]   r   )ri   rj   rm   rn   int32rl   zerosro   ranger   rk   r   dotrp   r*   rq   )0r   r   r   r    AAqkr!   r$   chunk_indicesr.   r0   r2   rT   r   r4   r   r:   r}   i_ci_bhi_bi_hi_ii_ji_nboseosp_br   b_Ab_Aqki_kp_qp_kr   b_ktp_gko_km_kb_gnr   b_kb_gkb_ktgb_qb_qgp_Ap_Aqkr   r   r   3chunk_kda_scaled_dot_kkt_fwd_kernel_intra_sub_inter  s   "
,00000,& 00r   c                 C   s   g | ]	}t ji |d qS ))r<   r   )r   r<   r   r   r   r   T  s    r   )r4   rT   r:   c           ,      C   s  t dt dt d}}}||
 ||
 }}|rWt ||d  t jt ||d  d t j}}t || t jt || d t j}}|| }	n||	 ||	 |	 }}|| ||  |	krnd S t d|}t d|}||k }|| ||  | |	k }|||  ||  | |
 | ||  ||  }t | ||
 | |  |	|f|
| df|| ||  df||fd}t |||
 | |  |	|f|
| df|| ||  df||fd}t |||
 | |  |	|f|
| df|| ||  df||fd}t j|dd} t j|dd}!t j|dd}"||||  ||  | |
  | }#|!t j|#|ddd d d f  }!||||  ||  |
 |  ||  | }$||||  ||  |
 |  ||  | }%tdt||	||  ||  D ]w}&t j|$|ddt j	}'t j|%|ddt j	}(|'d d d f t
|"|(d d d f   })t |!|) d}*t ||&k|*d}*t | |) d}+t ||&k|+| d}+t j|| |& |*|d	 t j|| |& |+|d	 |$|
| 7 }$|%|
| 7 }%q{d S )
Nr   r   r   r\   r]   r^   r   rc   rd   )ri   rj   rm   rn   r   rk   rl   r   r@   ro   r   rr   rs   rp   ),r   r   r   r    r   r   r!   r$   r   r.   r0   r2   rT   r   r4   r:   r}   r   r   r   r   r   r   r   o_ir   r   m_Ao_Ar   r   r   r   r   r   r   p_ktr   jr   r   r   r   r   r   r   r   3chunk_kda_scaled_dot_kkt_fwd_kernel_intra_sub_intraS  s|   "
0


$ ,,$(r   r   gk
chunk_sizeoutput_dtypec                 C   s  |j \}}	}
}|dksJ |}|durt||nd}|du r#t|	|nt|}td|}t||}tt|d}tj||	|
||j	|d}tj||	|
||j	|d}||| ||
 f}t
| di d| d|d|d|d	|d
|d|d|d|d|	d|
d|d|d|d|d|du ||||
 f}t| di d| d|d|d|d	|d
|d|d|d|d|	d|
d|d|d|d|d|du ||fS )a  
    Compute beta * K * K^T.

    Args:
        k (torch.Tensor):
            The key tensor of shape `[B, T, H, K]`.
        beta (torch.Tensor):
            The beta tensor of shape `[B, T, H]`.
        gk (torch.Tensor):
            The cumulative sum of the gate tensor of shape `[B, T, H, K]` applied to the key tensor. Default: `None`.
        cu_seqlens (torch.LongTensor):
            The cumulative sequence lengths of the input tensor.
            Default: None
        chunk_size (int):
            The chunk size. Default: 64.
        output_dtype (torch.dtype):
            The dtype of the output tensor. Default: `torch.float32`

    Returns:
        beta * K * K^T of shape `[B, T, H, BT]` where `BT` is the chunk size.
       Nr   r   r   r   r   r    r   r   r!   r$   r   r.   r0   r2   rT   r   r   r:   r4   r   )r>   r   r   r?   r@   maxr   rA   r   r   r   r   )r   r   r   r    r!   r$   r   r   r/   r.   r0   r2   rT   r   NTr   r   r4   r   r   rL   r   r   r   chunk_kda_scaled_dot_kkt_fwd  s   

	
	
r   c                 C   s&   g | ]}d D ]
}t ji ||dqqS )r   r   r   )r   r<   r=   r   r   r   r   	  s    r   r   r   )r0   r2   r3   rT   r4   r5   r:   r3   r5   STORE_QGSTORE_KGDOT_PRECISIONc           :   	   C   s  t dt d}}|| || }}|rQt ||d  t jt ||d  d t j}}t |
| t jt |
| d t j}}|| }n|| || | }}t |||  | |f|f|| f|fd}t j|dd}t ||| | |  ||f|| df|| df||fd} t j| dd}!tt ||D ]l}"t ||| | |  ||f|| df|| |"| f||fd}#t ||| | |  ||f|| df|| |"| f||fd}$t j|#dd}%|%|d d d f  |%j}&t j	|!|&|d}'t j
|$|'|$jjdd qtt ||D ]Z}(t ||| | |  ||f|| df|| |(| f||fd})t ||| | |  ||f|| df|| |(| f||fd}*t j|*dd}+|+|d d d f  },t |	|| | |  ||f|| df|| |(| f||fd}-t j|-dd}.|,t|.9 },|rt | || | |  ||f|| df|| |(| f||fd}/t ||| | |  ||f|| df|| |(| f||fd}0t j|/dd}1|1t|. }2t j
|0|2|0jjdd |r^t|| | |d }3|(| t d| }4|4|k }5t j|	||3 | | |  |4 |5d	d
}6|+t|6|.  }7t ||| | |  ||f|| df|| |(| f||fd}8t j
|8|7|8jjdd t 	|!|,|+j}9t j
|)|9|)jjdd qd S )Nr   r   r   rb   r^   r\   r]   )input_precisionrc   r   )ri   rj   rm   rn   r   rl   r   r   r*   r   rp   rq   r   r@   rk   ):r   r   qgkgr   r    rw   ur   r   r$   r   r.   r0   r2   r3   rT   r4   r5   r   r   r:   r  r}   r   r   r   r   r   r   r   r   r   r   i_vp_vp_ub_vb_vbb_ur   p_wr   r   b_kbr   r   r   p_qgr   r   last_idxr   r   r   b_kgp_kgr   r   r   r   recompute_w_u_fwd_kernel  s   "
,4







r  r   c                 C   s@  g | j |j d R \}}}	}
}|j d }d}d}|d ur"t||nd }|d u r-t||nt|}t| }t|}|d urDt| nd }t|||	 f di d|d| dd d|d|d|d	|d
|d|d|d|d|d|d|	d|
d|d|d|d|ddd|d ud|d udd ||d |fS )Nr'   r   r   r   r  r  r   r    rw   r  r   r   r$   r   r.   r0   r2   r3   rT   r4   r5   r   Fr   r:   r  ieeer   )r>   r   r   r?   rA   rB   r  )r   r   r    r   r   r   r$   r/   r.   r0   r2   r3   rT   r4   r5   r   r   rw   r  r  r   r   r   recompute_w_u_fwd  st    	


	
r  c              
   C   s@   g | ]}d D ]}dD ]}dD ]}t j||d||dqq
qqS ))r   r   r   r   )r4   r5   r   r   )r   r4   r5   r<   r=   r   r   r   r     s    c           *   	   C   s  t dt dt d}}}||
 ||
 }}|r_|}t ||d  t jt ||d  d t j}}t || t jt || d t j}}|| }	t |	|}nt |	|}|| | }||	 ||	 |	 }}t d|d d d f t d|d d d f k}t j||gt jd}t	t ||D ]}t 
| ||
 | |  |	|f|
| df|| || f||fd}t 
|||
 | |  |	|f|
| df|| || f||fd}t 
|||
 | | |  ||f|df|| || f||fd} t j|dd}!|!| |!j}!t j|dd}"|!t|" |!j}#t j| dd}$|dkr:|t |#|$|#j7 }qt 
|||
 | |  |	|f|
| df|| || f||fd}%t 
|||
 | |  |	|f|
| df|| || f||fd}&t 
|||
 | |  |	|f|
| df|| df||fd}'t j|%dd}(t j|'dd})t ||)d|(j})|t j|)|(d	d
7 }t j|&||&jjdd d S )Nr   r   r   r)   r\   r]   r^   rc   F)
allow_tf32)ri   rj   rm   rn   r   r   rk   r   ro   r   rl   r*   r   r   rs   rp   rq   )*r   r   r   hr+   r   r$   r   r!   r.   r0   r2   r3   rT   r4   r5   r:   r  r}   r   r   r   i_tgr   r   r   r   m_sb_or   r   r   p_hr   r   r   b_hr  p_or   r	  r   r   r   r   chunk_gla_fwd_kernel_o  s   "0





4r  r  r+   c	                    s   g | j |j d R \ }	}
|}|d urt||nd }|d u r&t|	|nt| fdd}t| | |||||||||	|
||d ud |S )Nr'   c                    s   t | d   fS )Nr5   r   metar/   r0   r   r3   r   r   rL   ^  s   z chunk_gla_fwd_o_gk.<locals>.grid)r   r   r   r  r+   r   r$   r   r!   r.   r0   r2   r3   rT   r:   )r>   r   r   r?   r  )r   r   r   r   r  r+   r!   r$   r   r.   r2   rT   r   rL   r   r!  r   chunk_gla_fwd_o_gkI  s4    
r"  initial_state_indicesc	                 C   s   d}	t ||	|d}t| |||||tjd\}
}t|
||jd}
t||||
||d\}}}}~
t|||||||d\}}~~~t| ||||||||	d	}~~~|S )Nr   )r   r$   )r   r   r   r    r!   r$   r   )r   r$   r   )r   r   r    r   r   r$   )r   rw   r  r   r"   r#  r$   )	r   r   r   r   r  r+   r!   r$   r   )	r   r   rA   ro   r	   r*   r  r   r"  )r   r   r   r   r    r!   r"   r#  r$   r   r   r   rw   r  r   r  r  v_newr+   r   r   r   chunk_kda_fwdu  sX   
	
	r%  c
                 K   s\   |d u r|j d d }|rt|  } t| }t| || | | ||||	d	}|S )Nr'   rN   )	r   r   r   r   r    r!   r"   r#  r$   )r>   r   rP   r%  )r   r   r   r   r    r!   r"   r#  r&   r$   rQ   r+   r   r   r   	chunk_kda  s"   r&  c              	   C   s4   g | ]}t D ]}d D ]}tjd|i||dq
qqS ))r   r(   rT   r   )NUM_WARPS_AUTOTUNEr   r   )r   btnwnsr   r   r   r     r   	thresholdc              
   C   sh  t dt d}}||	 }t || t j}t | }|| }d}t j| ||  ||f||f|df|	|
fdd}t j|||  ||f||f|df|	|
fdd}t j|ddt j}|rt d|
}||k }t j|||  | |ddt j}||d d d f  }|| }||k}t ||d	| t	d	t |  }|| }t j
|||jjdd d S )
Nr   r   r\   )baser>   stridesoffsetsblock_shapeorderr]   r^   rc   r         ?)ri   rj   rm   rn   ro   r   rl   rk   rs   r   rp   r*   rq   )r   r   rv   g_biasr    r+  r.   r0   rS   rT   rU   r[   r}   r   n_tb_a
stride_row
stride_colg_ptry_ptrr   n_d	bias_maskb_biasg_scaled
use_linearspr   r   r   r   kda_gate_fwd_kernel  sF   
	
	$r?  r1        4@
head_k_dimr2  c           
         s   | j dd }| d| j d } | j d | j d }|   | |ks&J tj| tjd} fdd}	t|	 | ||||| |t||dud |jg | |R  }|S )	z
    Forward pass for KDA gate:
      input g: [..., H*D]
      param A: [H] or [1, 1, H, 1]
      beta: softplus beta parameter
      threshold: softplus threshold parameter
      return  : [..., H, D]
    Nr'   r   r   r)   c                    s   t | d  fS )NrT   r  r  r0   r.   r   r   rL   4  s   zfused_kda_gate.<locals>.grid)rU   r[   )r>   viewnumelrA   rB   ro   r?  r   )
r   r   rA  r2  r    r+  
orig_shapeHDrv   rL   r   rB  r   fused_kda_gate  s.   

rG  )TNNF)NNNTTN)rf   r   NNNF)rf   NFFr   )NNN)Nr   )N)NNNFN)Nr1  r@  );rA   torch.nnr   r   triton.languagelanguageri   -sglang.srt.layers.attention.fla.chunk_delta_hr   &sglang.srt.layers.attention.fla.cumsumr   /sglang.srt.layers.attention.fla.fused_recurrentr   %sglang.srt.layers.attention.fla.indexr   &sglang.srt.layers.attention.fla.l2normr   "sglang.srt.layers.attention.fla.opr   r   *sglang.srt.layers.attention.fla.solve_trilr	   %sglang.srt.layers.attention.fla.utilsr
   BT_LIST_AUTOTUNEr'  r   r   r   r   r   r   
LongTensortuplerM   rR   jit	constexprr   r   r   r*   r   r   Moduler   autotuner   r   ro   r   r  r  r  r"  r%  r&  r?  rG  r   r   r   r   <module>   s*  	
U	

'OH	

b	
&0	d]	
V 
2
t	
5	
A	

"		
B