U
    •Ê¹jw{  ã                   @   s(  d dl Z d dlZd dlmZ dd„ Zdd„ Zdd„ Zd	d
„ Zdd„ Zdd„ Z	dd„ Z
dd„ Zdd„ Zdd„ Zdd„ Zdd„ Zdd„ ZdCdd„Zdd „ Zd!d"„ Zd#d$„ Zeƒ �rd dlZd dlmZ d d%lmZmZ ejejejejejejejd&œd'd(„ƒZejejejejejejd)œd*d+„ƒZd,d-„ Z d.d/„ Z!d0d0dd1dd2œej"ej"ej"eej" e#eeee$ ee$ ee$ f  d3œd4d5„Z%dd1dd6œej"ej"eej" e#eeee$ ee$ ee$ f  d7œd8d9„Z&ejejejd:œd;d<„ƒZ'dDd=d>„Z(dEej"ej"ej"eej" e)e#ee) d@œdAdB„Z*ndZ(dZ&dZ%dZ*dS )Fé    N)Úget_device_capabilityc                  C   sF   t j ¡ sdS zdd l} | d k	o(tƒ dkW S  tk
r@   Y dS X d S )NFr   )é   r   )ÚtorchÚcudaZis_availableÚtritonr   ÚImportError)r   © r   úW/var/www/html/bwave/python/venv/lib/python3.8/site-packages/torch/sparse/_triton_ops.pyÚ_has_triton   s    
r
   c                 C   s   | st |ƒ‚d S ©N)Ú
ValueError)ZcondÚmsgr   r   r	   Úcheck   s    r   c                 C   s   t |jtjk| › d�ƒ d S )Nz@(): only BSR sparse format is supported for the sparse argument.)r   Úlayoutr   Ú
sparse_bsr)Úf_nameÚtr   r   r	   Úcheck_bsr_layout   s    
þr   c                 C   s&   t |j|ko|jjdk| › d�ƒ d S )Nr   z9(): all inputs are expected to be on the same GPU device.)r   ÚdeviceÚtype)r   r   r   r   r   r	   Úcheck_device   s    þr   c                 C   s€   t | ¡ dko| ¡ dk| › d| ¡ › d| ¡ › d�ƒ |jdd … \}}|jdd … \}}t ||k| › d|› d|› d�ƒ d S )Né   zc(): all inputs involved in the matrix product are expected to be at least 2D, but got lhs.dim() == z and rhs.dim() == Ú.éþÿÿÿzw(): arguments' sizes involved in the matrix product are not compatible for matrix multiplication, got lhs.shape[-1] == z( which is not equal to rhs.shape[-2] == )r   ÚdimÚshape)r   ÚlhsÚrhsÚmÚklÚkrÚnr   r   r	   Úcheck_mm_compatible_shapes%   s    þþr"   c                 G   sF   t |j|ko(|jtjtjtjft|Ž  k| › d|› d|j› d�ƒ d S )Nz\(): all inputs are expected to be of the same dtype and one of (half, bfloat16, float32) or z, but got dtype == r   )r   Údtyper   ÚhalfÚbfloat16ÚfloatÚtuple)r   r   r#   Zadditional_dtypesr   r   r	   Úcheck_dtype6   s
    
ýr(   c                    sP   t |ƒdkst‚dd„ ‰ ‡ fdd„}t||ƒ| › d|d › d|d	 › d
�ƒ d S )Nr   c                 S   s   | | d @  S ©Né   r   )Úvr   r   r	   Úis_power_of_twoC   s    z(check_blocksize.<locals>.is_power_of_twoc                    s&   d}| D ]}|dkoˆ |ƒo|}q|S )NTé   r   )ÚbÚresÚ	blocksize©r,   r   r	   Úis_compatible_blocksizeF   s    z0check_blocksize.<locals>.is_compatible_blocksizez(): sparse inputs' blocksize (r   z, r*   z;) should be at least 16 and a power of 2 in each dimension.)ÚlenÚAssertionErrorr   )r   r0   r2   r   r1   r	   Úcheck_blocksize@   s    þr5   c                 C   s,   |   d¡dkr$|   d¡dkr$|  ¡ S | S d S )Nr   r*   éÿÿÿÿ)ÚstrideÚ
contiguous)r   r   r   r	   Úmake_triton_contiguousT   s    r9   c                 G   s@   zt jdd„ |D ƒŽ W S  tk
r:   td| › d�ƒ Y nX d S )Nc                 s   s   | ]}|j d d… V  qd S ©Nr   ©r   ©Ú.0r   r   r   r	   Ú	<genexpr>]   s     z'broadcast_batch_dims.<locals>.<genexpr>Fz3(): inputs' batch dimensions are not broadcastable!)r   Úbroadcast_shapesÚ	Exceptionr   )r   Útensorsr   r   r	   Úbroadcast_batch_dims[   s    rB   c                 g   s2   |D ](}t d ƒg| ¡  }||| < || V  qd S r   )Úslicer   )r   Zslice_rangerA   r   Úslicesr   r   r	   Úslicerb   s    rE   c                 g   sN   |D ]D}t d ƒg| ¡  }t| |ƒD ]\}}|d k	r$|||< q$|| V  qd S r   )rC   r   Úzip)ZdimsrD   rA   r   ÚsÚdZd_slicer   r   r	   Úmultidim_sliceri   s    
rI   c                  g   s"   | D ]}|V  |  ¡ E d H  qd S r   )r7   )rA   r   r   r   r	   Úptr_stride_extractorr   s    rJ   c           	      #   s¸   dt ˆ ƒ  krdksn t‚dt ˆƒ  kr6dks<n t‚dd l}‡ ‡fdd„}‡fdd„}|j|ƒ Ž D ]H}dd„ tˆ |ˆƒD ƒ}d	d„ t||ƒD ƒ}|d d d
… f||ƒ˜V  qjd S )Nr   é   c                  3   s&   t ˆ ˆƒD ]\} }td| |ƒV  q
d S )Nr   )rF   Úrange)ÚfgÚmg)Ú	full_gridÚgrid_blocksr   r	   Úgenerate_grid_points~   s    z.grid_partitioner.<locals>.generate_grid_pointsc                 3   s(   ˆ   ¡ D ]\}}tt|| |ƒƒV  qd S r   )ÚitemsÚnextrI   )rD   r   Zt_dims)Útensor_dims_mapr   r	   Úgenerate_sliced_tensors‚   s    z1grid_partitioner.<locals>.generate_sliced_tensorsc                 S   s    g | ]\}}}t || |ƒ‘qS r   )Úmin)r=   rM   ÚgprN   r   r   r	   Ú
<listcomp>‡   s     z$grid_partitioner.<locals>.<listcomp>c                 S   s   g | ]\}}t ||| ƒ‘qS r   )rC   )r=   rW   Úgr   r   r	   rX   ˆ   s     r6   )r3   r4   Ú	itertoolsÚproductrF   )	rO   rP   rT   rZ   rQ   rU   Z
grid_pointÚgridrD   r   )rO   rP   rT   r	   Úgrid_partitionerx   s    r]   c                    sh   dd d d… }|d kr|}n$dd„ ‰ t ‡ fdd„t||ƒD ƒƒ}t|||ƒD ]^}}| |f|žŽ  qLd S )N)iÿÿÿéÿÿ  r^   r6   c                 S   s    | d kr|S t dt| |ƒƒS d S r)   )ÚmaxrV   )rY   rN   r   r   r	   Úvalid_grid_dim–   s    z%launch_kernel.<locals>.valid_grid_dimc                 3   s   | ]\}}ˆ ||ƒV  qd S r   r   )r=   rY   rN   ©r`   r   r	   r>   �   s    z launch_kernel.<locals>.<genexpr>)r'   rF   r]   )ÚkernelrT   rO   rP   Zcuda_max_gridr\   Úsliced_tensorsr   ra   r	   Úlaunch_kernel�   s    ÿrd   c                    s¸   |   ¡  d¡}|  ¡  d¡}t|  ¡  d¡ƒ}dd„ |D ƒ}tj|jd d… fdd„ |D ƒžŽ ‰dd„ ‰ ˆ |ˆd	ƒ}ˆ |ˆd	ƒ}ˆ |ˆ|jdd … ƒ}‡ ‡fd
d„|D ƒ}|||f|˜S )Nr   c                 S   s   g | ]}t | d ¡ƒ‘qS )r   )r9   Ú	unsqueezer<   r   r   r	   rX   ª   s     z"prepare_inputs.<locals>.<listcomp>éýÿÿÿc                 s   s   | ]}|j d d… V  qd S r:   r;   r<   r   r   r	   r>   ­   s     z!prepare_inputs.<locals>.<genexpr>c                 S   s   |   || ¡ dt|ƒd ¡S )Nr   r*   )Úbroadcast_toÚflattenr3   )r   Z
batch_dimsZinvariant_dimsr   r   r	   Úbatch_broadcast_and_squash°   s     
ÿz2prepare_inputs.<locals>.batch_broadcast_and_squash©r6   c                    s"   g | ]}ˆ |ˆ|j d d… ƒ‘qS )r   Nr;   r<   ©ri   Zbatch_dims_broadcastedr   r	   rX   ¿   s    )Úcrow_indicesre   Úcol_indicesr9   Úvaluesr   r?   r   )ÚbsrZdense_tensorsrl   rm   rn   rA   r   rk   r	   Úprepare_inputs¥   s2    $  ÿ  ÿ  ÿÿrp   c                 G   s|   t | |f|žŽ }| ¡  |d ¡}| ¡  |d ¡}| ¡  || ¡ jdd …  ¡}||jdd …  }tj|||||jd�S )Nrj   rf   r   ©Úsizer   )	rB   rl   rg   rm   rn   r   r   Úsparse_compressed_tensorr   )r   ro   rA   Zbatch_shaperl   rm   rn   rr   r   r   r	   Úbroadcast_batch_dims_bsrÆ   s     rt   c                 C   sH   | j �^ }}}|||d  |d ||d  |d g }|  |¡ dd¡S )Nr   r*   rf   r   )r   ÚreshapeÚ	transpose)r   r0   Úrestr   r!   Z	new_shaper   r   r	   Útile_to_blocksizeÑ   s    

ürx   )ÚOptionalÚTuple)ÚIS_BETA_ZEROÚBLOCKSIZE_ROWÚBLOCKSIZE_COLÚTILE_KÚ	acc_dtypeÚ
allow_tf32c            5   	   C   s2  t jdd�} t jdd�}!|||   ||!  }"t  |"¡}#t  |"| ¡}$|$|# }%|%dkrXd S t  d|¡}&t  d|¡}'|||   |	|#  |
|&d d …d f   ||'d d d …f   }(|||   ||#  })|||   ||!  ||&d d …d f   }*|||   ||'d d d …f   }+t  d|¡},t|%ƒD �]}-t j||f|d�}.t  |)¡}/td||ƒD ]’}0|0|, }1|1|k }2t j|*||1d d d …f   |2d d d …f dd�}3t j|+||/  ||1d d …d f   |2d d …d f dd�}4|.t j|3|4|d�7 }.�qJ|�rî|.| 9 }.n| |. |t  |(¡  }.t  |(|. |j	j
¡¡ |(|	7 }(|)|7 })�qd S )Nr*   ©Zaxisr   ©r#   ç        ©ÚmaskÚother©r€   )ÚtlÚ
program_idÚloadÚarangerL   ÚzerosÚdotÚstoreÚtor#   Ú
element_ty)5ÚalphaÚbetar{   r|   r}   Úkr~   Ú
values_ptrÚvalues_batch_strideÚvalues_nnz_strideÚvalues_row_block_strideÚvalues_col_block_strideÚcrow_indices_ptrÚcrow_indices_batch_strideÚcrow_indices_strideÚcol_indices_ptrÚcol_indices_batch_strideÚcol_indices_strideZmat1_ptrZmat1_batch_strideZmat1_tiled_row_strideZmat1_tiled_col_strideZmat1_row_block_strideZmat1_col_block_strideZmat2_ptrZmat2_batch_strideZmat2_tiled_row_strideZmat2_tiled_col_strideZmat2_row_block_strideZmat2_col_block_strider   r€   Ú	batch_pidÚrow_block_pidÚcrow_indices_offset_ptrÚ
nnz_offsetÚnnz_offset_nextÚrow_nnzÚrow_block_arangeÚcol_block_arangeÚvalues_block_ptrsÚcol_index_nnz_ptrZmat1_block_ptrsZmat2_block_ptrsZk_tile_arangeÚ_Z	acc_blockÚ	col_blockZk_tileZ	k_offsetsZmask_kZ
mat1_blockZ
mat2_blockr   r   r	   Ú_sampled_addmm_kernelá   s–    #ÿþÿ
ÿþýüÿ	ÿþÿÿþýÿ	ÿþÿ
ÿ ýÿþ ü
r«   )r|   r}   r   r€   ÚGROUP_SIZE_ROWc           0      C   s   t jdd�}t jdd�}t jdd�}t jdd�}t jdd�} t  |||| |¡\}}|||  |	|  }!t  |!¡}"t  |!|	 ¡}#|#|" }$|$dkr’d S t  d| ¡}%t  d|¡}&|||  ||"  ||%d d …d f   ||&d d d …f   }'|||  ||  ||&d d …d f   ||%d d d …f   }(|||  ||  ||  ||%d d …d f   ||%d d d …f   })|
||  ||"  }*t j| | f|d�}+t|$ƒD ]P},t  |'¡}-t  |*¡}.t  |(||.  ¡}/|+t j|-|/|d�7 }+|'|7 }'|*|7 }*�q”t  	|)|+ 
|jj¡¡ d S )Nr   r�   r   r*   r‚   r‡   )rˆ   r‰   Znum_programsZ	swizzle2drŠ   r‹   rŒ   rL   r�   rŽ   r�   r#   r�   )0r|   r}   r”   r•   r–   r—   r˜   r™   rš   r›   rœ   r�   rž   Z	dense_ptrZdense_batch_strideZdense_tiled_row_strideZdense_tiled_col_strideZdense_row_block_strideZdense_col_block_strideZ
output_ptrZoutput_batch_strideZoutput_tiled_row_strideZoutput_tiled_col_strideZoutput_row_block_strideZoutput_col_block_strider   r€   r¬   rŸ   r    Zcol_block_pidÚn_block_rowsÚn_block_colsr¡   r¢   r£   r¤   r¥   r¦   r§   Zdense_block_ptrsZoutput_ptrsr¨   Zoutput_acc_blockr©   Zvalues_blockZdense_row_idxZdense_blockr   r   r	   Ú"_bsr_strided_dense_rowspace_kernel\  s�    )    ÿÿþÿ
ÿþýüÿÿþýüÿ
ÿþýüûÿÿþÿ

r¯   c              
      sÎ   |  d¡}|  d¡d }|  d¡}	||	|f}
|d k	rht|d d… d d d… ƒddt|d d… ƒ   }nd }|d|d|d|d	|d
i}|jtjtjfkr¢tj‰ d‰n
tj	‰ d‰‡ ‡‡fdd„}t
|||
|ƒ d S )Nr   r6   r*   rf   rK   r   ©r   NN©r   Nr6   )r   rf   N)r   rf   éüÿÿÿTFc                    s&   t |  ˆt|Ž žˆ ˆddddœŽ d S )Né   r*   )r   r€   r¬   Ú
num_stagesÚ	num_warps)r¯   rJ   ©r\   rc   ©r   r€   r0   r   r	   rb   ò  s    þùz*_run_dense_rowspace_kernel.<locals>.kernel©rr   r'   r3   r#   r   r$   r%   rˆ   Úfloat32Zfloat64rd   )r0   rn   rl   rm   ÚdenseÚoutputÚmax_gridÚ	n_batchesr­   r®   rO   rP   rT   rb   r   r·   r	   Ú_run_dense_rowspace_kernelØ  s2    


4     ûr¾   c              
      sÌ   |  d¡}|  d¡d }||f}|d k	r\t|d d… d d d… ƒddt|d d… ƒ   }nd }|d|d|d|	d|
di}|jtjtjfkr–tj‰ d	‰n
tj	‰ d
‰‡ ‡‡‡‡‡‡‡fdd„}t
||||ƒ d S )Nr   r6   r*   r   r   )r   N)r   r6   )r   r²   TFc                    s2   t |  ˆˆˆfˆˆˆft|Ž žˆ ˆdddœŽ d S )Nr*   r³   )r   r€   r´   rµ   )r«   rJ   r¶   ©r   r€   r‘   r’   r0   Úis_beta_zeror“   Útile_kr   r	   rb     s       ÿ  þýùz)_run_sampled_addmm_kernel.<locals>.kernelr¸   )r‘   r’   rÀ   r0   r“   rÁ   rn   rl   rm   Úmat1Úmat2r¼   r½   r­   rO   rP   rT   rb   r   r¿   r	   Ú_run_sampled_addmm_kernel   s0    
4     ûrÄ   g      ð?F)r’   r‘   ÚoutÚskip_checksr¼   )ÚinputrÂ   rÃ   rÅ   rÆ   r¼   c                C   sn  d}t || ƒ t|| ||ƒ}	|�s&t||| jƒ t||| jƒ |dkrh| jtjkrhtd|› d|› d�ƒ | jtjk	r’t||| jƒ t||| jƒ nt|||jƒ t	|||ƒ |d k	�r&t ||ƒ t|||jƒ t||| jƒ t|j
|	j
koø| ¡ |  ¡ k|› d|	j
› d|	 ¡ › d|j
› d	| ¡ › �	ƒ |d k�rB|	j|jd
d�}n
| |	¡ | ¡ dk�sh| ¡ dk�rl|S | ¡ j
dd … }
| d¡}| d¡}| d¡}|dk�s°|dk�rÂ| ¡  |¡ |S |}t|||ƒ\}}}}}t||
d |fƒ}t|||
d fƒ}t|
Ž }t|||dk|
||||||||ƒ | ¡  ¡ dd … | ¡ dd … k�rj| ¡  | | ¡ j
¡¡ |S )NÚsampled_addmmrƒ   Fz(): having beta == z3 not equal to 0.0 with boolean mask is not allowed.z!(): Expects `out` to be of shape z and with nnz equal to z but got out.shape = z and out.nnz = T)Úcopyr   r   r6   r*   rf   )r   rt   r   r   r#   r   Úboolr   r(   r"   r   Ú_nnzr�   Zcopy_Únumelrn   rr   Úmul_rp   rx   r_   rÄ   r7   ru   )rÇ   rÂ   rÃ   r’   r‘   rÅ   rÆ   r¼   r   Zinput_broadcastedr0   r   r!   r“   Ú
out_backuprl   rm   rn   rÁ   r   r   r	   rÈ   +  sv    
þ

*ý




       û&rÈ   )rÅ   rÆ   r¼   )ro   rº   rÅ   rÆ   r¼   c                C   s¾  d}|s’t || ƒ t|| |jƒ t|| |jƒ t|| |ƒ |  d¡}| d¡}|  ¡ jdd … \}}	t	||  d|› d|› d�ƒ t
|||	fƒ n$| jdd … \}}
|jdd … \}}t|| |ƒ}|d k	�r|�s|||f }t	|j|kd|› d|j› d�ƒ t	| ¡ �p| dd¡ ¡ d	ƒ |d k�r:| |||f ¡}|  ¡ d
k�rP| ¡ S |  ¡ jdd … }|}t| ||ƒ\}}}}}t||d d d… ƒ}t||d
 |d
 fƒ}t|||||||ƒ |S )NÚbsr_dense_mmr   r6   z"bsr_dense_mm(): dense.size(-1) == z( should be divisible by blocksize[0] == r   z9bsr_dense_mm(): `out` argument has wrong shape, expected z
, but got z›bsr_dense_mm(): only row-major/col-major `out` arguments are supported, i.e. (out.is_contiguous() or out.transpose(-2, -1).is_contiguous()) should be True.r   )r   r   r   r(   r#   r"   rr   rn   r   r   r5   rB   Úis_contiguousrv   Z	new_emptyrË   Zzero_rp   rx   r¾   )ro   rº   rÅ   rÆ   r¼   r   r   r!   Ú	row_blockrª   r   r    Zoriginal_batch_dims_broadcastedZexpected_out_shaper0   rÎ   rl   rm   rn   r   r   r	   rÏ     sL    


þþþ
rÏ   )ÚMAX_ROW_NNZÚTILEc                 C   sV  t jdd�}t jdd�}t jdd�}| ||  ||  }t  |¡}t  || ¡}|| }|dkrdd S t  d|
¡}||| k }|||  ||  ||  }t j|| |tdƒ d� t j¡}t j|dd�}t|
|	|
ƒD ]Z}||
7 }||| k }t j|| |tdƒ d� t j¡}t j|dd�}t  	||k||¡}qÔt  
|| ¡}t j|dd�}t|
|	|
ƒD ]\}||
8 }||| k }t j|| |tdƒ d� t j¡}t  
|| ¡}|t j|dd�7 }�qXt j|| ||  |jj¡|d� t|
|	|
ƒD ]l}||
7 }||| k }t j|| |tdƒ d� t j¡}t  
|| ¡}t j|| ||  |jj¡|d� �qäd S )Nr   r�   r*   r   Úinfr„   )r…   )rˆ   r‰   rŠ   r‹   r&   r�   r¹   r_   rL   ÚwhereÚexpÚsumrŽ   r#   r�   )r™   rš   r›   r”   r•   r—   Zvalues_nnz_col_block_striderÑ   rª   rÒ   rÓ   rŸ   Zrow_block_offset_pidr    r¡   r¢   r£   r¤   Z
row_aranger…   Zcurr_row_values_ptrsZrow_tileZmax_row_valuer©   Zcurr_max_row_valueÚnumÚdenomr   r   r	   Ú_bsr_softmax_kernelÎ  s^    ÿþÿ
ÿþýÿ"""""rÚ   c                    s�  d}t || ƒ t|| | jƒ |  ¡ dks4|  ¡ dkr<|  ¡ S | jdd … \}}|  ¡ }|  ¡ jdd … \‰‰ ˆd kr€t 	|¡‰n
t 	ˆ¡‰|  
¡  d¡ dd¡}|  ¡  dd¡ ¡ rÂ|  ¡  ¡ }n|  ¡ }| dd¡ ¡  d¡ dd¡ dˆ|ˆ  ¡}|jd ˆ|ˆ f}d }	|dd d…f d|d	i}
‡ ‡‡fd
d„}t||
||	ƒ | dˆ|ˆ ¡ dd¡j|  ¡ jŽ }tj|  
¡  ¡ |  ¡  ¡ || j| jd�S )NÚbsr_softmaxr   r   rf   r²   r6   .r±   r°   c                    s&   t |  t|Ž ˆˆ ˆtdˆƒfžŽ  d S )Ni   )rÚ   rJ   rV   r¶   ©rª   Úmax_row_nnzrÑ   r   r	   rb   <  s     ûzbsr_softmax.<locals>.kernelrq   )r   r(   r#   rË   rÌ   Úcloner   rn   r   Znext_power_of_2rl   re   rh   rv   rÐ   r8   ru   rd   r   rs   rm   r   )rÇ   rÝ   r   r   r!   Znnzrl   rn   rO   rP   rT   rb   r   rÜ   r	   rÛ     sB    

,  ü	$

ûrÛ   rƒ   )ÚqueryÚkeyÚvalueÚ	attn_maskÚ	dropout_pÚ	is_causalÚscalec           
      C   sl  d}t | |› d�ƒ t |d k	|› d�ƒ |d k	s6t‚t |jtjk|› dtj› d|j› d�ƒ t||| jƒ t||| jƒ t||| jƒ t||| jƒ t||| jƒ |jtj	k	r¾t||| jƒ t
|| | dd¡d	d
d�}|d krî|  d¡dksø|d	k�rt d
|› d|› d�ƒ |d k�r,dt |  d¡¡ n|}	| ¡  |	¡ t|ƒ}tjjj| ¡ |dd� t||ƒ}|S )NÚ_scaled_dot_product_attentionz'(): is_causal == True is not supported.z'(): attn_mask == None is not supported.z(): attn_mask.layout must be z, but got attn_mask.layout == r   r   r6   rƒ   F)r’   rÆ   r   z(): current value of scale == z results in division by zero.r*   T)ÚpZinplace)r   r4   r   r   r   r   r   r(   r#   rÊ   rÈ   rv   rr   ÚmathÚsqrtrn   rÍ   rÛ   ÚnnZ
functionalZdropoutrÏ   )
rß   rà   rá   râ   rã   rä   rå   r   ZsdpaZscale_factorr   r   r	   ræ   Q  sB    	þþ
þ þ"
ræ   )N)N)rƒ   FN)+rè   r   Ztorch._inductor.cuda_propertiesr   r
   r   r   r   r"   r(   r5   r9   rB   rE   rI   rJ   r]   rd   rp   rt   rx   r   Ztriton.languageÚlanguagerˆ   Útypingry   rz   ZjitZ	constexprr«   r¯   r¾   rÄ   ZTensorrÊ   ÚintrÈ   rÏ   rÚ   rÛ   r&   ræ   r   r   r   r	   Ú<module>   s¨   
	
!àz"Ú{(0÷÷XúúO
öF
A   ùù2