Ë
    [^(hP ã            !       ó¶  — d dl Z d dlZd dlZd dlm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  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drd„Zd„ Zd„ Zd„ Z d„ Z!ddœd„Z"	 	 	 	 	 	 dsd„Z#	 	 	 	 	 	 	 	 dtd„Z$ G d„ d «      Z% ee¬!«      d"„ «       Z&	 dud#„Z'dvd$„Z(dddddd%ddd&œd'ejR                  d(ejR                  d)ejR                  d*eejR                     d+eejR                     d,eejR                     d-e*d.ee+ee   ee   ee   f      d/ee,   fd0„Z-dddddd%ddd&œd'ejR                  d(ejR                  d)ejR                  d*eejR                     d+eejR                     d,eejR                     d-e*d.ee+ee   ee   ee   f      d/ee,   fd1„Z. e«       �rd dl/Z/d dl0m1Z2 e/jf                  d2e2jh                  d3e2jh                  d4e2jh                  d5e2jh                  d6e2jh                  d7e2jh                  fd8„«       Z5e/jf                  d3e2jh                  d4e2jh                  d6e2jh                  d7e2jh                  d9e2jh                  f
d:„«       Z6d;„ Z7d<d<dd%dd=œd'ejR                  d>ejR                  d?ejR                  d,eejR                     d-e*d.ee+ee   ee   ee   f      fd@„Z8dd%dddAœd(ejR                  d)ejR                  d,eejR                     d-e*d.ee+ee   ee   ee   f      d/ee,   fdB„Z9e/jf                  dCe2jh                  dDe2jh                  fdE„«       Z:drdF„Z;	 	 	 dwdGejR                  dHejR                  dIejR                  dJeejR                     dKe<dLe*dMee<   fdN„Z=e/jf                  dOe2jh                  dPe2jh                  dQe2jh                  dRe2jh                  dSe2jh                  dTe2jh                  d7e2jh                  fdU„«       Z>dVejR                  dWejR                  dXejR                  dYejR                  dZejR                  f
d[„Z?e/jf                  d\e2jh                  d]e2jh                  dRe2jh                  d^e2jh                  dSe2jh                  dTe2jh                  d_e2jh                  d7e2jh                  fd`„«       Z@	 dxdVejR                  dWejR                  daejR                  dbejR                  dcejR                  ddejR                  d/e,dZejR                  dee*fdf„ZAe/jf                  dge2jh                  dhe2jh                  die2jh                  dje2jh                  dke2jh                  dle2jh                  dme2jh                  dne2jh                  doe2jh                  d3e2jh                  d4e2jh                  dpe2jh                  d6e2jh                  d7e2jh                  d9e2jh                  d^e2jh                  f dq„«       ZBydZ;dZ9dZ8dZ=dZ?dZAdZBy)yé    N)Ú	lru_cache)ÚOptional)Ú	warn_once)Ú
has_tritoné   )Úget_metaÚ*TORCH_SPARSE_BSR_SCATTER_MM_LRU_CACHE_SIZEé   c                 ó   — | st        |«      ‚y ©N)Ú
ValueError)ÚcondÚmsgs     úV/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/torch/sparse/_triton_ops.pyÚcheckr      s   € ÙÜ˜‹oÐð ó    c                 óX   — t        |j                  t        j                  k(  | › d�«       y )Nz@(): only BSR sparse format is supported for the sparse argument.)r   ÚlayoutÚtorchÚ
sparse_bsr)Úf_nameÚts     r   Úcheck_bsr_layoutr      s'   € Ü	Ø	�‰”E×$Ñ$Ñ$Øˆ(ÐRÐSõr   c                 ór   — t        |j                  |k(  xr |j                  j                  dk(  | › d�«       y )NÚcudaz9(): all inputs are expected to be on the same GPU device.)r   ÚdeviceÚtype)r   r   r   s      r   Úcheck_devicer   !   s3   € Ü	Ø	�‰�FÑÒ6˜qŸx™xŸ}™}°Ñ6Øˆ(ÐKÐLõr   c           	      ó*  — t        |j                  «       dk\  xr |j                  «       dk\  | › d|j                  «       › d|j                  «       › d�«       |j                  dd  \  }}|j                  dd  \  }}t        ||k(  | › d|› d|› d�«       y )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¨™>Øˆ(ð  Ø #§¡£	˜{Ð*<¸S¿W¹W»Y¸KÀqð	Jôð �Y‰Y�r�sˆ^�F€BˆØ�Y‰Y�r�sˆ^�F€Bˆä	Ø
ˆb‰Øˆ(ð  Ø "˜tÐ#KÈBÈ4Èqð	Rõr   c           	      óè   — t        |j                  |k(  xrD |j                  t        j                  t        j                  t        j
                  ft        |Ž z   v | › d|› d|j                  › d�«       y )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_dtyper2   9   sp   € Ü	Ø	�‰�5Ñò 	SØ�G‰GÜ�Z‰ZœŸ™¬¯©Ð5¼Ð?PÐ8QÑQðSàˆ(ð 3Ø3DÐ2Eð FØŸG™G˜9 Að	'õ	r   c           	      óx   ‡— t        |«      dk(  sJ ‚d„ Šˆfd„}t         ||«      | › d|d   › d|d   › d�«       y )	Nr
   c                 ó   — | | dz
  z   S ©Nr   © )Úvs    r   Úis_power_of_twoz(check_blocksize.<locals>.is_power_of_twoG   s   € Ø˜˜Q™‘KÐ Ð r   c                 ó@   •— d}| D ]  }|dk\  xr  ‰|«      xr |}Œ |S )NTé   r6   )ÚbÚresÚ	blocksizer8   s      €r   Úis_compatible_blocksizez0check_blocksize.<locals>.is_compatible_blocksizeJ   s8   ø€ ØˆØò 	KˆIà ‘?ÒA¡°yÓ'AÒJÀs‰Cð	Kð ˆ
r   z(): sparse inputs' blocksize (r   z, r   z;) should be at least 16 and a power of 2 in each dimension.)Úlenr   )r   r=   r>   r8   s      @r   Úcheck_blocksizer@   D   sX   ø€ Üˆy‹>˜QÒÐÐò!ôô 
Ù 	Ó*Øˆ(Ð0°¸1±°¸bÀÈ1ÁÀð ODð 	Dõr   c                 ó^   — t        | j                  «       «      dkD  r| j                  «       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_contiguousrE   X   s)   € ô ˆ1�8‰8‹:ƒ˜Òð �|‰|‹~Ðàˆr   c                 ór   — 	 t        j                  d„ |D «       Ž S # t        $ r t        d| › d�«       Y y w xY w)Nc              3   ó:   K  — | ]  }|j                   d d –— Œ y ­w©Nr!   ©r#   ©Ú.0r   s     r   ú	<genexpr>z'broadcast_batch_dims.<locals>.<genexpr>m   s   è ø€ Ò'F¸¨¯©°°¬Ñ'Fùó   ‚Fz3(): inputs' batch dimensions are not broadcastable!)r   Úbroadcast_shapesÚ	Exceptionr   )r   Útensorss     r   Úbroadcast_batch_dimsrQ   k   sB   € ðUÜ×%Ñ%Ñ'F¸gÔ'FÐGÐGøÜò UÜˆe˜�xÐRÐSÖTðUús   ‚ ›6µ6c              '   ól   K  — |D ]+  }t        d «      g|j                  «       z  }||| <   ||   –— Œ- y ­wr   )Úslicer"   )r"   Úslice_rangerP   r   Úslicess        r   ÚslicerrV   r   s=   è ø€ Øò ˆÜ˜“+� §¡£Ñ(ˆØ!ˆˆs‰Ø�‰i‹ñùs   ‚24c              '   óš   K  — |D ]B  }t        d «      g|j                  «       z  }t        | |«      D ]  \  }}|€Œ	|||<   Œ ||   –— ŒD y ­wr   )rS   r"   Úzip)ÚdimsrU   rP   r   ÚsÚdÚd_slices          r   Úmultidim_slicerr]   y   s^   è ø€ Øò ˆÜ�4‹[ˆM˜AŸE™E›GÑ#ˆÜ˜d FÓ+ò 	‰JˆAˆwØ‰}Ø��!’ð	ð �‰d‹
ñùs
   ‚7AºAc               '   óV   K  — | D ]  }|–— |j                  «       E d {  –—†  Œ  y 7 Œ­wr   )rC   )rP   r   s     r   Úptr_stride_extractorr_   ‚   s,   è ø€ Øò ˆØŠØ—8‘8“:×Ññàús   ‚)Ÿ' )c           
   #   ó¼  ‡ ‡‡K  — dt        ‰ «      cxk  rdk  sJ ‚ J ‚dt        ‰«      cxk  rdk  sJ ‚ J ‚dd l}ˆ ˆfd„}ˆfd„} |j                   |«       Ž D ]p  }t        ‰ |‰«      D ���	cg c]  \  }}}	t	        ||z
  |	«      ‘Œ }
}}}	t        ||
«      D ��cg c]  \  }}t        |||z   «      ‘Œ }}}|
d d d…   g ||«      ¢­–— Œr y c c}	}}w c c}}w ­w)Nr   é   c               3   óT   •K  — t        ‰‰«      D ]  \  } }t        d| |«      –— Œ y ­w)Nr   )rX   Úrange)ÚfgÚmgÚ	full_gridÚgrid_blockss     €€r   Úgenerate_grid_pointsz.grid_partitioner.<locals>.generate_grid_pointsŽ   s0   øè ø€ Ü˜) [Ó1ò 	#‰FˆB�Ü˜˜2˜rÓ"Ó"ñ	#ùs   ƒ%(c              3   ón   •K  — ‰j                  «       D ]  \  }}t        t        || |«      «      –— Œ y ­wr   )ÚitemsÚnextr]   )rU   r   Út_dimsÚtensor_dims_maps      €r   Úgenerate_sliced_tensorsz1grid_partitioner.<locals>.generate_sliced_tensors’   s7   øè ø€ Ø(×.Ñ.Ó0ò 	;‰IˆAˆvÜ” v¨v°qÓ9Ó:Ó:ñ	;ùs   ƒ25éÿÿÿÿ)r?   Ú	itertoolsÚproductrX   rB   rS   )rf   rg   rm   rp   rh   rn   Ú
grid_pointrd   Úgpre   ÚgridÚgrU   s   ```          r   Úgrid_partitionerrv   ˆ   sþ   úè ø€ Ø”�I“Ô# !Ò#Ð#Ñ#Ð#Ð#Ø”�KÓ Ô% AÒ%Ð%Ñ%Ð%Ð%ãõ#ô;ð (�i×'Ñ'Ñ)=Ó)?Ð@ò ;ˆ
ä/2°9¸jÈ+Ó/V÷
ð 
Ù!+  R¨ŒC��R‘˜Õð
ˆò 
ô 25°ZÀÓ1F×G©¨¨A”%˜˜B ™FÕ#ÐGˆÑGð ‘4�R�4‰jÐ:Ñ2°6Ó:Ñ:Ó:ñ;ùô
ùó Hùs   …A*CÁ/C
Â	CÂCÂ5'Cc                 óœ   ‡— dd d d…   }|€|}n!d„ Št        ˆfd„t        ||«      D «       «      }t        |||«      D ]  ^}} | |g|¢­Ž  Œ y )N)iÿÿÿéÿÿ  rx   ro   c                 ó6   — | €|S t        dt        | |«      «      S r5   )ÚmaxrB   )ru   re   s     r   Úvalid_grid_dimz%launch_kernel.<locals>.valid_grid_dim¨   s!   € ØˆyØ�	ô ˜1œc ! R›jÓ)Ð)r   c              3   ó6   •K  — | ]  \  }} ‰||«      –— Œ y ­wr   r6   )rK   ru   re   r{   s      €r   rL   z launch_kernel.<locals>.<genexpr>¯   s!   øè ø€ ò 
Ù&+ a¨‰N˜1˜b×!ñ
ùs   ƒ)r0   rX   rv   )Úkernelrm   rf   rg   Úcuda_max_gridrt   Úsliced_tensorsr{   s          @r   Úlaunch_kernelr€   ¡   sr   ø€ à.©t°¨tÑ4€MØÐØ#‰ò	*ô ó 
Ü/2°;ÀÓ/Nô
ó 
ˆô "2Ø�; ó"ò &Ðˆˆ~ñ 	ˆtÐ%�nÔ%ñ&r   c           
      ó8  — | j                  «       j                  d«      }| j                  «       j                  d«      }t        | j	                  «       j                  d«      «      }|D �cg c]  }t        |j                  d«      «      ‘Œ }}t        j                  |j                  d d gd„ |D «       ¢­Ž }d„ } |||d«      } |||d«      } ||||j                  dd  «      }|D �cg c]  } ||||j                  dd  «      ‘Œ }}|||g|¢­S c c}w c c}w )Nr   éýÿÿÿc              3   ó:   K  — | ]  }|j                   d d –— Œ y ­wrH   rI   rJ   s     r   rL   z!prepare_inputs.<locals>.<genexpr>Â   s   è ø€ Ò;¨a˜QŸW™W S bœ\Ñ;ùrM   c                 ób   — | j                  ||z   «      j                  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Ç   s1   € Ø�~‰~˜j¨>Ñ9Ó:×BÑBØŒs�:‹ Ñ"ó
ð 	
r   ©ro   r!   )Úcrow_indicesÚ	unsqueezeÚcol_indicesrE   Úvaluesr   rN   r#   )	ÚbsrÚdense_tensorsr‹   r�   rŽ   r   rP   Úbatch_dims_broadcastedr‰   s	            r   Úprepare_inputsr’   ¹   s6  € à×#Ñ#Ó%×/Ñ/°Ó2€LØ—/‘/Ó#×-Ñ-¨aÓ0€KÜ# C§J¡J£L×$:Ñ$:¸1Ó$=Ó>€FØ?LÖM¸!Ô% a§k¡k°!£nÕ5ÐM€GÐMô #×3Ñ3Ø�‰�S�bÐðÙ;°7Ô;òÐò
ñ
 .ØÐ,¨eó€Lñ -¨[Ð:PÐRWÓX€KÙ'ØÐ&¨¯©°R°SÐ(9ó€Fð
 öàñ 	# 1Ð&<¸a¿g¹gÀbÀc¸lÕKð€Gð ð
 ˜ fÐ6¨wÑ6Ð6ùò7 Nùò,s   Á+!DÃ*Dc                 óŠ  — t        | |g|¢­Ž }|j                  «       j                  |dz   «      }|j                  «       j                  |dz   «      }|j	                  «       j                  ||j	                  «       j
                  dd  z   «      }||j
                  dd  z   }t        j                  |||||j                  ¬«      S )NrŠ   r‚   r!   ©Úsizer   )	rQ   r‹   r…   r�   rŽ   r#   r   Úsparse_compressed_tensorr   )r   r�   rP   Úbatch_shaper‹   r�   rŽ   r•   s           r   Úbroadcast_batch_dims_bsrr˜   Ü   s±   € Ü& v¨sÐ=°WÒ=€Kà×#Ñ#Ó%×2Ñ2°;ÀÑ3FÓG€LØ—/‘/Ó#×0Ñ0°¸uÑ1DÓE€KØ�Z‰Z‹\×&Ñ& {°S·Z±Z³\×5GÑ5GÈÈÐ5LÑ'LÓM€FØ˜Ÿ™ 2 3˜Ñ'€DÜ×)Ñ)Ø�k 6°¸S¿Z¹Zôð r   c                 óš   — | j                   �^ }}}|||d   z  |d   ||d   z  |d   gz   }| j                  |«      j                  dd«      S )Nr   r   r‚   r!   )r#   ÚviewÚ	transpose)r   r=   ÚrestÚmÚnÚ	new_shapes         r   Útile_to_blocksizer    é   sd   € Ø—'‘'�K€Tˆ1ˆaØØ	ˆY�q‰\ÑØ�!‰Ø	ˆY�q‰\ÑØ�!‰ð	ñ €Ið �6‰6�)Ó×&Ñ& r¨2Ó.Ð.r   c                 ó  — | j                   dk  r!| j                  d«      } | j                   dk  rŒ!| j                   dkD  r| j                  d| j                   dz
  «      } | j                   dk(  sJ | j                  «       ‚| S )zÙReturn tensor as 3D tensor by either prepending new dimensions to
    the tensor shape (when ``tensor.ndim < 3``), or by collapsing
    starting dimensions into the first dimension (when ``tensor.ndim >
    3``).
    ra   r   )ÚndimrŒ   r†   r#   )Útensors    r   Ú	as1Dbatchr¤   ö   so   € ð �+‰+˜Š/Ø×!Ñ! !Ó$ˆð �+‰+˜‹/à‡{�{�Q‚Ø—‘  6§;¡;°¡?Ó3ˆØ�;‰;˜!ÒÐ)˜VŸ\™\Ó)ÐØ€Mr   ©Úaccumulatorsc                ó  — |d   }| j                   dk(  sJ ‚| j                  \  }}}|dk(  �r|dd \  }}	|j                   dk(  sJ ‚|j                  \  }
}}||k(  sJ ‚|€B|j                  d   dz
  }t        j                  |||f| j                  | j
                  ¬«      }n|j                  \  }}}||k(  sJ ‚||k(  sJ ‚|dz  s|dz  s|dz  st        €^t        |j                  d   dz
  «      D ]>  }||   }||dz      }t        ||«      D ]   }|	|   \  }}||xx   | |   ||   z  z  cc<   Œ" Œ@ |S t        | |||	|«       |S |dk(  �rü|j                  }t        |«      }|j                  \  }}}||z  dk(  sJ ‚|dd \  }}}}}|d	   }|€^||j                  «       j                  «       dz   |z  z   } t        j                  g |dd
 ¢| ‘|‘­| j                  | j
                  ¬«      }n|j                  d
d \  } }!|!|k(  sJ ‚|j                  }"t        |«      }||z  }|dz  s|dz  s|dz  st        €î|j                  «        t        |«      D ]Ï  }#t        |j                  d   «      D ]²  }||   j                  «       }$||   j                  «       }||dz      j                  «       }t        |$|«      \  }%}&||#|%|%|z   …|&|&|z   …f   }'t        ||«      D ]D  }||   ||   }}t        |j                  «       |«      \  }(})|'| |   ||#|(|(|z   …|)|)|z   …f   z  z  }'ŒF Œ´ ŒÑ nt        | |||||||«       |j                  |"«      S |dk(  �r1|j                  }t        |«      }|j                  \  }}}||z  dk(  sJ ‚|dd \  }}}}|d	   }|€^||j                  «       j                  «       dz   |z  z   } t        j                  g |dd
 ¢| ‘|‘­| j                  | j
                  ¬«      }n|j                  d
d \  } }!|!|k(  sJ ‚|j                  }"t        |«      }||z  }|dz  s|dz  s|dz  st        €øt        |«      D ]é  }#t        t        |«      «      D ]Ð  }*t        ||*   j                  «       |«      \  }%}&|%|z  }+|&|z  },||+   j                  «       }-||+dz      j                  «       }.||#|%|%|z   …|&|&|z   …f   }'t!        t        |-|.«      «      D ]Q  \  }/}||,|.z  ||,z
  |-z  z   |/z      j                  «       }t        ||«      \  }(})|'| |   ||#|(|(|z   …|)|)|z   …f   z  z  }'ŒS ŒÒ Œë n>t        j"                  d|j                  |j
                  ¬«      }t        | |||||||«       |j                  |"«      S t%        |«      ‚)ad  Scattered matrix multiplication of tensors.

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

    accumulators (Tensor, optional): a tensor of matrix product
      accumulators. If ``indices_data[0]=="scatter_mm"``, the tensor
      is a 1-D batch tensor of output matrices. Otherwise, output
      matrices are slices of the :attr:`accumulators` tensor.
    r   ra   Ú
scatter_mmr   N©r,   r   r:   Úbsr_strided_mmÚSPLIT_Nr!   Úbsr_strided_mm_compressed)r   )r¢   r#   r   Úzerosr,   r   Ú_scatter_mm2rc   r¤   rz   Ú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Úg1ru   ÚpÚqÚothers_shapeÚBÚKÚNÚ	c_indicesÚ	r_offsetsÚ	p_offsetsÚ	q_offsetsÚmetar«   ÚMÚN_Úaccumulators_shaper;   Úr_Úr0Úr1ÚaccÚq0Úq1Újr�   rž   Úc0Úc1Úis0                                                   r   r¨   r¨     s‹  € ðH " !‘_€Nà�;‰;˜!ÒÐÐØ—‘�J€BˆˆBà˜Ó%Ø$ Q RÐ(‰ˆ	�2à�{‰{˜aÒÐÐØ—l‘l‰ˆˆC�Ø�SŠyÐˆyàÐØ—‘ Ñ" QÑ&ˆAÜ Ÿ;™;Ø�B˜� 6§<¡<¸¿¹ô‰Lð '×,Ñ,‰KˆAˆs�CØ˜"’9Ð�9Ø˜"’9Ð�9à�Š7�b˜2’g  b¢¬LÐ,@Ü˜9Ÿ?™?¨1Ñ-°Ñ1Ó2ò =�Ø˜q‘\�Ø˜q 1™uÑ%�Ü˜r 2›ò =�AØ˜a™5‘D�A�qØ  “O v¨a¡y°6¸!±9Ñ'<Ñ<”Oñ=ð=ð Ðô ˜ ¨°B¸ÔEØÐà	Ð+Ó	+Ø—|‘|ˆÜ˜6Ó"ˆà—,‘,‰ˆˆ1ˆaØ�2‰v˜Š{Ðˆ{à;GÈÈÐ;KÑ8ˆ	�9˜i¨°DØ�y‘/ˆàÐØ�i—m‘m“o×*Ñ*Ó,¨qÑ0°QÑ6Ñ6ˆAÜ Ÿ;™;Ø*�,˜s Ð#Ð* QÐ*¨Ñ*°&·,±,ÀvÇ}Á}ô‰Lð !×&Ñ& r sÐ+‰EˆAˆrØ˜’7ˆN�7à)×/Ñ/ÐÜ  Ó.ˆà�'‰\ˆà�Š7�b˜2’g  b¢¬LÐ,@Ø×ÑÔ Ü˜1“Xò 
Q�Ü˜yŸ™¨qÑ1Ó2ò 	Q�AØ" 1™×*Ñ*Ó,�BØ" 1™×*Ñ*Ó,�BØ" 1 q¡5Ñ)×.Ñ.Ó0�BÜ# B¨›]‘F�B˜Ø& q¨"¨r°B©w¨,¸¸RÀ"¹W¸Ð'DÑE�CÜ" 2 r›]ò Q˜Ø(¨™|¨Y°q©\˜1˜Ü!'¨¯©«°!Ó!4™˜˜BØ˜v a™y¨6°!°R¸"¸r¹'°\À2ÈÈRÉÀ<Ð2OÑ+PÑPÑP™ñQñ	Qñ
Qô ØØØØØØØØô	ð × Ñ Ð!3Ó4Ð4à	Ð6Ó	6Ø—|‘|ˆÜ˜6Ó"ˆà—,‘,‰ˆˆ1ˆaØ�2‰v˜Š{Ðˆ{à0<¸Q¸RÐ0@Ñ-ˆ	�9˜i¨Ø�y‘/ˆàÐØ�i—m‘m“o×*Ñ*Ó,¨qÑ0°QÑ6Ñ6ˆAÜ Ÿ;™;Ø*�,˜s Ð#Ð* QÐ*¨Ñ*°&·,±,ÀvÇ}Á}ô‰Lð !×&Ñ& r sÐ+‰EˆAˆrØ˜’7ˆN�7à)×/Ñ/ÐÜ  Ó.ˆà�'‰\ˆà�Š7�b˜2’g  b¢¬LÐ,@Ü˜1“Xò Q�Üœs 9›~Ó.ò 
Q�AÜ# I¨a¡L×$5Ñ$5Ó$7¸Ó;‘F�B˜Ø˜b™�AØ˜b™�AØ" 1™×*Ñ*Ó,�BØ" 1 q¡5Ñ)×.Ñ.Ó0�BØ& q¨"¨r°B©w¨,¸¸RÀ"¹W¸Ð'DÑE�CÜ )¬%°°B«-Ó 8ò Q™˜˜1Ø% a¨"¡f°¸!±¸rÑ/AÑ&AÀAÑ&EÑF×KÑKÓM˜Ü!'¨¨1£™˜˜BØ˜v a™y¨6°!°R¸"¸r¹'°\À2ÈÈRÉÀ<Ð2OÑ+PÑPÑP™ñQñ
QñQô Ÿ™Ø˜IŸO™O°I×4DÑ4DôˆIô ØØØØØØØØô	ð × Ñ Ð!3Ó4Ð4ô " .Ó1Ð1r   c           
      ó\  — ||||	|
|hd hk(  �r.t         j                  j                  «       }t        d| ||||f|dt         j                  df¬«      }|� |j
                  d$i |¤Ž |S | ||fdk(  rU||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	}	�nx| ||fdk(  rU||fdk(  rd}d}d}d}d}
d}	�n[||fd
k(  rd}d}d}d	}d}
d}	�nF||fdk(  rd	}d}d}d	}d}
d	}	�n1||fdk(  �r)d}d}d}d	}d}
d	}	�n| ||fdk(  rd||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(  rd||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	}	nk||fdk(  rd}d}d}d	}d}
d	}	nW||fdk(  rPd	}d}d}d}d}
d	}	nC| ||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œ	j                  |d«      }|dk\  r|dk\  rd}||z  }|€t        |dk  rdnd|«      }|€t        |dk  rdnd|«      }|
xs d}
|	€†t        | |«      dkD  rddddœj                  |d	«      }	n`t        | |«      dk(  rddddœj                  |d	«      }	n:t        | |«      dk(  rdd	dœj                  |d	«      }	ndddœj                  |d	«      }	|xs d	}||k  sJ t        ||¬«      «       ‚||k  sJ t        ||¬«      «       ‚|| k  sJ t        | |¬ «      «       ‚||k  sJ t        ||¬!«      «       ‚||k  sJ t        ||¬"«      «       ‚t        d$||||
|	|d#œ|¤ŽS )%Nr¨   r   ç      à?©Úversion)é   rä   rä   )r:   r:   r   r:   é   )é    ræ   r
   ræ   )é@   rç   )é€   rè   )é   ré   ré   é   rç   rè   )é   rë   rë   )rä   rä   )é   rì   rì   )é   rí   rí   rä   )	r:   ræ   rç   rè   rä   ré   rë   rí   i    ré   rì   rë   )r:   ræ   rç   )r:   ræ   )ÚTILE_Mr»   )ÚTILE_NrÁ   )rÓ   r»   )rÍ   rÁ   )rÌ   r¼   )rî   rï   Ú
GROUP_SIZEÚ
num_stagesÚ	num_warpsr«   r6   )	r   r   Úget_device_namer   Úfloat16ÚupdateÚgetrB   Ú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Ü—j‘j×0Ñ0Ó2ˆÜØØ��1�b˜"ÐØØœŸ™ sÐ+ô	
ˆð ÐØˆD�K‰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˜˜dšØˆGØ	
ˆg‰€BØ€~Ü˜2 š8‘R¨¨RÓ0ˆØ€~Ü˜2 š8‘R¨¨RÓ0ˆØ’˜q€JØÐÜˆq�!‹9�tÒØ A¨1Ñ-×1Ñ1°"°aÓ8‰IÜ��A‹Y˜$ÒØ A¨1Ñ-×1Ñ1°"°aÓ8‰IÜ��A‹Y˜#ÒØ A™×*Ñ*¨2¨qÓ1‰Ià A™×*Ñ*¨2¨qÓ1ˆIØ’˜q€Jà�RŠ<Ð3œ V°Ô3Ó3ˆ<Ø�RŠ<Ð3œ V°Ô3Ó3ˆ<Ø�Š7Ð$”D˜1 Ô$Ó$ˆ7Ø�Š7Ð$”D˜1 Ô$Ó$ˆ7Ø�Š7Ð$”D˜1 Ô$Ó$ˆ7äð ØØØØØØñð ñð r   c                 óL  — |€t         j                  }|€|}|€d}||	|
|hd hk(  �rKt         j                  j                  «       }| |||||dk(  |dk(  |dk(  f}||u r|}n||f}t	        d|||||f¬«      }|€|dk7  rt	        d||||df¬«      }|€||urt	        d||||df¬«      }|€Žt	        dg |d d ¢d‘|dd  ¢­|||df¬«      }|€#||urt	        dg |d d ¢d‘|dd  ¢­|||df¬«      }t        |xs i «      D ]8  }||   }|d   }|d	   }||z  }||z  dk(  sŒ ||k  sŒ&t        |«      }||z  |d	<   Œ: |� |j                  di |¤Ž |S t        d
| ›d|›d|›d|›d|›d|›d|›d|›d|›�«       |xs t        ||z  d«      }|xs d}|
xs d}
|	xs d}	t        d|||
|	dœ|¤ŽS )Nrá   r   r   Úbsr_dense_addmmrâ   r
   Ú*ra   r«   z@bsr_dense_addmm uses non-optimal triton kernel parameters for M=z K=z N=z Ms=z, Ks=z beta=z alpha=z dtype=z out_dtype=rå   )r«   ÚGROUP_SIZE_ROWrñ   rò   r6   )
r   rô   r   ró   r   Úsortedr÷   rõ   r   rz   )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Ü—j‘j×0Ñ0Ó2ˆØ�!�Q˜˜B ¨¡	¨4°1©9°e¸q±jÐAˆØ�IÑØ!‰Mà! 9Ð,ˆMÜØØØØ˜}¨hÐ7ô	
ˆð ˆ<˜H¨šOÜØ!ØØØ! =°#Ð6ô	ˆDð ˆ<˜E¨Ñ2ÜØ! 3¨¸hÈÈsÐ=SôˆDð ˆ<ä$Ø!Ø)�#�b�q�'Ð)˜3Ð)  Q R Ñ)ØØ! =°#Ð6ô	ˆMð Ð$¨°iÑ)?Ü (Ø%Ø-�c˜"˜1�gÐ-˜sÐ- S¨¨ WÑ-ØØ% u¨cÐ2ô	!�ô ˜}Ò2°Ó3ò -�Ø% dÑ+�Ø˜‘G�Ø 	Ñ*�Ø˜‘L�Ø�q‘5˜A“: ! q£&Ü ›;�DØ&'¨1¡f�D˜’Oð-ð ÐØˆD�K‰KÑ ˜%Ò ØˆKô
 ðØ�t˜4˜Q˜D  !  U r e¨6¨b¨U°'°D°7¸(¸E¸8À8ÀUÀHÈLÈiÈ\ð[ôð
 Ò(œ˜Q "™W a›€GØ#Ò( q€NØ’˜q€JØ’˜Q€IÜð ØØ%ØØñ	ð
 ñð r   c                   ó2   — e Zd ZdZd„ Zd„ Zd„ Zed„ «       Zy)ÚTensorAsKeyaS  A light-weight wrapper of a tensor that enables storing tensors as
    keys with efficient memory reference based comparision as an
    approximation to data equality based keys.

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

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

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

    where left_alpha, right_alpha are (* + 1)-D tensors when
    specified, otherwise, these are treated as tensors filled with
    ones.
    rü   r   r
   ra   ro   Nr   FTr6   r!   )r  r,   r  r«   r‚   r   ©r   NN©r   Nro   )r   r‚   éüÿÿÿ)r   r‚   Nc                 ó†   •	— t        |    g t        |Ž ¢‰‘‰‘­‰dk(  ‰dk7  ‰dk(  ‰‰
‰‰‰‰t        j                  k(  ‰dœ
‰	¤Ž y )Nr   r   )
Úbeta_is_oneÚbeta_is_nonzeroÚalpha_is_oneÚleft_alpha_is_oneÚright_alpha_is_oneÚBLOCKSIZE_ROWÚBLOCKSIZE_INNERÚBLOCKSIZE_COLrL  Ú	acc_dtype)Ú_bsr_strided_addmm_kernelr_   ÚtlÚfloat32)rt   r   ÚBKÚBMÚBNr  r   Údot_out_dtyperq  rÒ   rr  s     €€€€€€€€€r   r}   zbsr_dense_addmm.<locals>.kernel  ss   ø€ Ü! $Ñ'ð 	
Ü! >Ð2ð	
àð	
ð ñ	
ð  ™	Ø  A™IØ !™Ø/Ø1ØØØØ$¬¯
©
Ñ2Ø#ñ	
ð ó	
r   )!rŽ   r‹   r�   r"   r#   rQ   Ú	new_emptyrU  r±   rW  Úmul_Úexpandrš   rC   Úroundr  r,   r’   rö   r    r   rô   rx  ry  r.   Úfloat64rf  r5  r•   r0   r?   r€   r  )&rb  r�   rc  r   r  r^  r_  rZ  r`  ra  rÒ   r   rŽ   r‹   r�   rg  rÓ   rÌ   r=   rÍ   rh  r  Ú
out_backupr«   Úout_untiledÚ	n_batchesÚn_block_rowsÚn_block_colsrf   rg   rm   r}   rz  r{  r|  r}  rq  rr  s&      ``     `                     @@@@@@r   rü   rü   …  s   ÿø€ ð, €FØ�Z‰Z‹\€FØ×#Ñ#Ó%€LØ—/‘/Ó#€KØ×!Ñ!Ó# aÑ'€JØ�9‰9�Z *¨q¡.Ð1�D€A€qØ—‘˜Z¨!™^¨j¸1©nÐ=€IØ�‰�B‰€Aô ';¸6À3ÈÓ&NÐ#Ø
€{Ø�o‰oÐ=ÀÀAÀÑFÓGˆà
‡x�xƒz�Q‚˜% 1š*¨¨Qª°!°q²&¸AÀºFØ�1Š9Ø�I‰IŒKð
 ˆ
ð �I‰I�eÔØ�qŠyØ—‘˜”Øˆ
àÐØÐØÐØ ÐØ/�U—_‘_ RÓ(×/Ñ/ð 
Ø,ð
Ø./ð
Ø12ò
‰
ð T�_�Z—_‘_ÐLÐ&EÐLÀqÐLÈ!ÒL×SÑSð 
Ø,ð
Ø./ð
Ø12ò
ˆ
ð ÐØ!ÐØ0�e—o‘o bÓ)×0Ñ0ð 
Ø,ð
Ø./ð
Ø12ò
‰ð VÐ&�k×&Ñ&ÐNÐ(GÐNÈÐNÈAÒN×UÑUð 
Ø,ð
Ø./ð
Ø12ò
ˆð ×ÑÓ˜rÑ" aÒ'Ð'Ð'Ø×ÑÓ Ñ# qÒ(Ð(Ð(à€|Ü˜˜SŸX™X›Z¨)°A©,Ñ6¸À1¹ÑEÈÈQÉÑOÑOÐQRÓSˆÜ#ØØØØ�a‰LØ�a‰LØØØØ—+‘+Ø—i‘iô
ˆð €Jô 	�s˜E 5¨*°kÀ3Ó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€JÜ# K°"°b°Ó:€Kô 	�‰”r—z‘zÜ�‰œŸ
™
Ü�‰”r—z‘zÜ�‰”r—z‘zÜ�
‰
”B—H‘HÜ�‰”R—X‘Xðð 
‡i�iñ€Mð —
‘
˜1“€IØ×$Ñ$ RÓ(¨1Ñ,€LØ—:‘:˜b“>€Là˜L¨,Ð7€IØÐÜ˜H R a˜L©¨2¨Ñ.Ó/°'¸QÄÀXÈbÈqÀ\ÓARÑ=RÑ2SÑS‰àˆð 	�Ø�mØ�_Øˆ{Øˆ}Ø�KØ�[Øˆ[ð	€Oð �AŠ:Ðˆ:÷
ô 
ô$ �&˜/¨9°kÔBà
‡|�|ƒ~˜×,Ñ,Ó.Ò.ð 	×Ñ˜×)Ñ)¨*×*:Ñ*:Ó;Ô<àÐr   ÚIS_BETA_ZEROrs  ru  ÚTILE_Krv  rL  c            
      ó¶  — t        j                  d¬«      } t        j                  d¬«      }!||| z  z   ||!z  z   }"t        j                  |"«      }#t        j                  |"|z   «      }$|$|#z
  }%|%dk(  ry 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 �]0  }-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                  j                  «      «       |(|	z  }(|)|z  })�Œ3 y )Nr   ©Úaxisr   ©r,   ç        ©Úmaskr)  ©rL  r  )rx  Ú
program_idÚloadr6  rc   r­   ÚdotÚstoreÚtor,   Ú
element_ty)5r  r   rˆ  rs  ru  Ú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_striderv  rL  Ú	batch_pidÚrow_block_pidÚcrow_indices_offset_ptrÚ
nnz_offsetÚnnz_offset_nextÚrow_nnzÚrow_block_arangeÚcol_block_arangeÚvalues_block_ptrsÚcol_index_nnz_ptrÚmat1_block_ptrsÚmat2_block_ptrsÚk_tile_arangeÚ_Ú	acc_blockÚ	col_blockÚk_tileÚ	k_offsetsÚmask_kÚ
mat1_blockÚ
mat2_blocks5                                                        r   Ú_sampled_addmm_kernelrÅ  2  s
  € ôF —M‘M qÔ)ˆ	ÜŸ™¨1Ô-ˆð Ø'¨)Ñ3ñ4à! MÑ1ñ2ð 	 ô
 —W‘WÐ4Ó5ˆ
ÜŸ'™'Ð"9Ð<OÑ"OÓPˆð " JÑ.ˆØ�aŠ<ØäŸ9™9 Q¨Ó6ÐÜŸ9™9 Q¨Ó6Ðð Ø! IÑ-ñ.à *Ñ,ñ-ð &Ð(8º¸D¸Ñ(AÑAñBð &Ð(8¸ºq¸Ñ(AÑAñ	Bð 	ð Ø&¨Ñ2ñ3à  :Ñ-ñ.ð 	ð Ø )Ñ+ñ,à# mÑ3ñ4ð $Ð&6²q¸$°wÑ&?Ñ?ñ@ð 	ð Ø )Ñ+ñ,à#Ð&6°tºQ°wÑ&?Ñ?ñ@ð 	ô Ÿ	™	 ! VÓ,ˆÜ�w“ó &	4ˆAÜŸ™ -°Ð!?ÀyÔQˆIô Ÿ™Ð 1Ó2ˆIä  1 fÓ-ò �Ø" ]Ñ2�	Ø" Q™�äŸW™WØ#Ð&;¸iÈÊaÈÑ>PÑ&PÑPØ ¢a ™Øô�
ô  ŸW™WØ#Ø+¨iÑ7ñ8à+¨iº¸4¸Ñ.@Ñ@ñAð  ¢ 4 ™Øô�
ð œRŸV™VØ 
°zÈYôñ ‘	ð%ñ, Ø˜UÑ"‘	à! IÑ-°´r·w±wÐ?PÓ7QÑ0QÑQ�	ô �H‰HÐ&¨	¯©°Z×5EÑ5E×5PÑ5PÓ(QÔRð Ð!2Ñ2ÐØÐ!3Ñ3ÒñM&	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y 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 ]m  },t        j                  |'«      }-t        j                  |*«      }.t        j                  |(||.z  z   «      }/|+t        j                  |-|/||¬«      z  }+|'|z  }'|*|
z  }*Œo t        j                  |)|+j                  |j                  j                  «      «       y )Nr
   r‹  r   r   r�  r‘  ©rx  r’  Únum_programsÚ	swizzle2dr“  r6  r­   rc   r”  r•  r–  r,   r—  )0r™  rš  r›  rœ  r�  rž  rŸ  r   r¡  r¢  r£  Ú	dense_ptrÚdense_batch_strideÚdense_tiled_row_strideÚdense_tiled_col_strideÚdense_row_block_strideÚdense_col_block_strideÚ
output_ptrÚoutput_batch_strideÚoutput_tiled_row_strideÚoutput_tiled_col_strideÚoutput_row_block_strideÚoutput_col_block_striders  ru  rv  rL  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Ý  °  sð  € ô\ —M‘M qÔ)ˆ	ÜŸ™¨1Ô-ˆÜŸ™¨1Ô-ˆÜ—‘¨AÔ.ˆÜ—‘¨AÔ.ˆä')§|¡|Ø˜=¨,¸Ànó(
Ñ$ˆ�}ð
 Ø'¨)Ñ3ñ4à! MÑ1ñ2ð 	 ô
 —W‘WÐ4Ó5ˆ
ÜŸ'™'Ð"9Ð<OÑ"OÓPˆð " JÑ.ˆØ�aŠ<ØäŸ9™9 Q¨Ó6ÐÜŸ9™9 Q¨Ó6Ðð Ø! IÑ-ñ.à *Ñ,ñ-ð &Ð(8º¸D¸Ñ(AÑAñBð &Ð(8¸ºq¸Ñ(AÑAñ	Bð 	ð Ø  9Ñ,ñ-à$ }Ñ4ñ5ð %Ð'7º¸4¸Ñ'@Ñ@ñAð %Ð'7¸ºa¸Ñ'@Ñ@ñ	Að 	ð Ø! IÑ-ñ.à%¨Ñ5ñ6ð &¨Ñ5ñ6ð &Ð(8º¸D¸Ñ(AÑAñ	Bð
 &Ð(8¸ºq¸Ñ(AÑAñBð 	ð Ø&¨Ñ2ñ3à  :Ñ-ñ.ð 	ô Ÿ8™8 ]°MÐ$BÈ)ÔTÐÜ�w“ò 	4ˆAÜŸ7™7Ð#4Ó5ˆLô ŸG™GÐ$5Ó6ˆMÜŸ'™'Ø Ð#9¸MÑ#IÑIóˆKð
 ¤§¡Ø˜k°jÈIô!ñ Ðð
 Ð!2Ñ2ÐØÐ!3Ñ3Ñð#	4ô( 	�‰�Ð.×1Ñ1°*×2BÑ2B×2MÑ2MÓNÕOr   c           
      ó°  ‡ ‡‡‡‡‡‡‡— |j                  d«      }|j                  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        ||||«       y )Nr   ro   r   r
   r   )r   N)r   ro   )r   rl  TFc                 óL   •— t        |    ‰‰‰g‰¢‰‘‰	‘t        |Ž ¢­‰‰dddœŽ y )Nr   rå   )rv  rL  rñ   rò   )rÅ  r_   )
rt   r   rv  rL  r  r   r=   Úis_beta_zeror˜  Útile_ks
     €€€€€€€€r   r}   z)_run_sampled_addmm_kernel.<locals>.kernelX  sY   ø€ Ü! $Ñ'ØØØðð ð	ð
 ðð ðô & ~Ð6ñð $Ø%ØØôr   )r•   r0   r?   r,   r   r-   r.   rx  ry  r‚  r€   )r  r   rà  r=   r˜  rá  rŽ   r‹   r�   Úmat1Úmat2ra  r…  r†  rf   rg   rm   r}   rv  rL  s   ``````            @@r   Ú_run_sampled_addmm_kernelrä  4  sâ   ÿ€ ð —K‘K “Nˆ	Ø#×(Ñ(¨Ó,¨qÑ0ˆà Ð-ˆ	ØÐÜ ¨¨! ©T¨r¨TÑ 2Ó3°gÀÄSÈÐRTÐSTÈÓEVÑAVÑ6WÑW‰KàˆKà�IØ˜'Ø˜Ø�'Ø�)ð
ˆð �<‰<œEŸJ™J¬¯©Ð7Ñ7ÜŸ
™
ˆIØ‰JäŸ
™
ˆIØˆJ÷	ó 	ô 	�f˜o¨y¸+ÕFr   g      ð?)r   r  rZ  r`  ra  râ  rã  c                ón  — d}t        || «       t        || ||«      }	|�s�t        ||| j                  «       t        ||| j                  «       |dk7  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(  xr! |j                  «       | j                  «       k(  |› d|	j                  › d|	j                  «       › d|j                  › d	|j                  «       › �	«       |€|	j                  |j                  d
¬«      }n|j                  |	«       |j                  «       dk(  s|j                  «       dk(  r|S |j                  «       j                  dd  }
|j!                  d«      }|dk(  s|dk(  r!|j                  «       j#                  |«       |S |}t%        |||«      \  }}}}}t'        ||
d   |f«      }t'        |||
d   f«      }t)        |
Ž }t+        |||dk(  |
||||||||«       |j                  «       j-                  «       dd  |j-                  «       dd  k7  rF|j                  «       j                  |j/                  |j                  «       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!   ro   r   r‚   )r   r˜   r   r   r,   r   Úboolr   r2   r*   r#   rU  r–  rW  rP  rŽ   r•   r  r’   r    rz   rä  rC   rY  )rb  râ  rã  r   r  rZ  r`  ra  r   Úinput_broadcastedr=   r˜  rƒ  r‹   r�   rŽ   rá  s                    r   ræ  ræ  i  sô  € ð !ˆä˜ Ô'Ü4°V¸UÀDÈ$ÓOÐâÜ˜  u§|¡|Ô4Ü˜  u§|¡|Ô4Ø�sŠ{˜uŸ{™{¬e¯j©jÑ8ÜØØ�hÐ1°$°Ð7jÐkôð �{‰{¤%§*¡*Ñ,Ü˜F D¨%¯+©+Ô6Ü˜F D¨%¯+©+Õ6ä˜F D¨$¯*©*Ô5Ü& v¨t°TÔ:ØˆÜ  ¨Ô-Ü˜V S¨$¯+©+Ô6Ü˜F C¨¯©Ô5ÜØ—I‘IÐ!2×!8Ñ!8Ñ8ÒW¸S¿X¹X»ZÈ5Ï:É:Ë<Ñ=WØ�hÐ?Ð@Q×@WÑ@WÐ?Xð Y-Ø->×-CÑ-CÓ-EÐ,Fð G+Ø+.¯9©9¨+°_ÀSÇXÁXÃZÀLðRôð ˆ;Ø#×&Ñ& t§z¡z¸Ð&Ó=‰Cà�I‰IÐ'Ô(à�9‰9‹;˜!Ò˜sŸx™x›z¨QšØˆJà—J‘J“L×&Ñ& r sÐ+ˆ	Ø�I‰I�b‹Mˆð �CŠ<˜1 š6Ø�J‰J‹L×Ñ˜dÔ#ØˆJð ˆ
Ü8FÀsÈDÐRVÓ8WÑ5ˆ�k 6¨4°ä  ¨	°!©°aÐ'8Ó9ˆÜ  ¨¨9°Q©<Ð'8Ó9ˆÜ�i�ˆä!ØØØ�C‰KØØØØØØØØØô	
ð$ ×ÑÓ×%Ñ%Ó'¨¨Ð,°·±³ÀÀÐ0DÒDØ×ÑÓ×%Ñ% f§n¡n°Z×5FÑ5FÓ5H×5NÑ5NÓ&OÔPØÐr   )rZ  r`  ra  rÒ   c                ó@  — d}| j                   dd  \  }}|s­t        || «       t        || |j                  «       t	        || |j
                  t        j                  f«       t        || |«       |j                  d«      }	| j                  «       j                   dd  \  }
}t        ||
|f«       t        |	dz   |› d|	› d�«       n|j                   dd  \  }}	t        || |«      }|�o|sm|||	fz   }t        |j                   |k(  d|› d|j                   › d	�«       t        |j                  «       xs  |j                  dd«      j                  «       d
«       |€|j!                  |||	fz   «      }| j#                  «       dk(  r|j%                  «       S t'        || |dd|¬«      S )NÚbsr_dense_mmr!   ro   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   rZ  )r#   r   r   r   r2   r,   r   rf  r*   r•   rŽ   r@   r   rQ   Úis_contiguousr›   r~  rU  r±   rü   )r�   rc  rZ  r`  ra  rÒ   r   r�   Ú_klrž   Ú	row_blockr¿  Ú_krrh  Úexpected_out_shapes                  r   rë  rë  À  s«  € ð  ˆØ—‘˜2˜3�‰ˆˆ3ÙÜ˜V SÔ)Ü˜  e§l¡lÔ3Ü˜  U§[¡[´5·:±:°-Ô@Ü& v¨s°EÔ:à—
‘
˜2“ˆAØ#&§:¡:£<×#5Ñ#5°b°cÐ#:Ñ ˆI�yÜ˜F Y°	Ð$:Ô;ÜØ˜‘F�
Ø�(Ð0°°Ð3MÐNõð
 —[‘[  Ð%‰FˆC�ä*>¸vÀsÈEÓ*RÐ'àˆ?¡;Ø!@ÀAÀqÀ6Ñ!IÐÜØ—	‘	Ð/Ñ/ðØ.Ð/¨z¸#¿)¹)¸ÀAðGôô
 Ø×!Ñ!Ó#ÒL s§}¡}°R¸Ó'<×'JÑ'JÓ'Lð"ôð ˆ;Ø—/‘/Ð"AÀQÈÀFÑ"JÓKˆCð �8‰8‹:˜Š?Ø—9‘9“;Ðô ˜s C¨°a¸aÀSÔ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y t        j                  d|
«      }|||z  k  }|||z  z   ||z  z   ||z  z   }t        j                  ||z   |t	        d«       ¬«      j                  t         j                  «      }t        j                  |d¬«      }t        |
|	|
«      D ]‚  }||
z  }|||z  k  }t        j                  ||z   |t	        d«       ¬«      j                  t         j                  «      }t        j                  |d¬«      }t        j                  ||kD  ||«      }Œ„ t        j                  ||z
  «      }t        j                  |d¬«      }t        |
|	|
«      D ]ƒ  }||
z  }|||z  k  }t        j                  ||z   |t	        d«       ¬«      j                  t         j                  «      }t        j                  ||z
  «      }|t        j                  |d¬«      z  }Œ… t        j                  ||z   ||z  j                  |j                  j                  «      |¬«       t        |
|	|
«      D ]ª  }||
z  }|||z  k  }t        j                  ||z   |t	        d«       ¬«      j                  t         j                  «      }t        j                  ||z
  «      }t        j                  ||z   ||z  j                  |j                  j                  «      |¬«       Œ¬ y )Nr
   r‹  r   r   Úinfr�  )r�  )rx  r’  r“  r6  r/   r–  ry  rz   rc   ÚwhereÚexpÚsumr•  r,   r—  )rž  rŸ  r   r™  rš  rœ  Úvalues_nnz_col_block_striderî  r¿  rñ  rò  r°  Úrow_block_offset_pidr±  r²  r³  r´  rµ  Ú
row_aranger�  Úcurr_row_values_ptrsÚrow_tileÚmax_row_valuer½  Úcurr_max_row_valueÚnumÚdenoms                              r   Ú_bsr_softmax_kernelr  ÷  sJ  € ô —M‘M qÔ)ˆ	Ü!Ÿ}™}°!Ô4ÐÜŸ™¨1Ô-ˆð Ø'¨)Ñ3ñ4à! MÑ1ñ2ð 	 ô
 —W‘WÐ4Ó5ˆ
ÜŸ'™'Ð"9Ð<OÑ"OÓPˆð " JÑ.ˆØ�aŠ<Øä—Y‘Y˜q $Ó'ˆ
Ø˜G iÑ/Ñ/ˆð Ø! IÑ-ñ.à%Ð(<Ñ<ñ=ð ˜9Ñ$ñ%ð 	ô —7‘7Ø  :Ñ-°DÄÀuÃÀô
ç
‰"ŒR�Z‰Z‹.ð 	ô Ÿ™˜x¨aÔ0ˆÜ�t˜[¨$Ó/ò 		ˆAØ˜$ÑˆJØ ¨)Ñ 3Ñ3ˆDÜ—w‘wØ$ zÑ1¸ÄUÈ5Ã\ÀMôç‰b”—‘‹nð ô "$§¡¨°qÔ!9ÐÜŸH™HØÐ 2Ñ2°MÐCUó‰Mð		ô �f‰f�X Ñ-Ó.ˆÜ—‘�s Ô#ˆÜ�t˜[¨$Ó/ò 	)ˆAØ˜$ÑˆJØ ¨)Ñ 3Ñ3ˆDÜ—w‘wØ$ zÑ1¸ÄUÈ5Ã\ÀMôç‰b”—‘‹nð ô —&‘&˜ MÑ1Ó2ˆCØ”R—V‘V˜C aÔ(Ñ(‰Eð	)ô 	�‰Ø  :Ñ-Ø�5‰[×Ñ˜Z×-Ñ-×8Ñ8Ó9Øõ	
ô
 �t˜[¨$Ó/ò 	ˆAØ˜$ÑˆJØ ¨)Ñ 3Ñ3ˆDÜ—w‘wØ$ zÑ1¸ÄUÈ5Ã\ÀMôç‰b”—‘‹nð ô —&‘&˜ MÑ1Ó2ˆCÜ�H‰HØ$ zÑ1Ø�u‘× Ñ  ×!1Ñ!1×!<Ñ!<Ó=Øöñ	r   c                 ó*  ‡‡‡— d}t        || «       t        || | j                  «       | j                  «       dk(  s| j	                  «       dk(  r| j                  «       S | j                  dd  \  }}| j                  «       }| j                  «       j                  dd  \  ŠŠ‰€t        j                  |«      Šnt        j                  ‰«      Š| j                  «       j                  d«      j                  dd«      }| j                  «       j                  dd«      j                  «       r| j                  «       j                  «       }n| j                  «       }|j                  dd«      j                  «       j                  d«      j                  dd«      j!                  d‰|‰z  «      }|j                  d   ‰|‰z  f}d }	|dd d…f   d|d	i}
ˆˆˆfd
„}t#        ||
||	«        |j!                  d‰|‰«      j                  dd«      j                   | j                  «       j                  Ž }t%        j&                  | j                  «       j                  «       | j)                  «       j                  «       || j                  | j*                  ¬«      S )NÚbsr_softmaxr   r!   r‚   rl  ro   .rk  rj  c                 óR   •— t        |    g t        |Ž ¢‰‘‰‘‰‘t        d‰«      ‘­Ž  y )Ni   )r  r_   rB   )rt   r   r¿  Úmax_row_nnzrî  s     €€€r   r}   zbsr_softmax.<locals>.kernel}  sJ   ø€ Ü Ñ%ð Ü% ~Ð6ðàðð ðð ð	ô �E˜;Ó'ôr   r”   )r   r2   r,   rU  rP  Úcloner#   rŽ   ÚtritonÚnext_power_of_2r‹   rŒ   r†   r›   rì  rD   rY  r€   r   r–   r�   r   )rb  r  r   r�   rž   Únnzr‹   rŽ   rf   rg   rm   r}   r¿  rî  s    `          @@r   r  r  P  s6  ú€ Øˆä˜ Ô'Ü�F˜E 5§;¡;Ô/à�:‰:‹<˜1Ò §¡£°Ò 2Ø—;‘;“=Ð à�{‰{˜2˜3Ð‰ˆˆ1Ø�j‰j‹lˆØ$Ÿ|™|›~×3Ñ3°B°CÐ8Ñˆ	�9àÐÜ ×0Ñ0°Ó3‰Kä ×0Ñ0°Ó=ˆKà×)Ñ)Ó+×5Ñ5°aÓ8×@Ñ@ÀÀBÓGˆð �<‰<‹>×#Ñ# B¨Ó+×9Ñ9Ô;à—\‘\“^×)Ñ)Ó+‰Fà—\‘\“^ˆFà×Ñ˜R Ó$ß‰Z‹\ß‰Y�q‹\ß‰W�Q˜‹^ß‰W�R˜ C¨)¡OÓ4ð 	ð —\‘\ !‘_ i°°i±Ð@ˆ	Øˆð ˜˜c˜r˜c˜Ñ" MØ�Oð	
ˆö	ô 	�f˜o¨y¸+ÔFðˆF�N‰N˜2˜y¨#¨yÓ9ß‰Y�r˜2Óß‰W�e—l‘l“n×*Ñ*ð,ð 	ô ×-Ñ-Ø×ÑÓ ×&Ñ&Ó(Ø×ÑÓ×%Ñ%Ó'ØØ—‘Ø—<‘<ô
ð 	
r   Úqueryr  ÚvalueÚ	attn_maskÚ	dropout_pÚ	is_causalÚscalec           	      ó   — d}t        | |› d�«       t        |d u|› d�«       |€J ‚t        |j                  t        j                  k(  |› dt        j                  › d|j                  › d�«       t	        ||| j
                  «       t	        ||| j
                  «       t	        ||| j
                  «       t        ||| j                  «       t        ||| j                  «       |j                  t        j                  urt        ||| j                  «       t        || |j                  dd«      d	d
¬«      }|€| j                  d«      dk(  s|d	k(  rt        d
|› d|› d�«       |€'dt        j                  | j                  d«      «      z  n|}	|j                  «       j                  |	«       t!        |«      }t        j"                  j$                  j'                  |j                  «       |d¬«       t)        ||«      }|S )NÚ_scaled_dot_product_attentionz'(): is_causal == True is not supported.z'(): attn_mask == None is not supported.z(): attn_mask.layout must be z, but got attn_mask.layout == r    r!   ro   rŽ  F)r   r`  r   z(): current value of scale == z results in division by zero.r   T)rÈ   Úinplace)r   r   r   r   r   r   r2   r,   rè  ræ  r›   r•   ÚmathÚsqrtrŽ   r  r  ÚnnÚ
functionalÚdropoutrë  )
r
  r  r  r  r  r  r  r   ÚsdpaÚscale_factors
             r   r  r  —  sÑ  € ð 1ˆÜ�)ˆm ˜xÐ'NÐOÔPÜˆi˜tÐ#¨ xÐ/VÐ%WÔXØÐ$Ð$Ð$äØ×Ñ¤× 0Ñ 0Ñ0Øˆhð (Ü(-×(8Ñ(8Ð'9ð :#Ø#,×#3Ñ#3Ð"4°Að7ô	
ô 	�V˜S %§,¡,Ô/Ü�V˜U E§L¡LÔ1Ü�V˜Y¨¯©Ô5ä�F˜C §¡Ô-Ü�F˜E 5§;¡;Ô/Ø�?‰?¤%§*¡*Ñ,Ü˜ 	¨5¯;©;Ô7äØ�u˜cŸm™m¨B°Ó3¸#È5ô
ˆð ˆ=˜UŸZ™Z¨›^¨qÒ0°E¸S²LÜØØ�(Ð8¸¸ð @/ð /ôð
 9>¸�qœ4Ÿ9™9 U§Z¡Z°£^Ó4Ò4È5ˆØ�‰‹×Ñ˜<Ô(Ü˜4Ó ˆÜ�‰×Ñ×#Ñ# D§K¡K£M°YÈÐ#Ô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y 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                  j                  «      «       y ©Nr   r‹  r   r�  )r  rL  )rx  r’  r6  r“  r­   rc   r”  r•  r–  r,   r—  )+rÓ   rÌ   rÍ   Ú
blocks_ptrÚblocks_stride_PÚblocks_stride_MÚblocks_stride_KÚ
others_ptrÚothers_stride_QÚothers_stride_KÚothers_stride_NÚaccumulators_ptrÚaccumulators_stride_RÚaccumulators_stride_MÚaccumulators_stride_NÚpq_offsets_ptrÚpq_offsets_strideÚpq_ptrÚpq_stride_TÚpq_stride_1r}  rî   rï   rL  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_kernelr8  Å  s,  € ð6 �&‰[ˆä—‘ 1Ô%ˆä�m‰m Ô#ˆØ�r‘	ˆØ�b‘ˆà�V‰^œbŸi™i¨¨6Ó2Ñ2ˆØ�V‰^œbŸi™i¨¨6Ó2Ñ2ˆÜ�Y‰Y�q˜!‹_ˆàØŠq�$ˆw‰K˜/Ñ)¨B¨t²Q¨w©K¸/Ñ,IÑIñ
ˆð ØŠq�$ˆw‰K˜/Ñ)¨B¨t²Q¨w©K¸/Ñ,IÑIñ
ˆô �W‰W�^ eÐ.?Ñ&?Ñ?Ó@ˆÜ�W‰W�^ u¨q¡yÐ4EÑ&EÑEÓFˆà�Š8Øä—H‘H˜f fÐ-°]ÔCˆ	ä�r˜2“ò 	VˆAÜ—‘˜  [¡Ñ0Ó1ˆAÜ—‘˜  [¡Ñ0°;Ñ>Ó?ˆAÜ—‘˜  OÑ 3Ñ3Ó4ˆAÜ—‘˜  OÑ 3Ñ3Ó4ˆAØœŸ™  1°È*ÔUÑU‰Ið	Vð ØÐ+Ñ+ñ,ð ’1�d�7‘Ð3Ñ3Ø�Tš1�W‘+Ð 5Ñ5ñ6ñð 	ô 	�‰�˜	Ÿ™Ð%5×%;Ñ%;×%FÑ%FÓGÕHr   r¶   r·   rI  Ú
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#|	j                  |t
        j                  k(  ¬«       t        |
   ‰|‰| | j                  d	«      | j                  d«      | j                  d«      ||j                  d	«      |j                  d«      |j                  d«      ||j                  d	«      |j                  d«      |j                  d«      ‰‰j                  d	«      ||j                  d	«      |j                  d«      fd
|i|	¤Ž y )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Í   rI  s    €€€r   rt   z_scatter_mm2.<locals>.grid  sI   ø€ à× Ñ  Ñ# aÑ'Ü—‘˜A˜t H™~Ó.´·±¸QÀÀXÁÓ1OÑOØðð r   rL  rM  r   r}  )r#   r÷   rz   r   rô   rx  ry  r.   r‚  r,   rõ   r8  rC   )r¶   r·   rI  r9  r¦   rº   rÌ   r¿   r½  rÒ   rt   r}  rÓ   rÍ   s     `         @@r   r®   r®     s“  ú€ ð —<‘<‰ˆˆAˆqØ—<‘<‰ˆˆAˆqäÜ�r˜1 ™6“?¬3¨r°1¸±6«?ÀqÐTUô
ˆö	ô �M‰Mœ2Ÿ:™:Ü�N‰NœBŸJ™JÜ�M‰Mœ2Ÿ:™:Ü�M‰Mœ2Ÿ:™:ð	
ð
 ×
Ñ
ñˆð ˜tÑ#Ø�K‰K =´B·J±JÑ#>ˆKÔ?Ü˜DÑ!ØØØØØ�M‰M˜!ÓØ�M‰M˜!ÓØ�M‰M˜!ÓØØ�M‰M˜!ÓØ�M‰M˜!ÓØ�M‰M˜!ÓØØ×Ñ Ó"Ø×Ñ Ó"Ø×Ñ Ó"ØØ×Ñ˜aÓ ØØ×Ñ˜aÓ Ø×Ñ˜aÓ ñ)	
ð* (ð+	
ð, ó-	
r   r¼   rN  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   «      }-|rW|-|z  |z  }.|-|z  |z  }/t        j                  ||.z   «      }0t        j                  ||.z   dz   «      }1|/|1z  ||/z
  |0z  z   }2|1|0z
  }3n8t        j                  ||!z   «      }2t        j                  ||!z   dz   «      }4|4|2z
  }3||2z   }5t        j
                  ||f|¬«      }6|r�|+0|z  z  }+t        |3«      D ]j  }7t        j                  |5«      }8t        j                  |,|8z   «      }9t        j                  |+«      }:|6t        j                  |:|9||¬«      z  }6|+|z  }+|5dz  }5Œl n˜||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                  |=|6j                  |j                  j                  «      «       y r  )rx  r’  rB   r6  r“  r­   rc   r”  r•  r–  r,   r—  )>r?  r»   r¼   rÍ   r  r  r  r  r   Úothers_stride_Br"  r#  r$  Úaccumulators_stride_Br&  r'  Úc_indices_ptrÚr_offsets_ptrÚp_offsets_ptrÚq_offsets_ptrrN  r}  r«   rî   rï   rð   rL  rÁ   ÚBLOCKS_MÚBLOCKS_NÚpid_t_r.  Úpid_br-  Únum_pid_in_groupÚgroup_idÚfirst_pid_mÚgroup_size_mr/  r0  r1  r2  r3  r4  r5  rÅ   r�   rž   r×   rØ   rÆ   r	  rÇ   Úq_ptrr¾  r½  rÉ   rË   r6  Úp_ptrrÈ   r7  s>                                                                 r   Ú_scatter_mm6_kernelrP  B  s§  € ð< �'‰\ˆØ˜‘<ˆØ˜‘<ˆä—‘ AÔ&ˆÜ�m‰m Ô#ˆØ˜Ñ!ˆØ˜(Ñ"ˆà%¨Ñ0ÐØÐ*Ñ*ˆØ Ñ+ˆÜ˜8 kÑ1°:Ó>ˆØ˜s \Ñ1Ñ2ˆØÐ'Ñ'¨LÑ8ˆà�V‰^œbŸi™i¨¨6Ó2Ñ2ˆØ�V‰^œbŸi™i¨¨6Ó2Ñ2ˆÜ�Y‰Y�q˜"ÓˆØØŠq�$ˆw‰K˜/Ñ)¨B¨t²Q¨w©K¸/Ñ,IÑIñ
ˆð Ø�oÑ%ñ&à’!�T�'‰{˜_Ñ,¨r°$º°'©{¸_Ñ/LÑLñNð 	ô �G‰G�M EÑ)Ó*ˆáØ�a‘˜B‘ˆAØ�Q‘˜2‘ˆAÜ—‘˜¨Ñ*Ó+ˆBÜ—‘˜¨Ñ*¨QÑ.Ó/ˆBØ�R‘˜7 Q™;¨"Ñ,Ñ,ˆBØ�r‘'‰Cä—‘˜¨Ñ.Ó/ˆBÜ—‘˜¨Ñ.°Ñ2Ó3ˆBØ�r‘'ˆCà Ñ"ˆÜ—H‘H˜f fÐ-°]ÔCˆ	áØ�R˜/Ñ)Ñ)ˆEÜ˜3“Zò �Ü—G‘G˜E“N�Ü—G‘G˜E A™IÓ&�Ü—G‘G˜E“N�ØœRŸV™VØ�q M¸jôñ �	ð ˜Ñ(�Ø˜‘
‘ñð " BÑ&ˆEÜ˜3“Zò 	�Ü—G‘G˜E“N�Ü—G‘G˜E A™IÓ&�Ü—G‘G˜E“N�Ü—G‘G˜E A¨Ñ$7Ñ7Ó8�Ø˜‘
�Ø˜‘
�ØœRŸV™VØ�q M¸jôñ ‘	ð	ð ØñàÐ+Ñ+ñ,ð ’1�d�7‘Ð3Ñ3Ø�Tš1�W‘+Ð 5Ñ5ñ6ñ	ð 	ô 	�‰�˜	Ÿ™Ð%5×%;Ñ%;×%FÑ%FÓGÕHr   rÎ   rÏ   rÐ   rÑ   Úforce_contiguousc	                 óè  ‡‡‡‡— |d   }	| j                   \  }
Š}|j                   \  Š}}|j                   \  }}}||k(  sJ ‚||	z  Š|‰k(  sJ ‚ˆˆˆˆfd„}t        j                  t        j                  t        j
                  t        j                  t        j                  t        j                  t        j                  t        j                  i|j                     }d|vr#|j                  |t        j                  k(  ¬«       |j                  d«      dk(  sJ ‚‰j                  d«      dk(  sJ ‚|j                  d«      dk(  sJ ‚|j                  d«      dk(  sJ ‚|rD| j                  «       } |j                  «       }|j                  «       s|j                  «       }n|}n|}t        |   ‰‰||| | j                  d«      | j                  d«      | j                  d«      ||j                  d«      |j                  d«      |j                  d«      ||j                  d«      |j                  d«      |j                  d«      |‰||fd|i|¤Ž |r#|j                  «       s|j                  |«       y y y )	Nr«   c                 óŽ   •— ‰j                   d   ‰z  t        j                  ‰| d   «      t        j                  ‰| d   «      z  fS )Nr   rî   rï   r<  )r>  rË   r»   rÁ   rÏ   s    €€€€r   rt   z_scatter_mm6.<locals>.gridÇ  sD   ø€ à—‘ Ñ" QÑ&Ü—‘˜B  X¡Ó/´&·+±+¸bÀ$ÀxÁ.Ó2QÑQðð r   rL  rM  r   r   r
   r}  )r#   r   rô   rx  ry  r.   r‚  r,   rõ   rC   rD   rì  rP  rW  )r¶   r·   rÎ   rÏ   rÐ   rÑ   rÒ   r¦   rQ  r«   rº   r¼   Ú_KrÍ   ÚB_Ú_MrÔ   rt   r}  Úaccumulators_rË   r»   rÁ   s      `                @@@r   r°   r°   ´  sS  û€ ð �y‘/ˆØ—\‘\‰
ˆˆB�Ø—<‘<‰ˆˆ2ˆqØ!×'Ñ'‰
ˆˆB�Ø�QŠwˆˆwØ�'‰\ˆØ�QŠwˆˆw÷	ô �M‰Mœ2Ÿ:™:Ü�N‰NœBŸJ™JÜ�M‰Mœ2Ÿ:™:Ü�M‰Mœ2Ÿ:™:ð	
ð
 ×
Ñ
ñˆð ˜tÑ#Ø�K‰K =´B·J±JÑ#>ˆKÔ?à×Ñ Ó" aÒ'Ð'Ð'Ø×Ñ Ó" aÒ'Ð'Ð'Ø×Ñ Ó" aÒ'Ð'Ð'Ø×Ñ Ó" aÒ'Ð'Ð'ñ Ø×&Ñ&Ó(ˆFØ×&Ñ&Ó(ˆFØ×-Ñ-Ô/Ø ,× 7Ñ 7Ó 9‘à ,‘à(ˆMä˜DÑ!ØØØØØØ�M‰M˜!ÓØ�M‰M˜!ÓØ�M‰M˜!ÓØØ�M‰M˜!ÓØ�M‰M˜!ÓØ�M‰M˜!ÓØØ× Ñ  Ó#Ø× Ñ  Ó#Ø× Ñ  Ó#ØØØØñ)	
ð* (ð+	
ð, ò-	
ñ2  L×$>Ñ$>Ô$@Ø×Ñ˜}Õ-ð %AÐr   Úleft_alpha_tiled_col_strideÚleft_alpha_col_block_strideÚright_alpha_tiled_row_strideÚright_alpha_row_block_stridern  ro  rp  rq  rr  rt  c7                 ó   — |dk(  sJ ‚|dk(  sJ ‚|dk(  sJ ‚|!dk(  sJ ‚t        j                  d¬«      }7t        j                  d¬«      }8t        j                  d¬«      }9t        j                  d¬«      }:t        j                  d¬«      };t        j                  |8|9|:|;|5«      \  }8}9|||7z  z   ||8z  z   }<t        j                  |<«      }=t        j                  |<|z   «      }>|>|=z
  }?t        j
                  d|0«      }@t        j
                  d|2«      }At        j
                  d|1«      }B| ||7z  z   ||=z  z   ||@d d …d f   z  z   ||Ad d d …f   z  z   }C|||7z  z   ||9z  z   ||Ad d …d f   z  z   ||Bd d d …f   z  z   }D|#|$|7z  z   |%|8z  z   |&|9z  z   |'|@d d …d f   z  z   |(|Bd d d …f   z  z   }E||	|7z  z   |
|=z  z   }Ft        j                  |0|1f|3¬«      }Gt        |?«      D ]m  }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Œo |-sG|*z  }G|.sF|||7z  z   ||8z  z   ||9z  z   |@d d …d f   z  z   |Bd d d …f   z  z   }LGt        j                  |L«      z  }G|/sF|||7z  z   ||8z  z   | |9z  z   |!@d d …d f   z  z   |"Bd d d …f   z  z   }MGt        j                  |M«      z  }G|,rd|||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Gj                  |#j                  j                  «      «       y )Nr   r
   r‹  r   r�  r‘  rÇ  )Or™  rš  r›  rœ  r�  rž  rŸ  r   r¡  r¢  r£  Ú	input_ptrÚinput_batch_strideÚinput_tiled_row_strideÚinput_tiled_col_strideÚinput_row_block_strideÚinput_col_block_striderÊ  rË  rÌ  rÍ  rÎ  rÏ  Úleft_alpha_ptrÚleft_alpha_batch_strideÚleft_alpha_tiled_row_striderX  Úleft_alpha_row_block_striderY  Úright_alpha_ptrÚright_alpha_batch_striderZ  Úright_alpha_tiled_col_strider[  Úright_alpha_col_block_striderÐ  rÑ  rÒ  rÓ  rÔ  rÕ  r   r  rn  ro  rp  rq  rr  rs  ru  rt  rv  rL  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   rw  rw  	  sÖ  € ðV +¨aÒ/Ð/Ð/Ø*¨aÒ/Ð/Ð/Ø+¨qÒ0Ð0Ð0Ø+¨qÒ0Ð0Ð0ä—M‘M qÔ)ˆ	ÜŸ™¨1Ô-ˆÜŸ™¨1Ô-ˆÜ—‘¨AÔ.ˆÜ—‘¨AÔ.ˆä')§|¡|Ø˜=¨,¸Ànó(
Ñ$ˆ�}ð
 Ø'¨)Ñ3ñ4à! MÑ1ñ2ð 	 ô
 —W‘WÐ4Ó5ˆ
ÜŸ'™'Ð"9Ð<OÑ"OÓPˆð " JÑ.ˆäŸ9™9 Q¨Ó6ÐÜŸY™Y q¨/Ó:ÐÜŸ9™9 Q¨Ó6Ðð Ø! IÑ-ñ.à *Ñ,ñ-ð &Ð(8º¸D¸Ñ(AÑAñBð &Ð(:¸4Â¸7Ñ(CÑCñ	Dð 	ð Ø  9Ñ,ñ-à$ }Ñ4ñ5ð %Ð'9º!¸T¸'Ñ'BÑBñCð %Ð'7¸ºa¸Ñ'@Ñ@ñ	Að 	ð Ø! IÑ-ñ.à%¨Ñ5ñ6ð &¨Ñ5ñ6ð &Ð(8º¸D¸Ñ(AÑAñ	Bð
 &Ð(8¸ºq¸Ñ(AÑAñBð 	ð Ø&¨Ñ2ñ3à  :Ñ-ñ.ð 	ô Ÿ8™8 ]°MÐ$BÈ)ÔTÐä�w“ò 	4ˆAÜŸ7™7Ð#4Ó5ˆLô ŸG™GÐ$5Ó6ˆMÜŸ'™'Ø Ð#9¸MÑ#IÑIóˆKð
 ¤§¡Ø˜k°jÈIô!ñ Ðð
 Ð!2Ñ2ÐØÐ!3Ñ3Ñð#	4ñ& Ø Ñ%Ðá àØ)¨IÑ5ñ6à-°Ñ=ñ>ð .°Ñ=ñ>ð .Ð0@ÂÀDÀÑ0IÑIñ	Jð
 .Ð0@ÀÂqÀÑ0IÑIñJð ð ¤§¡¨Ó 8Ñ8Ðá!àØ*¨YÑ6ñ7à.°Ñ>ñ?ð /°Ñ>ñ?ð /Ð1AÂ!ÀTÀ'Ñ1JÑJñ	Kð
 /Ð1AÀ$ÊÀ'Ñ1JÑJñKð ð ¤§¡Ð(8Ó 9Ñ9ÐáàØ$ yÑ0ñ1à(¨=Ñ8ñ9ð )¨=Ñ8ñ9ð )Ð+;ºA¸t¸GÑ+DÑDñ	Eð
 )Ð+;¸DÂ!¸GÑ+DÑDñEð ñ Ø ¤B§G¡G¨JÓ$7Ñ7Ñ à  D¬2¯7©7°:Ó+>Ñ$>Ñ>Ð ô 	�‰�Ð.×1Ñ1°*×2BÑ2B×2MÑ2MÓNÕOr   r   )NNNNNN)NNNNNNNr   )r¬   )NN)rŽ  FN)T)Cr  Úosr  Ú	functoolsr   Útypingr   r   Útorch._dynamo.utilsr   Útorch.utils._tritonr   Ú_triton_ops_metar   ÚintÚgetenvr	   r   r   r   r*   r2   r@   rE   rQ   rV   r]   r_   rv   r€   r’   r˜   r    r¤   r¨   rú   r  r  rJ  rS  r\  ÚTensorrè  r0   r÷   re  rü   r  Útriton.languageÚlanguagerx  ÚjitÚ	constexprrÅ  rÝ  rä  ræ  rë  r  r  r/   r  r8  r®   rP  r°   rw  r6   r   r   ú<module>r|     s´  ðó Û 	Û Ý Ý ã Ý )Ý *å &ñ .1Ø€B‡I�IÐ:¸AÓ>ó.Ð *ò
ò
òòò"òò(ò&Uòòòò;ó2&ò0 7òF	ò
/òð >Bô m2ðl ØØØØØókðl ØØØØØ
ØØó`÷FKñ Kñ\ Ð=Ô>ñ]
ó ?ð]
ðB  ;óóBDðX 
Ø
Ø)-Ø*.Ø"&ØØMQØò&Ø�<‰<ð&à	�‰ð&ð �<‰<ð&ð ˜Ÿ™Ñ&ð&ð ˜%Ÿ,™,Ñ'ð&ð 
�%—,‘,Ñ	ð&ð ð&ð �u˜X c™]¨H°S©M¸8ÀC¹=ÐHÑIÑJð&ð �4‰.ó&ð\ 
Ø
Ø)-Ø*.Ø"&ØØMQØòfØ�<‰<ðfà	�‰ðfð �<‰<ðfð ˜Ÿ™Ñ&ðfð ˜%Ÿ,™,Ñ'ðfð 
�%—,‘,Ñ	ðfð ðfð �u˜X c™]¨H°S©M¸8ÀC¹=ÐHÑIÑJðfð �4‰.ófñR …<ÛÝ à‡Z�Zð{4ð —l‘lð{4ð —|‘|ð	{4ð
 —|‘|ð{4ð —‘ð{4ð> —<‘<ð?{4ð@ —L‘LòA{4ó ð{4ðz ‡Z�ZðAPðN —|‘|ðOAPðP —|‘|ðQAPðR —<‘<ðSAPðT —L‘LðUAPðV Ÿ™òWAPó ðAPòF3Gðt ØØ&*Ø!ØQUòUØ�|‰|ðUà�l‰lðUð �l‰lðUð �e—l‘lÑ#ðUð ðUð ˜5 ¨#¡°¸±¸xÈ¹}Ð!LÑMÑNóUðv '+Ø!ØQUØ#ò5JØ�\‰\ð5Jà�|‰|ð5Jð �e—l‘lÑ#ð	5Jð
 ð5Jð ˜5 ¨#¡°¸±¸xÈ¹}Ð!LÑMÑNð5Jð �t‰nó5Jðn ‡Z�ZðVð —\‘\ðVð �l‰lòVó ðVópE
ðX ØØ!%ñ,Ø�|‰|ð,à�\‰\ð,ð �|‰|ð,ð ˜EŸL™LÑ)ð	,ð
 ð,ð ð,ð ˜‰ó,ð\ ‡Z�ZðDIØ�<‰<ðDIà�<‰<ðDIð �<‰<ðDIð* —|‘|ð+DIð, —‘ð-DIð. —‘ð/DIð0 —L‘Lò1DIó ðDIðL4
Ø—‘ð4
à—‘ð4
ð —L‘Lð4
ð —L‘Lð	4
ð
 —l‘ló4
ðl ‡Z�ZðoIð �L‰LðoIð* —|‘|ð+oIð, —|‘|ð-oIð. —‘ð/oIð0 —‘ð1oIð2 —‘ð3oIð4 —L‘Lð5oIð6 —L‘Lò7oIó ðoIðt "&ñX.Ø—‘ðX.à—‘ðX.ð —<‘<ðX.ð —<‘<ð	X.ð
 —<‘<ðX.ð —<‘<ðX.ð ðX.ð —l‘lðX.ð óX.ðt ‡Z�ZðIPðL &(§\¡\ðMIPðP &(§\¡\ðQIPðZ ')§l¡lð[IPð^ ')§l¡lð_IPðx —\‘\ðyIPðz Ÿ™ð{IPð| —l‘lð}IPð~ Ÿ<™<ðIPð@ ŸL™LðAIPðB —|‘|ðCIPðD —|‘|ðEIPðF Ÿ™ðGIPðH —<‘<ðIIPðJ —L‘LðKIPðL Ÿ™ðMIPðN —‘òOIPó ñIPðX €KØ€LØ€MØ$(Ð!Ø€LØ€LØ $Ñr   