§
    Ë²ÅjgZ ã            !       óp	  — d dl Z d dlZd dlZd dlmZ d dlmZmZ d dlZd dl	m
Z
 d dlmZ ddlmZ  e ej        dd	¦  «        ¦  «        Zd
„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zdud„Zd„ Zd„ Z d„ Z!d„ Z"ddœd„Z#	 	 	 	 	 	 dvd„Z$	 	 	 	 	 	 	 	 dwd„Z% G d„ d ¦  «        Z& ee¬!¦  «        d"„ ¦   «         Z'	 dxd$„Z(dyd%„Z)dddddd&ddd'œd(ej*        d)ej*        d*ej*        d+eej*                 d,eej*                 d-eej*                 d.e+d/eeee         ee         ee         f                  d0ee,         fd1„Z-dddddd&ddd'œd(ej*        d)ej*        d*ej*        d+eej*                 d,eej*                 d-eej*                 d.e+d/eeee         ee         ee         f                  d0ee,         fd2„Z. e¦   «         �réd dl/Z/d dl0m1Z2 e/j3        d3e2j4        d4e2j4        d5e2j4        d6e2j4        d7e2j4        d8e2j4        fd9„¦   «         Z5e/j3        d4e2j4        d5e2j4        d7e2j4        d8e2j4        d:e2j4        f
d;„¦   «         Z6d<„ Z7d=d=dd&dd>œd(ej*        d?ej*        d@ej*        d-eej*                 d.e+d/eeee         ee         ee         f                  fdA„Z8dd&dddBœd)ej*        d*ej*        d-eej*                 d.e+d/eeee         ee         ee         f                  d0ee,         fdC„Z9e/j3        dDe2j4        dEe2j4        fdF„¦   «         Z:dudG„Z;	 	 	 dzdIej*        dJej*        dKej*        dLeej*                 dMe<dNe+dOee<         fdP„Z=e/j3        dQe2j4        dRe2j4        dSe2j4        dTe2j4        dUe2j4        dVe2j4        d8e2j4        fdW„¦   «         Z>dXej*        dYej*        dZej*        d[ej*        d\ej*        f
d]„Z?e/j3        d^e2j4        d_e2j4        dTe2j4        d`e2j4        dUe2j4        dVe2j4        dae2j4        d8e2j4        fdb„¦   «         Z@	 d{dXej*        dYej*        ddej*        deej*        dfej*        dgej*        d0e,d\ej*        dhe+fdi„ZAe/j3        dje2j4        dke2j4        dle2j4        dme2j4        dne2j4        doe2j4        dpe2j4        dqe2j4        dre2j4        d4e2j4        d5e2j4        dse2j4        d7e2j4        d8e2j4        d:e2j4        d`e2j4        f dt„¦   «         ZBdS dZ;dZ9dZ8dZ=dZ?dZAdZBdS )|é    N)Ú	lru_cache)ÚOptionalÚTuple)Ú	warn_once)Ú
has_tritoné   )Úget_metaÚ*TORCH_SPARSE_BSR_SCATTER_MM_LRU_CACHE_SIZEé   c                 ó(   — | st          |¦  «        ‚d S ©N)Ú
ValueError)ÚcondÚmsgs     úIC:\Projects\Historical\venv\Lib\site-packages\torch/sparse/_triton_ops.pyÚcheckr      s    € Øð Ý˜‰oŒoÐðð ó    c                 óR   — t          |j        t          j        k    | › d�¦  «         d S )Nz@(): only BSR sparse format is supported for the sparse argument.)r   ÚlayoutÚtorchÚ
sparse_bsr)Úf_nameÚts     r   Úcheck_bsr_layoutr      s7   € Ý	Ø	Œ•EÔ$Ò$ØÐSÐSÐSñô ð ð ð r   c                 ó^   — t          |j        |k    o|j        j        dk    | › d�¦  «         d S )NÚcudaz9(): all inputs are expected to be on the same GPU device.)r   ÚdeviceÚtype)r   r   r   s      r   Úcheck_devicer   !   sB   € Ý	Ø	Œ�FÒÐ6˜qœxœ}°Ò6ØÐLÐLÐLñô ð ð ð r   c           	      ó^  — t          |                     ¦   «         dk    o|                     ¦   «         dk    | › d|                     ¦   «         › d|                     ¦   «         › d�¦  «         |j        dd …         \  }}|j        dd …         \  }}t          ||k    | › d|› d|› d�¦  «         d S )Nr   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Ú_ns          r   Úcheck_mm_compatible_shapesr+   (   sõ   € Ý	Ø�Š‰	Œ	�QŠÐ)˜3Ÿ7š7™9œ9¨š>Øð 	Jð 	JØ #§¢¡	¤	ð	Jð 	JØ=@¿WºW¹Y¼Yð	Jð 	Jð 	Jñô ð ð ŒY�r�s�sŒ^�F€BˆØŒY�r�s�sŒ^�F€Bˆå	Ø
ˆbŠØð 	Rð 	RØ "ð	Rð 	RØLNð	Rð 	Rð 	Rñô ð ð ð r   c           	      ó¼   — t          |j        |k    o3|j        t          j        t          j        t          j        ft          |Ž z   v | › 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-   Úadditional_dtypess       r   Úcheck_dtyper3   9   s‚   € Ý	Ø	Œ�5Òð 	SØŒGÝŒZ�œ­¬Ð5½Ð?PÐ8QÑQðSàð 	'ð 	'Ø3Dð	'ð 	'àœGð	'ð 	'ð 	'ñ	ô ð ð ð r   c           	      óž   ‡— t          |¦  «        dk    sJ ‚d„ Šˆfd„}t           ||¦  «        | › d|d         › d|d         › d�¦  «         d S )	Nr   c                 ó   — | | dz
  z   S ©Nr   © )Úvs    r   Úis_power_of_twoz(check_blocksize.<locals>.is_power_of_twoG   s   € Ø˜˜Q™‘KÐ Ð r   c                 ó<   •— d}| D ]}|dk    o ‰|¦  «        o|}Œ|S )NTé   r7   )ÚbÚresÚ	blocksizer9   s      €r   Úis_compatible_blocksizez0check_blocksize.<locals>.is_compatible_blocksizeJ   s=   ø€ ØˆØð 	Kð 	KˆIà ’?ÐA  °yÑ'AÔ'AÐJÀsˆCˆCØˆ
r   z(): sparse inputs' blocksize (r   z, r   z;) should be at least 16 and a power of 2 in each dimension.)Úlenr   )r   r>   r?   r9   s      @r   Úcheck_blocksizerA   D   s§   ø€ Ýˆy‰>Œ>˜QÒÐÐÐð!ð !ð !ðð ð ð ð õ 
ØÐ 	Ñ*Ô*Øð 	Dð 	D°¸1´ð 	Dð 	DÀÈ1Äð 	Dð 	Dð 	Dñô ð ð ð r   c                 óx   — t          |                      ¦   «         ¦  «        dk    r|                      ¦   «         S | S )a¹  Return input as a triton-contiguous tensor.

    A triton-contiguous tensor is defined as a tensor that has strides
    with minimal value smaller than or equal to 1.

    While triton kernels support triton-non-contiguous tensors (all
    strides being greater than 1) arguments, a considerable slow-down
    occurs because tensor data is copied element-wise rather than
    chunk-wise. Zero strides is assumed to not have this defect.
    r   )ÚminÚstrideÚ
contiguous)r   s    r   Úmake_triton_contiguousrF   X   s2   € õ ˆ1�8Š8‰:Œ:�„˜ÒÐð �|Š|‰~Œ~Ðàˆr   c                 ó|   — 	 t          j        d„ |D ¦   «         Ž S # t          $ r t          d| › d�¦  «         Y d S w xY w)Nc              3   ó4   K  — | ]}|j         d d…         V — Œd S ©Nr"   ©r$   ©Ú.0r   s     r   ú	<genexpr>z'broadcast_batch_dims.<locals>.<genexpr>m   s,   è è € Ð'FÐ'F¸¨¬°°°¬Ð'FÐ'FÐ'FÐ'FÐ'FÐ'Fr   Fz3(): inputs' batch dimensions are not broadcastable!)r   Úbroadcast_shapesÚ	Exceptionr   )r   Útensorss     r   Úbroadcast_batch_dimsrQ   k   sf   € ðUÝÔ%Ð'FÐ'F¸gÐ'FÑ'FÔ'FÐGÐGøÝð Uð Uð UÝˆe˜ÐSÐSÐSÑTÔTÐTÐTÐTÐTðUøøøs   ‚ š;º;c              '   ó|   K  — |D ]6}t          d ¦  «        g|                     ¦   «         z  }||| <   ||         V — Œ7d S r   )Úslicer#   )r#   Úslice_rangerP   r   Úslicess        r   ÚslicerrV   r   sQ   è è € Øð ð ˆÝ˜‘+”+� §¢¡¤Ñ(ˆØ!ˆˆs‰Ø�Œiˆˆˆˆðð r   c              '   ó¬   K  — |D ]N}t          d ¦  «        g|                     ¦   «         z  }t          | |¦  «        D ]\  }}|�|||<   Œ||         V — ŒOd S r   )rS   r#   Úzip)ÚdimsrU   rP   r   ÚsÚdÚd_slices          r   Úmultidim_slicerr]   y   st   è è € Øð ð ˆÝ�4‰[Œ[ˆM˜AŸEšE™GœGÑ#ˆÝ˜d FÑ+Ô+ð 	ð 	‰JˆAˆwØˆ}Ø��!‘øØ�Œdˆ
ˆ
ˆ
ˆ
ðð r   c               '   óP   K  — | D ] }|V — |                      ¦   «         E d {V —† Œ!d S r   )rD   )rP   r   s     r   Úptr_stride_extractorr_   ‚   sK   è è € Øð ð ˆØˆˆˆØ—8’8‘:”:ÐÐÐÐÐÐÐÐðð r   c              #   ón  ‡ ‡‡K  — dt          ‰ ¦  «        cxk    rdk    sn J ‚dt          ‰¦  «        cxk    rdk    sn J ‚dd l}ˆ ˆfd„}ˆfd„} |j         |¦   «         Ž D ]Q}d„ t          ‰ |‰¦  «        D ¦   «         }d„ t          ||¦  «        D ¦   «         }|d d d…         g ||¦  «        ¢R V — ŒRd S )Nr   é   c               3   ó^   •K  — t          ‰‰¦  «        D ]\  } }t          d| |¦  «        V — Œd S )Nr   )rX   Úrange)ÚfgÚmgÚ	full_gridÚgrid_blockss     €€r   Úgenerate_grid_pointsz.grid_partitioner.<locals>.generate_grid_pointsŽ   sI   øè è € Ý˜) [Ñ1Ô1ð 	#ð 	#‰FˆB�Ý˜˜2˜rÑ"Ô"Ð"Ð"Ð"Ð"ð	#ð 	#r   c              3   ó€   •K  — ‰                      ¦   «         D ]%\  }}t          t          || |¦  «        ¦  «        V — Œ&d S r   )ÚitemsÚnextr]   )rU   r   Út_dimsÚtensor_dims_maps      €r   Úgenerate_sliced_tensorsz1grid_partitioner.<locals>.generate_sliced_tensors’   sS   øè è € Ø(×.Ò.Ñ0Ô0ð 	;ð 	;‰IˆAˆvÝ• v¨v°qÑ9Ô9Ñ:Ô:Ð:Ð:Ð:Ð:ð	;ð 	;r   c                 ó<   — g | ]\  }}}t          ||z
  |¦  «        ‘ŒS r7   )rC   )rL   rd   Úgpre   s       r   ú
<listcomp>z$grid_partitioner.<locals>.<listcomp>—   s9   € ð 
ð 
ð 
Ù!+  R¨�C��R‘˜ÑÔð
ð 
ð 
r   c                 ó:   — g | ]\  }}t          |||z   ¦  «        ‘ŒS r7   )rS   )rL   rp   Úgs      r   rq   z$grid_partitioner.<locals>.<listcomp>š   s*   € ÐGÐGÐG©¨¨A•%˜˜B ™FÑ#Ô#ÐGÐGÐGr   éÿÿÿÿ)r@   Ú	itertoolsÚproductrX   )	rf   rg   rm   ru   rh   rn   Ú
grid_pointÚgridrU   s	   ```      r   Úgrid_partitionerry   ˆ   sO  øøøè è € Ø•�I‘”Ð#Ð#Ò#Ð# !Ò#Ð#Ð#Ð#Ð#Ð#Ø•�KÑ Ô Ð%Ð%Ò%Ð% AÒ%Ð%Ð%Ð%Ð%Ð%àÐÐÐð#ð #ð #ð #ð #ð #ð;ð ;ð ;ð ;ð ;ð (�iÔ'Ð)=Ð)=Ñ)?Ô)?Ð@ð ;ð ;ˆ
ð
ð 
Ý/2°9¸jÈ+Ñ/VÔ/Vð
ñ 
ô 
ˆð HÐGµ°ZÀÑ1FÔ1FÐGÑGÔGˆð �4�4�R�4ŒjÐ:Ð2Ð2°6Ñ:Ô:Ð:Ð:Ð:Ð:Ð:Ð:ð;ð ;r   c                 ó¾   ‡— dd d d…         }|€|}n,d„ Št          ˆfd„t          ||¦  «        D ¦   «         ¦  «        }t          |||¦  «        D ]^}} | |g|¢R Ž  Œd S )N)iÿÿÿéÿÿ  r{   rt   c                 óF   — | €|S t          dt          | |¦  «        ¦  «        S r6   )ÚmaxrC   )rs   re   s     r   Úvalid_grid_dimz%launch_kernel.<locals>.valid_grid_dim¨   s&   € ØˆyØ�	õ ˜1�c ! R™jœjÑ)Ô)Ð)r   c              3   ó6   •K  — | ]\  }} ‰||¦  «        V — Œd S r   r7   )rL   rs   re   r~   s      €r   rM   z launch_kernel.<locals>.<genexpr>¯   sD   øè è € ð 
ð 
Ù&+ a¨ˆNˆN˜1˜bÑ!Ô!ð
ð 
ð 
ð 
ð 
ð 
r   )r1   rX   ry   )Úkernelrm   rf   rg   Úcuda_max_gridrx   Úsliced_tensorsr~   s          @r   Úlaunch_kernelrƒ   ¡   s¾   ø€ à.¨t¨t°¨tÔ4€MØÐØ#ˆˆð	*ð 	*ð 	*õ ð 
ð 
ð 
ð 
Ý/2°;ÀÑ/NÔ/Nð
ñ 
ô 
ñ 
ô 
ˆõ "2Ø�; ñ"ô "ð &ð &Ðˆˆ~ð 	ˆˆtÐ%�nÐ%Ð%Ð%Ð%Ð%ð&ð &r   c                 ó  ‡‡— |                       ¦   «                              d¦  «        }|                      ¦   «                              d¦  «        }t          |                      ¦   «                              d¦  «        ¦  «        }d„ |D ¦   «         }t          j        |j        d d…         gd„ |D ¦   «         ¢R Ž Šd„ Š ‰|‰d¦  «        } ‰|‰d¦  «        } ‰|‰|j        dd …         ¦  «        }ˆˆfd„|D ¦   «         }|||g|¢R S )Nr   c                 óR   — g | ]$}t          |                     d ¦  «        ¦  «        ‘Œ%S ©r   )rF   Ú	unsqueezerK   s     r   rq   z"prepare_inputs.<locals>.<listcomp>¾   s+   € ÐMÐMÐM¸!Õ% a§k¢k°!¡n¤nÑ5Ô5ÐMÐMÐMr   éýÿÿÿc              3   ó4   K  — | ]}|j         d d…         V — Œd S rI   rJ   rK   s     r   rM   z!prepare_inputs.<locals>.<genexpr>Â   s,   è è € Ð;Ð;¨a˜QœW S b Sœ\Ð;Ð;Ð;Ð;Ð;Ð;r   c                 óz   — |                       ||z   ¦  «                             dt          |¦  «        dz
  ¦  «        S )Nr   r   )Úbroadcast_toÚflattenr@   )r   Ú
batch_dimsÚinvariant_dimss      r   Úbatch_broadcast_and_squashz2prepare_inputs.<locals>.batch_broadcast_and_squashÇ   s;   € Ø�~Š~˜j¨>Ñ9Ñ:Ô:×BÒBØ�s�:‰Œ Ñ"ñ
ô 
ð 	
r   ©rt   c           	      óD   •— g | ]} ‰|‰|j         d d…         ¦  «        ‘ŒS )r"   NrJ   )rL   r   r�   Úbatch_dims_broadcasteds     €€r   rq   z"prepare_inputs.<locals>.<listcomp>Ô   sC   ø€ ð ð ð àð 	#Ð" 1Ð&<¸a¼gÀbÀcÀc¼lÑKÔKðð ð r   )Úcrow_indicesr‡   Úcol_indicesrF   Úvaluesr   rN   r$   )ÚbsrÚdense_tensorsr“   r”   r•   rP   r�   r’   s         @@r   Úprepare_inputsr˜   ¹   s`  øø€ à×#Ò#Ñ%Ô%×/Ò/°Ñ2Ô2€LØ—/’/Ñ#Ô#×-Ò-¨aÑ0Ô0€KÝ# C§J¢J¡L¤L×$:Ò$:¸1Ñ$=Ô$=Ñ>Ô>€FØMÐM¸}ÐMÑMÔM€Gõ #Ô3ØŒ�S�b�SÔðØ;Ð;°7Ð;Ñ;Ô;ðð ð Ðð
ð 
ð 
ð
 .Ð-ØÐ,¨eñô €Lð -Ð,¨[Ð:PÐRWÑXÔX€KØ'Ð'ØÐ&¨¬°R°S°SÔ(9ñô €Fðð ð ð ð àðñ ô €Gð
 ˜ fÐ6¨wÐ6Ð6Ð6r   c                 ó¸  — t          | |g|¢R Ž }|                     ¦   «                              |dz   ¦  «        }|                     ¦   «                              |dz   ¦  «        }|                     ¦   «                              ||                     ¦   «         j        dd …         z   ¦  «        }||j        dd …         z   }t          j        |||||j        ¬¦  «        S )Nr�   rˆ   r"   ©Úsizer   )	rQ   r“   r‹   r”   r•   r$   r   Úsparse_compressed_tensorr   )r   r–   rP   Úbatch_shaper“   r”   r•   r›   s           r   Úbroadcast_batch_dims_bsrrž   Ü   sÎ   € Ý& v¨sÐ=°WÐ=Ð=Ð=€Kà×#Ò#Ñ%Ô%×2Ò2°;ÀÑ3FÑGÔG€LØ—/’/Ñ#Ô#×0Ò0°¸uÑ1DÑEÔE€KØ�ZŠZ‰\Œ\×&Ò& {°S·Z²Z±\´\Ô5GÈÈÈÔ5LÑ'LÑMÔM€FØ˜œ 2 3 3œÑ'€DÝÔ)Ø�k 6°¸S¼Zðñ ô ð r   c                 ó¸   — | j         �^ }}}|||d         z  |d         ||d         z  |d         gz   }|                      |¦  «                             dd¦  «        S )Nr   r   rˆ   r"   )r$   ÚviewÚ	transpose)r   r>   ÚrestÚmÚnÚ	new_shapes         r   Útile_to_blocksizer¦   é   sh   € Ø”'�K€Tˆ1ˆaØØ	ˆY�qŒ\ÑØ�!ŒØ	ˆY�qŒ\ÑØ�!Œð	ñ €Ið �6Š6�)ÑÔ×&Ò& r¨2Ñ.Ô.Ð.r   c                 óâ   — | j         dk     r |                      d¦  «        } | j         dk     ° | j         dk    r|                      d| j         dz
  ¦  «        } | j         dk    sJ | j        ¦   «         ‚| S )zÙReturn tensor as 3D tensor by either prepending new dimensions to
    the tensor shape (when ``tensor.ndim < 3``), or by collapsing
    starting dimensions into the first dimension (when ``tensor.ndim >
    3``).
    ra   r   )Úndimr‡   rŒ   r$   )Útensors    r   Ú	as1Dbatchrª   ö   su   € ð Œ+˜Š/ˆ/Ø×!Ò! !Ñ$Ô$ˆð Œ+˜Š/ˆ/à„{�Q‚€Ø—’  6¤;°¡?Ñ3Ô3ˆØŒ;˜!ÒÐÐ˜Vœ\ÑÔÐØ€Mr   ©Úaccumulatorsc                ó2  — |d         }| j         dk    sJ ‚| j        \  }}}|dk    �r|dd…         \  }}	|j         dk    sJ ‚|j        \  }
}}||k    sJ ‚|€5|j        d         dz
  }t          j        |||f| j        | j        ¬¦  «        }n|j        \  }}}||k    sJ ‚||k    sJ ‚|dz  s|dz  s|dz  st          €qt          |j        d         dz
  ¦  «        D ]R}||         }||dz            }t          ||¦  «        D ],}|	|         \  }}||xx         | |         ||         z  z  cc<   Œ-ŒSnt          | |||	|¦  «         |S |dk    �rE|j        }t          |¦  «        }|j        \  }}}||z  dk    sJ ‚|dd…         \  }}}}}|d	         }|€`|| 	                    ¦   «          
                    ¦   «         dz   |z  z   } t          j        g |dd
…         ¢| ‘|‘R | j        | j        ¬¦  «        }n|j        d
d…         \  } }!|!|k    sJ ‚|j        }"t          |¦  «        }||z  }|dz  s|dz  s|dz  st          �€(|                     ¦   «          t          |¦  «        D �]}#t          |j        d         ¦  «        D ]ä}||          
                    ¦   «         }$||          
                    ¦   «         }||dz             
                    ¦   «         }t          |$|¦  «        \  }%}&||#|%|%|z   …|&|&|z   …f         }'t          ||¦  «        D ]X}||         ||         }}t          | 
                    ¦   «         |¦  «        \  }(})|'| |         ||#|(|(|z   …|)|)|z   …f         z  z  }'ŒYŒå�Œnt          | |||||||¦  «         |                     |"¦  «        S |dk    �rt|j        }t          |¦  «        }|j        \  }}}||z  dk    sJ ‚|dd…         \  }}}}|d	         }|€`|| 	                    ¦   «          
                    ¦   «         dz   |z  z   } t          j        g |dd
…         ¢| ‘|‘R | j        | j        ¬¦  «        }n|j        d
d…         \  } }!|!|k    sJ ‚|j        }"t          |¦  «        }||z  }|dz  s|dz  s|dz  st          �€7t          |¦  «        D �]%}#t          t          |¦  «        ¦  «        D �]}*t          ||*          
                    ¦   «         |¦  «        \  }%}&|%|z  }+|&|z  },||+          
                    ¦   «         }-||+dz             
                    ¦   «         }.||#|%|%|z   …|&|&|z   …f         }'t!          t          |-|.¦  «        ¦  «        D ]b\  }/}||,|.z  ||,z
  |-z  z   |/z             
                    ¦   «         }t          ||¦  «        \  }(})|'| |         ||#|(|(|z   …|)|)|z   …f         z  z  }'Œc�Œ�Œ'n7t          j        d|j        |j        ¬¦  «        }t          | |||||||¦  «         |                     |"¦  «        S t%          |¦  «        ‚)ad  Scattered matrix multiplication of tensors.

    A scattered matrix multiplication is defined as a series of matrix
    multiplications applied to input tensors according to the input
    and output mappings specified by indices data.

    The following indices data formats are supported for defining a
    scattered matrix multiplication operation (:attr:`indices_data[0]`
    holds the name of the indices data format as specified below):

    - ``"scatter_mm"`` - matrix multiplications scattered in batches
      of tensors.

      If :attr:`blocks` is a :math:`(* 	imes M 	imes K) tensor,
      :attr:`others` is a :math:`(* 	imes K 	imes N)` tensor,
      :attr:`accumulators` is a :math:`(* 	imes M 	imes N)` tensor,
      and :attr:`indices = indices_data['indices']` is a :math:`(*
      	imes 3)` tensor, then the operation is equivalent to the
      following code::

        c_offsets, pq = indices_data[1:]
        for r in range(len(c_offsets) - 1):
            for g in range(c_offsets[r], c_offsets[r + 1]):
                p, q = pq[g]
                accumulators[r] += blocks[p] @ others[q]

    - ``"bsr_strided_mm"`` - matrix multiplications scattered in
      batches of tensors and a tensor.

      If :attr:`blocks` is a :math:`(Ms 	imes Ks) tensor,
      :attr:`others` is a :math:`(* 	imes K 	imes N)` tensor,
      :attr:`accumulators` is a :math:`(* 	imes M 	imes N)` tensor, then
      the operation is equivalent to the following code::

        c_indices, r_offsets, p_offsets, q_offsets, meta = indices_data[1:]
        for b in range(nbatches):
            for i, r in enumerate(r_offsets):
                r0, r1 = divmod(r, N)
                acc = accumulators[b, r0:r0 + Ms, r1:r1 + Ns]
                for g in range(c_indices[i], c_indices[i+1]):
                    p = p_offsets[g]
                    q0, q1 = divmod(q_offsets[g], N)
                    acc += blocks[p] @ others[b, q0:q0 + Ks, q1:q1 + Ns]

      where ``Ns = N // meta['SPLIT_N']``, and ``M`` and ``K`` are
      integer multiples of ``Ms`` and ``Ks``, respectively.

    - ``"bsr_strided_mm_compressed"`` - matrix multiplications
      scattered in batches of tensors and a tensor. A memory and
      processor efficient version of ``"bsr_strided_mm"`` format.  If
      :attr:`blocks` is a :math:`(Ms 	imes Ks) tensor, :attr:`others`
      is a :math:`(* 	imes K 	imes N)` tensor, :attr:`accumulators`
      is a :math:`(* 	imes M 	imes N)` tensor, then the operation is
      equivalent to the following code::

        c_indices, r_offsets, q_offsets, meta = indices_data[1:]
        for b in range(nbatches):
            for r in r_offsets:
                m = (r // N) // Ms
                n = (r % N) // Ns
                r0, r1 = divmod(r, N)
                c0, c1 = c_indices[m], c_indices[m + 1]
                acc = accumulators[b, r0:r0 + Ms, r1:r1 + Ns]
                for i, p in enumerate(range(c0, c1)):
                    q = q_offsets[n * c1 + (SPLIT_N - n) * c0 + i]
                    q0, q1 = divmod(q, N)
                    acc += blocks[p] @ others[b, q0:q0 + Ks, q1:q1 + Ns]

      where ``Ns = N // meta['SPLIT_N']``, and ``M`` and ``K`` are
      integer multiples of ``Ms`` and ``Ks``, respectively.

      Notice that the order of ``r_offsets`` items can be arbitrary;
      this property enables defining swizzle operators via
      rearrangements of ``r_offsets`` items..

    Auxilary functions are provided for pre-computing
    :attr:`indices_data`. For example,
    :func:`bsr_scatter_mm_indices_data` is used to define indices data
    for matrix multiplication of BSR and strided tensors.

    Parameters
    ----------
    blocks (Tensor): a 3-D tensor of first matrices to be multiplied

    others (Tensor): a tensor of second matrices to be multiplied. If
      ``indices_data[0]=="scatter_mm"``, the tensor is a 1-D batch
      tensor of second input matrices to be multiplied. Otherwise, the
      second input matrices are slices of the :attr:`others` tensor.
    indices_data (tuple): a format data that defines the inputs and
      outputs of scattered matrix multiplications.

    Keyword arguments
    -----------------

    accumulators (Tensor, optional): a tensor of matrix product
      accumulators. If ``indices_data[0]=="scatter_mm"``, the tensor
      is a 1-D batch tensor of output matrices. Otherwise, output
      matrices are slices of the :attr:`accumulators` tensor.
    r   ra   Ú
scatter_mmr   N©r-   r   r;   Úbsr_strided_mmÚSPLIT_Nr"   Úbsr_strided_mm_compressedr†   )r¨   r$   r   Úzerosr-   r   Ú_scatter_mm2rc   rª   r}   ÚitemÚ_scatter_mm6Úzero_Údivmodr    r@   Ú	enumerateÚemptyÚNotImplementedError)0ÚblocksÚothersÚindices_datar¬   Úindices_formatÚ_PÚMsÚKsÚ	c_offsetsÚpqÚ_QÚKs_ÚNsÚRÚMs_ÚNs_ÚrÚg0Úg1rs   ÚpÚqÚothers_shapeÚBÚKÚNÚ	c_indicesÚ	r_offsetsÚ	p_offsetsÚ	q_offsetsÚmetar±   ÚMÚN_Úaccumulators_shaper<   Úr_Úr0Úr1ÚaccÚq0Úq1Újr£   r¤   Úc0Úc1Úis0                                                   r   r®   r®     sƒ  € ðH " !”_€NàŒ;˜!ÒÐÐÐØ”�J€BˆˆBà˜Ò%Ñ%Ø$ Q R RÔ(‰ˆ	�2àŒ{˜aÒÐÐÐØ”l‰ˆˆC�Ø�SŠyˆyˆyˆyàÐØ” Ô" QÑ&ˆAÝ œ;Ø�B˜� 6¤<¸¼ðñ ô ˆLˆLð 'Ô,‰KˆAˆs�CØ˜"’9�9�9�9Ø˜"’9�9�9�9à�‰7ð 	F�b˜2‘gð 	F  b¡ð 	F­LÐ,@Ý˜9œ?¨1Ô-°Ñ1Ñ2Ô2ð =ð =�Ø˜q”\�Ø˜q 1™uÔ%�Ý˜r 2™œð =ð =�AØ˜aœ5‘D�A�qØ  �O�O”O v¨a¤y°6¸!´9Ñ'<Ñ<�O�O‘O�Oð=ð=õ ˜ ¨°B¸ÑEÔEÐEØÐà	Ð+Ò	+Ñ	+Ø”|ˆÝ˜6Ñ"Ô"ˆà”,‰ˆˆ1ˆaØ�2‰v˜Š{ˆ{ˆ{ˆ{à;GÈÈÈÔ;KÑ8ˆ	�9˜i¨°DØ�y”/ˆàÐØ�i—m’m‘o”o×*Ò*Ñ,Ô,¨qÑ0°QÑ6Ñ6ˆAÝ œ;Ø*�,˜s ˜sÔ#Ð* QÐ*¨Ð*Ð*°&´,ÀvÄ}ðñ ô ˆLˆLð !Ô& r s sÔ+‰EˆAˆrØ˜’7�7�7�7à)Ô/ÐÝ  Ñ.Ô.ˆà�'‰\ˆà�‰7ð 	�b˜2‘gð 	  b¡ð 	­LÑ,@Ø×ÒÑ Ô Ð Ý˜1‘X”Xð 
Qñ 
Q�Ý˜yœ¨qÔ1Ñ2Ô2ð 	Qð 	Q�AØ" 1œ×*Ò*Ñ,Ô,�BØ" 1œ×*Ò*Ñ,Ô,�BØ" 1 q¡5Ô)×.Ò.Ñ0Ô0�BÝ# B¨™]œ]‘F�B˜Ø& q¨"¨r°B©w¨,¸¸RÀ"¹W¸Ð'DÔE�CÝ" 2 r™]œ]ð Qð Q˜Ø(¨œ|¨Y°q¬\˜1˜Ý!'¨¯ª©¬°!Ñ!4Ô!4™˜˜BØ˜v aœy¨6°!°R¸"¸r¹'°\À2ÈÈRÉÀ<Ð2OÔ+PÑPÑP˜˜ðQñ	Qð
Qõ ØØØØØØØØñ	ô 	ð 	ð × Ò Ð!3Ñ4Ô4Ð4à	Ð6Ò	6Ñ	6Ø”|ˆÝ˜6Ñ"Ô"ˆà”,‰ˆˆ1ˆaØ�2‰v˜Š{ˆ{ˆ{ˆ{à0<¸Q¸R¸RÔ0@Ñ-ˆ	�9˜i¨Ø�y”/ˆàÐØ�i—m’m‘o”o×*Ò*Ñ,Ô,¨qÑ0°QÑ6Ñ6ˆAÝ œ;Ø*�,˜s ˜sÔ#Ð* QÐ*¨Ð*Ð*°&´,ÀvÄ}ðñ ô ˆLˆLð !Ô& r s sÔ+‰EˆAˆrØ˜’7�7�7�7à)Ô/ÐÝ  Ñ.Ô.ˆà�'‰\ˆà�‰7ð 	�b˜2‘gð 	  b¡ð 	­LÑ,@Ý˜1‘X”Xð Qñ Q�Ý�s 9™~œ~Ñ.Ô.ð 
Qñ 
Q�AÝ# I¨a¤L×$5Ò$5Ñ$7Ô$7¸Ñ;Ô;‘F�B˜Ø˜b™�AØ˜b™�AØ" 1œ×*Ò*Ñ,Ô,�BØ" 1 q¡5Ô)×.Ò.Ñ0Ô0�BØ& q¨"¨r°B©w¨,¸¸RÀ"¹W¸Ð'DÔE�CÝ )­%°°B©-¬-Ñ 8Ô 8ð Qð Q™˜˜1Ø% a¨"¡f°¸!±¸rÑ/AÑ&AÀAÑ&EÔF×KÒKÑMÔM˜Ý!'¨¨1¡¤™˜˜BØ˜v aœy¨6°!°R¸"¸r¹'°\À2ÈÈRÉÀ<Ð2OÔ+PÑPÑP˜˜ñQñ
QðQõ œØ˜IœO°IÔ4Dðñ ô ˆIõ ØØØØØØØØñ	ô 	ð 	ð × Ò Ð!3Ñ4Ô4Ð4õ " .Ñ1Ô1Ð1r   c           
      ó.  — ||||	|
|hd hk    �rBt           j                             ¦   «         }t          d| ||||f|dt           j        df¬¦  «        }|� |j        d$i |¤Ž |S | ||fdk    rX||fdk    rd}d}d}d	}d}
d	}	�nÏ||fd
k    rd}d}d}d	}d}
d	}	�n¹||fdk    rd}d}d}d	}d}
d	}	�n£||fdk    rd}d}d}d}d}
d	}	�n�| ||fdk    rX||fdk    rd}d}d}d}d}
d}	�nn||fd
k    rd}d}d}d	}d}
d}	�nX||fdk    rd	}d}d}d	}d}
d	}	�nB||fdk    rd}d}d}d	}d}
d	}	�n,| ||fdk    rj||fdk    rd	}d}d}d}d}
d}	�n||fd
k    rd}d}d}d}d}
d}	nø||fdk    rd}d}d}d	}d}
d}	nã||fdk    rd}d}d}d	}d}
d	}	nÎ||fdk    rd}d}d}d}d}
d	}	n¹| ||fdk    ri||fdk    rd	}d}d}d}d}
d}	n›||fd
k    rd	}d}d}d	}d}
d}	n†||fdk    rd	}d}d}d	}d}
d	}	nq||fdk    rd}d}d}d	}d}
d	}	n\||fdk    rd	}d}d}d}d}
d	}	nG| ||fdk    r>||fdk    rd}d}d}d}d}
d}	n)||fd
k    rd}d}d}d}d}
d}	n||fdk    rd}d}d}d}d}
d	}	|€.ddd	dddddddœ	                     |d¦  «        }|dk    r|dk    rd}||z  }|€t          |dk     rdnd|¦  «        }|€t          |dk     rdnd|¦  «        }|
pd}
|	€¥t          | |¦  «        dk    rddddœ                     |d	¦  «        }	nvt          | |¦  «        dk    rddddœ                     |d	¦  «        }	nGt          | |¦  «        dk    rdd	dœ                     |d	¦  «        }	ndddœ                     |d	¦  «        }	|pd	}||k    sJ t          ||¬¦  «        ¦   «         ‚||k    sJ t          ||¬¦  «        ¦   «         ‚|| k    sJ t          | |¬ ¦  «        ¦   «         ‚||k    sJ t          ||¬!¦  «        ¦   «         ‚||k    sJ t          ||¬"¦  «        ¦   «         ‚t          d$||||
|	|d#œ|¤ŽS )%Nr®   r   ç      à?©Úversion)é   rê   rê   )r;   r;   r   r;   é   )é    rì   r   rì   )é@   rí   )é€   rî   )é   rï   rï   é   rí   rî   )é   rñ   rñ   )rê   rê   )é   rò   rò   )é   ró   ró   rê   )	r;   rì   rí   rî   rê   rï   rñ   ró   i    rï   rò   rñ   )r;   rì   rí   )r;   rì   )ÚTILE_MrÁ   )ÚTILE_NrÇ   )rÙ   rÁ   )rÓ   rÇ   )rÒ   rÂ   )rô   rõ   Ú
GROUP_SIZEÚ
num_stagesÚ	num_warpsr±   r7   )	r   r   Úget_device_namer	   Úfloat16ÚupdateÚgetrC   Údict)rÙ   rÒ   rÓ   rÁ   rÂ   rö   rô   rõ   r±   rø   r÷   ÚextraÚdevice_namerØ   rÇ   s                  r   Úscatter_mm_metar   ô  sú  € ð 	�˜ ¨J¸
ÐCÈÀvÒMÑMÝ”j×0Ò0Ñ2Ô2ˆÝØØ��1�b˜"ÐØØ�œ sÐ+ð	
ñ 
ô 
ˆð ÐØˆDŒKÐ Ð ˜%Ð Ð Ð ØˆKð ˆq�!ˆ9˜
Ò"Ð"Ø�Bˆx˜8Ò#Ð#Ø�Ø�Ø�Ø�
Ø�
Ø�	‘	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	‘	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	‘	Ø�b�˜ZÒ'Ð'Ø�Ø�Ø�Ø�
Ø�
Ø�	ùØ��AˆY˜*Ò$Ð$Ø�Bˆx˜8Ò#Ð#Ø�Ø�Ø�Ø�
Ø�
Ø�	‘	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	‘	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	‘	Ø�b�˜ZÒ'Ð'Ø�Ø�Ø�Ø�
Ø�
Ø�	ùØ��AˆY˜+Ò%Ð%Ø�Bˆx˜8Ò#Ð#Ø�Ø�Ø�Ø�
Ø�
Ø�	‘	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	�	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	�	Ø�b�˜ZÒ'Ð'Ø�Ø�Ø�Ø�
Ø�
Ø�	�	Ø�b�˜ZÒ'Ð'Ø�Ø�Ø�Ø�
Ø�
Ø�	øØ��AˆY˜+Ò%Ð%Ø�Bˆx˜8Ò#Ð#Ø�Ø�Ø�Ø�
Ø�
Ø�	�	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	�	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	�	Ø�b�˜ZÒ'Ð'Ø�Ø�Ø�Ø�
Ø�
Ø�	�	Ø�b�˜ZÒ'Ð'Ø�Ø�Ø�Ø�
Ø�
Ø�	øØ��AˆY˜+Ò%Ð%Ø�Bˆx˜8Ò#Ð#Ø�Ø�Ø�Ø�
Ø�
Ø�	�	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	�	Ø�b�˜XÒ%Ð%Ø�Ø�Ø�Ø�
Ø�
Ø�	à€ð ØØØØØØØØð

ð 

÷ Š#ˆa�‰*Œ*ð 	ð �Š9ˆ9˜˜dš˜ØˆGØ	
ˆg‰€BØ€~Ý˜2 š8˜8�R�R¨¨RÑ0Ô0ˆØ€~Ý˜2 š8˜8�R�R¨¨RÑ0Ô0ˆØ�˜q€JØÐÝˆq�!‰9Œ9�tÒÐØ A¨1Ð-Ð-×1Ò1°"°aÑ8Ô8ˆIˆIÝ��A‰YŒY˜$ÒÐØ A¨1Ð-Ð-×1Ò1°"°aÑ8Ô8ˆIˆIÝ��A‰YŒY˜#ÒÐØ A˜˜×*Ò*¨2¨qÑ1Ô1ˆIˆIà A˜˜×*Ò*¨2¨qÑ1Ô1ˆIØ�˜q€Jà�RŠ<ˆ<ˆ<� V°Ð3Ñ3Ô3‰<Œ<ˆ<Ø�RŠ<ˆ<ˆ<� V°Ð3Ñ3Ô3‰<Œ<ˆ<Ø�Š7ˆ7ˆ7•D˜1 Ð$Ñ$Ô$‰7Œ7ˆ7Ø�Š7ˆ7ˆ7•D˜1 Ð$Ñ$Ô$‰7Œ7ˆ7Ø�Š7ˆ7ˆ7•D˜1 Ð$Ñ$Ô$‰7Œ7ˆ7åð ØØØØØØðð ð ðð ð r   c                 ó°  — |€t           j        }|€|}|€d}||	|
|hd hk    �rƒt           j                             ¦   «         }| |||||dk    |dk    |dk    f}||u r|}n||f}t	          d|||||f¬¦  «        }|€|dk    rt	          d||||df¬¦  «        }|€||urt	          d||||df¬¦  «        }|€·t	          dg |d d…         ¢d‘|dd …         ¢R |||df¬¦  «        }|€1||ur-t	          dg |d d…         ¢d‘|dd …         ¢R |||df¬¦  «        }t          |pi ¦  «        D ]E}||         }|d         }|d	         }||z  }||z  dk    r||k    rt          |¦  «        }||z  |d	<   ŒF|� |j        di |¤Ž |S t          d
| ›d|›d|›d|›d|›d|›d|›d|›d|›�¦  «         |pt          ||z  d¦  «        }|pd}|
pd}
|	pd}	t          d|||
|	dœ|¤ŽS )Nrç   r   r   Úbsr_dense_addmmrè   r   Ú*ra   r±   z@bsr_dense_addmm uses non-optimal triton kernel parameters for M=z K=z N=z Ms=z, Ks=z beta=z alpha=z dtype=z out_dtype=rë   )r±   ÚGROUP_SIZE_ROWr÷   rø   r7   )
r   rú   r   rù   r	   Úsortedrý   rû   r   r}   )rÙ   rÒ   rÓ   rÁ   rÂ   ÚbetaÚalphar±   r  rø   r÷   Úsparsityr-   Ú	out_dtypeÚ_versionrþ   rÿ   ÚkeyÚversion_dtyperØ   Úmatching_metaÚmkeyÚmeta_r¤   Úsplit_nÚcs                             r   Úbsr_dense_addmm_metar  â  sf  € ð* €}Ý”ˆØÐØˆ	ØÐØˆØ�˜J¨Ð7¸D¸6ÒAÑAÝ”j×0Ò0Ñ2Ô2ˆØ�!�Q˜˜B ¨¢	¨4°1ª9°e¸q²jÐAˆØ�IÐÐØ!ˆMˆMà! 9Ð,ˆMÝØØØØ˜}¨hÐ7ð	
ñ 
ô 
ˆð ˆ<˜H¨šO˜OÝØ!ØØØ! =°#Ð6ð	ñ ô ˆDð ˆ<˜E¨Ð2Ð2ÝØ! 3¨¸hÈÈsÐ=Sðñ ô ˆDð ˆ<å$Ø!Ø)�#�b�q�b”'Ð)˜3Ð)  Q R R¤Ð)Ð)ØØ! =°#Ð6ð	ñ ô ˆMð Ð$¨°iÐ)?Ð)?Ý (Ø%Ø-�c˜"˜1˜"”gÐ-˜sÐ- S¨¨¨¤WÐ-Ð-ØØ% u¨cÐ2ð	!ñ !ô !�õ ˜}Ð2°Ñ3Ô3ð -ð -�Ø% dÔ+�Ø˜”G�Ø 	Ô*�Ø˜‘L�Ø�q‘5˜A’:�: ! q¢& &Ý ™;œ;�DØ&'¨1¡f�D˜‘OøØÐØˆDŒKÐ Ð ˜%Ð Ð Ð ØˆKõ
 ð[Øð[ð [Øð[ð [Ø"#ð[ð [Ø')ð[ð [Ø.0ð[ð [Ø48ð[ð [Ø<Að[ð [ØEJð[ð [ØNWð[ð [ñô ð ð
 Ð(�˜Q "™W a™œ€GØ#Ð( q€NØ�˜q€JØ�˜Q€IÝð ØØ%ØØð	ð ð
 ðð ð r   c                   ó:   — e Zd ZdZd„ Zd„ Zd„ Zed„ ¦   «         ZdS )ÚTensorAsKeyaS  A light-weight wrapper of a tensor that enables storing tensors as
    keys with efficient memory reference based comparision as an
    approximation to data equality based keys.

    Motivation: the hash value of a torch tensor is tensor instance
    based that does not use data equality and makes the usage of
    tensors as keys less useful. For instance, the result of
    ``len({a.crow_indices(), a.crow_indices()})`` is `2`, although,
    the tensor results from `crow_indices` method call are equal, in
    fact, these share the same data storage.
    On the other hand, for efficient caching of tensors we want to
    avoid calling torch.equal that compares tensors item-wise.

    TensorAsKey offers a compromise in that it guarantees key equality
    of tensors that references data in the same storage in the same
    manner and without accessing underlying data. However, this
    approach does not always guarantee correctness. For instance, for
    a complex tensor ``x``, we have ``TensorAsKey(x) ==
    TensorAsKey(x.conj())`` while ``torch.equal(x, x.conj())`` would
    return False.
    c                 ó\  — d„ }t          j        |¦  «        | _        |j        t          j        u r ||¦  «        | _        nÒ|j        t          j        t          j        hv r@ || 	                    ¦   «         ¦  «         || 
                    ¦   «         ¦  «        f| _        ns|j        t          j        t          j        hv r@ ||                     ¦   «         ¦  «         ||                     ¦   «         ¦  «        f| _        nt          |j        ¦  «        ‚t!          | j        ¦  «        | _        d S )Nc                 óÞ   — | j         j        s| j         j        rJ | j         ¦   «         ‚|                      ¦   «         |                      ¦   «         | j        |                      ¦   «         | j         fS r   )r-   Úis_floating_pointÚ
is_complexÚdata_ptrÚstorage_offsetr$   rD   )Úobjs    r   Úget_tensor_keyz,TensorAsKey.__init__.<locals>.get_tensor_key]  sb   € ð œ	Ô3ÐW°s´yÔ7KÐWÐWÈcÌiÑWÔWÐLà—’‘”Ø×"Ò"Ñ$Ô$Ø”	Ø—
’
‘”Ø”	ðð r   )ÚweakrefÚrefÚ_obj_refr   r   Ústridedr  Ú
sparse_csrr   r“   r”   Ú
sparse_cscÚ
sparse_bscÚccol_indicesÚrow_indicesr»   ÚhashÚ_hash)Úselfr  r  s      r   Ú__init__zTensorAsKey.__init__\  s  € ð	ð 	ð 	õ&  œ CÑ(Ô(ˆŒØŒ:�œÐ&Ð&Ø%�~ cÑ*Ô*ˆDŒHˆHØŒZ�EÔ,­eÔ.>Ð?Ð?Ð?à�˜s×/Ò/Ñ1Ô1Ñ2Ô2Ø�˜sŸšÑ0Ô0Ñ1Ô1ðˆDŒHˆHð ŒZ�EÔ,­eÔ.>Ð?Ð?Ð?à�˜s×/Ò/Ñ1Ô1Ñ2Ô2Ø�˜sŸšÑ0Ô0Ñ1Ô1ðˆDŒHˆHõ
 & c¤jÑ1Ô1Ð1Ý˜$œ(‘^”^ˆŒ
ˆ
ˆ
r   c                 ó   — | j         S r   )r'  ©r(  s    r   Ú__hash__zTensorAsKey.__hash__�  s
   € ØŒzÐr   c                 ót   — t          |t          ¦  «        sdS | j        �|j        €| |u S | j        |j        k    S )NF)Ú
isinstancer  r  r  )r(  Úothers     r   Ú__eq__zTensorAsKey.__eq__„  sC   € Ý˜%¥Ñ-Ô-ð 	Ø�5ØŒ8Ð˜uœyÐ0ð ˜5�=Ð ØŒx˜5œ9Ò$Ð$r   c                 ó*   — |                       ¦   «         S )z'Return object if alive, otherwise None.)r  r+  s    r   r  zTensorAsKey.obj�  s   € ð �}Š}‰ŒÐr   N)	Ú__name__Ú
__module__Ú__qualname__Ú__doc__r)  r,  r0  Úpropertyr  r7   r   r   r  r  E  sg   € € € € € ðð ð,#$ð #$ð #$ðJð ð ð%ð %ð %ð ðð ñ „Xðð ð r   r  )Úmaxsizec	           	      ó|	  — |j         }	|	€J ‚|	                     ¦   «         |	                     ¦   «         }}
|
j        }t          j        }| dk    �rd||z  }g }t	          j        |||¬¦  «        |z  }t          ||z  ¦  «        D ]�}|
|                              ¦   «         }|
|dz                                 ¦   «         }||k    rŒ@| 	                    |||…         ||z  z   
                    |¦  «        |                     ||z
  ¦  «        z   ¦  «         Œ�t	          j        |¦  «        }|
                     ¦   «         }|                     ¦   «         }|||z  z  }||z                        d¦  «        }|
}||                              |¦  «        }|                     dd¬¦  «        \  }}||         }| |||fS | dk    �rÄ||z  }g }g }t	          j        |||¬¦  «        |z  }t          ||z  ¦  «        D ]Í}|
|                              ¦   «         }|
|dz                                 ¦   «         }||k    rŒ@| 	                    t	          j        ||||¬¦  «         
                    |¦  «        ¦  «         | 	                    |||…         ||z  z   
                    |¦  «        |                     ||z
  ¦  «        z   ¦  «         ŒÎt	          j        |¦  «        }|
                     ¦   «         }|                     ¦   «         }|||z  z  }||z                        d¦  «        }t	          j        |
d d…         t	          j        ||                              |¦  «        d¦  «        f¦  «        }t	          j        |¦  «        }| ||||fS | d	k    �r'|}dg}g }t          |¦  «        D ]á}t          ||z  ¦  «        D ]Ì}|
|                              ¦   «         }|
|dz                                 ¦   «         }t          ||z  ¦  «        D ]€}| 	                    |d         |z   |z
  ¦  «         t          ||z
  ¦  «        D ]J} || z   }!||!                              ¦   «         |||z  z  z   ||z  z  |z   }"| 	                    |!|"g¦  «         ŒKŒ�ŒÍŒâ| t	          j        |||¬¦  «        t	          j        |||¬¦  «        fS t'          d
| ›d�¦  «        ‚)Nr²   r¯   r   rt   T)Ú
descendingÚstabler°   r   r®   zInvalid indices_format=z>. Expected bsr_strided_mm_compressed|bsr_strided_mm|scatter_mm)r  r“   r”   r   r   Úint32Úarangerc   rµ   ÚappendÚrepeatÚrepeat_interleaveÚcatÚdiffÚnonzeror    ÚsortÚcumsumr©   r   )#r¿   rÙ   rÒ   rÓ   rÁ   rÂ   Únbatchesr±   Úcompressed_sparse_tensor_as_keyr–   r“   r”   r   Úindices_dtyperÇ   Úq_offsets_lstr<   r£   rÝ   rÞ   r×   Úcrow_indices_diffÚnon_zero_row_indicesÚarÕ   rÔ   Únnz_per_rowÚindicesÚp_offsets_lstrÖ   Ú
pq_offsetsr¤   r   rÎ   rÏ   s#                                      r   Ú_bsr_scatter_mm_indices_datarP  “  sI  € ð *Ô
-€CØˆ?ˆ?ˆ?Ø #× 0Ò 0Ñ 2Ô 2°C·O²OÑ4EÔ4E�+€LØÔ €FÝ”K€MàÐ4Ò4Ñ4Ø�'‰\ˆØˆÝŒL˜¨¸fÐEÑEÔEÈÑJˆÝ�q˜B‘w‘”ð 	ð 	ˆAØ˜a”×%Ò%Ñ'Ô'ˆBØ˜a !™eÔ$×)Ò)Ñ+Ô+ˆBØ�RŠxˆxØØ× Ò Ø˜R ˜UÔ# r¨A¡vÑ.×6Ò6°wÑ?Ô?Ø×%Ò% b¨2¡gÑ.Ô.ñ/ñô ð ð õ ”I˜mÑ,Ô,ˆ	Ø(×-Ò-Ñ/Ô/ÐØ0×8Ò8Ñ:Ô:ÐØ  B¨¡FÑ+ˆØ˜‘U—L’L Ñ$Ô$ˆ	Ø ˆ	à'Ð(<Ô=×OÒOÐPWÑXÔXˆØ*×/Ò/¸4ÈÐ/ÑMÔMÑˆ�WØ˜gÔ&ˆ	Ø 	¨9°iÐ@Ð@à	Ð+Ò	+Ñ	+Ø�'‰\ˆØˆØˆÝŒL˜¨¸fÐEÑEÔEÈÑJˆÝ�q˜B‘w‘”ð 	ð 	ˆAØ˜a”×%Ò%Ñ'Ô'ˆBØ˜a !™eÔ$×)Ò)Ñ+Ô+ˆBØ�RŠxˆxØØ× Ò Ý”˜R ¨=ÀÐHÑHÔH×OÒOÐPWÑXÔXñô ð ð × Ò Ø˜R ˜UÔ# r¨A¡vÑ.×6Ò6°wÑ?Ô?Ø×%Ò% b¨2¡gÑ.Ô.ñ/ñô ð ð õ ”I˜mÑ,Ô,ˆ	Ø(×-Ò-Ñ/Ô/ÐØ0×8Ò8Ñ:Ô:ÐØ  B¨¡FÑ+ˆØ˜‘U—L’L Ñ$Ô$ˆ	Ý”Ià˜R˜a˜RÔ Ý”Ø%Ð&:Ô;×MÒMÈgÑVÔVØñô ðñ
ô 
ˆ	õ ”I˜mÑ,Ô,ˆ	Ø 	¨9°iÀÐKÐKà	˜<Ò	'Ñ	'ØˆØ�Cˆ	Øˆ
å�x‘”ð 		2ð 		2ˆAÝ˜1 ™7‘^”^ð 2ð 2�Ø! !”_×)Ò)Ñ+Ô+�Ø! ! a¡%Ô(×-Ò-Ñ/Ô/�Ý˜q B™w™œð 2ð 2�AØ×$Ò$ Y¨r¤]°RÑ%7¸"Ñ%<Ñ=Ô=Ð=Ý" 2¨¡7™^œ^ð 2ð 2˜Ø ™F˜Ø(¨œ^×0Ò0Ñ2Ô2°Q¸!¸r¹'±]ÑBÀqÈBÁwÑOÐRSÑS˜Ø"×)Ò)¨1¨a¨&Ñ1Ô1Ð1Ð1ð2ð2ð2ð ÝŒL˜¨-ÀÐGÑGÔGÝŒL˜¨=ÀÐHÑHÔHð
ð 	
õ Øf�~ÐfÐfÐfñ
ô 
ð 	
r   r²   c                 óš  — |                       ¦   «         dk    sJ ‚| j        dk    sJ ‚|                      ¦   «         j        dd…         }| j        \  }}|\  }}|j        dd…         \  }	}
|	|k    sJ ‚|j        dd…                              ¦   «         }t          |||
||fi |¤Ž}d|vr3|                     | j        t          j	        t          j
        hv ¬¦  «         |d         }t          ||||
||||t          | ¦  «        ¦	  «	        }|dk    r|                     d	¬
¦  «         ||fz   S |dk    r|                     d¬
¦  «         ||fz   S |S )zkComputes indices data for :func:`scatter_mm` used in BSR and
    strided tensor matrix multiplication.
    r   r   r"   NÚ
allow_tf32©rR  r±   r²   T)Úis_compressedr°   F)Ú	dense_dimr¨   r•   r$   Únumelr   rû   r-   r   rú   r/   rP  r  )r–   r/  r¿   Ú
meta_inputr>   rÙ   rÒ   rÁ   rÂ   ÚK_rÓ   rE  rØ   r±   r¾   s                  r   Úbsr_scatter_mm_indices_datarY  ô  sw  € ð �=Š=‰?Œ?˜aÒÐÐÐØŒ8�qŠ=ˆ=ˆ=ˆ=Ø—
’
‘”Ô" 2 3 3Ô'€IØŒ9�D€A€qØ�F€BˆØŒK˜˜˜Ô�E€BˆØ�Š7ˆ7ˆ7ˆ7ØŒ{˜3˜B˜3Ô×%Ò%Ñ'Ô'€Hå˜1˜a  B¨Ð9Ð9¨jÐ9Ð9€DØ˜:Ð%Ð%Ø�Š˜sœy­U¬]½E¼NÐ,KÐKˆÑLÔLÐLØ�9Œo€GÝ/Ø˜˜1˜a  R¨°7½KÈÑ<LÔ<Lñô €Lð Ð4Ò4Ð4Ø�Š $ˆÑ'Ô'Ð'Ø˜t˜gÑ%Ð%Ø	Ð+Ò	+Ð	+Ø�Š %ˆÑ(Ô(Ð(Ø˜t˜gÑ%Ð%àÐr   c           
      óÜ  — | j         dk    sJ ‚|j         dk    sJ ‚| j        d         | j        d         |j        d         }}}|                      ¦   «         j        dd…         }|€t          | |d¬¦  «        }|d         }|€5t	          j        g |j        dd…         ¢|‘|‘R | j        | j        ¬¦  «        }|j        }	t          |¦  «        }|  	                    ¦   «         dk    r| 
                    ¦   «          �nß|d	v r;| 
                    ¦   «          t          |                      ¦   «         |||¬
¦  «         �n |dk    �rŠ|j        dd…                              ¦   «         }
t	          j        |
|z  |d         z  |z  |d         z  |d         |d         f| j        | j        ¬¦  «        }t          |¦  «                             dd¦  «                             |
||d         z  |d         ||d         z  |d         ¦  «                             dd¦  «                             dd¦  «        }t          |                      ¦   «         |||¬
¦  «         |                     |                     d|
||d         z  ||d         z  f¦  «                             dd¦  «                             |
||¦  «                             dd¦  «        ¦  «         nt+          |¦  «        ‚|                     |	¦  «        S )zBSR @ strided -> stridedr   r"   rt   Nr²   )r¿   r   r¯   >   r°   r²   r«   r®   r   )ra   r   rë   r   )r   r   ra   rë   )r¨   r$   r•   rY  r   rº   r-   r   rª   Ú_nnzr·   r®   rV  r³   r¡   r    ÚmovedimrŒ   Úcopy_Ú	unflattenÚreshaper»   )r–   r/  r¾   ÚoutrÁ   rÂ   rÇ   r>   r¿   Ú	out_shaperE  r¬   r½   s                r   Úbsr_scatter_mmrb    sø  € ð Œ8�qŠ=ˆ=ˆ=ˆ=ØŒ:˜Š?ˆ?ˆ?ˆ?à”˜2” ¤	¨"¤¨u¬{¸2¬ˆBˆ€BØ—
’
‘”Ô" 2 3 3Ô'€IàÐÝ2Ø�Ð'Bð
ñ 
ô 
ˆð " !”_€Nà
€{ÝŒkØ'ˆeŒk˜#˜2˜#ÔÐ' Ð' BÐ'Ð'¨s¬yÀÄð
ñ 
ô 
ˆð ”	€IÝ
�C‰.Œ.€Cà
‡x‚x�z„z�Q‚€Ø�	Š	‰Œˆ‰Ø	ÐJÐ	JÐ	JØ�	Š	‰ŒˆÝ�3—:’:‘<”< ¨À3ÐGÑGÔGÐGÑGØ	˜<Ò	'Ñ	'Ø”;˜s ˜sÔ#×)Ò)Ñ+Ô+ˆÝ”{à˜2‘ ¨1¤Ñ-°Ñ2°iÀ´lÑBØ˜!”Ø˜!”ðð
 ”)Ø”:ð
ñ 
ô 
ˆõ �eÑÔßŠY�r˜2ÑÔßŠTØØ�i ”lÑ"Ø˜!”Ø�i ”lÑ"Ø˜!”ñô ÷ ŠWØ˜lñô ÷ ŠW�Q˜‰]Œ]ð 	õ 	�3—:’:‘<”< ¨ÀLÐQÑQÔQÐQØ�	Š	Ø×"Ò"Ø�H˜b I¨a¤LÑ0°"¸	À!¼Ñ2DÐEñô ÷ ŠWØ˜lñô ÷ ŠW�X˜r 2Ñ&Ô&ßŠY�r˜2ÑÔñ		
ô 		
ð 		
ð 		
õ " .Ñ1Ô1Ð1à�8Š8�IÑÔÐr   F©r  r  Ú
left_alphaÚright_alphar`  Úskip_checksÚmax_gridrØ   Úinputr–   Údenserd  re  r`  rf  rg  rØ   c                óf  — |€–|j         t          j        u rƒd}|                     ¦   «         }|                     ¦   «         dz
  }|j        |         }|j        d         }t          |||¦  «        }t          j        |||fz   t          j        |j	        ¬¦  «        }t          | |||||||||	|
¬¦  «        S )NÚ_int_bsr_dense_addmmr   rt   r¯   rc  )r-   r   Úint8r“   r#   r$   rQ   rº   r;  r   r  )rh  r–   ri  r  r  rd  re  r`  rf  rg  rØ   r   r“   Ú
batch_ndimrÙ   rÓ   Úoriginal_batch_dims_broadcasteds                    r   rk  rk  \  sÒ   € ð €{�u”{¥e¤jÐ0Ð0Ø'ˆØ×'Ò'Ñ)Ô)ˆØ!×%Ò%Ñ'Ô'¨!Ñ+ˆ
ØŒI�jÔ!ˆØŒK˜ŒOˆÝ*>¸vÀsÈEÑ*RÔ*RÐ'ÝŒkØ+¨q°!¨fÑ4Ý”+Ø”<ð
ñ 
ô 
ˆõ
 ØØØØØØØØØØØðñ ô ð r   c                óž	  ‡‡‡
‡ ‡!‡"‡#‡$‡%— d}|                      ¦   «         }|                     ¦   «         }|                     ¦   «         }|                     ¦   «         dz
  }|j        ||dz   …         \  }}|j        |dz   |dz   …         }|j        d         }t          |||¦  «        }|€|                     |||fz   ¦  «        }|                     ¦   «         dk    s‰dk    s|dk    s|dk    s|dk    rM‰dk    r|                     ¦   «          n0| 	                    | ¦  «         ‰dk    r| 
                    ‰¦  «         |S dŠ$dŠ%|€'d	Š$ |                     d
¦  «        j        g |¢|‘|‘R Ž }n   |j        g |¢|‘d‘R Ž j        g |¢|‘|‘R Ž }|€'d	Š% |                     d
¦  «        j        g |¢|‘|‘R Ž }n   |j        g |¢d‘|‘R Ž j        g |¢|‘|‘R Ž }|                     ¦   «         d         dk    sJ ‚|                     ¦   «         d         dk    sJ ‚‰
€lt          d|                     ¦   «         |d         z  |d         z  ||z  z  z
  d¦  «        }t          ||||d         |d         ‰‰||j        |j        ¬¦
  «
        Š
|}t#          || ||||¦  «        \  }}}} }}}}|\  Š!Š ‰
                     d|‰!z  ¦  «        }||z  Š"|}t'          |‰!‰"f¦  «        }t'          |‰ ‰"f¦  «        }t'          | ‰!‰"f¦  «        } t'          |‰!‰"f¦  «        }t'          |‰!‰"f¦  «        }t(          j        t,          j        t(          j        t,          j        t(          j        t,          j        t(          j        t,          j        t(          j        t,          j        t(          j        t,          j        i|j                 Š#|                     d¦  «        }|                     d¦  «        dz
  }|                     d¦  «        }|||f}|	�?t;          |	dd…         ddd…         ¦  «        ddt=          |	dd…         ¦  «        z
  z  z   }nd}|d|d|d| d|d|d|d|di}‰dk    sJ ‚ˆ ˆ!ˆ"ˆˆˆ#ˆ$ˆ
ˆ%f	d„}t?          ||||¦  «         |                      ¦   «         |                      ¦   «         k    r-| 	                    |                     |j        ¦  «        ¦  «         |S )a  Compute

      out = beta * input + left_alpha.reshape(-1, 1) * (alpha * (bsr @ dense)) * right_alpha.reshape(1, -1)

    where left_alpha, right_alpha are (* + 1)-D tensors when
    specified, otherwise, these are treated as tensors filled with
    ones.
    r  r   r   ra   rt   Nr   FTr7   r"   )r  r-   r	  r±   rˆ   r   ©r   NN©r   Nrt   )r   rˆ   éüÿÿÿ)r   rˆ   Nc                 ó’   •	— t          |          g t          |Ž ¢‰‘‰‘R ‰dk    ‰dk    ‰dk    ‰‰
‰‰‰‰t          j        k    ‰dœ
‰	¤Ž d S )Nr   r   )
Úbeta_is_oneÚbeta_is_nonzeroÚalpha_is_oneÚleft_alpha_is_oneÚright_alpha_is_oneÚBLOCKSIZE_ROWÚBLOCKSIZE_INNERÚBLOCKSIZE_COLrR  Ú	acc_dtype)Ú_bsr_strided_addmm_kernelr_   ÚtlÚfloat32)rx   r‚   ÚBKÚBMÚBNr  r  Údot_out_dtyperw  rØ   rx  s     €€€€€€€€€r   r€   zbsr_dense_addmm.<locals>.kernel  s�   ø€ Ý! $Ô'ð 	
Ý! >Ð2ð	
àð	
ð ð	
ð 	
ð  š	Ø  AšIØ !šØ/Ø1ØØØØ$­¬
Ò2Ø#ð	
ð 	
ð ð	
ð 	
ð 	
ð 	
ð 	
r   )!r•   r“   r”   r#   r$   rQ   Ú	new_emptyr[  r·   r]  Úmul_Úexpandr    rD   Úroundr  r-   r˜   rü   r¦   r   rú   r~  r  r/   Úfloat64rl  r;  r›   r1   r@   rƒ   r  )&rh  r–   ri  r  r  rd  re  r`  rf  rg  rØ   r   r•   r“   r”   rm  rÙ   rÒ   r>   rÓ   rn  r  Ú
out_backupr±   Úout_untiledÚ	n_batchesÚn_block_rowsÚn_block_colsrf   rg   rm   r€   r€  r�  r‚  rƒ  rw  rx  s&      ``     `                     @@@@@@r   r  r  …  sÏ  øøøøøøøøø€ ð, €FØ�ZŠZ‰\Œ\€FØ×#Ò#Ñ%Ô%€LØ—/’/Ñ#Ô#€KØ×!Ò!Ñ#Ô# aÑ'€JØŒ9�Z *¨q¡.Ð0Ô1�D€A€qØ”˜Z¨!™^¨j¸1©nÐ<Ô=€IØŒ�BŒ€Aõ ';¸6À3ÈÑ&NÔ&NÐ#Ø
€{Ø�oŠoÐ=ÀÀAÀÑFÑGÔGˆà
‡x‚x�z„z�Q‚€˜% 1š*˜*¨¨Qª¨°!°q²&°&¸AÀºF¸FØ�1Š9ˆ9Ø�IŠI‰KŒKˆKˆKà�IŠI�eÑÔÐØ�qŠyˆyØ—’˜‘”�Øˆ
àÐØÐØÐØ ÐØ/�U—_’_ RÑ(Ô(Ô/ð 
Ø,ð
Ø./ð
Ø12ð
ð 
ð 
ˆ
ˆ
ð T�_�Z”_ÐLÐ&EÐLÀqÐLÈ!ÐLÐLÐLÔSð 
Ø,ð
Ø./ð
Ø12ð
ð 
ð 
ˆ
ð ÐØ!ÐØ0�e—o’o bÑ)Ô)Ô0ð 
Ø,ð
Ø./ð
Ø12ð
ð 
ð 
ˆˆð VÐ&�kÔ&ÐNÐ(GÐNÈÐNÈAÐNÐNÐNÔUð 
Ø,ð
Ø./ð
Ø12ð
ð 
ð 
ˆð ×ÒÑÔ˜rÔ" aÒ'Ð'Ð'Ð'Ø×ÒÑÔ Ô# qÒ(Ð(Ð(Ð(à€|Ý˜˜SŸXšX™ZœZ¨)°A¬,Ñ6¸À1¼ÑEÈÈQÉÑOÑOÐQRÑSÔSˆÝ#ØØØØ�aŒLØ�aŒLØØØØ”+Ø”ið
ñ 
ô 
ˆð €Jõ 	�s˜E 5¨*°kÀ3ÑGÔGñ	ØØØØØØØØð �F€BˆØ�hŠh�y ! r¡'Ñ*Ô*€GØ	
ˆg‰€Bà€KÝ
˜C " b Ñ
*Ô
*€CÝ˜e b¨" XÑ.Ô.€EÝ˜e b¨" XÑ.Ô.€EÝ" :°°B¨xÑ8Ô8€JÝ# K°"°b°Ñ:Ô:€Kõ 	Œ•r”zÝŒ�œ
ÝŒ•r”zÝŒ•r”zÝŒ
•B”HÝŒ•R”Xðð 
„iô€Mð —
’
˜1‘”€IØ×$Ò$ RÑ(Ô(¨1Ñ,€LØ—:’:˜b‘>”>€Là˜L¨,Ð7€IØÐÝ˜H R a RœL¨¨¨2¨Ô.Ñ/Ô/°'¸QÅÀXÈbÈqÈbÄ\ÑARÔARÑ=RÑ2SÑSˆˆàˆð 	�Ø�mØ�_Øˆ{Øˆ}Ø�KØ�[Øˆ[ð	€Oð �AŠ:ˆ:ˆ:ˆ:ð
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
õ$ �&˜/¨9°kÑBÔBÐBà
‡|‚|�~„~˜×,Ò,Ñ.Ô.Ò.Ð.ð 	×Ò˜×)Ò)¨*Ô*:Ñ;Ô;Ñ<Ô<Ð<àÐr   ÚIS_BETA_ZEROry  r{  ÚTILE_Kr|  rR  c            
      óÎ  — t          j        d¬¦  «        } t          j        d¬¦  «        }!||| z  z   ||!z  z   }"t          j        |"¦  «        }#t          j        |"|z   ¦  «        }$|$|#z
  }%|%dk    rd S t          j        d|¦  «        }&t          j        d|¦  «        }'||| z  z   |	|#z  z   |
|&d d …d f         z  z   ||'d d d …f         z  z   }(||| z  z   ||#z  z   })||| z  z   ||!z  z   ||&d d …d f         z  z   }*||| z  z   ||'d d d …f         z  z   }+t          j        d|¦  «        },t	          |%¦  «        D �]2}-t          j        ||f|¬¦  «        }.t          j        |)¦  «        }/t	          d||¦  «        D ]�}0|0|,z   }1|1|k     }2t          j        |*||1d d d …f         z  z   |2d d d …f         d¬¦  «        }3t          j        |+||/z  z   ||1d d …d f         z  z   |2d d …d f         d¬¦  «        }4|.t          j        |3|4||¬¦  «        z  }.Œ‘|r|.| z  }.n| |.z  |t          j        |(¦  «        z  z   }.t          j        |(|.                     |j	        j
        ¦  «        ¦  «         |(|	z  }(|)|z  })�Œ4d S )Nr   ©Úaxisr   ©r-   ç        ©Úmaskr/  ©rR  r	  )r~  Ú
program_idÚloadr<  rc   r³   ÚdotÚstoreÚtor-   Ú
element_ty)5r  r  rŽ  ry  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_strideÚmat1_ptrÚmat1_batch_strideÚmat1_tiled_row_strideÚmat1_tiled_col_strideÚmat1_row_block_strideÚmat1_col_block_strideÚmat2_ptrÚmat2_batch_strideÚmat2_tiled_row_strideÚmat2_tiled_col_strideÚmat2_row_block_strideÚmat2_col_block_strider|  rR  Ú	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_ptrÚmat1_block_ptrsÚmat2_block_ptrsÚk_tile_arangeÚ_Ú	acc_blockÚ	col_blockÚk_tileÚ	k_offsetsÚmask_kÚ
mat1_blockÚ
mat2_blocks5                                                        r   Ú_sampled_addmm_kernelrË  2  sv  € õF ”M qÐ)Ñ)Ô)ˆ	Ýœ¨1Ð-Ñ-Ô-ˆð Ø'¨)Ñ3ñ4à! MÑ1ñ2ð 	 õ
 ”WÐ4Ñ5Ô5ˆ
Ýœ'Ð"9Ð<OÑ"OÑPÔPˆð " JÑ.ˆØ�aŠ<ˆ<ØˆFåœ9 Q¨Ñ6Ô6ÐÝœ9 Q¨Ñ6Ô6Ðð Ø! IÑ-ñ.à *Ñ,ñ-ð &Ð(8¸¸¸¸D¸Ô(AÑAñBð &Ð(8¸¸q¸q¸q¸Ô(AÑAñ	Bð 	ð Ø&¨Ñ2ñ3à  :Ñ-ñ.ð 	ð Ø )Ñ+ñ,à# mÑ3ñ4ð $Ð&6°q°q°q¸$°wÔ&?Ñ?ñ@ð 	ð Ø )Ñ+ñ,à#Ð&6°t¸Q¸Q¸Q°wÔ&?Ñ?ñ@ð 	õ œ	 ! VÑ,Ô,ˆÝ�w‘”ð &	4ñ &	4ˆAÝœ -°Ð!?ÀyÐQÑQÔQˆIõ œÐ 1Ñ2Ô2ˆIå  1 fÑ-Ô-ð ð �Ø" ]Ñ2�	Ø" Qš�åœWØ#Ð&;¸iÈÈaÈaÈaÈÔ>PÑ&PÑPØ  a a a œØðñ ô �
õ  œWØ#Ø+¨iÑ7ñ8à+¨i¸¸¸¸4¸Ô.@Ñ@ñAð      4 œØðñ ô �
ð �RœVØ 
°zÈYðñ ô ñ �	�	ð ð RØ˜UÑ"�	�	à! IÑ-°µr´wÐ?PÑ7QÔ7QÑ0QÑQ�	õ ŒHÐ&¨	¯ª°ZÔ5EÔ5PÑ(QÔ(QÑRÔRÐRð Ð!2Ñ2ÐØÐ!3Ñ3ÐÑðM&	4ð &	4r   r  c                 ó°  — t          j        d¬¦  «        }t          j        d¬¦  «        }t          j        d¬¦  «        }t          j        d¬¦  «        }t          j        d¬¦  «        } t          j        |||| |¦  «        \  }}|||z  z   ||z  z   }!t          j        |!¦  «        }"t          j        |!|z   ¦  «        }#|#|"z
  }$|$dk    rd S t          j        d|¦  «        }%t          j        d|¦  «        }&| ||z  z   ||"z  z   ||%d d …d f         z  z   ||&d d d …f         z  z   }'|||z  z   ||z  z   ||&d d …d f         z  z   ||%d d d …f         z  z   }(|||z  z   ||z  z   ||z  z   ||%d d …d f         z  z   ||%d d d …f         z  z   })||	|z  z   |
|"z  z   }*t          j        ||f|¬¦  «        }+t          |$¦  «        D ]i},t          j        |'¦  «        }-t          j        |*¦  «        }.t          j        |(||.z  z   ¦  «        }/|+t          j        |-|/||¬¦  «        z  }+|'|z  }'|*|
z  }*Œjt          j	        |)|+ 
                    |j        j        ¦  «        ¦  «         d S )Nr   r‘  r   r   r“  r—  ©r~  r˜  Únum_programsÚ	swizzle2dr™  r<  r³   rc   rš  r›  rœ  r-   r�  )0rŸ  r   r¡  r¢  r£  r¤  r¥  r¦  r§  r¨  r©  Ú	dense_ptrÚdense_batch_strideÚdense_tiled_row_strideÚdense_tiled_col_strideÚdense_row_block_strideÚdense_col_block_strideÚ
output_ptrÚoutput_batch_strideÚoutput_tiled_row_strideÚoutput_tiled_col_strideÚoutput_row_block_strideÚoutput_col_block_stridery  r{  r|  rR  r  r¶  r·  Úcol_block_pidrŒ  r�  r¸  r¹  rº  r»  r¼  r½  r¾  Údense_block_ptrsÚoutput_ptrsr¿  Úoutput_acc_blockrÃ  Úvalues_blockÚdense_row_idxÚdense_blocks0                                                   r   Ú"_bsr_strided_dense_rowspace_kernelrã  °  sD  € õ\ ”M qÐ)Ñ)Ô)ˆ	Ýœ¨1Ð-Ñ-Ô-ˆÝœ¨1Ð-Ñ-Ô-ˆÝ”¨AÐ.Ñ.Ô.ˆÝ”¨AÐ.Ñ.Ô.ˆå')¤|Ø˜=¨,¸Ànñ(
ô (
Ñ$ˆ�}ð
 Ø'¨)Ñ3ñ4à! MÑ1ñ2ð 	 õ
 ”WÐ4Ñ5Ô5ˆ
Ýœ'Ð"9Ð<OÑ"OÑPÔPˆð " JÑ.ˆØ�aŠ<ˆ<ØˆFåœ9 Q¨Ñ6Ô6ÐÝœ9 Q¨Ñ6Ô6Ðð Ø! IÑ-ñ.à *Ñ,ñ-ð &Ð(8¸¸¸¸D¸Ô(AÑAñBð &Ð(8¸¸q¸q¸q¸Ô(AÑAñ	Bð 	ð Ø  9Ñ,ñ-à$ }Ñ4ñ5ð %Ð'7¸¸¸¸4¸Ô'@Ñ@ñAð %Ð'7¸¸a¸a¸a¸Ô'@Ñ@ñ	Að 	ð Ø! IÑ-ñ.à%¨Ñ5ñ6ð &¨Ñ5ñ6ð &Ð(8¸¸¸¸D¸Ô(AÑAñ	Bð
 &Ð(8¸¸q¸q¸q¸Ô(AÑAñBð 	ð Ø&¨Ñ2ñ3à  :Ñ-ñ.ð 	õ œ8 ]°MÐ$BÈ)ÐTÑTÔTÐÝ�w‘”ð 	4ð 	4ˆAÝœ7Ð#4Ñ5Ô5ˆLõ œGÐ$5Ñ6Ô6ˆMÝœ'Ø Ð#9¸MÑ#IÑIñô ˆKð
 ¥¤Ø˜k°jÈIð!ñ !ô !ñ Ðð
 Ð!2Ñ2ÐØÐ!3Ñ3ÐÐõ 	Œ�Ð.×1Ò1°*Ô2BÔ2MÑNÔNÑOÔOÐOÐOÐOr   c           
      óÊ  ‡ ‡‡‡‡‡‡‡— |                      d¦  «        }|                      d¦  «        dz
  }||f}|�?t          |d d…         d d d…         ¦  «        ddt          |d d…         ¦  «        z
  z  z   }nd }|d|d|d|	d|
di}|j        t          j        t          j        fv rt          j        Šd	Šnt          j	        Šd
Šˆˆˆ ˆˆˆˆˆfd„}t          ||||¦  «         d S )Nr   rt   r   r   r   )r   N)r   rt   )r   rr  TFc                 óX   •— t          |          ‰‰‰g‰¢‰‘‰	‘t          |Ž ¢R ‰‰dddœŽ d S )Nr   rë   )r|  rR  r÷   rø   )rË  r_   )
rx   r‚   r|  rR  r  r  r>   Úis_beta_zerorž  Útile_ks
     €€€€€€€€r   r€   z)_run_sampled_addmm_kernel.<locals>.kernelX  sw   ø€ Ý! $Ô'ØØØðð ð	ð
 ðð ðõ & ~Ð6ðð ð $Ø%ØØðð ð ð ð ð r   )r›   r1   r@   r-   r   r.   r/   r~  r  rˆ  rƒ   )r  r  ræ  r>   rž  rç  r•   r“   r”   Úmat1Úmat2rg  r‹  rŒ  rf   rg   rm   r€   r|  rR  s   ``````            @@r   Ú_run_sampled_addmm_kernelrê  4  s6  øøøøøøøø€ ð —K’K ‘N”Nˆ	Ø#×(Ò(¨Ñ,Ô,¨qÑ0ˆà Ð-ˆ	ØÐÝ ¨¨!¨¤¨T¨T¨r¨TÔ 2Ñ3Ô3°gÀÅSÈÐRTÐSTÐRTÌÑEVÔEVÑAVÑ6WÑWˆKˆKàˆKà�IØ˜'Ø˜Ø�'Ø�)ð
ˆð Œ<�EœJ­¬Ð7Ð7Ð7Ýœ
ˆIØˆJˆJåœ
ˆIØˆJð	ð 	ð 	ð 	ð 	ð 	ð 	ð 	ð 	ð 	ð 	ð 	õ 	�f˜o¨y¸+ÑFÔFÐFÐFÐFr   g      ð?)r  r  r`  rf  rg  rè  ré  c                ó  — d}t          || ¦  «         t          || ||¦  «        }	|�s„t          ||| j        ¦  «         t          ||| j        ¦  «         |dk    r)| j        t
          j        u rt          d|› d|› d�¦  «         | j        t
          j        ur-t          ||| j        ¦  «         t          ||| j        ¦  «         nt          |||j        ¦  «         t          |||¦  «         |�Àt          ||¦  «         t          |||j        ¦  «         t          ||| j        ¦  «         t          |j
        |	j
        k    o)|                     ¦   «         |                      ¦   «         k    |› d|	j
        › d|	                     ¦   «         › d|j
        › d	|                     ¦   «         › �	¦  «         |€|	                     |j        d
¬¦  «        }n|                     |	¦  «         |                     ¦   «         dk    s|                     ¦   «         dk    r|S |                     ¦   «         j
        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    rQ|                     ¦   «                              |                     |                     ¦   «         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"   rt   r   rˆ   )r   rž   r   r   r-   r   Úboolr   r3   r+   r$   r[  rœ  r]  rV  r•   r›   r…  r˜   r¦   r}   rê  rD   r_  )rh  rè  ré  r  r  r`  rf  rg  r   Úinput_broadcastedr>   rž  r‰  r“   r”   r•   rç  s                    r   rì  rì  i  s«  € ð !ˆå˜ Ñ'Ô'Ð'Ý4°V¸UÀDÈ$ÑOÔOÐàñ 	Ý˜  u¤|Ñ4Ô4Ð4Ý˜  u¤|Ñ4Ô4Ð4Ø�sŠ{ˆ{˜uœ{­e¬jÐ8Ð8ÝØØÐkÐk°$ÐkÐkÐkñô ð ð Œ{¥%¤*Ð,Ð,Ý˜F D¨%¬+Ñ6Ô6Ð6Ý˜F D¨%¬+Ñ6Ô6Ð6Ð6å˜F D¨$¬*Ñ5Ô5Ð5Ý& v¨t°TÑ:Ô:Ð:ØˆÝ  ¨Ñ-Ô-Ð-Ý˜V S¨$¬+Ñ6Ô6Ð6Ý˜F C¨¬Ñ5Ô5Ð5ÝØ”IÐ!2Ô!8Ò8ÐW¸S¿XºX¹Z¼ZÈ5Ï:Ê:É<Ì<Ò=WØð Rð RÐ@QÔ@Wð Rð RØ->×-CÒ-CÑ-EÔ-EðRð Rà+.¬9ðRð RàEHÇXÂXÁZÄZðRð Rñô ð ð ˆ;Ø#×&Ò& t¤z¸Ð&Ñ=Ô=ˆCˆCà�IŠIÐ'Ñ(Ô(Ð(à�9Š9‰;Œ;˜!ÒÐ˜sŸxšx™zœz¨Qš˜ØˆJà—J’J‘L”LÔ& r s sÔ+ˆ	Ø�IŠI�b‰MŒMˆð �CŠ<ˆ<˜1 š6˜6Ø�JŠJ‰LŒL×Ò˜dÑ#Ô#Ð#ØˆJð ˆ
Ý8FÀsÈDÐRVÑ8WÔ8WÑ5ˆ�k 6¨4°å  ¨	°!¬°aÐ'8Ñ9Ô9ˆÝ  ¨¨9°Q¬<Ð'8Ñ9Ô9ˆÝ�i�ˆå!ØØØ�CŠKØØØØØØØØØñ	
ô 	
ð 	
ð$ ×ÒÑÔ×%Ò%Ñ'Ô'¨¨¨Ô,°·²±´ÀÀÀÔ0DÒDÐDØ×ÒÑÔ×%Ò% f§n¢n°Z×5FÒ5FÑ5HÔ5HÔ5NÑ&OÔ&OÑPÔPÐPØÐr   )r`  rf  rg  rØ   c                ó   — d}| j         dd …         \  }}|s¿t          || ¦  «         t          || |j        ¦  «         t	          || |j        t          j        f¦  «         t          || |¦  «         | 	                    d¦  «        }	|  
                    ¦   «         j         dd …         \  }
}t          ||
|f¦  «         t          |	dz   |› d|	› d�¦  «         n|j         dd …         \  }}	t          || |¦  «        }|�x|sv|||	fz   }t          |j         |k    d|› d|j         › d	�¦  «         t          |                     ¦   «         p'|                     dd¦  «                             ¦   «         d
¦  «         |€|                     |||	fz   ¦  «        }|                      ¦   «         dk    r|                     ¦   «         S t'          || |dd|¬¦  «        S )NÚbsr_dense_mmr"   rt   r;   z(): dense.size(-1) == z should be divisible by 16z9bsr_dense_mm(): `out` argument has wrong shape, expected z
, but got r!   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   r   r   r3   r-   r   rl  r+   r›   r•   rA   r   rQ   Úis_contiguousr¡   r„  r[  r·   r  )r–   ri  r`  rf  rg  rØ   r   r£   Ú_klr¤   Ú	row_blockrÅ  Ú_krrn  Úexpected_out_shapes                  r   rñ  rñ  À  s  € ð  ˆØ”˜2˜3˜3”‰ˆˆ3Øð 	&Ý˜V SÑ)Ô)Ð)Ý˜  e¤lÑ3Ô3Ð3Ý˜  U¤[µ5´:°-Ñ@Ô@Ð@Ý& v¨s°EÑ:Ô:Ð:à—
’
˜2‘”ˆAØ#&§:¢:¡<¤<Ô#5°b°c°cÔ#:Ñ ˆI�yÝ˜F Y°	Ð$:Ñ;Ô;Ð;ÝØ˜‘F�
ØÐNÐN°ÐNÐNÐNñô ð ð ð
 ”[   Ô%‰FˆC�å*>¸vÀsÈEÑ*RÔ*RÐ'àˆ? ;ˆ?Ø!@ÀAÀqÀ6Ñ!IÐÝØ”	Ð/Ò/ðGØ.ðGð GØ:=¼)ðGð Gð Gñô ð õ
 Ø×!Ò!Ñ#Ô#ÐL s§}¢}°R¸Ñ'<Ô'<×'JÒ'JÑ'LÔ'Lð"ñô ð ð ˆ;Ø—/’/Ð"AÀQÈÀFÑ"JÑKÔKˆCð �8Š8‰:Œ:˜Š?ˆ?Ø—9’9‘;”;Ðõ ˜s C¨°a¸aÀSÐIÑIÔIÐIr   ÚMAX_ROW_NNZÚTILEc                 óÄ  — t          j        d¬¦  «        }t          j        d¬¦  «        }t          j        d¬¦  «        }| ||z  z   ||z  z   }t          j        |¦  «        }t          j        ||z   ¦  «        }||z
  }|dk    rd S t          j        d|
¦  «        }|||z  k     }|||z  z   ||z  z   ||z  z   }t          j        ||z   |t	          d¦  «         ¬¦  «                             t           j        ¦  «        }t          j        |d¬¦  «        }t          |
|	|
¦  «        D ]…}||
z  }|||z  k     }t          j        ||z   |t	          d¦  «         ¬¦  «                             t           j        ¦  «        }t          j        |d¬¦  «        }t          j	        ||k    ||¦  «        }Œ†t          j
        ||z
  ¦  «        }t          j        |d¬¦  «        }t          |
|	|
¦  «        D ]…}||
z  }|||z  k     }t          j        ||z   |t	          d¦  «         ¬¦  «                             t           j        ¦  «        }t          j
        ||z
  ¦  «        }|t          j        |d¬¦  «        z  }Œ†t          j        ||z   ||z                       |j        j        ¦  «        |¬¦  «         t          |
|	|
¦  «        D ]¦}||
z  }|||z  k     }t          j        ||z   |t	          d¦  «         ¬¦  «                             t           j        ¦  «        }t          j
        ||z
  ¦  «        }t          j        ||z   ||z                       |j        j        ¦  «        |¬¦  «         Œ§d S )Nr   r‘  r   r   Úinfr•  )r–  )r~  r˜  r™  r<  r0   rœ  r  r}   rc   ÚwhereÚexpÚsumr›  r-   r�  )r¤  r¥  r¦  rŸ  r   r¢  Úvalues_nnz_col_block_striderô  rÅ  r÷  rø  r¶  Úrow_block_offset_pidr·  r¸  r¹  rº  r»  Ú
row_aranger–  Úcurr_row_values_ptrsÚrow_tileÚmax_row_valuerÃ  Úcurr_max_row_valueÚnumÚdenoms                              r   Ú_bsr_softmax_kernelr  ÷  sË  € õ ”M qÐ)Ñ)Ô)ˆ	Ý!œ}°!Ð4Ñ4Ô4ÐÝœ¨1Ð-Ñ-Ô-ˆð Ø'¨)Ñ3ñ4à! MÑ1ñ2ð 	 õ
 ”WÐ4Ñ5Ô5ˆ
Ýœ'Ð"9Ð<OÑ"OÑPÔPˆð " JÑ.ˆØ�aŠ<ˆ<ØˆFå”Y˜q $Ñ'Ô'ˆ
Ø˜G iÑ/Ò/ˆð Ø! IÑ-ñ.à%Ð(<Ñ<ñ=ð ˜9Ñ$ñ%ð 	õ ”7Ø  :Ñ-°DÅÀuÁÄÀð
ñ 
ô 
ç
Š"�RŒZ‰.Œ.ð 	õ œ˜x¨aÐ0Ñ0Ô0ˆÝ�t˜[¨$Ñ/Ô/ð 		ð 		ˆAØ˜$ÑˆJØ ¨)Ñ 3Ò3ˆDÝ”wØ$ zÑ1¸ÅUÈ5Á\Ä\ÀMðñ ô çŠb•”‰nŒnð õ "$¤¨°qÐ!9Ñ!9Ô!9ÐÝœHØÐ 2Ò2°MÐCUñô ˆMˆMõ
 Œf�X Ñ-Ñ.Ô.ˆÝ”�s Ð#Ñ#Ô#ˆÝ�t˜[¨$Ñ/Ô/ð 	)ð 	)ˆAØ˜$ÑˆJØ ¨)Ñ 3Ò3ˆDÝ”wØ$ zÑ1¸ÅUÈ5Á\Ä\ÀMðñ ô çŠb•”‰nŒnð õ ”&˜ MÑ1Ñ2Ô2ˆCØ•R”V˜C aÐ(Ñ(Ô(Ñ(ˆEˆEõ 	ŒØ  :Ñ-Ø�5‰[×Ò˜ZÔ-Ô8Ñ9Ô9Øð	
ñ 	
ô 	
ð 	
õ
 �t˜[¨$Ñ/Ô/ð 	ð 	ˆAØ˜$ÑˆJØ ¨)Ñ 3Ò3ˆDÝ”wØ$ zÑ1¸ÅUÈ5Á\Ä\ÀMðñ ô çŠb•”‰nŒnð õ ”&˜ MÑ1Ñ2Ô2ˆCÝŒHØ$ zÑ1Ø�u‘× Ò  Ô!1Ô!<Ñ=Ô=Øðñ ô ð ð ð	ð 	r   c                 óà  ‡‡‡— d}t          || ¦  «         t          || | j        ¦  «         |                      ¦   «         dk    s|                      ¦   «         dk    r|                      ¦   «         S | j        dd …         \  }}|                      ¦   «         }|                      ¦   «         j        dd …         \  ŠŠ‰€t          j	        |¦  «        Šnt          j	        ‰¦  «        Š|  
                    ¦   «                              d¦  «                             dd¦  «        }|                      ¦   «                              dd¦  «                             ¦   «         r'|                      ¦   «                              ¦   «         }n|                      ¦   «         }|                     dd¦  «                             ¦   «                              d¦  «                             dd¦  «                             d‰|‰z  ¦  «        }|j        d         ‰|‰z  f}d }	|dd d…f         d|d	i}
ˆˆˆfd
„}t#          ||
||	¦  «          |                     d‰|‰¦  «                             dd¦  «        j        |                      ¦   «         j        Ž }t%          j        |  
                    ¦   «                              ¦   «         |                      ¦   «                              ¦   «         || j        | j        ¬¦  «        S )NÚbsr_softmaxr   r"   rˆ   rr  rt   .rq  rp  c                 óf   •— t          |          g t          |Ž ¢‰‘‰‘‰‘t          d‰¦  «        ‘R Ž  d S )Ni   )r  r_   rC   )rx   r‚   rÅ  Úmax_row_nnzrô  s     €€€r   r€   zbsr_softmax.<locals>.kernel}  sf   ø€ Ý Ô%ð Ý% ~Ð6ðàðð ðð ð	õ �E˜;Ñ'Ô'ðð ð ð ð ð r   rš   )r   r3   r-   r[  rV  Úcloner$   r•   ÚtritonÚnext_power_of_2r“   r‡   rŒ   r¡   rò  rE   r_  rƒ   r   rœ   r”   r   )rh  r  r   r£   r¤   Únnzr“   r•   rf   rg   rm   r€   rÅ  rô  s    `          @@r   r	  r	  P  s²  øøø€ Øˆå˜ Ñ'Ô'Ð'Ý�F˜E 5¤;Ñ/Ô/Ð/à�:Š:‰<Œ<˜1ÒÐ §¢¡¤°Ò 2Ð 2Ø—;’;‘=”=Ð àŒ{˜2˜3˜3Ô‰ˆˆ1Ø�jŠj‰lŒlˆØ$Ÿ|š|™~œ~Ô3°B°C°CÔ8Ñˆ	�9àÐÝ Ô0°Ñ3Ô3ˆKˆKå Ô0°Ñ=Ô=ˆKà×)Ò)Ñ+Ô+×5Ò5°aÑ8Ô8×@Ò@ÀÀBÑGÔGˆð �<Š<‰>Œ>×#Ò# B¨Ñ+Ô+×9Ò9Ñ;Ô;ð 	$à—\’\‘^”^×)Ò)Ñ+Ô+ˆFˆFà—\’\‘^”^ˆFà×Ò˜R Ñ$Ô$ßŠZ‰\Œ\ßŠY�q‰\Œ\ßŠW�Q˜‰^Œ^ßŠW�R˜ C¨)¡OÑ4Ô4ð 	ð ”\ !”_ i°°i±Ð@ˆ	Øˆð ˜˜c˜r˜c˜Ô" MØ�Oð	
ˆð	ð 	ð 	ð 	ð 	ð 	ð 	õ 	�f˜o¨y¸+ÑFÔFÐFðˆF�NŠN˜2˜y¨#¨yÑ9Ô9ßŠY�r˜2ÑÔÜ�e—l’l‘n”nÔ*ð,ð 	õ Ô-Ø×ÒÑ Ô ×&Ò&Ñ(Ô(Ø×ÒÑÔ×%Ò%Ñ'Ô'ØØ”Ø”<ð
ñ 
ô 
ð 	
r   r”  Úqueryr  ÚvalueÚ	attn_maskÚ	dropout_pÚ	is_causalÚscalec           	      ó   — d}t          | |› d�¦  «         t          |d u|› d�¦  «         |€J ‚t          |j        t          j        k    |› dt          j        › d|j        › d�¦  «         t	          ||| j        ¦  «         t	          ||| j        ¦  «         t	          ||| j        ¦  «         t          ||| j        ¦  «         t          ||| j        ¦  «         |j        t          j        urt          ||| j        ¦  «         t          || | 
                    dd¦  «        d	d
¬¦  «        }|€|                      d¦  «        dk    s|d	k    rt          d
|› d|› d�¦  «         |€*dt          j        |                      d¦  «        ¦  «        z  n|}	|                     ¦   «                              |	¦  «         t!          |¦  «        }t          j        j                             |                     ¦   «         |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"   rt   r”  F)r  rf  r   z(): current value of scale == z results in division by zero.r   T)rÎ   Úinplace)r   r   r   r   r   r   r3   r-   rî  rì  r¡   r›   ÚmathÚsqrtr•   r…  r	  ÚnnÚ
functionalÚdropoutrñ  )
r  r  r  r  r  r  r  r   ÚsdpaÚscale_factors
             r   r  r  —  s:  € ð 1ˆÝ�)ˆm ÐOÐOÐOÑPÔPÐPÝˆi˜tÐ#¨Ð%WÐ%WÐ%WÑXÔXÐXØÐ$Ð$Ð$åØÔ¥Ô 0Ò0Øð 7ð 7Ý(-Ô(8ð7ð 7à#,Ô#3ð7ð 7ð 7ñ	
ô 	
ð 	
õ 	�V˜S %¤,Ñ/Ô/Ð/Ý�V˜U E¤LÑ1Ô1Ð1Ý�V˜Y¨¬Ñ5Ô5Ð5å�F˜C ¤Ñ-Ô-Ð-Ý�F˜E 5¤;Ñ/Ô/Ð/ØŒ?¥%¤*Ð,Ð,Ý˜ 	¨5¬;Ñ7Ô7Ð7åØ�u˜cŸmšm¨B°Ñ3Ô3¸#È5ð
ñ 
ô 
ˆð ˆ=˜UŸZšZ¨™^œ^¨qÒ0Ð0°E¸S²L°LÝØØð /ð /¸ð /ð /ð /ñô ð ð
 9>¸�q�4œ9 U§Z¢Z°¡^¤^Ñ4Ô4Ñ4Ð4È5ˆØ�Š‰Œ×Ò˜<Ñ(Ô(Ð(Ý˜4Ñ Ô ˆÝŒÔ×#Ò# D§K¢K¡M¤M°YÈÐ#ÑMÔMÐMÝ˜D %Ñ(Ô(ˆØˆr   rÙ   rÒ   rÓ   rƒ  rô   rõ   c                 ó*  — | |z  }t          j        d¬¦  «        }t          j        d¬¦  «        }||z  }||z  }||z  t          j        d|¦  «        z   }||z  t          j        d|¦  «        z   }t          j        d|¦  «        }||d d …d f         |z  |d d d …f         |z  z   z   } ||d d …d f         |	z  |d d d …f         |
z  z   z   }!t          j        |||z  z   ¦  «        }"t          j        ||dz   |z  z   ¦  «        }#|"|#k    rd S t          j        ||f|¬¦  «        }$t          |"|#¦  «        D ]ˆ}%t          j        ||%|z  z   ¦  «        }&t          j        ||%|z  z   |z   ¦  «        }'t          j        | |&|z  z   ¦  «        }(t          j        |!|'|z  z   ¦  «        })|$t          j        |(|)||¬¦  «        z  }$Œ‰|||z  z   |d d …d f         |z  |d d d …f         |z  z   z   }*t          j        |*|$                     |j	        j
        ¦  «        ¦  «         d S ©Nr   r‘  r   r“  )r	  rR  )r~  r˜  r<  r™  r³   rc   rš  r›  rœ  r-   r�  )+rÙ   rÒ   rÓ   Ú
blocks_ptrÚblocks_stride_PÚblocks_stride_MÚblocks_stride_KÚ
others_ptrÚothers_stride_QÚothers_stride_KÚothers_stride_NÚaccumulators_ptrÚaccumulators_stride_RÚaccumulators_stride_MÚaccumulators_stride_NÚpq_offsets_ptrÚpq_offsets_strideÚpq_ptrÚpq_stride_TÚpq_stride_1rƒ  rô   rõ   rR  rÁ   Úpid_tÚpidÚpid_mÚpid_nÚrmÚrnÚrkÚA_ptrÚB_ptrrÌ   rÍ   rÄ  rå   rÎ   rÏ   ÚArÑ   ÚC_ptrs+                                              r   Ú_scatter_mm2_kernelr>  Å  sk  € ð6 �&‰[ˆå” 1Ð%Ñ%Ô%ˆåŒm Ð#Ñ#Ô#ˆØ�r‘	ˆØ�b‘ˆà�V‰^�bœi¨¨6Ñ2Ô2Ñ2ˆØ�V‰^�bœi¨¨6Ñ2Ô2Ñ2ˆÝŒY�q˜!‰_Œ_ˆàØˆqˆqˆq�$ˆwŒK˜/Ñ)¨B¨t°Q°Q°Q¨w¬K¸/Ñ,IÑIñ
ˆð Øˆqˆqˆq�$ˆwŒK˜/Ñ)¨B¨t°Q°Q°Q¨w¬K¸/Ñ,IÑIñ
ˆõ ŒW�^ eÐ.?Ñ&?Ñ?Ñ@Ô@ˆÝŒW�^ u¨q¡yÐ4EÑ&EÑEÑFÔFˆà�Š8ˆ8ØˆFå”H˜f fÐ-°]ÐCÑCÔCˆ	å�r˜2‘”ð 	Vð 	VˆAÝ”˜  [¡Ñ0Ñ1Ô1ˆAÝ”˜  [¡Ñ0°;Ñ>Ñ?Ô?ˆAÝ”˜  OÑ 3Ñ3Ñ4Ô4ˆAÝ”˜  OÑ 3Ñ3Ñ4Ô4ˆAØ�œ  1°È*ÐUÑUÔUÑUˆIˆIð ØÐ+Ñ+ñ,ð �1�1�1�d�7”Ð3Ñ3Ø�T˜1˜1˜1�W”+Ð 5Ñ5ñ6ñð 	õ 	Œ�˜	ŸšÐ%5Ô%;Ô%FÑGÔGÑHÔHÐHÐHÐHr   r¼   r½   rO  Ú
pq_indicesr¬   c                 óâ  ‡‡‡— | j         \  }Š}|j         \  }}Št          t          d‰dz  ¦  «        t          d‰dz  ¦  «        dd¬¦  «        }	ˆˆˆfd„}
t          j        t
          j        t          j        t
          j        t          j        t
          j        t          j        t
          j        i|j	                 }d|	vr$|	 
                    |t
          j        k    ¬¦  «         t          |
         ‰|‰| |                      d	¦  «        |                      d¦  «        |                      d¦  «        ||                     d	¦  «        |                     d¦  «        |                     d¦  «        ||                     d	¦  «        |                     d¦  «        |                     d¦  «        ‰‰                     d	¦  «        ||                     d	¦  «        |                     d¦  «        fd
|i|	¤Ž d S )Nr;   rë   r   r   )rô   rõ   r÷   rø   c                 ó”   •— ‰j         d         dz
  t          j        ‰| d         ¦  «        t          j        ‰| d         ¦  «        z  dfS )Nr   r   rô   rõ   ©r$   r  Úcdiv)ÚMETArÙ   rÓ   rO  s    €€€r   rx   z_scatter_mm2.<locals>.grid  sH   ø€ àÔ  Ô# aÑ'Ý”˜A˜t Hœ~Ñ.Ô.µ´¸QÀÀXÄÑ1OÔ1OÑOØðð r   rR  rS  r   rƒ  )r$   rý   r}   r   rú   r~  r  r/   rˆ  r-   rû   r>  rD   )r¼   r½   rO  r?  r¬   rÀ   rÒ   rÅ   rÃ  rØ   rx   rƒ  rÙ   rÓ   s     `         @@r   r´   r´     sî  øøø€ ð ”<‰ˆˆAˆqØ”<‰ˆˆAˆqåÝ�r˜1 ™6‘?”?­3¨r°1¸±6©?¬?ÀqÐTUð
ñ 
ô 
ˆð	ð 	ð 	ð 	ð 	ð 	ð 	õ ŒM�2œ:ÝŒN�BœJÝŒM�2œ:ÝŒM�2œ:ð	
ð
 Ô
ôˆð ˜tÐ#Ð#Ø�KŠK =µB´JÒ#>ˆKÑ?Ô?Ð?Ý˜DÔ!ØØØØØ�MŠM˜!ÑÔØ�MŠM˜!ÑÔØ�MŠM˜!ÑÔØØ�MŠM˜!ÑÔØ�MŠM˜!ÑÔØ�MŠM˜!ÑÔØØ×Ò Ñ"Ô"Ø×Ò Ñ"Ô"Ø×Ò Ñ"Ô"ØØ×Ò˜aÑ Ô ØØ×Ò˜aÑ Ô Ø×Ò˜aÑ Ô ð)	
ð 	
ð* (ð+	
ð, ð-	
ð 	
ð 	
ð 	
ð 	
r   rÂ   rT  r±   rö   c                 ó‚  — ||z  }||z  }||z  }t          j        d¬¦  «        }t          j        d¬¦  «        }|| z  } || z  }!||z  }"||"z  }#|#|z  }$t          ||$z
  |¦  «        }%|$||%z  z   }&||"z  |%z  }'|&|z  t          j        d|¦  «        z   }(|'|z  t          j        d|¦  «        z   })t          j        d|¦  «        }*||(d d …d f         |z  |*d d d …f         |z  z   z   }+|| |	z  z   |*d d …d f         |
z  |)d d d …f         |z  z   z   },t          j        ||!z   ¦  «        }-|rU|-|z  |z  }.|-|z  |z  }/t          j        ||.z   ¦  «        }0t          j        ||.z   dz   ¦  «        }1|/|1z  ||/z
  |0z  z   }2|1|0z
  }3n6t          j        ||!z   ¦  «        }2t          j        ||!z   dz   ¦  «        }4|4|2z
  }3||2z   }5t          j        ||f|¬¦  «        }6|r|+|0|z  z  }+t          |3¦  «        D ]f}7t          j        |5¦  «        }8t          j        |,|8z   ¦  «        }9t          j        |+¦  «        }:|6t          j        |:|9||¬¦  «        z  }6|+|z  }+|5dz  }5Œgn•||2z   };t          |3¦  «        D ]€}7t          j        |5¦  «        }8t          j        |,|8z   ¦  «        }9t          j        |;¦  «        }<t          j        |+|<|z  z   ¦  «        }:|;dz  };|5dz  }5|6t          j        |:|9||¬¦  «        z  }6Œ�||-z   | |z  z   |(d d …d f         |z  |)d d d …f         |z  z   z   }=t          j        |=|6 	                    |j
        j        ¦  «        ¦  «         d S r!  )r~  r˜  rC   r<  r™  r³   rc   rš  r›  rœ  r-   r�  )>rE  rÁ   rÂ   rÓ   r"  r#  r$  r%  r&  Úothers_stride_Br(  r)  r*  Úaccumulators_stride_Br,  r-  Úc_indices_ptrÚr_offsets_ptrÚp_offsets_ptrÚq_offsets_ptrrT  rƒ  r±   rô   rõ   rö   rR  rÇ   ÚBLOCKS_MÚBLOCKS_NÚpid_t_r4  Úpid_br3  Únum_pid_in_groupÚgroup_idÚfirst_pid_mÚgroup_size_mr5  r6  r7  r8  r9  r:  r;  rË   r£   r¤   rÝ   rÞ   rÌ   r  rÍ   Úq_ptrrÄ  rÃ  rÏ   rÑ   r<  Úp_ptrrÎ   r=  s>                                                                 r   Ú_scatter_mm6_kernelrV  B  s	  € ð< �'‰\ˆØ˜‘<ˆØ˜‘<ˆå” AÐ&Ñ&Ô&ˆÝŒm Ð#Ñ#Ô#ˆØ˜Ñ!ˆØ˜(Ñ"ˆà%¨Ñ0ÐØÐ*Ñ*ˆØ Ñ+ˆÝ˜8 kÑ1°:Ñ>Ô>ˆØ˜s \Ñ1Ñ2ˆØÐ'Ñ'¨LÑ8ˆà�V‰^�bœi¨¨6Ñ2Ô2Ñ2ˆØ�V‰^�bœi¨¨6Ñ2Ô2Ñ2ˆÝŒY�q˜"ÑÔˆØØˆqˆqˆq�$ˆwŒK˜/Ñ)¨B¨t°Q°Q°Q¨w¬K¸/Ñ,IÑIñ
ˆð Ø�oÑ%ñ&à�!�!�!�T�'Œ{˜_Ñ,¨r°$¸¸¸°'¬{¸_Ñ/LÑLñNð 	õ ŒG�M EÑ)Ñ*Ô*ˆàð 
	Ø�a‘˜B‘ˆAØ�Q‘˜2‘ˆAÝ”˜¨Ñ*Ñ+Ô+ˆBÝ”˜¨Ñ*¨QÑ.Ñ/Ô/ˆBØ�R‘˜7 Q™;¨"Ñ,Ñ,ˆBØ�r‘'ˆCˆCå”˜¨Ñ.Ñ/Ô/ˆBÝ”˜¨Ñ.°Ñ2Ñ3Ô3ˆBØ�r‘'ˆCà Ñ"ˆÝ”H˜f fÐ-°]ÐCÑCÔCˆ	àð 	Ø�R˜/Ñ)Ñ)ˆEÝ˜3‘Z”Zð ð �Ý”G˜E‘N”N�Ý”G˜E A™IÑ&Ô&�Ý”G˜E‘N”N�Ø�RœVØ�q M¸jðñ ô ñ �	ð ˜Ñ(�Ø˜‘
��ðð " BÑ&ˆEÝ˜3‘Z”Zð 	ð 	�Ý”G˜E‘N”N�Ý”G˜E A™IÑ&Ô&�Ý”G˜E‘N”N�Ý”G˜E A¨Ñ$7Ñ7Ñ8Ô8�Ø˜‘
�Ø˜‘
�Ø�RœVØ�q M¸jðñ ô ñ �	�	ð
 ØñàÐ+Ñ+ñ,ð �1�1�1�d�7”Ð3Ñ3Ø�T˜1˜1˜1�W”+Ð 5Ñ5ñ6ñ	ð 	õ 	Œ�˜	ŸšÐ%5Ô%;Ô%FÑGÔGÑHÔHÐHÐHÐHr   TrÔ   rÕ   rÖ   r×   Úforce_contiguousc	                 ó@  ‡‡‡‡— |d         }	| j         \  }
Š}|j         \  Š}}|j         \  }}}||k    sJ ‚||	z  Š|‰k    sJ ‚ˆˆˆˆfd„}t          j        t          j        t          j        t          j        t          j        t          j        t          j        t          j        i|j                 }d|vr$|                     |t          j        k    ¬¦  «         | 	                    d¦  «        dk    sJ ‚‰ 	                    d¦  «        dk    sJ ‚| 	                    d¦  «        dk    sJ ‚| 	                    d¦  «        dk    sJ ‚|rT|  
                    ¦   «         } | 
                    ¦   «         }|                     ¦   «         s| 
                    ¦   «         }n|}n|}t          |         ‰‰||| |  	                    d¦  «        |  	                    d¦  «        |  	                    d¦  «        || 	                    d¦  «        | 	                    d¦  «        | 	                    d¦  «        || 	                    d¦  «        | 	                    d¦  «        | 	                    d¦  «        |‰||fd|i|¤Ž |r+|                     ¦   «         s|                     |¦  «         d S d S d S )	Nr±   c                 ó’   •— ‰j         d         ‰z  t          j        ‰| d         ¦  «        t          j        ‰| d         ¦  «        z  fS )Nr   rô   rõ   rB  )rD  rÑ   rÁ   rÇ   rÕ   s    €€€€r   rx   z_scatter_mm6.<locals>.gridÇ  sD   ø€ à” Ô" QÑ&Ý”˜B  X¤Ñ/Ô/µ&´+¸bÀ$ÀxÄ.Ñ2QÔ2QÑQðð r   rR  rS  r   r   r   rƒ  )r$   r   rú   r~  r  r/   rˆ  r-   rû   rD   rE   rò  rV  r]  )r¼   r½   rÔ   rÕ   rÖ   r×   rØ   r¬   rW  r±   rÀ   rÂ   Ú_KrÓ   ÚB_Ú_MrÚ   rx   rƒ  Úaccumulators_rÑ   rÁ   rÇ   s      `                @@@r   r¶   r¶   ´  sê  øøøø€ ð �y”/ˆØ”\‰
ˆˆB�Ø”<‰ˆˆ2ˆqØ!Ô'‰
ˆˆB�Ø�QŠwˆwˆwˆwØ�'‰\ˆØ�QŠwˆwˆwˆwð	ð 	ð 	ð 	ð 	ð 	ð 	ð 	õ ŒM�2œ:ÝŒN�BœJÝŒM�2œ:ÝŒM�2œ:ð	
ð
 Ô
ôˆð ˜tÐ#Ð#Ø�KŠK =µB´JÒ#>ˆKÑ?Ô?Ð?à×Ò Ñ"Ô" aÒ'Ð'Ð'Ð'Ø×Ò Ñ"Ô" aÒ'Ð'Ð'Ð'Ø×Ò Ñ"Ô" aÒ'Ð'Ð'Ð'Ø×Ò Ñ"Ô" aÒ'Ð'Ð'Ð'ð ð 	)Ø×&Ò&Ñ(Ô(ˆFØ×&Ò&Ñ(Ô(ˆFØ×-Ò-Ñ/Ô/ð -Ø ,× 7Ò 7Ñ 9Ô 9��à ,��à(ˆMå˜DÔ!ØØØØØØ�MŠM˜!ÑÔØ�MŠM˜!ÑÔØ�MŠM˜!ÑÔØØ�MŠM˜!ÑÔØ�MŠM˜!ÑÔØ�MŠM˜!ÑÔØØ× Ò  Ñ#Ô#Ø× Ò  Ñ#Ô#Ø× Ò  Ñ#Ô#ØØØØð)	
ð 	
ð* (ð+	
ð, ð-	
ð 	
ð 	
ð2 ð 	. L×$>Ò$>Ñ$@Ô$@ð 	.Ø×Ò˜}Ñ-Ô-Ð-Ð-Ð-ð	.ð 	.ð 	.ð 	.r   Úleft_alpha_tiled_col_strideÚleft_alpha_col_block_strideÚright_alpha_tiled_row_strideÚright_alpha_row_block_stridert  ru  rv  rw  rx  rz  c7                 ó   — |dk    sJ ‚|dk    sJ ‚|dk    sJ ‚|!dk    sJ ‚t          j        d¬¦  «        }7t          j        d¬¦  «        }8t          j        d¬¦  «        }9t          j        d¬¦  «        }:t          j        d¬¦  «        };t          j        |8|9|:|;|5¦  «        \  }8}9|||7z  z   ||8z  z   }<t          j        |<¦  «        }=t          j        |<|z   ¦  «        }>|>|=z
  }?t          j        d|0¦  «        }@t          j        d|2¦  «        }At          j        d|1¦  «        }B| ||7z  z   ||=z  z   ||@d d …d f         z  z   ||Ad d d …f         z  z   }C|||7z  z   ||9z  z   ||Ad d …d f         z  z   ||Bd d d …f         z  z   }D|#|$|7z  z   |%|8z  z   |&|9z  z   |'|@d d …d f         z  z   |(|Bd d d …f         z  z   }E||	|7z  z   |
|=z  z   }Ft          j        |0|1f|3¬¦  «        }Gt          |?¦  «        D ]i}Ht          j        |C¦  «        }It          j        |F¦  «        }Jt          j        |D||Jz  z   ¦  «        }K|Gt          j        |I|K|4|3¬¦  «        z  }G|C|z  }C|F|
z  }FŒj|-s|G|*z  }G|.sK|||7z  z   ||8z  z   ||9z  z   ||@d d …d f         z  z   ||Bd d d …f         z  z   }L|Gt          j        |L¦  «        z  }G|/sK|||7z  z   ||8z  z   | |9z  z   |!|@d d …d f         z  z   |"|Bd d d …f         z  z   }M|Gt          j        |M¦  «        z  }G|,rh|||7z  z   ||8z  z   ||9z  z   ||@d d …d f         z  z   ||Bd d d …f         z  z   }N|+r|Gt          j        |N¦  «        z  }Gn|G|)t          j        |N¦  «        z  z  }Gt          j	        |E|G 
                    |#j        j        ¦  «        ¦  «         d S )Nr   r   r‘  r   r“  r—  rÍ  )OrŸ  r   r¡  r¢  r£  r¤  r¥  r¦  r§  r¨  r©  Ú	input_ptrÚinput_batch_strideÚinput_tiled_row_strideÚinput_tiled_col_strideÚinput_row_block_strideÚinput_col_block_striderÐ  rÑ  rÒ  rÓ  rÔ  rÕ  Úleft_alpha_ptrÚleft_alpha_batch_strideÚleft_alpha_tiled_row_strider^  Úleft_alpha_row_block_strider_  Úright_alpha_ptrÚright_alpha_batch_strider`  Úright_alpha_tiled_col_stridera  Úright_alpha_col_block_striderÖ  r×  rØ  rÙ  rÚ  rÛ  r  r  rt  ru  rv  rw  rx  ry  r{  rz  r|  rR  r  r±   r¶  r·  rÜ  rŒ  r�  r¸  r¹  rº  r»  r¼  Úinner_block_aranger½  r¾  rÝ  rÞ  r¿  rß  rÃ  rà  rá  râ  Úleft_alpha_ptrsÚright_alpha_ptrsÚ
input_ptrssO                                                                                  r   r}  r}  	  si  € ðV +¨aÒ/Ð/Ð/Ð/Ø*¨aÒ/Ð/Ð/Ð/Ø+¨qÒ0Ð0Ð0Ð0Ø+¨qÒ0Ð0Ð0Ð0å”M qÐ)Ñ)Ô)ˆ	Ýœ¨1Ð-Ñ-Ô-ˆÝœ¨1Ð-Ñ-Ô-ˆÝ”¨AÐ.Ñ.Ô.ˆÝ”¨AÐ.Ñ.Ô.ˆå')¤|Ø˜=¨,¸Ànñ(
ô (
Ñ$ˆ�}ð
 Ø'¨)Ñ3ñ4à! MÑ1ñ2ð 	 õ
 ”WÐ4Ñ5Ô5ˆ
Ýœ'Ð"9Ð<OÑ"OÑPÔPˆð " JÑ.ˆåœ9 Q¨Ñ6Ô6ÐÝœY q¨/Ñ:Ô:ÐÝœ9 Q¨Ñ6Ô6Ðð Ø! IÑ-ñ.à *Ñ,ñ-ð &Ð(8¸¸¸¸D¸Ô(AÑAñBð &Ð(:¸4ÀÀÀ¸7Ô(CÑCñ	Dð 	ð Ø  9Ñ,ñ-à$ }Ñ4ñ5ð %Ð'9¸!¸!¸!¸T¸'Ô'BÑBñCð %Ð'7¸¸a¸a¸a¸Ô'@Ñ@ñ	Að 	ð Ø! IÑ-ñ.à%¨Ñ5ñ6ð &¨Ñ5ñ6ð &Ð(8¸¸¸¸D¸Ô(AÑAñ	Bð
 &Ð(8¸¸q¸q¸q¸Ô(AÑAñBð 	ð Ø&¨Ñ2ñ3à  :Ñ-ñ.ð 	õ œ8 ]°MÐ$BÈ)ÐTÑTÔTÐå�w‘”ð 	4ð 	4ˆAÝœ7Ð#4Ñ5Ô5ˆLõ œGÐ$5Ñ6Ô6ˆMÝœ'Ø Ð#9¸MÑ#IÑIñô ˆKð
 ¥¤Ø˜k°jÈIð!ñ !ô !ñ Ðð
 Ð!2Ñ2ÐØÐ!3Ñ3ÐÐàð 	&Ø Ñ%Ðà ð 		9àØ)¨IÑ5ñ6à-°Ñ=ñ>ð .°Ñ=ñ>ð .Ð0@ÀÀÀÀDÀÔ0IÑIñ	Jð
 .Ð0@ÀÀqÀqÀqÀÔ0IÑIñJð ð ¥¤¨Ñ 8Ô 8Ñ8Ðà!ð 		:àØ*¨YÑ6ñ7à.°Ñ>ñ?ð /°Ñ>ñ?ð /Ð1AÀ!À!À!ÀTÀ'Ô1JÑJñ	Kð
 /Ð1AÀ$ÈÈÈÀ'Ô1JÑJñKð ð ¥¤Ð(8Ñ 9Ô 9Ñ9Ðàð 	?àØ$ yÑ0ñ1à(¨=Ñ8ñ9ð )¨=Ñ8ñ9ð )Ð+;¸A¸A¸A¸t¸GÔ+DÑDñ	Eð
 )Ð+;¸DÀ!À!À!¸GÔ+DÑDñEð ð ð ?Ø ¥B¤G¨JÑ$7Ô$7Ñ7Ð Ð à  D­2¬7°:Ñ+>Ô+>Ñ$>Ñ>Ð õ 	Œ�Ð.×1Ò1°*Ô2BÔ2MÑNÔNÑOÔOÐOÐOÐOr   r   )NNNNNN)NNNNNNNr   )r²   )NN)r”  FN)T)Cr  Úosr  Ú	functoolsr   Útypingr   r   r   Útorch._dynamo.utilsr   Útorch.utils._tritonr   Ú_triton_ops_metar	   ÚintÚgetenvr
   r   r   r   r+   r3   rA   rF   rQ   rV   r]   r_   ry   rƒ   r˜   rž   r¦   rª   r®   r   r  r  rP  rY  rb  ÚTensorrî  rý   rk  r  r  Útriton.languageÚlanguager~  ÚjitÚ	constexprrË  rã  rê  rì  rñ  r  r	  r0   r  r>  r´   rV  r¶   r}  r7   r   r   ú<module>r‚     s  ðð €€€Ø 	€	€	€	Ø €€€Ø Ð Ð Ð Ð Ð Ø "Ð "Ð "Ð "Ð "Ð "Ð "Ð "à €€€Ø )Ð )Ð )Ð )Ð )Ð )Ø *Ð *Ð *Ð *Ð *Ð *à &Ð &Ð &Ð &Ð &Ð &ð .1¨SØ€B„IÐ:¸AÑ>Ô>ñ.ô .Ð *ð
ð ð ð
ð ð ðð ð ðð ð ð"ð ð ðð ð ð(ð ð ð&Uð Uð Uðð ð ðð ð ðð ð ð;ð ;ð ;ð2&ð &ð &ð &ð0 7ð  7ð  7ðF	ð 	ð 	ð
/ð 
/ð 
/ðð ð ð >Bð m2ð m2ð m2ð m2ð m2ðl ØØØØØðkð kð kð kðl ØØØØØ
ØØð`ð `ð `ð `ðFKð Kð Kð Kð Kñ Kô Kð Kð\ €Ð=Ð>Ñ>Ô>ð]
ð ]
ñ ?Ô>ð]
ðB  ;ðð ð ð ðBDð Dð Dð DðX 
Ø
Ø)-Ø*.Ø"&ØØMQØð&ð &ð &ØŒ<ð&à	Œð&ð Œ<ð&ð ˜œÔ&ð&ð ˜%œ,Ô'ð&ð 
�%”,Ô	ð&ð ð&ð �u˜X cœ]¨H°S¬M¸8ÀC¼=ÐHÔIÔJð&ð �4Œ.ð&ð &ð &ð &ð\ 
Ø
Ø)-Ø*.Ø"&ØØMQØðfð fð fØŒ<ðfà	Œðfð Œ<ðfð ˜œÔ&ðfð ˜%œ,Ô'ðfð 
�%”,Ô	ðfð ðfð �u˜X cœ]¨H°S¬M¸8ÀC¼=ÐHÔIÔJðfð �4Œ.ðfð fð fð fðR €:�<„<ñ s%Ø€M€M€MØ Ð Ð Ð Ð Ð à„Zð{4ð ”lð{4ð ”|ð	{4ð
 ”|ð{4ð ”ð{4ð> ”<ð?{4ð@ ”LðA{4ð {4ð {4ñ „Zð{4ðz „ZðAPðN ”|ðOAPðP ”|ðQAPðR ”<ðSAPðT ”LðUAPðV œðWAPð APð APñ „ZðAPðF3Gð 3Gð 3Gðt ØØ&*Ø!ØQUðUð Uð UØŒ|ðUàŒlðUð ŒlðUð �e”lÔ#ðUð ðUð ˜5 ¨#¤°¸´¸xÈ¼}Ð!LÔMÔNðUð Uð Uð Uðv '+Ø!ØQUØ#ð5Jð 5Jð 5JØŒ\ð5JàŒ|ð5Jð �e”lÔ#ð	5Jð
 ð5Jð ˜5 ¨#¤°¸´¸xÈ¼}Ð!LÔMÔNð5Jð �tŒnð5Jð 5Jð 5Jð 5Jðn „ZðVð ”\ðVð ŒlðVð Vð Vñ „ZðVðpE
ð E
ð E
ð E
ðX ØØ!%ð,ð ,ØŒ|ð,àŒ\ð,ð Œ|ð,ð ˜EœLÔ)ð	,ð
 ð,ð ð,ð ˜Œð,ð ,ð ,ð ,ð\ „ZðDIØŒ<ðDIàŒ<ðDIð Œ<ðDIð* ”|ð+DIð, ”ð-DIð. ”ð/DIð0 ”Lð1DIð DIð DIñ „ZðDIðL4
Ø”ð4
à”ð4
ð ”Lð4
ð ”Lð	4
ð
 ”lð4
ð 4
ð 4
ð 4
ðl „ZðoIð ŒLðoIð* ”|ð+oIð, ”|ð-oIð. ”ð/oIð0 ”ð1oIð2 ”ð3oIð4 ”Lð5oIð6 ”Lð7oIð oIð oIñ „ZðoIðt "&ðX.ð X.Ø”ðX.à”ðX.ð ”<ðX.ð ”<ð	X.ð
 ”<ðX.ð ”<ðX.ð ðX.ð ”lðX.ð ðX.ð X.ð X.ð X.ðt „ZðIPðL &(¤\ðMIPðP &(¤\ðQIPðZ ')¤lð[IPð^ ')¤lð_IPðx ”\ðyIPðz œð{IPð| ”lð}IPð~ œ<ðIPð@ œLðAIPðB ”|ðCIPðD ”|ðEIPðF œðGIPðH ”<ðIIPðJ ”LðKIPðL œðMIPðN ”ðOIPð IPð IPñ „ZðIPð IPð IPðX €KØ€LØ€MØ$(Ð!Ø€LØ€LØ $ÐÐÐr   