§
    �Štjf^ ã            !       ó¶  — d dl Z d dlZd dlZd dlmZ d dlZd dlmZ d dlm	Z	 ddl
mZ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dtd„Zd„ Zd„ Zd„ Zd„ Z ddœd„Z!	 	 	 	 	 	 dud„Z"	 	 	 	 	 	 	 	 dvd„Z# G d„ d¦  «        Z$ ee¬ ¦  «        d!„ ¦   «         Z%	 dwd#„Z&dxd$„Z'dddddd%ddd&œd'ej(        d(ej(        d)ej(        d*ej(        dz  d+ej(        dz  d,ej(        dz  d-e)d.e*edz  edz  edz  f         dz  d/e+dz  fd0„Z,dddddd%ddd&œd'ej(        d(ej(        d)ej(        d*ej(        dz  d+ej(        dz  d,ej(        dz  d-e)d.e*edz  edz  edz  f         dz  d/e+dz  fd1„Z- e	¦   «         �rÂd dl.Z.d dl/m0Z1 e.j2        d2e1j3        d3e1j3        d4e1j3        d5e1j3        d6e1j3        d7e1j3        fd8„¦   «         Z4e.j2        d3e1j3        d4e1j3        d6e1j3        d7e1j3        d9e1j3        f
d:„¦   «         Z5d;„ Z6d<d<dd%dd=œd'ej(        d>ej(        d?ej(        d,ej(        dz  d-e)d.e*edz  edz  edz  f         dz  fd@„Z7dd%dddAœd(ej(        d)ej(        d,ej(        dz  d-e)d.e*edz  edz  edz  f         dz  d/e+dz  fdB„Z8e.j2        dCe1j3        dDe1j3        fdE„¦   «         Z9dtdF„Z:	 	 	 dydHej(        dIej(        dJej(        dKej(        dz  dLe;dMe)dNe;dz  fdO„Z<e.j2        dPe1j3        dQe1j3        dRe1j3        dSe1j3        dTe1j3        dUe1j3        d7e1j3        fdV„¦   «         Z=dWej(        dXej(        dYej(        dZej(        d[ej(        f
d\„Z>e.j2        d]e1j3        d^e1j3        dSe1j3        d_e1j3        dTe1j3        dUe1j3        d`e1j3        d7e1j3        fda„¦   «         Z?	 dzdWej(        dXej(        dcej(        ddej(        deej(        dfej(        d/e+d[ej(        dge)fdh„Z@e.j2        die1j3        dje1j3        dke1j3        dle1j3        dme1j3        dne1j3        doe1j3        dpe1j3        dqe1j3        d3e1j3        d4e1j3        dre1j3        d6e1j3        d7e1j3        d9e1j3        d_e1j3        f ds„¦   «         ZAdS dZ:dZ8dZ7dZ<dZ>dZ@dZAdS ){é    N)Ú	lru_cache)Ú	warn_once)Ú
has_tritoné   )Ú_get_device_nameÚget_metaÚ*TORCH_SPARSE_BSR_SCATTER_MM_LRU_CACHE_SIZEé   c                 ó(   — | st          |¦  «        ‚d S ©N)Ú
ValueError)ÚcondÚmsgs     úV/var/www/html/CA-Chatbot/venv/lib/python3.11/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                 óZ   — t          |j        |k    o|j        j        dv | › d�¦  «         d S )N)ÚcudaÚxpuz9(): all inputs are expected to be on the same GPU device.)r   ÚdeviceÚtype)r   r   r   s      r   Úcheck_devicer       sB   € Ý	Ø	Œ�FÒÐ?˜qœxœ}°Ð?ØÐ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   8   s‚   € Ý	Ø	Œ�5Òð 	SØŒGÝŒZ�œ­¬Ð5½Ð?PÐ8QÑQðSàð 	'ð 	'Ø3Dð	'ð 	'àœGð	'ð 	'ð 	'ñ	ô ð ð ð r   c           	      óØ   ‡— t          |¦  «        dk    rt          dt          |¦  «        › �¦  «        ‚d„ Šˆfd„}t           ||¦  «        | › d|d         › d|d         › d	�¦  «         d S )
Nr
   z"blocksize must have length 2, got 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.)ÚlenÚAssertionErrorr   )r   r>   r?   r9   s      @r   Úcheck_blocksizerB   C   s¾   ø€ Ý
ˆ9�~„~˜ÒÐÝÐRÅ#ÀiÁ.Ä.ÐRÐRÑSÔSÐSð!ð !ð !ðð ð ð ð õ 
ØÐ 	Ñ*Ô*Øð 	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_contiguousrG   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_dimsrR   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_rangerQ   r   Úslicess        r   ÚslicerrW   r   sQ   è è € Øð ð ˆÝ˜‘+”+� §¢¡¤Ñ(ˆØ!ˆˆs‰Ø�Œiˆˆˆˆðð r   c              '   óÊ   K  — |D ]]}t          d ¦  «        g|                     ¦   «         z  }t          | |d¬¦  «        D ]\  }}|�|||<   Œ|t          |¦  «                 V — Œ^d S )NF©Ústrict)rT   r#   Úzipr1   )ÚdimsrV   rQ   r   ÚsÚdÚd_slices          r   Úmultidim_slicerr`   y   sƒ   è è € Øð ð ˆÝ�4‰[Œ[ˆM˜AŸEšE™GœGÑ#ˆÝ˜d F°5Ð9Ñ9Ô9ð 	ð 	‰JˆAˆwØˆ}Ø��!‘øØ•�a‘”ŒkÐÐÐÐðð r   c               '   óP   K  — | D ] }|V — |                      ¦   «         E d {V —† Œ!d S r   )rE   )rQ   r   s     r   Úptr_stride_extractorrb   ‚   sK   è è € Øð ð ˆØˆˆˆØ—8’8‘:”:ÐÐÐÐÐÐÐÐðð r   c              #   ó  ‡ ‡‡K  — t          ‰ ¦  «        dk     st          ‰ ¦  «        dk    rt          dt          ‰ ¦  «        › �¦  «        ‚t          ‰¦  «        dk     st          ‰¦  «        dk    rt          dt          ‰¦  «        › �¦  «        ‚dd l}ˆ ˆfd„}ˆfd„} |j         |¦   «         Ž D ]U}d„ t	          ‰ |‰d¬	¦  «        D ¦   «         }d
„ t	          ||d¬	¦  «        D ¦   «         }|d d d…         g ||¦  «        ¢R V — ŒVd S )Nr   é   z"full_grid length must be 0-3, got z$grid_blocks length must be 0-3, got c               3   ób   •K  — t          ‰‰d¬¦  «        D ]\  } }t          d| |¦  «        V — Œd S )NFrY   r   )r[   Úrange)ÚfgÚmgÚ	full_gridÚgrid_blockss     €€r   Úgenerate_grid_pointsz.grid_partitioner.<locals>.generate_grid_points�   sN   øè è € Ý˜) [¸Ð?Ñ?Ô?ð 	#ð 	#‰FˆB�Ý˜˜2˜rÑ"Ô"Ð"Ð"Ð"Ð"ð	#ð 	#r   c              3   ó€   •K  — ‰                      ¦   «         D ]%\  }}t          t          || |¦  «        ¦  «        V — Œ&d S r   )ÚitemsÚnextr`   )rV   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   )rD   )rM   rg   Úgprh   s       r   ú
<listcomp>z$grid_partitioner.<locals>.<listcomp>™   s<   € ð 
ð 
ð 
á��B˜õ ��R‘˜ÑÔð
ð 
ð 
r   FrY   c                 ó:   — g | ]\  }}t          |||z   ¦  «        ‘ŒS r7   )rT   )rM   rs   Úgs      r   rt   z$grid_partitioner.<locals>.<listcomp>�   s*   € ÐUÐUÐU©¨¨A•%˜˜B ™FÑ#Ô#ÐUÐUÐUr   éÿÿÿÿ)r@   rA   Ú	itertoolsÚproductr[   )	ri   rj   rp   rx   rk   rq   Ú
grid_pointÚgridrV   s	   ```      r   Úgrid_partitionerr|   ˆ   sƒ  øøøè è € Ý
ˆ9�~„~˜ÒÐ�S ™^œ^¨aÒ/Ð/ÝÐRÅ#ÀiÁ.Ä.ÐRÐRÑSÔSÐSÝ
ˆ;ÑÔ˜!ÒÐ�s ;Ñ/Ô/°!Ò3Ð3ÝÐVÅCÈÑDTÔDTÐVÐVÑWÔWÐWàÐÐÐð#ð #ð #ð #ð #ð #ð;ð ;ð ;ð ;ð ;ð (�iÔ'Ð)=Ð)=Ñ)?Ô)?Ð@ð 	;ð 	;ˆ
ð
ð 
å! )¨Z¸ÈUÐSÑSÔSð
ñ 
ô 
ˆð VÐUµ°ZÀÈeÐ1TÑ1TÔ1TÐUÑUÔUˆð �4�4�R�4ŒjÐ:Ð2Ð2°6Ñ:Ô:Ð:Ð:Ð:Ð:Ð:Ð:ð	;ð 	;r   c                 óÂ   ‡— dd d d…         }|€|}n.d„ Št          ˆfd„t          ||d¬¦  «        D ¦   «         ¦  «        }t          |||¦  «        D ]^}} | |g|¢R Ž  Œd S )N)iÿÿÿéÿÿ  r~   rw   c                 óF   — | €|S t          dt          | |¦  «        ¦  «        S r6   )ÚmaxrD   )rv   rh   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   )rM   rv   rh   r�   s      €r   rN   z launch_kernel.<locals>.<genexpr>²   sG   øè è € ð 
ð 
á��2ð ˆN˜1˜bÑ!Ô!ð
ð 
ð 
ð 
ð 
ð 
r   FrY   )r1   r[   r|   )Úkernelrp   ri   rj   Úcuda_max_gridr{   Úsliced_tensorsr�   s          @r   Úlaunch_kernelr†   ¤   sÃ   ø€ à.¨t¨t°¨tÔ4€MØÐØ#ˆˆð	*ð 	*ð 	*õ ð 
ð 
ð 
ð 
å˜[¨-ÀÐFÑFÔFð
ñ 
ô 
ñ 
ô 
ˆõ
 "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   )rG   Ú	unsqueezerL   s     r   rt   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 rJ   rK   rL   s     r   rN   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   ©rw   c           	      óD   •— g | ]} ‰|‰|j         d d…         ¦  «        ‘ŒS )r"   NrK   )rM   r   r’   Úbatch_dims_broadcasteds     €€r   rt   z"prepare_inputs.<locals>.<listcomp>Ø   sC   ø€ ð ð ð àð 	#Ð" 1Ð&<¸a¼gÀbÀcÀc¼lÑKÔKðð ð r   )Úcrow_indicesrŠ   Úcol_indicesrG   Úvaluesr   rO   r$   )ÚbsrÚdense_tensorsr–   r—   r˜   rQ   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   )	rR   r–   rŽ   r—   r˜   r$   r   Úsparse_compressed_tensorr   )r   r™   rQ   Ú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    rt          d| 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``).
    rd   r   z3tensor should have 3 dimensions after reshape, got )ÚndimrŠ   r�   rA   r$   )Útensors    r   Ú	as1Dbatchr­   ú   s†   € ð Œ+˜Š/ˆ/Ø×!Ò! !Ñ$Ô$ˆð Œ+˜Š/ˆ/à„{�Q‚€Ø—’  6¤;°¡?Ñ3Ô3ˆØ„{�aÒÐÝØPÀ&Ä,ÐPÐPñ
ô 
ð 	
ð €Mr   ©Úaccumulatorsc                ó¢  — |d         }| j         dk    rt          d| j         › d�¦  «        ‚| j        \  }}}|dk    �rm|dd…         \  }}	|j         dk    rt          d|j         › d�¦  «        ‚|j        \  }
}}||k    rt          d	|› d
|› d�¦  «        ‚|€5|j        d         dz
  }t          j        |||f| j        | j        ¬¦  «        }nC|j        \  }}}||k    rt          d|› d|› d�¦  «        ‚||k    rt          d|› d|› d�¦  «        ‚|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    �rm|j        }t          |¦  «        }|j        \  }}}||z  dk    rt          d|› d|› d�¦  «        ‚|dd…         \  }}}}}|d         }|€`|| 
                    ¦   «                              ¦   «         dz   |z  z   } t          j        g |dd…         ¢| ‘|‘R | j        | j        ¬¦  «        }n.|j        dd…         \  } }!|!|k    rt          d|!› d|› d�¦  «        ‚|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    �rœ|j        }t          |¦  «        }|j        \  }}}||z  dk    rt          d|› d|› d�¦  «        ‚|dd…         \  }}}}|d         }|€`|| 
                    ¦   «                              ¦   «         dz   |z  z   } t          j        g |dd…         ¢| ‘|‘R | j        | j        ¬¦  «        }n.|j        dd…         \  } }!|!|k    rt          d|!› d|› d�¦  «        ‚|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'          |¦  «        ‚)aw  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..

    Auxiliary 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   rd   zblocks must be 3D, got ÚDÚ
scatter_mmr   Nzothers must be 3D, got z
blocks K (z) != others K (ú)©r-   r   zaccumulators Ms (z) != blocks Ms (zaccumulators Ns (z) != others Ns (r;   Úbsr_strided_mmzK (z) must be divisible by Ks (ÚSPLIT_Nr"   úaccumulators N (ú) != others N (Úbsr_strided_mm_compressedr‰   )r«   rA   r$   r   Úzerosr-   r   Ú_scatter_mm2rf   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Úg1rv   Ú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à„{�aÒÐÝÐE°v´{ÐEÐEÐEÑFÔFÐFØ”�J€BˆˆBà˜Ò%Ñ%Ø$ Q R RÔ(‰ˆ	�2àŒ;˜!ÒÐÝ Ð!I¸6¼;Ð!IÐ!IÐ!IÑJÔJÐJØ”l‰ˆˆC�Ø�Š9ˆ9Ý Ð!G¨bÐ!GÐ!GÀÐ!GÐ!GÐ!GÑHÔHÐHàÐØ” Ô" QÑ&ˆAÝ œ;Ø�B˜� 6¤<¸¼ðñ ô ˆLˆLð 'Ô,‰KˆAˆs�CØ�bŠyˆyÝ$Ð%S¸Ð%SÐ%SÈbÐ%SÐ%SÐ%SÑTÔTÐTØ�bŠyˆyÝ$Ð%S¸Ð%SÐ%SÈbÐ%SÐ%SÐ%SÑTÔTÐTà�‰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Øˆr‰6�QŠ;ˆ;Ý Ð!J qÐ!JÐ!JÀRÐ!JÐ!JÐ!JÑKÔKÐKà;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Ø�QŠwˆwÝ$Ð%O¸Ð%OÐ%OÈ1Ð%OÐ%OÐ%OÑPÔPÐPà)Ô/ÐÝ  Ñ.Ô.ˆà�'‰\ˆà�‰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Øˆr‰6�QŠ;ˆ;Ý Ð!J qÐ!JÐ!JÀRÐ!JÐ!JÐ!JÑKÔKÐKà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Ø�QŠwˆwÝ$Ð%O¸Ð%OÐ%OÈ1Ð%OÐ%OÐ%OÑPÔPÐPà)Ô/ÐÝ  Ñ.Ô.ˆà�'‰\ˆà�‰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    �r2t          ¦   «         }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    rt          d|› d|› d �¦  «        ‚||k    rt          d!|› d"|› d �¦  «        ‚|| k    rt          d#|› d$| › d �¦  «        ‚||k    rt          d%|› d&|› d �¦  «        ‚||k    rt          d'|› d(|› d �¦  «        ‚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ó   zTILE_M (z) must be <= Ms (r³   zTILE_N (z) must be <= Ns (zMs (z) must be <= M (zNs (z) must be <= N (zKs (z) must be <= K ()ÚTILE_MÚTILE_NÚ
GROUP_SIZEÚ
num_stagesÚ	num_warpsr¶   r7   )	r   r   r   Úfloat16ÚupdateÚgetrD   rA   Ú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Ý&Ñ(Ô(ˆÝØØ��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à�‚{€{ÝÐF¨ÐFÐFÀÐFÐFÐFÑGÔGÐGØ�‚{€{ÝÐF¨ÐFÐFÀÐFÐFÐFÑGÔGÐGØ	ˆA‚v€vÝÐ< BÐ<Ð<¸Ð<Ð<Ð<Ñ=Ô=Ð=Ø	ˆA‚v€vÝÐ< BÐ<Ð<¸Ð<Ð<Ð<Ñ=Ô=Ð=Ø	ˆA‚v€vÝÐ< BÐ<Ð<¸Ð<Ð<Ð<Ñ=Ô=Ð=åð ØØØØØØðð ð ðð ð r   c                 ó�  — |€t           j        }|€|}|€d}||	|
|hd hk    �rst          ¦   «         }| |||||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
   Ú*rd   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   Ú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  ø  s^  € ð* €}Ý”ˆØÐØˆ	ØÐØˆØ�˜J¨Ð7¸D¸6ÒAÑAÝ&Ñ(Ô(ˆØ�!�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 )ÚTensorAsKeyaR  A light-weight wrapper of a tensor that enables storing tensors as
    keys with efficient memory reference based comparison 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t          d| j         › �¦  «        ‚|                      ¦   «         |                      ¦   «         | j        |                      ¦   «         | j         fS )Nz>TensorAsKey does not support floating point or complex dtype: )r-   Úis_floating_pointÚ
is_complexrA   Údata_ptrÚstorage_offsetr$   rE   )Úobjs    r   Úget_tensor_keyz,TensorAsKey.__init__.<locals>.get_tensor_keys  sy   € ð ŒyÔ*ð ¨c¬iÔ.Bð Ý$Ø`ÐUXÔU^Ð`Ð`ñô ð ð —’‘”Ø×"Ò"Ñ$Ô$Ø”	Ø—
’
‘”Ø”	ðð 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__r  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%  r1  s    r   r!  zTensorAsKey.obj¦  s   € ð �}Š}‰ŒÐr   N)	Ú__name__Ú
__module__Ú__qualname__Ú__doc__r/  r2  r6  Úpropertyr!  r7   r   r   r  r  [  sg   € € € € € ðð ð,&$ð &$ð &$ðPð ð ð%ð %ð %ð ðð ñ „Xðð ð r   r  )Úmaxsizec	           	      ó–	  — |j         }	|	€t          d¦  «        ‚|	                     ¦   «         |	                     ¦   «         }}
|
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�¦  «        ‚)Nz+compressed_sparse_tensor_as_key.obj is Noner¹   r´   r   rw   T)Ú
descendingÚstablerµ   r   r²   zInvalid indices_format=z>. Expected bsr_strided_mm_compressed|bsr_strided_mm|scatter_mm)r!  rA   r–   r—   r   r   Úint32Úarangerf   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_datarV  ¬  sT  € ð *Ô
-€CØ
€{ÝÐJÑKÔKÐKØ #× 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                 ó2  — |                       ¦   «         dk    r$t          d|                       ¦   «         › �¦  «        ‚| j        dk    rt          d| j        › d�¦  «        ‚|                      ¦   «         j        dd…         }| j        \  }}|\  }}|j        dd…         \  }	}
|	|k    rt          d|	› d	|› d
�¦  «        ‚|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   zbsr.dense_dim() must be 0, got r
   z$bsr must be 2D (no batch dims), got r±   r"   Nz	other K (z) != bsr K (r³   Ú
allow_tf32©rX  r¶   r¹   T)Úis_compressedrµ   F)Ú	dense_dimrA   r«   r˜   r$   Únumelr  r  r-   r   r   r/   rV  r  )r™   r5  rÆ   Ú
meta_inputr>   rà   rÙ   rÈ   rÉ   ÚK_rÚ   rK  rß   r¶   rÅ   s                  r   Úbsr_scatter_mm_indices_datar_    sÆ  € ð ‡}‚}�„˜!ÒÐÝÐP¸s¿}º}¹¼ÐPÐPÑQÔQÐQØ
„x�1‚}€}ÝÐOÀCÄHÐOÐOÐOÑPÔPÐPØ—
’
‘”Ô" 2 3 3Ô'€IØŒ9�D€A€qØ�F€BˆØŒK˜˜˜Ô�E€BˆØ	ˆQ‚w€wÝÐ=¨Ð=Ð=¸Ð=Ð=Ð=Ñ>Ô>Ð>ØŒ{˜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           
      ó4  — | j         dk    rt          d| j         › d�¦  «        ‚|j         dk     rt          d|j         › d�¦  «        ‚| 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
   zbsr must be 2D, got r±   z+other must have at least 2 dimensions, got r"   rw   Nr¹   )rÆ   r   r´   >   rµ   r¹   r®   r²   r   )rd   r   rò   r
   )r   r
   rd   rò   )r«   rA   r$   r˜   r_  r   rÁ   r-   r   r­   Ú_nnzr¾   r²   r\  rº   r¤   r£   Úmovedimr�   Úcopy_Ú	unflattenÚreshaperÂ   )r™   r5  rÅ   ÚoutrÈ   rÉ   rÎ   r>   rÆ   Ú	out_shaperK  r¯   rÄ   s                r   Úbsr_scatter_mmrh  2  s0  € ð „x�1‚}€}ÝÐ?°C´HÐ?Ð?Ð?Ñ@Ô@Ð@Ø„z�A‚~€~ÝØG¸%¼*ÐGÐGÐGñ
ô 
ð 	
ð ”˜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_alpharf  Úskip_checksÚmax_gridrß   Úinputr™   Údenserj  rk  rf  rl  rm  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   rw   r´   ri  )r-   r   Úint8r–   r#   r$   rR   rÁ   rA  r   r  )rn  r™   ro  r  r  rj  rk  rf  rl  rm  rß   r   r–   Ú
batch_ndimrà   rÚ   Úoriginal_batch_dims_broadcasteds                    r   rq  rq  }  sÒ   € ð €{�u”{¥e¤jÐ0Ð0Ø'ˆØ×'Ò'Ñ)Ô)ˆØ!×%Ò%Ñ'Ô'¨!Ñ+ˆ
ØŒI�jÔ!ˆØŒK˜ŒOˆÝ*>¸vÀsÈEÑ*RÔ*RÐ'ÝŒkØ+¨q°!¨fÑ4Ý”+Ø”<ð
ñ 
ô 
ˆõ
 ØØØØØØØØØØØðñ ô ð r   c                óX
  ‡‡‡
‡ ‡!‡"‡#‡$‡%— 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    r*t          d|                     ¦   «         d         › �¦  «        ‚|                     ¦   «         d         dk    r*t          d|                     ¦   «         d         › �¦  «        ‚‰
€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    rt          d¦  «        ‚ˆ ˆ!ˆ"ˆˆˆ#ˆ$ˆ
ˆ%f	d„}tA          ||||¦  «         | !                    ¦   «         | !                    ¦   «         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
   rd   rw   Nr   FTr7   z'left_alpha.stride()[-1] must be 0, got r"   z(right_alpha.stride()[-2] must be 0, got )r  r-   r  r¶   r‹   r   ©r   NN©r   Nrw   )r   r‹   éüÿÿÿ)r   r‹   Nzalpha must not be 0c                 ó’   •	— 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_COLrX  Ú	acc_dtype)Ú_bsr_strided_addmm_kernelrb   ÚtlÚfloat32)r{   r…   ÚBKÚBMÚBNr  r  Údot_out_dtyper}  rß   r~  s     €€€€€€€€€r   rƒ   zbsr_dense_addmm.<locals>.kernel:  s�   ø€ å! $Ô'ð 	
Ý! >Ð2ð	
àð	
ð ð	
ð 	
ð  š	Ø  AšIØ !šØ/Ø1ØØØØ$­¬
Ò2Ø#ð	
ð 	
ð ð	
ð 	
ð 	
ð 	
ð 	
r   )"r˜   r–   r—   r#   r$   rR   Ú	new_emptyra  r¾   rc  Úmul_Úexpandr£   rE   rA   Úroundr  r-   r›   r  r©   r   r   r„  r…  r/   Úfloat64rr  rA  rž   r1   r@   r†   r  )&rn  r™   ro  r  r  rj  rk  rf  rl  rm  rß   r   r˜   r–   r—   rs  rà   rÙ   r>   rÚ   rt  r  Ú
out_backupr¶   Úout_untiledÚ	n_batchesÚn_block_rowsÚn_block_colsri   rj   rp   rƒ   r†  r‡  rˆ  r‰  r}  r~  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ð
ð 
ð 
ˆð ×ÒÑÔ˜2Ô !Ò#Ð#ÝØO°j×6GÒ6GÑ6IÔ6IÈ"Ô6MÐOÐOñ
ô 
ð 	
ð ×ÒÑÔ˜BÔ 1Ò$Ð$ÝØQ°{×7IÒ7IÑ7KÔ7KÈBÔ7OÐQÐ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ð �‚z€zÝÐ2Ñ3Ô3Ð3ð
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
ð 
õ& �&˜/¨9°kÑBÔBÐBà
‡|‚|�~„~˜×,Ò,Ñ.Ô.Ò.Ð.ð 	×Ò˜×)Ò)¨*Ô*:Ñ;Ô;Ñ<Ô<Ð<àÐr   ÚIS_BETA_ZEROr  r�  ÚTILE_Kr‚  rX  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-   ç        ©Úmaskr5  ©rX  r  )r„  Ú
program_idÚloadrB  rf   rº   ÚdotÚstoreÚtor-   Ú
element_ty)5r  r  r”  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_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‚  rX  Ú	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Ñ  [  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ÐÑðQ(	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Ÿ  rB  rº   rf   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_strider  r�  r‚  rX  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   rw   r   r
   r   )r   N)r   rw   )r   rx  TFc                 óX   •— t          |          ‰‰‰g‰¢‰‘‰	‘t          |Ž ¢R ‰‰dddœŽ d S )Nr   rò   )r‚  rX  rþ   rÿ   )rÑ  rb   )
r{   r…   r‚  rX  r  r  r>   Úis_beta_zeror¤  Útile_ks
     €€€€€€€€r   rƒ   z)_run_sampled_addmm_kernel.<locals>.kernelƒ  sw   ø€ Ý! $Ô'ØØØðð ð	ð
 ðð ðõ & ~Ð6ðð ð $Ø%ØØðð ð ð ð ð r   )rž   r1   r@   r-   r   r.   r/   r„  r…  rŽ  r†   )r  r  rë  r>   r¤  rì  r˜   r–   r—   Úmat1Úmat2rm  r‘  r’  ri   rj   rp   rƒ   r‚  rX  s   ``````            @@r   Ú_run_sampled_addmm_kernelrï  _  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  rf  rl  rm  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"   rw   r   r‹   )r   r¡   r   r   r-   r   Úboolr   r3   r+   r$   ra  r¢  rc  r\  r˜   rž   r‹  r›   r©   r€   rï  rE   re  )rn  rí  rî  r  r  rf  rl  rm  r   Úinput_broadcastedr>   r¤  r�  r–   r—   r˜   rì  s                    r   rñ  rñ  ”  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   )rf  rl  rm  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"   rw   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  rf  )r$   r   r   r   r3   r-   r   rr  r+   rž   r˜   rB   r   rR   Úis_contiguousr¤   rŠ  ra  r¾   r  )r™   ro  rf  rl  rm  rß   r   r¦   Ú_klr§   Ú	row_blockrË  Ú_krrt  Ú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Ÿ  rB  r0   r¢  r…  r€   rf   Ú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‹   rx  rw   .rw  rv  c                 óf   •— t          |          g t          |Ž ¢‰‘‰‘‰‘t          d‰¦  «        ‘R Ž  d S )Ni   )r  rb   rD   )r{   r…   rË  Úmax_row_nnzrù  s     €€€r   rƒ   zbsr_softmax.<locals>.kernel¨  sf   ø€ Ý Ô%ð Ý% ~Ð6ðàðð ðð ð	õ �E˜;Ñ'Ô'ðð ð ð ð ð r   r�   )r   r3   r-   ra  r\  Úcloner$   r˜   ÚtritonÚnext_power_of_2r–   rŠ   r�   r¤   r÷  rF   re  r†   r   rŸ   r—   r   )rn  r  r   r¦   r§   Únnzr–   r˜   ri   rj   rp   rƒ   rË  rù  s    `          @@r   r  r  {  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�¦  «         |€t          d¦  «        ‚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 must not be Nonez(): attn_mask.layout must be z, but got attn_mask.layout == r!   r"   rw   rš  F)r  rl  r   z(): current value of scale == z results in division by zero.r   T)rÕ   Úinplace)r   rA   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  Â  sE  € ð 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  rX  )r„  rž  rB  rŸ  rº   rf   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ü   rX  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_kernelrC  ô  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Ä   rU  Ú
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Ú   rU  s    €€€r   r{   z_scatter_mm2.<locals>.gridI  sH   ø€ àÔ  Ô# aÑ'Ý”˜A˜t Hœ~Ñ.Ô.µ´¸QÀÀXÄÑ1OÔ1OÑOØðð r   rX  rY  r   r‰  )r$   r  r€   r   r   r„  r…  r/   rŽ  r-   r  rC  rE   )rÃ   rÄ   rU  rD  r¯   rÇ   rÙ   rÌ   rÉ  rß   r{   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É   rZ  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ž  rD   rB  rŸ  rº   rf   r   r¡  r¢  r-   r£  )>rK  rÈ   rÉ   rÚ   r'  r(  r)  r*  r+  Úothers_stride_Br-  r.  r/  Úaccumulators_stride_Br1  r2  Úc_indices_ptrÚr_offsets_ptrÚp_offsets_ptrÚq_offsets_ptrrZ  r‰  r¶   rû   rü   rý   rX  rÎ   ÚBLOCKS_MÚBLOCKS_NÚpid_t_r9  Úpid_br8  Únum_pid_in_groupÚgroup_idÚfirst_pid_mÚgroup_size_mr:  r;  r<  r=  r>  r?  r@  rÒ   r¦   r§   rä   rå   rÓ   r  rÔ   Úq_ptrrÊ  rÉ  rÖ   rØ   rA  Úp_ptrrÕ   rB  s>                                                                 r   Ú_scatter_mm6_kernelr[  q  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    rt          d|› d|› d�¦  «        ‚||	z  Š|‰k    rt          d|› 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    ¬	¦  «         | 
                    d
¦  «        dk    r%t          d| 
                    d
¦  «        › �¦  «        ‚‰ 
                    d
¦  «        dk    r%t          d‰ 
                    d
¦  «        › �¦  «        ‚| 
                    d
¦  «        dk    r%t          d| 
                    d
¦  «        › �¦  «        ‚| 
                    d
¦  «        dk    r%t          d| 
                    d
¦  «        › �¦  «        ‚|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¶   r·   r¸   r³   zaccumulators B (z) != others B (c                 ó’   •— ‰j         d         ‰z  t          j        ‰| d         ¦  «        t          j        ‰| d         ¦  «        z  fS )Nr   rû   rü   rG  )rI  rØ   rÈ   rÎ   rÜ   s    €€€€r   r{   z_scatter_mm6.<locals>.gridø  sD   ø€ à” Ô" QÑ&Ý”˜B  X¤Ñ/Ô/µ&´+¸bÀ$ÀxÄ.Ñ2QÔ2QÑQðð r   rX  rY  r   r   z#c_indices.stride(0) must be 1, got z#r_offsets.stride(0) must be 1, got z#p_offsets.stride(0) must be 1, got z#q_offsets.stride(0) must be 1, got r
   r‰  )r$   rA   r   r   r„  r…  r/   rŽ  r-   r  rE   rF   r÷  r[  rc  )rÃ   rÄ   rÛ   rÜ   rÝ   rÞ   rß   r¯   r\  r¶   rÇ   rÉ   Ú_KrÚ   ÚB_Ú_Mrá   r{   r‰  Úaccumulators_rØ   rÈ   rÎ   s      `                @@@r   r½   r½   ã  sÀ  øøøø€ ð �y”/ˆØ”\‰
ˆˆB�Ø”<‰ˆˆ2ˆqØ!Ô'‰
ˆˆB�Ø�Š7ˆ7Ý Ð!K°BÐ!KÐ!KÀqÐ!KÐ!KÐ!KÑLÔLÐLØ�'‰\ˆØ�Š7ˆ7Ý Ð!K°BÐ!KÐ!KÀqÐ!KÐ!KÐ!KÑLÔLÐLð	ð 	ð 	ð 	ð 	ð 	ð 	ð 	õ ŒM�2œ:ÝŒN�BœJÝŒM�2œ:ÝŒM�2œ:ð	
ð
 Ô
ôˆð ˜tÐ#Ð#Ø�KŠK =µB´JÒ#>ˆKÑ?Ô?Ð?à×Ò˜AÑÔ !Ò#Ð#Ý ØK°i×6FÒ6FÀqÑ6IÔ6IÐKÐKñô ð ð ×Ò˜AÑÔ !Ò#Ð#Ý ØK°i×6FÒ6FÀqÑ6IÔ6IÐKÐKñô ð ð ×Ò˜AÑÔ !Ò#Ð#Ý ØK°i×6FÒ6FÀqÑ6IÔ6IÐKÐKñô ð ð ×Ò˜AÑÔ !Ò#Ð#Ý ØK°i×6FÒ6FÀqÑ6IÔ6IÐKÐKñô ð ð" ð 	)Ø×&Ò&Ñ(Ô(ˆ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_striderz  r{  r|  r}  r~  r€  c7                 ó   — |dk    rt          d|› �¦  «        ‚|dk    rt          d|› �¦  «        ‚|dk    rt          d|› �¦  «        ‚|!dk    rt          d|!› �¦  «        ‚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   z+left_alpha_tiled_col_stride must be 0, got z+left_alpha_col_block_stride must be 0, got z,right_alpha_tiled_row_stride must be 0, got z,right_alpha_row_block_stride must be 0, got r
   r—  r   r™  r�  )rA   r„  rž  rÓ  rÔ  rŸ  rB  rº   rf   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_striderc  Úleft_alpha_row_block_striderd  Úright_alpha_ptrÚright_alpha_batch_stridere  Úright_alpha_tiled_col_striderf  Úright_alpha_col_block_striderÛ  rÜ  rÝ  rÞ  rß  rà  r  r  rz  r{  r|  r}  r~  r  r�  r€  r‚  rX  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ƒ  K	  sÑ  € ðV '¨!Ò+Ð+Ý Ø[Ð>YÐ[Ð[ñô ð ð '¨!Ò+Ð+Ý Ø[Ð>YÐ[Ð[ñô ð ð (¨1Ò,Ð,Ý Ø]Ð?[Ð]Ð]ñô ð ð (¨1Ò,Ð,Ý Ø]Ð?[Ð]Ð]ñô ð õ ”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)Br  Úosr#  Ú	functoolsr   r   Útorch._dynamo.utilsr   Útorch.utils._tritonr   Ú_triton_ops_metar   r   ÚintÚgetenvr	   r   r   r   r+   r3   rB   rG   rR   rW   r`   rb   r|   r†   r›   r¡   r©   r­   r²   r  r  r  rV  r_  rh  ÚTensorró  r1   r  rq  r  r  Útriton.languageÚlanguager„  ÚjitÚ	constexprrÑ  rè  rï  rñ  rö  r  r  r0   r  rC  r»   r[  r½   rƒ  r7   r   r   ú<module>r†     s  ðð €€€Ø 	€	€	€	Ø €€€Ø Ð Ð Ð Ð Ð à €€€Ø )Ð )Ð )Ð )Ð )Ð )Ø *Ð *Ð *Ð *Ð *Ð *à 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8ð .1¨SØ€B„IÐ:¸AÑ>Ô>ñ.ô .Ð *ð
ð ð ð
ð ð ðð ð ðð ð ð"ð ð ðð ð ð*ð ð ð&Uð Uð Uðð ð ðð ð ðð ð ð;ð ;ð ;ð8&ð &ð &ð &ð2 7ð  7ð  7ðF	ð 	ð 	ð
/ð 
/ð 
/ðð ð ð" >Bð w2ð w2ð w2ð w2ð w2ð@ ØØØØØðpð pð pð pðv ØØØØØ
ØØð`ð `ð `ð `ðFNð Nð Nð Nð Nñ Nô Nð Nðb €Ð=Ð>Ñ>Ô>ð^
ð ^
ñ ?Ô>ð^
ðD  ;ð!ð !ð !ð !ðHHð Hð Hð Hð` 
Ø
Ø&*Ø'+Ø#ØØAEØð&ð &ð &ØŒ<ð&à	Œð&ð Œ<ð&ð ”˜tÑ#ð&ð ” Ñ$ð&ð 
Œ˜Ñ	ð&ð ð&ð �C˜$‘J  d¡
¨C°$©JÐ6Ô7¸$Ñ>ð&ð �‰+ð&ð &ð &ð &ð\ 
Ø
Ø&*Ø'+Ø#ØØAEØðnð nð nØŒ<ðnà	Œðnð Œ<ðnð ”˜tÑ#ðnð ” Ñ$ðnð 
Œ˜Ñ	ðnð ðnð �C˜$‘J  d¡
¨C°$©JÐ6Ô7¸$Ñ>ðnð �‰+ðnð nð nð nðb €:�<„<ñ S%Ø€M€M€MØ Ð Ð Ð Ð Ð à„Zð}4ð ”lð}4ð ”|ð	}4ð
 ”|ð}4ð ”ð}4ð> ”<ð?}4ð@ ”LðA}4ð }4ð }4ñ „Zð}4ð~ „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 ØØ#'Ø!ØEIðUð Uð UØŒ|ðUàŒlðUð ŒlðUð Œ\˜DÑ ðUð ðUð ˜˜d™
 C¨$¡J°°d±
Ð:Ô;¸dÑBðUð Uð Uð Uðv $(Ø!ØEIØ ð5Jð 5Jð 5JØŒ\ð5JàŒ|ð5Jð Œ\˜DÑ ð	5Jð
 ð5Jð ˜˜d™
 C¨$¡J°°d±
Ð:Ô;¸dÑBð5Jð �T‰kð5Jð 5Jð 5Jð 5Jðn „ZðVð ”\ðVð ŒlðVð Vð Vñ „ZðVðpE
ð E
ð E
ð E
ðX ØØ"ð0ð 0ØŒ|ð0àŒ\ð0ð Œ|ð0ð ”< $Ñ&ð	0ð
 ð0ð ð0ð �t‰|ð0ð 0ð 0ð 0ðd „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 "&ðf.ð f.Ø”ðf.à”ðf.ð ”<ðf.ð ”<ð	f.ð
 ”<ðf.ð ”<ðf.ð ðf.ð ”lðf.ð ðf.ð f.ð f.ð f.ðP „ZðUPðL &(¤\ðMUPðP &(¤\ðQUPðZ ')¤lð[UPð^ ')¤lð_UPðx ”\ðyUPðz œð{UPð| ”lð}UPð~ œ<ðUPð@ œLðAUPðB ”|ðCUPðD ”|ðEUPðF œðGUPðH ”<ðIUPðJ ”LðKUPðL œðMUPðN ”ðOUPð UPð UPñ „ZðUPð UPð UPðp €KØ€LØ€MØ$(Ð!Ø€LØ€LØ $ÐÐÐr   