+
    ˜°³ji^ ã                   óˆ  € ^ RI t ^ RIt^ RIt^ RIHt ^ RIt^ RIHt ^ RIH	t	 ^RI
HtHt ]! ]P                  ! R^4      4      tR tR tR	 tR
 tR tR tR tR tR tR tR tR tRHR ltR tR tR tR t RR/R lt!RIR lt"RJR lt# ! R R4      t$]! ]R7      R 4       t%RKR  lt&RLR! lt'R"^R#^R$RR%RR&RR'R(R)RR*R/R+ R, llt(R"^R#^R$RR%RR&RR'R(R)RR*R/R- R. llt)]	! 4       '       dË   ^ RI*t*^ RI+H,t- ]*P\                  R/ R0 l4       t/]*P\                  R1 R2 l4       t0R3 t1R"R4R#R4R&RR'R(R)R/R5 R6 llt2R&RR'R(R)RR*R/R7 R8 llt3]*P\                  R9 R: l4       t4RHR; lt5RMR< R= llt6]*P\                  R> R? l4       t7R@ RA lt8]*P\                  RB RC l4       t9RNRD RE llt:]*P\                  RF RG l4       t;R# Rt5Rt3Rt2Rt6Rt8Rt:Rt;R# )Oé    N)Ú	lru_cache)Ú	warn_once)Ú
has_triton)Ú_get_device_nameÚget_metaÚ*TORCH_SPARSE_BSR_SCATTER_MM_LRU_CACHE_SIZEc                 ó,   € V '       g   \        V4      hR # ©N)Ú
ValueError)ÚcondÚmsgs   &&ÚY/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/sparse/_triton_ops.pyÚcheckr      s   € ßÜ˜‹oÐñ ó    c                 óZ   € \        VP                  \        P                  8H  V  R 24       R# )z@(): only BSR sparse format is supported for the sparse argument.N)r   ÚlayoutÚtorchÚ
sparse_bsr)Úf_nameÚts   &&r   Úcheck_bsr_layoutr      s'   € Ü	Ø	�‰”E×$Ñ$Ñ$Øˆ(ÐRÐSör   c                 ó€   € \        VP                  V8H  ;'       d    VP                  P                  R9   V  R24       R# )Úcudaz9(): all inputs are expected to be on the same GPU device.N)r   Úxpu)r   ÚdeviceÚtype)r   r   r   s   &&&r   Úcheck_devicer       s6   € Ü	Ø	�‰�FÑ×?Ð?˜qŸx™xŸ}™}°Ñ?Øˆ(ÐKÐLör   c           	      ó2  € \        VP                  4       ^8¬  ;'       d    VP                  4       ^8¬  V  RVP                  4        RVP                  4        R24       VP                  RR w  r4VP                  RR w  rV\        WE8H  V  RV RV R24       R# )é   zc(): all inputs involved in the matrix product are expected to be at least 2D, but got lhs.dim() == z and rhs.dim() == Ú.Nzw(): 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ä	Ø
‰Øˆ(ð  Ø "˜tÐ#KÈBÈ4Èqð	Rör   c           	      ó  € \        VP                  V8H  ;'       dK    VP                  \        P                  \        P                  \        P
                  3\        V!  ,           9   V  R V RVP                   R24       R# )z\(): all inputs are expected to be of the same dtype and one of (half, bfloat16, float32) or z, but got dtype == r    N)r   Údtyper   ÚhalfÚbfloat16ÚfloatÚtuple)r   r   r,   Úadditional_dtypess   &&&*r   Úcheck_dtyper2   8   sv   € Ü	Ø	�‰�5Ñ÷ 	Sð 	SØ�G‰GÜ�Z‰ZœŸ™¬¯©Ð5¼Ð?PÑ8QÕQñSàˆ(ð 3Ø3DÐ2Eð FØŸG™G˜9 Að	'ö	r   c           	      óº   a€ \        V4      ^8w  d   \        R\        V4       24      hR oV3R lp\        V! V4      V  RV^ ,           RV^,           R24       R# )r   z"blocksize must have length 2, got c                 ó*   € W ^,
          ,          '       * # )é   © )Úvs   &r   Úis_power_of_twoÚ(check_blocksize.<locals>.is_power_of_twoG   s   € Ø˜Q�•KÔ Ð r   c                 ó\   <€ R pV  F"  pV^8¬  ;'       d    S! V4      ;'       d    TpK$  	  V# ©Tr6   )ÚbÚresÚ	blocksizer8   s   &  €r   Úis_compatible_blocksizeÚ0check_blocksize.<locals>.is_compatible_blocksizeJ   s7   ø€ ØˆÛˆIà ‘?×JÐA¡°yÓ'A×JÐJÀsŠCñ ð ˆ
r   z(): sparse inputs' blocksize (z, z;) should be at least 16 and a power of 2 in each dimension.N)ÚlenÚAssertionErrorr   )r   r>   r?   r8   s   && @r   Úcheck_blocksizerC   C   sj   ø€ Ü
ˆ9ƒ~˜ÔÜÐAÄ#ÀiÃ.ÐAQÐRÓSÐSò!õô 
Ù 	Ó*Øˆ(Ð0°¸1µ°¸bÀÈ1ÅÀð ODð 	Dör   c                ób   € \        V P                  4       4      ^8”  d   V P                  4       # V # )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.
)ÚminÚstrideÚ
contiguous)r   s   &r   Úmake_triton_contiguousrH   X   s)   € ô ˆ1�8‰8‹:ƒ˜Ôð �|‰|‹~Ðàˆr   c                 ó|   €  \         P                  ! R  V 4       !  #   \         d    \        RT  R24        R# i ; i)c              3   ó>   "  € T F  qP                   R R x € K  	  R # 5i©Nr!   ©r#   ©Ú.0r   s   & r   Ú	<genexpr>Ú'broadcast_batch_dims.<locals>.<genexpr>m   s   é € Ð'F¹g¸¯©°°­»gùó   ‚Fz3(): inputs' batch dimensions are not broadcastable!N)r   Úbroadcast_shapesÚ	Exceptionr   )r   Útensorss   &*r   Úbroadcast_batch_dimsrU   k   sB   € ðUÜ×%Ò%Ñ'F¹gÓ'FÑGÐGøÜô UÜˆe˜�xÐRÐS×TðUús   ‚ �;º;c              '   ó~   "  € V F2  p\        R 4      .VP                  4       ,          pWV &   W4,          x € K4  	  R # 5ir
   )Úslicer"   )r"   Úslice_rangerT   r   Úslicess   &&*  r   ÚslicerrZ   r   s3   é € ÛˆÜ˜“+� §¡£Õ(ˆØ!ˆs‰Ø�iŒó ùs   ‚;=c              '   óÊ   "  € V FX  p\        R 4      .VP                  4       ,          p\        WRR7       F  w  rVVf   K  WdV&   K  	  V\        V4      ,          x € KZ  	  R # 5i)NF©Ústrict)rW   r"   Úzipr0   )ÚdimsrY   rT   r   ÚsÚdÚd_slices   &&*    r   Úmultidim_slicerrc   y   sT   é € ÛˆÜ�4‹[ˆM˜AŸE™E›GÕ#ˆÜ˜d°5×9‰JˆAØŒ}Ø�!“ñ :ð ”�a“�kÔó ùs   ‚:A#Á"A#c               '   ó\   "  € V  F  pVx € VP                  4        R j  x€L
  K!  	  R #  L
5ir
   )rF   )rT   r   s   * r   Úptr_stride_extractorre   ‚   s&   é € ÛˆØŠØ—8‘8“:×Òó áùs   ‚,Ÿ* ,c           
   #   ó^  a aa"  € \        S 4      ^ 8  g   \        S 4      ^8”  d   \        R\        S 4       24      h\        S4      ^ 8  g   \        S4      ^8”  d   \        R\        S4       24      h^ RIpV V3R lpV3R lpVP                  ! V! 4       !   F€  p\	        S VSRR7       UUU	u. uF  w  rxp	\        Wx,
          V	4      NK  	  p
ppp	\	        WjRR7       UUu. uF  w  r‹\        WˆV,           4      NK  	  pppV
RRR1,          .V! V4      O5x € K‚  	  R# u up	ppi u uppi 5i)	r   z"full_grid length must be 0-3, got z$grid_blocks length must be 0-3, got Nc               3   óZ   <"  € \        SSR R7       F  w  r\        ^ W4      x € K  	  R# 5i)Fr\   N)r^   Úrange)ÚfgÚmgÚ	full_gridÚgrid_blockss     €€r   Úgenerate_grid_pointsÚ.grid_partitioner.<locals>.generate_grid_points�   s)   øé € Ü˜) [¸×?‰FˆBÜ˜˜2Ó"Ô"ó @ùs   ƒ(+c              3   óp   <"  € SP                  4        F  w  r\        \        W V4      4      x € K  	  R # 5ir
   )ÚitemsÚnextrc   )rY   r   Út_dimsÚtensor_dims_maps   &  €r   Úgenerate_sliced_tensorsÚ1grid_partitioner.<locals>.generate_sliced_tensors”   s-   øé € Ø(×.Ñ.Ö0‰IˆAÜ” v°qÓ9Ó:Ô:ó 1ùs   ƒ36Fr\   éÿÿÿÿ)rA   rB   Ú	itertoolsÚproductr^   rE   rW   )rk   rl   rs   rw   rm   rt   Ú
grid_pointri   Úgprj   ÚgridÚgrY   s   fff          r   Úgrid_partitionerr}   ˆ   s  úé € Ü
ˆ9ƒ~˜ÔœS ›^¨aÔ/ÜÐAÄ#ÀiÃ.ÐAQÐRÓSÐSÜ
ˆ;Ó˜!Ôœs ;Ó/°!Ô3ÜÐCÄCÈÓDTÐCUÐVÓWÐWãö#õ;ð  ×'Ò'Ñ)=Ó)?Ô@ˆ
ô " )¨Z¸ÈUÕSõ
áS‘
�˜ô �•˜ÖÙSð 	ò 
ô 25°ZÈeÕ1TÔUÑ1T©¨”%˜ �FÖ#Ñ1TˆÑUð ‘4�R�4�jÐ:Ñ2°6Ó:Ñ:Ô:ó Aùô
ùó Vùs   …B+D-Â0D 
ÃD-Ã"D'Ã?.D-c           	      ó   a€ RRRR1,          pVf   TpMKR o\         ;QJ d"    . V3R l\        W4RR7       4       F  NK  	  5M! V3R l\        W4RR7       4       4      p\        W#V4       F  vrVV ! V.VO5!   K  	  R# )éÿÿÿNc                 ó8   € V f   V# \        ^\        W4      4      # r
   )ÚmaxrE   )r|   rj   s   &&r   Úvalid_grid_dimÚ%launch_kernel.<locals>.valid_grid_dim«   s   € ØŠyØ�	ô ˜1œc !›jÓ)Ð)r   c              3   ó8   <"  € T F  w  rS! W4      x € K  	  R # 5ir
   r6   )rN   r|   rj   r‚   s   &  €r   rO   Ú launch_kernel.<locals>.<genexpr>²   s"   øé € ð 
áF‘�ñ ˜1×!Ð!ÛFùs   ƒFr\   )r   éÿÿ  r†   rv   )r0   r^   r}   )Úkernelrs   rk   rl   Úcuda_max_gridr{   Úsliced_tensorsr‚   s   &&&&   @r   Úlaunch_kernelrŠ   ¤   sƒ   ø€ à.©t°¨tÕ4€MØÒØ#‰ò	*÷ ”eô 
ä˜[ÀÕFó
—e‘eô 
ä˜[ÀÕFó
ó 
ˆô
 "2Ø ö"Ðˆñ 	ˆtÐ%�nÖ%ó"r   c           
      ó:  € V P                  4       P                  ^ 4      pV P                  4       P                  ^ 4      p\        V P	                  4       P                  ^ 4      4      pV Uu. uF  p\        VP                  ^ 4      4      NK  	  pp\
        P                  ! VP                  RR .R V 4       O5!  pR pV! W'R4      pV! W7R4      pV! WGVP                  RR 4      pV Uu. uF  pV! WWVP                  RR 4      NK  	  ppW#V.VO5# u upi u upi )r   Nc              3   ó>   "  € T F  qP                   R R x € K  	  R # 5irK   rL   rM   s   & r   rO   Ú!prepare_inputs.<locals>.<genexpr>Æ   s   é € Ð;±7¨aŸW™W S b�\³7ùrQ   c                 óp   € V P                  W,           4      P                  ^ \        V4      ^,
          4      # ©r   )Úbroadcast_toÚflattenrA   )r   Ú
batch_dimsÚinvariant_dimss   &&&r   Úbatch_broadcast_and_squashÚ2prepare_inputs.<locals>.batch_broadcast_and_squashË   s/   € Ø�~‰~˜jÕ9Ó:×BÑBØŒs�:‹ Õ"ó
ð 	
r   éýÿÿÿ©rv   r!   )Úcrow_indicesÚ	unsqueezeÚcol_indicesrH   Úvaluesr   rR   r#   )	ÚbsrÚdense_tensorsr˜   rš   r›   r   rT   Úbatch_dims_broadcastedr”   s	   &*       r   Úprepare_inputsrŸ   ½   s,  € à×#Ñ#Ó%×/Ñ/°Ó2€LØ—/‘/Ó#×-Ñ-¨aÓ0€KÜ# C§J¡J£L×$:Ñ$:¸1Ó$=Ó>€FÙ?LÓM¹}¸!Ô% a§k¡k°!£nÖ5¹}€GÐMô #×3Ò3Ø�‰�S�bÐðÙ;±7Ó;óÐò
ñ
 .Ø¨eó€Lñ -¨[ÐRWÓX€KÙ'Ø¨¯©°R°SÐ(9ó€Fñ
 óáˆAñ 	# 1¸a¿g¹gÀbÀc¸lÖKÙð ð ð
  fÐ6¨wÑ6Ð6ùò7 Nùò,s   Á+#DÃ+Dc                 ó¤  € \        W.VO5!  pVP                  4       P                  VR,           4      pVP                  4       P                  VR,           4      pVP	                  4       P                  W1P	                  4       P
                  RR ,           4      pW1P
                  RR ,           p\        P                  ! WEWgVP                  R7      # )r5   N©Úsizer   r—   r–   r!   )	rU   r˜   r�   rš   r›   r#   r   Úsparse_compressed_tensorr   )r   rœ   rT   Úbatch_shaper˜   rš   r›   r¢   s   &&*     r   Úbroadcast_batch_dims_bsrr¥   à   s§   € Ü& vÐ=°WÓ=€Kà×#Ñ#Ó%×2Ñ2°;ÀÕ3FÓG€LØ—/‘/Ó#×0Ñ0°¸uÕ1DÓE€KØ�Z‰Z‹\×&Ñ& {·Z±Z³\×5GÑ5GÈÈÐ5LÕ'LÓM€FØŸ™ 2 3˜Õ'€DÜ×)Ò)Ø 6¸S¿Z¹Zôð r   c                 óÌ   € V P                   Ev r#pVW1^ ,          ,          V^ ,          WA^,          ,          V^,          .,           pV P                  V4      P                  RR4      # )r   r–   r!   )r#   ÚviewÚ	transpose)r   r>   ÚrestÚmÚnÚ	new_shapes   &&    r   Útile_to_blocksizer­   í   s^   € Ø—'‘'�K€TˆaØØ	�q�\ÕØ�!�Ø	�q�\ÕØ�!�ð	õ €Ið �6‰6�)Ó×&Ñ& r¨2Ó.Ð.r   c                ó  € V P                   ^8  d   V P                  ^ 4      p K$  V P                   ^8”  d$   V P                  ^ V P                   ^,
          4      p V P                   ^8w  d   \        RV P                   24      hV # )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``).
z3tensor should have 3 dimensions after reshape, got )Úndimr™   r‘   rB   r#   )Útensors   &r   Ú	as1Dbatchr±   ú   sp   € ð �+‰+˜Œ/Ø×!Ñ! !Ó$ŠØ‡{�{�Q„Ø—‘  6§;¡;°¥?Ó3ˆØ‡{�{�aÔÜØAÀ&Ç,Á,ÀÐPó
ð 	
ð €Mr   Úaccumulatorsc               óz  € V^ ,          pV P                   ^8w  d   \        RV P                    R24      hV P                  w  rVpVR8X  Ed¼   VR,          w  r‰VP                   ^8w  d   \        RVP                    R24      hVP                  w  r«pW{8w  d   \        RV RV R24      hVfK   VP                  ^ ,          ^,
          p\        P                  ! WÖV3V P
                  V P                  R
7      pM?VP                  w  rÞpWæ8w  d   \        RV RV R24      hWü8w  d   \        RV RV R24      hV^,          '       g'   V^,          '       g   V^,          '       g	   \        f�   \        VP                  ^ ,          ^,
          4       Fe  pVV,          pVV^,           ,          p\        VV4       F9  pV	V,          w  ppVV;;,          V V,          VV,          ,          ,          uu&   K;  	  Kg  	  V# \        WW‰V4       V# VR8X  Ed©   VP                  p\        V4      pVP                  w  pppVV,          ^ 8w  d   \        RV RV R24      hVR,          w  pppppVR,          pVfl   VVP                  4       P                  4       ^,           V,          ,           p \        P                  ! . VR	R OV NVN5V P
                  V P                  R
7      pM+VP                  RR	 w  p p!V!V8w  d   \        RV! RV R24      hVP                  p"\        V4      pVV,          pV^,          '       g(   V^,          '       g   V^,          '       g
   \        Ef7   VP                  4        \        V4       EF  p#\        VP                  ^ ,          4       Fñ  pVV,          P                  4       p$VV,          P                  4       pVV^,           ,          P                  4       p\        V$V4      w  p%p&VV#V%V%V,           1V&V&V,           13,          p'\        VV4       Fe  pVV,          VV,          pp\        VP                  4       V4      w  p(p)V'V V,          VV#V(V(V,           1V)V)V,           13,          ,          ,          p'Kg  	  Kó  	  EK  	  M\        V VVVVVVV4       VP                  V"4      # VR8X  Edõ   VP                  p\        V4      pVP                  w  pppVV,          ^ 8w  d   \        RV RV R24      hVR,          w  ppppVR,          pVfl   VVP                  4       P                  4       ^,           V,          ,           p \        P                  ! . VR	R OV NVN5V P
                  V P                  R
7      pM+VP                  RR	 w  p p!V!V8w  d   \        RV! RV R24      hVP                  p"\        V4      pVV,          pV^,          '       g(   V^,          '       g   V^,          '       g
   \        EfW   \        V4       EFE  p#\        \!        V4      4       EF(  p*\        VV*,          P                  4       V4      w  p%p&V%V,          p+V&V,          p,VV+,          P                  4       p-VV+^,           ,          P                  4       p.VV#V%V%V,           1V&V&V,           13,          p'\#        \        V-V.4      4       F‚  w  p/pVV,V.,          VV,,
          V-,          ,           V/,           ,          P                  4       p\        VV4      w  p(p)V'V V,          VV#V(V(V,           1V)V)V,           13,          ,          ,          p'K„  	  EK+  	  EKH  	  M?\        P$                  ! RVP
                  VP                  R
7      p\        V VVVVVVV4       VP                  V"4      # \'        V4      h)aU  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.
zblocks must be 3D, got ÚDÚ
scatter_mm:r5   NNzothers must be 3D, got z
blocks K (z) != others K (Ú)N©r,   r   zaccumulators Ms (z) != blocks Ms (zaccumulators Ns (z) != others Ns (Úbsr_strided_mmzK (z) must be divisible by Ks (ÚSPLIT_Núaccumulators N (ú) != others N (Úbsr_strided_mm_compressedr!   r�   )r¯   rB   r#   r   Úzerosr,   r   Ú_scatter_mm2rh   r±   r�   ÚitemÚ_scatter_mm6Úzero_Údivmodr§   rA   Ú	enumerateÚemptyÚNotImplementedError)0ÚblocksÚothersÚindices_datar²   Úindices_formatÚ_PÚMsÚKsÚ	c_offsetsÚpqÚ_QÚKs_ÚNsÚRÚMs_ÚNs_ÚrÚg0Úg1r|   Ú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ÔÜÐ6°v·{±{°mÀ1ÐEÓFÐFØ—‘�J€BˆBà˜Õ%Ø$ RÕ(‰ˆ	à�;‰;˜!ÔÜ Ð#:¸6¿;¹;¸-ÀqÐ!IÓJÐJØ—l‘l‰ˆ�ØŒ9Ü  :¨b¨T°ÀÀÀQÐ!GÓHÐHàÒØ—‘ Õ" QÕ&ˆAÜ Ÿ;š;Ø˜� 6§<¡<¸¿¹ô‰Lð '×,Ñ,‰KˆA�CØŒyÜ$Ð'8¸¸Ð=MÈbÈTÐQRÐ%SÓTÐTØŒyÜ$Ð'8¸¸Ð=MÈbÈTÐQRÐ%SÓTÐTà��7Œ7�b˜2—g”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ó 'ñ 3ð Ðô ˜¨¸ÔEØÐà	Ð+Õ	+Ø—|‘|ˆÜ˜6Ó"ˆà—,‘,‰ˆˆ1ˆaØˆr�6�QŒ;Ü  3 q cÐ)DÀRÀDÈÐ!JÓKÐKà;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Ø�QŒwÜ$Ð'7¸°t¸?È1È#ÈQÐ%OÓPÐPà)×/Ñ/ÐÜ  Ó.ˆà�'�\ˆà��7Œ7�b˜2—g”g  b§¤¬LÓ,@Ø×ÑÔ Ü˜1—X�Ü˜yŸ™¨qÕ1Ö2�AØ" 1�×*Ñ*Ó,�BØ" 1�×*Ñ*Ó,�BØ" 1 q¥5Õ)×.Ñ.Ó0�BÜ# B¨›]‘F�B˜Ø& q¨"¨r°B­w¨,¸¸RÀ"½W¸Ð'DÕE�CÜ" 2 rž]˜Ø(¨�|¨Y°q­\˜1˜Ü!'¨¯©«°!Ó!4™˜˜BØ˜v a�y¨6°!°R¸"¸r½'°\À2ÈÈRÍÀ<Ð2OÕ+PÕPÕPšó +ô 3ò ô ØØØØØØØØô	ð × Ñ Ð!3Ó4Ð4à	Ð6Õ	6Ø—|‘|ˆÜ˜6Ó"ˆà—,‘,‰ˆˆ1ˆaØˆr�6�QŒ;Ü  3 q cÐ)DÀRÀDÈÐ!JÓKÐKà0<¸RÕ0@Ñ-ˆ	�9˜i¨Ø�y•/ˆàÒØ�i—m‘m“o×*Ñ*Ó,¨qÕ0°QÕ6Õ6ˆAÜ Ÿ;š;Ø*�,˜s Ð#Ð* QÐ*¨Ñ*°&·,±,ÀvÇ}Á}ô‰Lð !×&Ñ& r sÐ+‰EˆAˆrØ�QŒwÜ$Ð'7¸°t¸?È1È#ÈQÐ%OÓPÐPà)×/Ñ/ÐÜ  Ó.ˆà�'�\ˆà��7Œ7�b˜2—g”g  b§¤¬LÓ,@Ü˜1—X�Üœs 9›~×.�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™˜˜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šô !9ô /ò ô ŸšØ˜IŸO™O°I×4DÑ4DôˆIô ØØØØØØØØô	ð × Ñ Ð!3Ó4Ð4ô " .Ó1Ð1r   c                 óØ  € WgW‰W¥0R 08X  Ed4   \        4       p\        RWW#V3V^ \        P                  R3R7      pVe   VP                  ! R/ VB  V# WV3R8X  dY   W43R8X  d   ^p^p^p^p^p
^p	EMÏW43R8X  d   ^p^ p^p^p^p
^p	EM¹W43R8X  d   ^p^ p^ p^p^p
^p	EM£W43R 8X  d   ^p^ p^ p^p^p
^p	EM�WV3R!8X  dY   W43R8X  d   ^p^p^@p^p^p
^p	EMnW43R8X  d   ^p^ p^@p^p^p
^p	EMXW43R8X  d   ^p^ p^€p^p^p
^p	EMBW43R 8X  d   ^p^@p^@p^p^p
^p	EM,WV3R"8X  dk   W43R8X  d   ^p^p^€p^p^p
^p	EMW43R8X  d   ^p^ p^@p^p^p
^p	MøW43R8X  d   ^p^@p^@p^p^p
^p	MãW43R 8X  d   ^p^@p^@p^p^p
^p	MÎW43R#8X  d   ^p^@p^@p^p^p
^p	M¹WV3R$8X  dj   W43R8X  d   ^p^p^€p^p^p
^p	M›W43R8X  d   ^p^ p^@p^p^p
^p	M†W43R8X  d   ^p^@p^€p^p^p
^p	MqW43R 8X  d   ^p^@p^@p^p^p
^p	M\W43R#8X  d   ^p^@p^@p^p^p
^p	MGWV3R%8X  d?   W43R8X  d   ^p^pRp^p^p
^p	M)W43R8X  d   ^p^ p^@p^p^p
^p	MW43R8X  d   ^p^@p^€p^p^p
^p	Vf5   ^^^ ^^@^^€^R^R^R^R^ R	^@/	P                  V^4      pVR8¼  d
   VR8¼  d   ^pW(,          pVf   \        VR8  d   ^@M^ V4      pVf   \        VR8  d   ^@M^ V4      pT
;'       g    ^p
V	f�   \        W4      R8”  d   ^^^ ^^@^/P                  V^4      p	Mf\        W4      R8X  d   ^^^ ^^@^/P                  V^4      p	M=\        W4      R8X  d   ^^^ ^/P                  V^4      p	M^^^ ^/P                  V^4      p	T;'       g    ^pWc8”  d   \        R
V RV R24      hW~8”  d   \        RV RV R24      hW08”  d   \        RV RV  R24      hWâ8”  d   \        RV RV R24      hWA8”  d   \        RV RV R24      h\        RRVRVRVRV
RV	RV/VB # )&Nrµ   ç      à?©Úversioné   é   é   é   é   i    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¹   r6   )rô   rô   rô   )é   rþ   )é    rÿ   )é@   r   )é€   r  )rõ   rõ   rõ   )rö   rö   rö   )rô   rô   )r÷   r÷   r÷   )rø   rø   rø   )	r   r   r   Úfloat16ÚupdateÚgetrE   rB   Ú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Ü&Ó(ˆÜØØ�1˜"ÐØØœŸ™ sÐ+ô	
ˆð ÒØ�KŠKÑ ˜%Ò ØˆKð �!ˆ9˜
Ô"Øˆx˜8Ô#Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�˜ZÔ'Ø�Ø�Ø�Ø�
Ø�
Ø�	ùØ�AˆY˜*Ô$Øˆx˜8Ô#Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�˜ZÔ'Ø�Ø�Ø�Ø�
Ø�
Ø�	ùØ�AˆY˜+Ô%Øˆx˜8Ô#Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�˜ZÔ'Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�˜ZÔ'Ø�Ø�Ø�Ø�
Ø�
Ø�	øØ�AˆY˜+Ô%Øˆx˜8Ô#Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�˜ZÔ'Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�˜ZÔ'Ø�Ø�Ø�Ø�
Ø�
Ø�	øØ�AˆY˜+Ô%Øˆx˜8Ô#Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�˜XÔ%Ø�Ø�Ø�Ø�
Ø�
Ø�	à‚ð �Ø�Ø�Ø�Ø�Ø�Ø�"Ø�"Ø�"ð

÷ ‰#ˆa�‹*ð 	ð �Œ9˜˜dœØˆGØ	
�€BØ‚~Ü˜2 œ8‘R¨¨RÓ0ˆØ‚~Ü˜2 œ8‘R¨¨RÓ0ˆØ—�˜q€JØÒÜˆq‹9�tÔØ˜Q  A r¨1Ð-×1Ñ1°"°aÓ8‰IÜ�‹Y˜$ÔØ˜Q  A r¨1Ð-×1Ñ1°"°aÓ8‰IÜ�‹Y˜#ÔØ˜Q  A˜×*Ñ*¨2¨qÓ1‰Ià˜Q  A˜×*Ñ*¨2¨qÓ1ˆIØ—�˜q€Jà„{Ü˜x¨ xÐ/@ÀÀÀAÐFÓGÐGØ„{Ü˜x¨ xÐ/@ÀÀÀAÐFÓGÐGØ	„vÜ˜t B 4Ð'7¸°s¸!Ð<Ó=Ð=Ø	„vÜ˜t B 4Ð'7¸°s¸!Ð<Ó=Ð=Ø	„vÜ˜t B 4Ð'7¸°s¸!Ð<Ó=Ð=äñ Øðàðð ðð ð	ð
 ðð ðð ñð r   c                 óü  € Vf   \         P                  pVf   TpVf   RpWyW¨0R 08X  Ed€   \        4       pWW#WE^ 8H  V^8H  V^8H  3pWÍJ d   TpMWÍ3p\        RVVVVV3R7      pVf   VR8w  d   \        RVVVVR3R7      pVf   WÍJd   \        RVVWìR3R7      pVfÇ   \        R. VR,          ORNVR,          O5VVVR3R7      pVf,   WÍJd'   \        R. VR,          ORNVR,          O5VWìR3R7      p\	        T;'       g    / 4       FW  pVV,          pV^,          pVR,          pVV,          pVV,          ^ 8X  g   K7  VV8:  g   K@  \        V4      pVV,          VR&   KY  	  Ve   VP                  ! R/ VB  V# \        RV : R	V: R
V: RV: RV: RV: RV: RV: RV: 24       T;'       g    \        W#,          ^4      pT;'       g    ^pT
;'       g    ^p
T	;'       g    ^p	\        RRVRVRV
RV	/VB # )Nrñ   Úbsr_dense_addmmrò   ºNr   NÚ*:é   NNr¹   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=ÚGROUP_SIZE_ROWrü   rý   r6   )	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Ü&Ó(ˆØ�Q˜B¨¡	¨4°1©9°e¸q±jÐAˆØÓØ!‰Mà!Ð,ˆMÜØØØØ˜}¨hÐ7ô	
ˆð Š<˜H¨œOÜØ!ØØØ! =°#Ð6ô	ˆDð Š<˜EÓ2ÜØ! 3¨¸hÈsÐ=SôˆDð Š<ä$Ø!Ø)�#�b•'Ð)˜3Ð)  R¥Ñ)ØØ! =°#Ð6ô	ˆMð Ò$¨Ó)?Ü (Ø%Ø-�c˜"•gÐ-˜sÐ- S¨¥WÑ-ØØ%¨cÐ2ô	!�ô ˜}×2Ð2°Ö3�Ø% dÕ+�Ø˜•G�Ø 	Õ*�Ø˜•L�Ø�q•5˜A–: ! q¦&Ü ›;�DØ&'¨1¥f�D˜“Oñ 4ð ÒØ�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                   óF   a € ] tR tRt o RtR tR tR t]R 4       t	Rt
V tR# )	ÚTensorAsKeyi[  a
  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                ó|  € R  p\         P                  ! V4      V n        VP                  \        P
                  J d   V! V4      V n        M×VP                  \        P                  \        P                  09   d3   V! VP                  4       4      V! VP                  4       4      3V n        MvVP                  \        P                  \        P                  09   d3   V! VP                  4       4      V! VP                  4       4      3V n        M\        VP                  4      h\!        V P                  4      V n        R# )c                 ó,  € V P                   P                  '       g   V P                   P                  '       d   \        R V P                    24      hV P	                  4       V P                  4       V P                  V P                  4       V P                   3# )z>TensorAsKey does not support floating point or complex dtype: )r,   Úis_floating_pointÚ
is_complexrB   Údata_ptrÚstorage_offsetr#   rF   )Úobjs   &r   Úget_tensor_keyÚ,TensorAsKey.__init__.<locals>.get_tensor_keys  sz   € ð �y‰y×*×*Ð*¨c¯i©i×.B×.BÐ.BÜ$ØTÐUX×U^ÑU^ÐT_Ð`óð ð —‘“Ø×"Ñ"Ó$Ø—	‘	Ø—
‘
“Ø—	‘	ðð r   N)Ú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__ÚTensorAsKey.__init__r  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                ó   € V P                   # r
   )r2  ©r3  s   &r   Ú__hash__ÚTensorAsKey.__hash__š  s   € Ø�z‰zÐr   c                ó¢   € \        V\        4      '       g   R # V P                  e   VP                  f   WJ # V P                  VP                  8H  # )F)Ú
isinstancer  r%  r  )r3  Úothers   &&r   Ú__eq__ÚTensorAsKey.__eq__�  sB   € Ü˜%¤×-Ò-ÙØ�8‰8Ò˜uŸy™yÒ0ð �=Ð Ø�x‰x˜5Ÿ9™9Ñ$Ð$r   c                ó"   € V P                  4       # )z'Return object if alive, otherwise None.)r*  r7  s   &r   r%  ÚTensorAsKey.obj¦  s   € ð �}‰}‹Ðr   )r2  r*  r  N)Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__r4  r8  r=  Úpropertyr%  Ú__static_attributes__Ú__classdictcell__)Ú__classdict__s   @r   r  r  [  s.   ø‡ € ñò,&$òPò%ð ñó ör   r  )Úmaxsizec	           
      óö	  € VP                   p	V	f   \        R4      hV	P                  4       V	P                  4       rºV
P                  p\
        P                  pV R8X  Edi   W7,          p. p\
        P                  ! W}VR7      V,          p\        W,          4       FŽ  pV
V,          P                  4       pV
V^,           ,          P                  4       pVV8X  d   KA  VP                  VVV WS,          ,          P                  V4      VP                  VV,
          4      ,           4       K�  	  \
        P                  ! V4      pV
P                  4       pVP                  4       pVWC,          ,          pVV,           P!                  R4      pT
pVV,          P                  V4      pVP#                  RRR7      w  ppVV,          pV VVV3# V R8X  EdÊ   W7,          p. p. p\
        P                  ! W}VR7      V,          p\        W,          4       FÅ  pV
V,          P                  4       pV
V^,           ,          P                  4       pVV8X  d   KA  VP                  \
        P                  ! VVWÜR7      P                  V4      4       VP                  VVV WS,          ,          P                  V4      VP                  VV,
          4      ,           4       KÇ  	  \
        P                  ! V4      pV
P                  4       pVP                  4       pVWC,          ,          pVV,           P!                  R4      p\
        P                  ! V
R,          \
        P$                  ! VV,          P                  V4      ^ 4      34      p\
        P                  ! V4      pV VVVV3# V R8X  EdL   Tp^ .p. p\        V4       EF  p\        W,          4       Fê  pV
V,          P                  4       pV
V^,           ,          P                  4       p\        W>,          4       F�  pVP                  VR,          V,           V,
          4       \        VV,
          4       F^  p VV ,           p!VV!,          P                  4       VW%,          ,          ,           W>,          ,          V,           p"VP                  V!V".4       K`  	  KŸ  	  Kì  	  EK  	  V \
        P&                  ! VWÜR7      \
        P&                  ! VWÜR7      3# \)        R	V : R
24      h)Nz+compressed_sparse_tensor_as_key.obj is Noner¼   r·   T)Ú
descendingÚstabler¸   :Nr5   Nrµ   zInvalid indices_format=z>. Expected bsr_strided_mm_compressed|bsr_strided_mm|scatter_mmrv   )r%  rB   r˜   rš   r   r   Úint32Úarangerh   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_datarc  ¬  sI  € ð *×
-Ñ
-€CØ
‚{ÜÐJÓKÐKØ #× 0Ñ 0Ó 2°C·O±OÓ4E�+Ø× Ñ €FÜ—K‘K€MàÐ4Õ4Ø�\ˆØˆÜ�LŠL˜¸fÔEÈÕJˆÜ�q•w–ˆAØ˜a•×%Ñ%Ó'ˆBØ˜a !�eÕ$×)Ñ)Ó+ˆBØ�RŒxÙØ× Ñ Ø˜R Ð# r¥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•w–ˆAØ˜a•×%Ñ%Ó'ˆBØ˜a !�eÕ$×)Ñ)Ó+ˆBØ�RŒxÙØ× Ñ Ü—’˜R ¨=ÔH×OÑOÐPWÓXôð × Ñ Ø˜R Ð# r¥vÕ.×6Ñ6°wÓ?Ø×%Ñ% b¨2¥gÓ.õ/öñ  ô —I’I˜mÓ,ˆ	Ø(×-Ñ-Ó/ÐØ0×8Ñ8Ó:ÐØ  B¥FÕ+ˆØ˜•U—L‘L Ó$ˆ	Ü—I’Ià˜RÕ Ü—’Ø%Ð&:Õ;×MÑMÈgÓVØóðó
ˆ	ô —I’I˜mÓ,ˆ	Ø 	¨9°iÀÐKÐKà	˜<Õ	'ØˆØ�Cˆ	Øˆ
ä�x—ˆAÜ˜1�7–^�Ø! !•_×)Ñ)Ó+�Ø! ! a¥%Õ(×-Ñ-Ó/�Ü˜q�wž�AØ×$Ñ$ Y¨r¥]°RÕ%7¸"Õ%<Ô=Ü" 2¨¥7ž^˜Ø �F˜Ø(¨�^×0Ñ0Ó2°Q¸!½'µ]ÕBÀqÅwÕOÐRSÕS˜Ø"×)Ñ)¨1¨a¨&Ö1ó ,ó (ô $ñ !ð Ü�LŠL˜¨-ÔGÜ�LŠL˜¨=ÔHð
ð 	
ô Ø&�~Ñ'Ð'eÐfó
ð 	
r   c                ó  € V P                  4       ^ 8w  d   \        RV P                  4        24      hV P                  ^8w  d   \        RV P                   R24      hV P                  4       P                  RR pV P                  w  rVVw  rxVP                  RR w  ršW–8w  d   \        RV	 RV R24      hVP                  RR P                  4       p\        WVW§V3/ VB pRV9  d>   VP                  V P                  \        P                  \        P                  09   R	7       VR
,          p\        W%WjWxW½\        V 4      4	      pVR8X  d   VP                  RR7       Wì3,           # VR8X  d   VP                  RR7       Wì3,           # V# )zcComputes indices data for :func:`scatter_mm` used in BSR and
strided tensor matrix multiplication.
zbsr.dense_dim() must be 0, got z$bsr must be 2D (no batch dims), got r´   Nz	other K (z) != bsr K (r¶   Ú
allow_tf32©re  r¹   r¼   T)Úis_compressedr¸   Fr!   )Ú	dense_dimrB   r¯   r›   r#   Únumelr  r  r,   r   r  r.   rc  r  )rœ   r<  rÉ   Ú
meta_inputr>   rã   rÜ   rË   rÌ   ÚK_rÝ   rX  râ   r¹   rÈ   s   &&&,           r   Úbsr_scatter_mm_indices_datarl    sr  € ð ‡}�}ƒ˜!ÔÜÐ>¸s¿}¹}»Ð>OÐPÓQÐQØ
‡x�x�1„}ÜÐCÀCÇHÁHÀ:ÈQÐOÓPÐPØ—
‘
“×"Ñ" 2 3Ð'€IØ�9‰9�D€AØ�F€BØ�K‰K˜˜Ð�E€BØ	„wÜ˜y¨¨¨L¸¸¸1Ð=Ó>Ð>Ø�{‰{˜3˜BÐ×%Ñ%Ó'€Hä˜1 ¨Ñ9¨jÑ9€DØ˜:Ô%Ø�‰˜sŸy™y¬U¯]©]¼E¿N¹NÐ,KÑKˆÔLØ�9�o€GÜ/Ø˜1 ¨¼KÈÓ<Ló€Lð Ð4Ô4Ø�‰ $ˆÔ'Ø˜gÕ%Ð%Ø	Ð+Ô	+Ø�‰ %ˆÔ(Ø˜gÕ%Ð%àÐr   c           
     óJ  € V P                   ^8w  d   \        RV P                    R24      hVP                   ^8  d   \        RVP                    R24      hV P                  R
,          V P                  R,          VP                  R,          repV P                  4       P                  R
R pVf   \	        WRR7      pV^ ,          pVfB   \
        P                  ! . VP                  RR
 OVNVN5V P                  V P                  R7      pVP                  p	\        V4      pV P                  4       ^ 8X  d   VP                  4        EMÍVR9   d/   VP                  4        \        V P                  4       WVR7       EM˜VR	8X  Ed†   VP                  RR
 P                  4       p
\
        P                  ! W¤,          V^ ,          ,          V,          V^ ,          ,          V^ ,          V^ ,          3V P                  V P                  R7      p\        V4      P                  R
R4      P!                  V
Wg^ ,          ,          V^ ,          WW^,          ,          V^,          4      P#                  RR4      P%                  ^ ^4      p\        V P                  4       WÂVR7       VP'                  VP)                  ^ W¤V^ ,          ,          Wg^ ,          ,          34      P#                  RR4      P+                  W¦V4      P                  R
R4      4       M\-        V4      hVP!                  V	4      # )zBSR @ strided -> stridedzbsr must be 2D, got r´   z+other must have at least 2 dimensions, got Nr¼   )rÉ   r·   )r²   rµ   r!   rv   >   r¸   r¼   )r  r5   é   r   )r5   r   r  rn  )r¯   rB   r#   r›   rl  r   rÄ   r,   r   r±   Ú_nnzrÁ   rµ   ri  r½   r¨   r§   Úmovedimr‘   Úcopy_Ú	unflattenÚreshaperÅ   )rœ   r<  rÈ   ÚoutrË   rÌ   rÑ   r>   rÉ   Ú	out_shaperX  r²   rÇ   s   &&&&         r   Úbsr_scatter_mmrv  2  s�  € ð ‡x�x�1„}ÜÐ3°C·H±H°:¸QÐ?Ó@Ð@Ø‡z�z�A„~ÜØ9¸%¿*¹*¸ÀQÐGó
ð 	
ð —‘˜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 Ð#×)Ñ)Ó+ˆÜ—{’{à• ¨1¥Õ-°Õ2°iÀµlÕBØ˜!•Ø˜!•ðð
 —)‘)Ø—:‘:ô
ˆô �eÓß‰Y�r˜2Óß‰TØØ •lÕ"Ø˜!•Ø •lÕ"Ø˜!•ó÷ ‰WØ˜ló÷ ‰W�Q˜‹]ð 	ô 	�3—:‘:“< ÀLÕQØ�	‰	Ø×"Ñ"Ø�H I¨a¥LÕ0°"À!½Õ2DÐEó÷ ‰WØ˜ló÷ ‰W�X 2Ó&ß‰Y�r˜2Óõ		
ô " .Ó1Ð1à�8‰8�IÓÐr   r  r  Ú
left_alphaÚright_alphart  Úskip_checksFÚmax_gridrâ   c                ó˜  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R,          R\         P                  R,          R\         P                  R,          R\        R	\        \        R,          \        R,          \        R,          3,          R,          R
\
        R,          /	# ©r   Úinputrœ   Údenserw  Nrx  rt  ry  rz  râ   ©r   ÚTensorÚboolr0   Úintr  )Úformats   "r   Ú__annotate__r„  }  sª   € ÷ &ñ &Ü�<‰<ð&ä	�‰ð&ô �<‰<ð&ô —‘˜tÕ#ð&ô —‘ Õ$ð&ô 
�‰˜Õ	ð&ô ð&ô ”C˜$•J¤ d¥
¬C°$­JÐ6Õ7¸$Õ>ð&ô ��+ñ&r   c                óš  € Vf³   VP                   \        P                  J d•   RpVP                  4       pVP	                  4       ^,
          pVP
                  V,          pVP
                  R,          p\        W±V4      p\        P                  ! VWï3,           \        P                  VP                  R7      p\        V VVVVVVVVV	V
R7      # )NÚ_int_bsr_dense_addmmr·   )r  r  rw  rx  rt  ry  rz  râ   rv   )r,   r   Úint8r˜   r"   r#   rU   rÄ   rN  r   r
  )r}  rœ   r~  r  r  rw  rx  rt  ry  rz  râ   r   r˜   Ú
batch_ndimrã   rÝ   Úoriginal_batch_dims_broadcasteds   &&&$$$$$$$$      r   r†  r†  }  s¼   € ð ‚{�u—{‘{¤e§j¡jÓ0Ø'ˆØ×'Ñ'Ó)ˆØ!×%Ñ%Ó'¨!Õ+ˆ
Ø�I‰I�jÕ!ˆØ�K‰K˜�OˆÜ*>¸vÈEÓ*RÐ'Ü�kŠkØ+¨q¨fÕ4Ü—+‘+Ø—<‘<ô
ˆô
 ØØØØØØØØØØØôð r   c                ó˜  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R,          R\         P                  R,          R\         P                  R,          R\        R	\        \        R,          \        R,          \        R,          3,          R,          R
\
        R,          /	# r|  r  )rƒ  s   "r   r„  r„  ¦  sµ   € ÷ nñ nÜ�<‰<ðnä	�‰ðnô �<‰<ðnô —‘˜tÕ#ðnô —‘ Õ$ðnô 
�‰˜Õ	ðnô ðnô ”C˜$•J¤ d¥
¬C°$­JÐ6Õ7¸$Õ>ðnô ��+ñnr   c               ó°
  aaa
a a!a"a#a$a%€ RpVP                  4       pVP                  4       pVP                  4       pVP                  4       ^,
          pVP                  Wÿ^,            w  ppVP                  V^,           V^,            pVP                  R,          p\        W±V4      pVf   VP                  VVV3,           4      pVP                  4       ^ 8X  g   S^ 8X  g   V^ 8X  g   V^ 8X  g   V^ 8X  dE   S^ 8X  d   VP                  4        V# VP                  V 4       S^8w  d   VP                  S4       V# Ro$Ro%Vf*   Ro$VP                  R4      P                  ! . VOVNVN5!  pM,VP                  ! . VOVN^N5!  P                  ! . VOVNVN5!  pVf*   Ro%VP                  R4      P                  ! . VOVNVN5!  pM,VP                  ! . VO^NVN5!  P                  ! . VOVNVN5!  pVP                  4       R,          ^ 8w  d$   \        RVP                  4       R,           24      hVP                  4       R,          ^ 8w  d$   \        RVP                  4       R,           24      hS
fƒ   \        ^VP                  4       V^ ,          ,          V^,          ,          VV,          ,          ,
          ^4      p\!        VVVV^ ,          V^,          SSVVP"                  VP"                  R7
      o
Tp\%        WW%Wg4      w  pppp ppppVw  o!o S
P'                  RVS!,          4      pVV,          o"Tp\)        VS!S"34      p\)        VS S"34      p\)        V S!S"34      p \)        VS!S"34      p\)        VS!S"34      p\*        P,                  \.        P0                  \*        P2                  \.        P0                  \*        P0                  \.        P4                  \*        P4                  \.        P4                  \*        P6                  \.        P8                  \*        P8                  \.        P8                  /VP"                  ,          o#VP;                  ^ 4      pVP;                  R4      ^,
          pVP;                  R4      pVVV3pV	eC   \=        V	R	,          RRR1,          4      R^\?        V	R	,          4      ,
          ,          ,           pMRpVRVRVRV RVRVRVRVR/pS^ 8X  d   \        R
4      hV V!V"VVV#V$V
V%3	R lp\A        VVVV4       VPC                  4       VPC                  4       8w  d+   VP                  VP                  VP                  4      4       V# )zð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
  NFTz'left_alpha.stride()[-1] must be 0, got z(right_alpha.stride()[-2] must be 0, got )r  r,   r  r¹   :Nr  Nzalpha must not be 0c                 ó¦   <	€ \         V ,          ! . \        V!  OSNSN5R S^8H  RS^ 8g  RS^8H  RSRS
RSRSRSRS\        P                  8H  R	S/
S	B  R
# )Úbeta_is_oneÚbeta_is_nonzeroÚalpha_is_oneÚleft_alpha_is_oneÚright_alpha_is_oneÚBLOCKSIZE_ROWÚBLOCKSIZE_INNERÚBLOCKSIZE_COLre  Ú	acc_dtypeN)Ú_bsr_strided_addmm_kernelre   ÚtlÚfloat32)r{   r‰   ÚBKÚBMÚBNr  r  Údot_out_dtyper�  râ   r‘  s   &*€€€€€€€€€r   r‡   Úbsr_dense_addmm.<locals>.kernel:  s²   ø€ ä! $Ö'ð 	
Ü! >Ñ2ð	
àð	
ð ò	
ð  ™	ð		
ð
 ! A™Ið	
ð  !™ð	
ð 0ð	
ð  2ð	
ð ð	
ð ð	
ð ð	
ð %¬¯
©
Ñ2ð	
ð $ð	
ð ô	
r   rv   r6   r!   r–   r
   ©r   NN©r   Nrv   )r   r–   éüÿÿÿ)r   r–   N)"r›   r˜   rš   r"   r#   rU   Ú	new_emptyro  rÁ   rq  Úmul_Úexpandr§   rF   rB   Úroundr  r,   rŸ   r  r­   r   r  r—  r˜  r.   Úfloat64r‡  rN  r¢   r0   rA   rŠ   r#  )&r}  rœ   r~  r  r  rw  rx  rt  ry  rz  râ   r   r›   r˜   rš   rˆ  rã   rÜ   r>   rÝ   r‰  r  Ú
out_backupr¹   Úout_untiledÚ	n_batchesÚn_block_rowsÚn_block_colsrk   rl   rs   r‡   r™  rš  r›  rœ  r�  r‘  s&   &&&dd$$$$$d                     @@@@@@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ÈÓ&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Ø—‘˜”Øˆ
àÐØÐØÒØ ÐØ—_‘_ RÓ(×/Ò/ð 
Ø,ð
Ø./ð
Ø12ó
‰
ð  —_’_ÐLÐ&EÐLÀqÐLÈ!ÓL×SÒSð 
Ø,ð
Ø./ð
Ø12ó
ˆ
ð ÒØ!ÐØ—o‘o bÓ)×0Ò0ð 
Ø,ð
Ø./ð
Ø12ó
‰ð "×&Ò&ÐNÐ(GÐNÈÐNÈAÓN×UÒUð 
Ø,ð
Ø./ð
Ø12ó
ˆð ×ÑÓ˜2Õ !Ô#ÜØ5°j×6GÑ6GÓ6IÈ"Õ6MÐ5NÐOó
ð 	
ð ×ÑÓ˜BÕ 1Ô$ÜØ6°{×7IÑ7IÓ7KÈBÕ7OÐ6PÐQó
ð 	
ð ‚|Ü˜˜SŸX™X›Z¨)°A­,Õ6¸À1½ÕEÈÈQÍÕOÕOÐQRÓSˆÜ#ØØØØ�a�LØ�a�LØØØØ—+‘+Ø—i‘iô
ˆð €Jô 	�s 5°kÓ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�L©¨2¨Õ.Ó/°'¸QÄÀXÈbÅ\ÓARÕ=RÕ2SÕS‰àˆð 	�Ø�mØ�_Øˆ{Øˆ}Ø�KØ�[Øˆ[ð	€Oð �„zÜÐ2Ó3Ð3÷
õ 
ô& �&˜/¨9°kÔBà
‡|�|ƒ~˜×,Ñ,Ó.Ô.ð 	×Ñ˜×)Ñ)¨*×*:Ñ*:Ó;Ô<àÐr   c                óØ   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  /# )r   ÚIS_BETA_ZEROr’  r”  ÚTILE_Kr•  re  ©r—  Ú	constexpr)rƒ  s   "r   r„  r„  \  sj   € ÷ }4ñ }4ô —l‘lð}4ô —|‘|ð	}4ô
 —|‘|ð}4ô —‘ð}4ô> —<‘<ð?}4ô@ —L‘LñA}4r   c            
      ó:  € \         P                  ! ^R7      p \         P                  ! ^ R7      p!VVV ,          ,           VV!,          ,           p"\         P                  ! V"4      p#\         P                  ! V"V,           4      p$V$V#,
          p%V%^ 8X  d   R# \         P                  ! ^ V4      p&\         P                  ! ^ V4      p'VVV ,          ,           V	V#,          ,           V
V&R,          ,          ,           VV'R,          ,          ,           p(VVV ,          ,           VV#,          ,           p)VVV ,          ,           VV!,          ,           VV&R,          ,          ,           p*VVV ,          ,           VV'R,          ,          ,           p+\         P                  ! ^ V4      p,\	        V%4       EFv  p-\         P
                  ! W43VR7      p.\         P                  ! V)4      p/\	        ^ WV4       Fª  p0V0V,,           p1V1V8  p2\         P                  ! V*VV1R,          ,          ,           V2R,          RR7      p3\         P                  ! V+VV/,          ,           VV1R,          ,          ,           V2R,          RR7      p4V.\         P                  ! V3V4VVR7      ,          p.K¬  	  V'       d   V.V ,          p.M+V V.,          V\         P                  ! V(4      ,          ,           p.\         P                  ! V(V.P                  VP                  P                  4      4       V(V	,          p(V)V,          p)EKy  	  R# )	r5   ©ÚaxisN©r,   ç        ©Úmaskr<  ©re  r  ©ºNNNN©Nr¹  )r—  Ú
program_idÚloadrO  rh   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•  re  Ú	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î  [  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¸Õ(AÕAõBð &Ð(8¸Õ(AÕAõ	Bð 	ð Ø&¨Õ2õ3à  :Õ-õ.ð 	ð Ø )Õ+õ,à# mÕ3õ4ð $Ð&6°wÕ&?Õ?õ@ð 	ð Ø )Õ+õ,à#Ð&6°wÕ&?Õ?õ@ð 	ô Ÿ	š	 ! VÓ,ˆÜ�w—ˆAÜŸš -Ð!?ÀyÔQˆIô ŸšÐ 1Ó2ˆIä  1Ö-�Ø" ]Õ2�	Ø" Q™�äŸWšWØ#Ð&;¸iÈÕ>PÕ&PÕPà �Øô	�
ô  ŸWšWØ#Ø+¨iÕ7õ8à+¨i¸Õ.@Õ@õAð   �Øô�
ð œRŸVšVØ 
°zÈYôõ ’	ñ) .÷0 Ø˜UÕ"‘	à! IÕ-°´r·w²wÐ?PÓ7QÕ0QÕQ�	ô �HŠHÐ&¨	¯©°Z×5EÑ5E×5PÑ5PÓ(QÔRð Ð!2Õ2ÐØÐ!3Õ3ÓóQ  r   c          
      ó¸   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  /# )r   r’  r”  r•  re  r  r®  )rƒ  s   "r   r„  r„  Ü  sj   € ÷ APñ APôN —|‘|ðOAPôP —|‘|ðQAPôR —<‘<ðSAPôT —L‘LðUAPôV Ÿ™ñWAPr   c           	      ó"  € \         P                  ! ^R7      p\         P                  ! ^ R7      p\         P                  ! ^R7      p\         P                  ! ^ R7      p\         P                  ! ^R7      p \         P                  ! VVVV V4      w  ppVVV,          ,           VV,          ,           p!\         P                  ! V!4      p"\         P                  ! V!V,           4      p#V#V",
          p$V$^ 8X  d   R# \         P
                  ! ^ V4      p%\         P
                  ! ^ V4      p&V VV,          ,           VV",          ,           VV%R,          ,          ,           VV&R,          ,          ,           p'VVV,          ,           VV,          ,           VV&R,          ,          ,           VV%R,          ,          ,           p(VVV,          ,           VV,          ,           VV,          ,           VV%R,          ,          ,           VV%R,          ,          ,           p)VV	V,          ,           V
V",          ,           p*\         P                  ! VV3VR7      p+\        V$4       F†  p,\         P                  ! V'4      p-\         P                  ! V*4      p.\         P                  ! V(VV.,          ,           4      p/V+\         P                  ! V-V/VVR7      ,          p+V'V,          p'V*V
,          p*Kˆ  	  \         P                  ! V)V+P                  VP                  P                  4      4       R# )r   r±  Nr³  r·  r¸  rº  )r—  r»  Únum_programsÚ	swizzle2dr¼  rO  r½   rh   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•  re  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¸Õ(AÕAõBð &Ð(8¸Õ(AÕAõ	Bð 	ð Ø  9Õ,õ-à$ }Õ4õ5ð %Ð'7¸Õ'@Õ@õAð %Ð'7¸Õ'@Õ@õ	Að 	ð Ø! IÕ-õ.à%¨Õ5õ6ð &¨Õ5õ6ð &Ð(8¸Õ(AÕAõ	Bð
 &Ð(8¸Õ(AÕAõBð 	ð Ø&¨Õ2õ3à  :Õ-õ.ð 	ô Ÿ8š8 ]°MÐ$BÈ)ÔTÐÜ�w–ˆAÜŸ7š7Ð#4Ó5ˆLô ŸGšGÐ$5Ó6ˆMÜŸ'š'Ø Ð#9¸MÕ#IÕIóˆKð
 ¤§¢Ø˜k°jÈIô!õ Ðð
 Ð!2Õ2ÐØÐ!3Õ3Òñ#  ô( 	�Š�Ð.×1Ñ1°*×2BÑ2B×2MÑ2MÓNÖOr   c           
      óò  a aaaaaaa€ VP                  ^ 4      pVP                  R4      ^,
          pWÍ3pVeC   \        VR,          RRR1,          4      R^\        VR,          4      ,
          ,          ,           pMRpVRVR	VRV	R
V
R/pVP                  \        P
                  \        P                  39   d   \        P                  oRoM\        P                  oRoVVV VVVVV3R lp\        VVWï4       R# )r   Nr  TFc                 ó`   <€ \         V ,          ! SSS.SOSNS	N\        V!  O5R SRSR^R^/  R# )r•  re  rü   rý   N)rî  re   )
r{   r‰   r•  re  r  r  r>   Úis_beta_zerorÁ  Útile_ks
   &*€€€€€€€€r   r‡   Ú)_run_sampled_addmm_kernel.<locals>.kernelƒ  sn   ø€ Ü! $Ö'ØØØðð ð	ð
 ðð ðô & ~Ñ6òð $ðð &ðð ðð ôr   rv   r
   )r   N)r   rv   )r   r   )r¢   r0   rA   r,   r   r-   r.   r—  r˜  r¥  rŠ   )r  r  r	  r>   rÁ  r
  r›   r˜   rš   Úmat1Úmat2rz  r¨  r©  rk   rl   rs   r‡   r•  re  s   ffffff&&&&&&      @@r   Ú_run_sampled_addmm_kernelr  _  sÙ   ÿ€ ð —K‘K “Nˆ	Ø#×(Ñ(¨Ó,¨qÕ0ˆàÐ-ˆ	ØÒÜ ¨¥©T¨r¨TÕ 2Ó3°gÀÄSÈÐRTÍÓEVÕAVÕ6WÕW‰KàˆKà�IØ˜'Ø˜Ø�'Ø�)ð
ˆð �<‰<œEŸJ™J¬¯©Ð7Ô7ÜŸ
™
ˆIØ‰JäŸ
™
ˆIØˆJ÷	ô 	ô 	�f˜o¨yÖFr   g      ð?c                ó"  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R,          R\        R\        \        R,          \        R,          \        R,          3,          R,          /# )r   r}  r  r  rt  Nry  rz  )r   r€  r�  r0   r‚  )rƒ  s   "r   r„  r„  ”  s~   € ÷ Uñ UÜ�|‰|ðUä�l‰lðUô �l‰lðUô �\‰\˜DÕ ðUô ðUô œ˜d�
¤C¨$¥J´°dµ
Ð:Õ;¸dÕBñUr   c                óž  € R p\        W€4       \        W€W4      p	V'       Eg–   \        W�V P                  4       \        W‚V P                  4       VR8w  d1   V P                  \
        P                  J d   \        RV RV R24       V P                  \
        P                  Jd.   \        W�V P                  4       \        W‚V P                  4       M\        W�VP                  4       \        W�V4       VeÂ   \        W…4       \        W…VP                  4       \        W…V P                  4       \        VP                  V	P                  8H  ;'       d"    VP                  4       V P                  4       8H  V RV	P                   RV	P                  4        RVP                   R	VP                  4        2	4       Vf   V	P                  VP                  R
R7      pMVP                  V	4       VP                  4       ^ 8X  g   VP                  4       ^ 8X  d   V# VP                  4       P                  RR p
VP!                  R4      pVR8X  g   V^ 8X  d"   VP                  4       P#                  V4       V# Tp\%        WQV4      w  rÞrñp\'        W^ ,          V34      p\'        W+V
^,          34      p\)        V
!  p\+        VVVR8H  V
VVVVVVVV4       VP                  4       P-                  4       RR VP-                  4       RR 8w  dG   VP                  4       P                  VP/                  VP                  4       P                  4      4       V# )Úsampled_addmmr´  Fz(): having beta == z3 not equal to 0.0 with boolean mask is not allowed.Nz!(): Expects `out` to be of shape z and with nnz equal to z but got out.shape = z and out.nnz = T)Úcopyr!   rv   r–   )r   r¥   r   r   r,   r   r�  r   r2   r*   r#   ro  r¿  rq  ri  r›   r¢   r¢  rŸ   r­   r�   r  rF   rs  )r}  r  r  r  r  rt  ry  rz  r   Úinput_broadcastedr>   rÁ  r¦  r˜   rš   r›   r
  s   &&&$$$$$         r   r  r  ”  s×  € ð !ˆä˜Ô'Ü4°VÀDÓOÐçˆ{Ü˜ u§|¡|Ô4Ü˜ u§|¡|Ô4Ø�sŒ{˜uŸ{™{¬e¯j©jÓ8ÜØØ�hÐ1°$°Ð7jÐkôð �{‰{¤%§*¡*Ó,Ü˜F¨%¯+©+Ô6Ü˜F¨%¯+©+Õ6ä˜F¨$¯*©*Ô5Ü& v°TÔ:ØŠÜ  Ô-Ü˜V¨$¯+©+Ô6Ü˜F¨¯©Ô5ÜØ—I‘IÐ!2×!8Ñ!8Ñ8×WÐ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ÐRVÓ8WÑ5ˆ 6°ä  °!­°aÐ'8Ó9ˆÜ  ¨9°Q­<Ð'8Ó9ˆÜ�i‘ˆä!ØØØ�C‰KØØØØØØØØØô	
ð$ ×ÑÓ×%Ñ%Ó'¨¨Ð,°·±³ÀÀÐ0DÔDØ×ÑÓ×%Ñ% f§n¡n°Z×5FÑ5FÓ5H×5NÑ5NÓ&OÔPØÐr   c                ó  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R,          R\        R\        \        R,          \        R,          \        R,          3,          R,          R\
        R,          /# )r   rœ   r~  rt  Nry  rz  râ   r  )rƒ  s   "r   r„  r„  ë  s~   € ÷ 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r   c          	      óx  € R pV P                   R	R w  rxV'       g°   \        W`4       \        W`VP                  4       \	        W`VP
                  \        P                  34       \        W`V4       VP                  R
4      p	V P                  4       P                   R	R w  r«\        WjV34       \        V	^,          '       * V RV	 R24       MVP                   R	R w  rÉ\        W`V4      pVe   V'       gw   W×V	3,           p\        VP                   V8H  RV RVP                    R24       \        VP                  4       ;'       g!    VP                  R	R
4      P                  4       R4       Vf   VP!                  W×V	3,           4      pV P#                  4       ^ 8X  d   VP%                  4       # \'        W V^^ VR7      # )Úbsr_dense_mmNz(): 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  rt  r!   rv   )r#   r   r   r   r2   r,   r   r‡  r*   r¢   r›   rC   r   rU   Úis_contiguousr¨   r¡  ro  rÁ   r
  )rœ   r~  rt  ry  rz  râ   r   rª   Ú_klr«   Ú	row_blockrè  Ú_krr‰  Úexpected_out_shapes   &&$$$$         r   r  r  ë  s–  € ð  ˆØ—‘˜2˜3�‰ˆßÜ˜VÔ)Ü˜ e§l¡lÔ3Ü˜ U§[¡[´5·:±:°-Ô@Ü& v°EÔ:à—
‘
˜2“ˆAØ#&§:¡:£<×#5Ñ#5°b°cÐ#:Ñ ˆIÜ˜F°	Ð$:Ô;ÜØ˜•F”
Ø�(Ð0°°Ð3MÐNõð
 —[‘[  Ð%‰FˆCä*>¸vÈEÓ*RÐ'àŠ?§;Ø!@ÀqÀ6Õ!IÐÜØ—	‘	Ð/Ñ/ðØ.Ð/¨z¸#¿)¹)¸ÀAðGôô
 Ø×!Ñ!Ó#×LÐL s§}¡}°R¸Ó'<×'JÑ'JÓ'Lð"ôð Š;Ø—/‘/Ð"AÈÀFÕ"JÓKˆCð �8‰8‹:˜Œ?Ø—9‘9“;Ðô ˜s¨°a¸aÀSÔIÐIr   c                óX   € V ^8„  d   QhR\         P                  R\         P                  /# )r   ÚMAX_ROW_NNZÚTILEr®  )rƒ  s   "r   r„  r„  #  s,   € ÷ Vñ Vô —\‘\ðVô �l‰lñVr   c                 óê  € \         P                  ! ^R7      p\         P                  ! ^R7      p\         P                  ! ^ R7      pV W,          ,           W-,          ,           p\         P                  ! V4      p\         P                  ! Wâ,           4      pVV,
          pV^ 8X  d   R# \         P                  ! ^ V
4      pVVV,          8  pVWK,          ,           W\,          ,           Wø,          ,           p\         P                  ! VV,           V\	        R4      ) R7      P                  \         P                  4      p\         P                  ! V^ R7      p\        W©V
4       F’  pVV
,          pVVV,          8  p\         P                  ! VV,           V\	        R4      ) R7      P                  \         P                  4      p\         P                  ! V^ R7      p\         P                  ! VV8„  VV4      pK”  	  \         P                  ! VV,
          4      p\         P                  ! V^ R7      p\        W©V
4       F›  pVV
,          pVVV,          8  p\         P                  ! VV,           V\	        R4      ) R7      P                  \         P                  4      p\         P                  ! VV,
          4      pV\         P                  ! V^ R7      ,          pK�  	  \         P                  ! VV,           VV,          P                  VP                  P                  4      VR7       \        W©V
4       FÆ  pVV
,          pVVV,          8  p\         P                  ! VV,           V\	        R4      ) R7      P                  \         P                  4      p\         P                  ! VV,
          4      p\         P                  ! VV,           VV,          P                  VP                  P                  4      VR7       KÈ  	  R# )r   r±  NÚinfrµ  )r¶  )r—  r»  r¼  rO  r/   r¿  r˜  r�   rh   Ú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’M qÔ)ˆ	Ü!Ÿ}š}°!Ô4ÐÜŸš¨1Ô-ˆð Ø'Õ3õ4à!Õ1õ2ð 	 ô
 —W’WÐ4Ó5ˆ
ÜŸ'š'Ð"9Õ"OÓPˆð " JÕ.ˆØ�aŒ<Ùä—Y’Y˜q $Ó'ˆ
Ø˜G iÕ/Ñ/ˆð Ø!Õ-õ.à%Õ<õ=ð Õ$õ%ð 	ô —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ñ 0ô �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ñ 0ô 	�ŠØ  :Õ-Ø�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×!<Ñ!<Ó=Ø÷ó 0r   c                 ób  aaa€ R p\        W 4       \        W V P                  4       V P                  4       ^ 8X  g   V P	                  4       ^ 8X  d   V P                  4       # V P                  RR w  r4V P                  4       pV P                  4       P                  RR w  ooSf   \        P                  ! V4      oM\        P                  ! S4      oV P                  4       P                  ^ 4      P                  ^ R4      pV P                  4       P                  RR4      P                  4       '       d    V P                  4       P                  4       pMV P                  4       pVP                  RR4      P                  4       P                  ^ 4      P                  ^ R4      P!                  RSVS,          4      pVP                  ^ ,          SVS,          3pRp	VRRR13,          R	VR
/p
VVV3R lp\#        WºW‰4       VP!                  RSVS4      P                  RR4      P                   ! V P                  4       P                  !  p\$        P&                  ! V P                  4       P                  4       V P)                  4       P                  4       VV P                  V P*                  R7      # )Úbsr_softmaxN.c                 ób   <€ \         V ,          ! . \        V!  OSNSNSN\        RS4      N5!   R# )r   Ni   )r-  re   rE   )r{   r‰   rè  Úmax_row_nnzr  s   &*€€€r   r‡   Úbsr_softmax.<locals>.kernel¨  sJ   ø€ Ü Ö%ð Ü% ~Ñ6ðàðð ðð ð	ô �E˜;Ó'ör   r¡   r!   r–   r   rv   rŸ  rž  )r   r2   r,   ro  ri  Úcloner#   r›   ÚtritonÚnext_power_of_2r˜   r™   r‘   r¨   r  rG   rs  rŠ   r   r£   rš   r   )r}  r1  r   rª   r«   Únnzr˜   r›   rk   rl   rs   r‡   rè  r  s   &f          @@r   r/  r/  {  s3  ú€ Øˆä˜Ô'Ü�F 5§;¡;Ô/à�:‰:‹<˜1Ô §¡£°Ô 2Ø—;‘;“=Ð à�{‰{˜2˜3Ð‰ˆØ�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¨yÔFð �N‰N˜2˜y¨#¨yÓ9ß‰Y�r˜2Óß‰Wðà—l‘l“n×*Ñ*ñ,ð 	ô ×-Ò-Ø×ÑÓ ×&Ñ&Ó(Ø×ÑÓ×%Ñ%Ó'ØØ—‘Ø—<‘<ô
ð 	
r   c                óØ   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R,          R\        R\        R\        R,          /# )	r   Úqueryr  ÚvalueÚ	attn_maskNÚ	dropout_pÚ	is_causalÚscale)r   r€  r/   r�  )rƒ  s   "r   r„  r„  Â  sg   € ÷ 0ñ 0Ü�|‰|ð0ä�\‰\ð0ô �|‰|ð0ô —<‘< $Õ&ð	0ô
 ð0ô ð0ô �t�|ñ0r   c           	      ó,  € R p\        V'       * V R24       \        VRJV R24       Vf   \        R4      h\        VP                  \        P                  8H  V R\        P                   RVP                   R24       \        WqV P                  4       \        WrV P                  4       \        WsV P                  4       \        WqV P                  4       \        WrV P                  4       VP                  \        P                  Jd   \        WsV P                  4       \        W0VP                  RR4      RR	R
7      pVf   V P                  R4      ^ 8X  g   VR8X  d   \        R	V RV R24       Vf-   ^\        P                  ! V P                  R4      4      ,          MTp	VP                  4       P!                  V	4       \#        V4      p\        P$                  P&                  P)                  VP                  4       VRR7       \+        W‚4      pV# )Ú_scaled_dot_product_attentionz'(): is_causal == True is not supported.Nz'(): 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´  F)r  ry  z(): current value of scale == z results in division by zero.T)rØ   Úinplacer!   rv   )r   rB   r   r   r   r   r   r2   r,   r�  r  r¨   r¢   ÚmathÚsqrtr›   r¢  r/  ÚnnÚ
functionalÚdropoutr  )
r8  r  r9  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 %§,¡,Ô/Ü�V E§L¡LÔ1Ü�V¨¯©Ô5ä�F §¡Ô-Ü�F 5§;¡;Ô/Ø�?‰?¤%§*¡*Ó,Ü˜¨5¯;©;Ô7ô Ø˜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   c                óø   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  /# )r   rã   rÜ   rÝ   rœ  rù   rú   re  r®  )rƒ  s   "r   r„  r„  õ  s|   € ÷ DIñ DIÜ�<‰<ðDIä�<‰<ðDIô �<‰<ðDIô* —|‘|ð+DIô, —‘ð-DIô. —‘ð/DIô0 —L‘Lñ1DIr   c           	      ól  € V V,          p\         P                  ! ^ R7      p\         P                  ! ^R7      pVV,          pVV,          pVV,          \         P                  ! ^ V4      ,           pVV,          \         P                  ! ^ V4      ,           p\         P                  ! ^ V4      pVVR,          V,          VR,          V,          ,           ,           p VVR,          V	,          VR,          V
,          ,           ,           p!\         P                  ! VVV,          ,           4      p"\         P                  ! VV^,           V,          ,           4      p#V"V#8X  d   R# \         P                  ! VV3VR7      p$\        V"V#4       F»  p%\         P                  ! VV%V,          ,           4      p&\         P                  ! VV%V,          ,           V,           4      p'\         P                  ! V V&V,          ,           4      p(\         P                  ! V!V'V,          ,           4      p)V$\         P                  ! V(V)VVR7      ,          p$K½  	  VVV,          ,           VR,          V,          VR,          V,          ,           ,           p*\         P                  ! V*V$P                  VP                  P                  4      4       R# ©r   r±  Nr³  )r  re  r¸  rº  )r—  r»  rO  r¼  r½   rh   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ú   re  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_kernelrg  ô  s  € ð6 �&�[ˆä—’ 1Ô%ˆä�mŠm Ô#ˆØ�r•	ˆØ�b•ˆà�V�^œbŸiši¨¨6Ó2Õ2ˆØ�V�^œbŸiši¨¨6Ó2Õ2ˆÜ�YŠY�q˜!‹_ˆàØˆw�K˜/Õ)¨B¨w­K¸/Õ,IÕIõ
ˆð Øˆw�K˜/Õ)¨B¨w­K¸/Õ,IÕIõ
ˆô �WŠW�^ eÐ.?Õ&?Õ?Ó@ˆÜ�WŠW�^ u¨q¥yÐ4EÕ&EÕEÓFˆà�Œ8Ùä—H’H˜f fÐ-°]ÔCˆ	ä�r˜2–ˆAÜ—’˜  [¥Õ0Ó1ˆAÜ—’˜  [¥Õ0°;Õ>Ó?ˆAÜ—’˜  OÕ 3Õ3Ó4ˆAÜ—’˜  OÕ 3Õ3Ó4ˆAØœŸš  1°È*ÔUÕUŠIñ ð ØÐ+Õ+õ,ð �7•Ð3Õ3Ø�W•+Ð 5Õ5õ6õð 	ô 	�Š�˜	Ÿ™Ð%5×%;Ñ%;×%FÑ%FÓGÖHr   c          
      ó¸   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  /# )r   rÆ   rÇ   rb  Ú
pq_indicesr²   )r   r€  )rƒ  s   "r   r„  r„  ;  sN   € ÷ 4
ñ 4
Ü—‘ð4
ä—‘ð4
ô —L‘Lð4
ô —L‘Lð	4
ô
 —l‘lñ4
r   c                 óà  aaa€ V P                   w  popVP                   w  rxo\        \        ^S^,          4      \        ^S^,          4      ^^R7      p	VVV3R lp
\        P                  \
        P                  \        P                  \
        P                  \        P                  \
        P                  \        P                  \
        P                  /VP                  ,          pRV	9  d$   V	P                  V\
        P                  8H  R7       \        V
,          ! SVSV V P                  ^ 4      V P                  ^4      V P                  ^4      VVP                  ^ 4      VP                  ^4      VP                  ^4      VVP                  ^ 4      VP                  ^4      VP                  ^4      SSP                  ^ 4      VVP                  ^ 4      VP                  ^4      3RV/V	B  R# )rþ   )rù   rú   rü   rý   c                 ó¼   <€ SP                   ^ ,          ^,
          \        P                  ! SV R,          4      \        P                  ! SV R,          4      ,          ^3# ©r   rù   rú   ©r#   r4  Úcdiv)ÚMETArã   rÝ   rb  s   &€€€r   r{   Ú_scatter_mm2.<locals>.gridI  sI   ø€ à× Ñ  Õ# aÕ'Ü—’˜A˜t H�~Ó.´·²¸QÀÀXÅÓ1OÕOØðð r   re  rf  rœ  N)r#   r  r�   r   r  r—  r˜  r.   r¥  r,   r  rg  rF   )rÆ   rÇ   rb  ri  r²   rÊ   rÜ   rÏ   ræ  râ   r{   rœ  rã   rÝ   s   &&f&&       @@r   r¾   r¾   ;  s‘  ú€ ð —<‘<‰ˆˆAˆqØ—<‘<‰ˆˆ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   c                ó  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  /# )	r   rÌ   rg  rœ  r¹   rù   rú   rû   re  r®  )rƒ  s   "r   r„  r„  r  s�   € ÷ oIñ 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r   c           	      óÖ  € VV,          pVV,          pVV,          p\         P                  ! ^ R7      p\         P                  ! ^R7      pVV ,          p VV ,          p!VV,          p"VV",          p#V#V,          p$\        VV$,
          V4      p%V$VV%,          ,           p&VV",          V%,          p'V&V,          \         P                  ! ^ V4      ,           p(V'V,          \         P                  ! ^ V4      ,           p)\         P                  ! ^ V4      p*VV(R,          V,          V*R,          V,          ,           ,           p+VV V	,          ,           V*R,          V
,          V)R,          V,          ,           ,           p,\         P                  ! VV!,           4      p-V'       dŠ   V-V,          V,          p.V-V,          V,          p/\         P                  ! VV.,           4      p0\         P                  ! VV.,           ^,           4      p1V/V1,          VV/,
          V0,          ,           p2V1V0,
          p3MJ\         P                  ! VV!,           4      p2\         P                  ! VV!,           ^,           4      p4V4V2,
          p3VV2,           p5\         P
                  ! VV3VR7      p6V'       d    V+X0V,          ,          p+\        V34       F  p7\         P                  ! V54      p8\         P                  ! V,V8,           4      p9\         P                  ! V+4      p:V6\         P                  ! V:V9VVR7      ,          p6V+V,          p+V5^,          p5K�  	  M»VV2,           p;\        V34       F£  p7\         P                  ! V54      p8\         P                  ! V,V8,           4      p9\         P                  ! V;4      p<\         P                  ! V+V<V,          ,           4      p:V;^,          p;V5^,          p5V6\         P                  ! V:V9VVR7      ,          p6K¥  	  VV-,           V V,          ,           V(R,          V,          V)R,          V,          ,           ,           p=\         P                  ! V=V6P                  VP                  P                  4      4       R# rJ  )r—  r»  rE   rO  r¼  r½   rh   r½  r¾  r¿  r,   rÀ  )>rX  rË   rÌ   rÝ   rK  rL  rM  rN  rO  Úothers_stride_BrQ  rR  rS  Úaccumulators_stride_BrU  rV  Úc_indices_ptrÚr_offsets_ptrÚp_offsets_ptrÚq_offsets_ptrrg  rœ  r¹   rù   rú   rû   re  rÑ   ÚBLOCKS_MÚBLOCKS_NÚpid_t_r]  Úpid_br\  Únum_pid_in_groupÚgroup_idÚfirst_pid_mÚgroup_size_mr^  r_  r`  ra  rb  rc  rd  rÕ   rª   r«   rç   rè   rÖ   r6  r×   Úq_ptrrç  ræ  rÙ   rÛ   re  Úp_ptrrØ   rf  s>   &&&&&&&&&&&&&&&&&&&&&&&&&&&                                   r   Ú_scatter_mm6_kernelrƒ  q  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˜"ÓˆØØˆw�K˜/Õ)¨B¨w­K¸/Õ,IÕIõ
ˆð Ø�oÕ%õ&à�'�{˜_Õ,¨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ôõ ’	ñ  ð ØõàÐ+Õ+õ,ð �7•Ð3Õ3Ø�W•+Ð 5Õ5õ6õ	ð 	ô 	�Š�˜	Ÿ™Ð%5×%;Ñ%;×%FÑ%FÓGÖHr   c                ó  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\        R\         P                  R	\        /	# )
r   rÆ   rÇ   rÞ   rß   rà   rá   râ   r²   Úforce_contiguous)r   r€  r  r�  )rƒ  s   "r   r„  r„  ã  s‰   € ÷ f.ñ f.Ü—‘ðf.ä—‘ðf.ô —<‘<ðf.ô —<‘<ð	f.ô
 —<‘<ðf.ô —<‘<ðf.ô ðf.ô —l‘lðf.ô ñf.r   c	                 ór  aaaa€ VR ,          p	V P                   w  p
opVP                   w  orÍVP                   w  rïpVV8w  d   \        RV RV R24      hWÙ,          oVS8w  d   \        RV RS R24      hVVVV3R lp\        P                  \        P
                  \        P                  \        P
                  \        P
                  \        P                  \        P                  \        P                  /VP                  ,          pRV9  d$   VP                  V\        P
                  8H  R7       VP                  ^ 4      ^8w  d   \        R	VP                  ^ 4       24      hSP                  ^ 4      ^8w  d   \        R
SP                  ^ 4       24      hVP                  ^ 4      ^8w  d   \        RVP                  ^ 4       24      hVP                  ^ 4      ^8w  d   \        RVP                  ^ 4       24      hV'       dK   V P                  4       p VP                  4       pVP                  4       '       g   VP                  4       pMTpMTp\        V,          ! SSVVV V P                  ^ 4      V P                  ^4      V P                  ^4      VVP                  ^ 4      VP                  ^4      VP                  ^4      VVP                  ^ 4      VP                  ^4      VP                  ^4      VSVV3RV/VB  V'       d,   VP                  4       '       g   VP                  V4       R# R# R# )r¹   rº   r»   r¶   zaccumulators B (z) != others B (c                 óº   <€ SP                   ^ ,          S,          \        P                  ! SV R,          4      \        P                  ! SV R,          4      ,          3# rl  rm  )ro  rÛ   rË   rÑ   rß   s   &€€€€r   r{   Ú_scatter_mm6.<locals>.gridø  sD   ø€ à—‘ Õ" QÕ&Ü—’˜B  X¥Ó/´&·+²+¸bÀ$ÀxÅ.Ó2QÕQðð r   re  rf  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œ  N)r#   rB   r   r  r—  r˜  r.   r¥  r,   r  rF   rG   r  rƒ  rq  )rÆ   rÇ   rÞ   rß   rà   rá   râ   r²   r…  r¹   rÊ   rÌ   Ú_KrÝ   ÚB_Ú_Mrä   r{   rœ  Úaccumulators_rÛ   rË   rÑ   s   &&&f&&&&&           @@@r   rÀ   rÀ   ã  s   û€ ð �y•/ˆØ—\‘\‰
ˆˆB�Ø—<‘<‰ˆˆ2Ø!×'Ñ'‰
ˆ�Ø�Œ7Ü Ð#3°B°4°ÀqÀcÈÐ!KÓLÐLØ�\ˆØ�Œ7Ü Ð#3°B°4°ÀqÀcÈÐ!KÓLÐL÷	ð 	ô �M‰Mœ2Ÿ:™:Ü�N‰NœBŸJ™JÜ�M‰Mœ2Ÿ:™:Ü�M‰Mœ2Ÿ:™:ð	
ð
 ×
Ñ
õˆð ˜tÔ#Ø�K‰K =´B·J±JÑ#>ˆKÔ?à×Ñ˜AÓ !Ô#Ü Ø5°i×6FÑ6FÀqÓ6IÐ5JÐKóð ð ×Ñ˜AÓ !Ô#Ü Ø5°i×6FÑ6FÀqÓ6IÐ5JÐKóð ð ×Ñ˜AÓ !Ô#Ü Ø5°i×6FÑ6FÀqÓ6IÐ5JÐKóð ð ×Ñ˜AÓ !Ô#Ü Ø5°i×6FÑ6FÀqÓ6IÐ5JÐKóð ÷" Ø×&Ñ&Ó(ˆFØ×&Ñ&Ó(ˆFØ×-Ñ-×/Ò/Ø ,× 7Ñ 7Ó 9‘à ,‘à(ˆMä˜DÖ!ØØØØØØ�M‰M˜!ÓØ�M‰M˜!ÓØ�M‰M˜!ÓØØ�M‰M˜!ÓØ�M‰M˜!ÓØ�M‰M˜!ÓØØ× Ñ  Ó#Ø× Ñ  Ó#Ø× Ñ  Ó#ØØØØñ)	
ð* (ð+	
ð, ò-	
÷2  L×$>Ñ$>×$@Ò$@Ø×Ñ˜}Ö-ñ %AÑr   c                 ó  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R	\         P                  R
\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  /# )r   Úleft_alpha_tiled_col_strideÚleft_alpha_col_block_strideÚright_alpha_tiled_row_strideÚright_alpha_row_block_strider�  rŽ  r�  r�  r‘  r’  r”  r“  r•  re  r  r¹   r®  )rƒ  s   "r   r„  r„  L	  s0  € ÷ UPñ UPôL &(§\¡\ðMUPôP &(§\¡\ðQUPôZ ')§l¡lð[UPô^ ')§l¡lð_UPôx —\‘\ðyUPôz Ÿ™ð{UPô| —l‘lð}UPô~ Ÿ<™<ðUPô@ ŸL™LðAUPôB —|‘|ðCUPôD —|‘|ðEUPôF Ÿ™ðGUPôH —<‘<ðIUPôJ —L‘LðKUPôL Ÿ™ðMUPôN —‘ñOUPr   c7           	      óD
  € V^ 8w  d   \        RV 24      hV^ 8w  d   \        RV 24      hV^ 8w  d   \        RV 24      hV!^ 8w  d   \        RV! 24      h\        P                  ! ^R7      p7\        P                  ! ^ R7      p8\        P                  ! ^R7      p9\        P                  ! ^ R7      p:\        P                  ! ^R7      p;\        P                  ! V8V9V:V;V54      w  p8p9VVV7,          ,           VV8,          ,           p<\        P
                  ! V<4      p=\        P
                  ! V<V,           4      p>V>V=,
          p?\        P                  ! ^ V04      p@\        P                  ! ^ V24      pA\        P                  ! ^ V14      pBV VV7,          ,           VV=,          ,           VV@R	,          ,          ,           VVAR
,          ,          ,           pCVVV7,          ,           VV9,          ,           VVAR	,          ,          ,           VVBR
,          ,          ,           pDV#V$V7,          ,           V%V8,          ,           V&V9,          ,           V'V@R	,          ,          ,           V(VBR
,          ,          ,           pEVV	V7,          ,           V
V=,          ,           pF\        P                  ! V0V13V3R7      pG\        V?4       F†  pH\        P
                  ! XC4      pI\        P
                  ! XF4      pJ\        P
                  ! XDVVJ,          ,           4      pKXG\        P                  ! VIVKV4V3R7      ,          pGVCV,          pCVFV
,          pFKˆ  	  V-'       g
   XGV*,          pGV.'       gt   VVV7,          ,           VV8,          ,           VV9,          ,           VX@R	,          ,          ,           VXBR
,          ,          ,           pLXG\        P
                  ! VL4      ,          pGV/'       gt   VVV7,          ,           VV8,          ,           V V9,          ,           V!X@R	,          ,          ,           V"XBR
,          ,          ,           pMXG\        P
                  ! VM4      ,          pGV,'       d¡   VVV7,          ,           VV8,          ,           VV9,          ,           VX@R	,          ,          ,           VXBR
,          ,          ,           pNV+'       d   XG\        P
                  ! XN4      ,          pGM$XGV)\        P
                  ! XN4      ,          ,          pG\        P                  ! XEXGP                  V#P                  P                  4      4       R# )r   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±  Nr³  r·  r¸  rº  )rB   r—  r»  rñ  rò  r¼  rO  r½   rh   r½  r¾  r¿  r,   rÀ  )OrÂ  rÃ  rÄ  rÅ  rÆ  rÇ  rÈ  rÉ  rÊ  rË  rÌ  Ú	input_ptrÚinput_batch_strideÚinput_tiled_row_strideÚinput_tiled_col_strideÚinput_row_block_strideÚinput_col_block_strideró  rô  rõ  rö  r÷  rø  Úleft_alpha_ptrÚleft_alpha_batch_strideÚleft_alpha_tiled_row_striderŽ  Úleft_alpha_row_block_strider�  Úright_alpha_ptrÚright_alpha_batch_strider�  Úright_alpha_tiled_col_strider‘  Úright_alpha_col_block_striderù  rú  rû  rü  rý  rþ  r  r  r�  rŽ  r�  r�  r‘  r’  r”  r“  r•  re  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Ð=ZÐ[óð ð '¨!Ô+Ü Ø=Ð>YÐ=ZÐ[óð ð (¨1Ô,Ü Ø>Ð?[Ð>\Ð]óð ð (¨1Ô,Ü Ø>Ð?[Ð>\Ð]óð ô —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¸Õ(AÕAõBð &Ð(:¸7Õ(CÕCõ	Dð 	ð Ø  9Õ,õ-à$ }Õ4õ5ð %Ð'9¸'Õ'BÕBõCð %Ð'7¸Õ'@Õ@õ	Að 	ð Ø! IÕ-õ.à%¨Õ5õ6ð &¨Õ5õ6ð &Ð(8¸Õ(AÕAõ	Bð
 &Ð(8¸Õ(AÕAõBð 	ð Ø&¨Õ2õ3à  :Õ-õ.ð 	ô Ÿ8š8 ]°MÐ$BÈ)ÔTÐä�w–ˆAÜŸ7š7Ð#4Ó5ˆLô ŸGšGÐ$5Ó6ˆMÜŸ'š'Ø Ð#9¸MÕ#IÕIóˆKð
 ¤§¢Ø˜k°jÈIô!õ Ðð
 Ð!2Õ2ÐØÐ!3Õ3Òñ#  ÷& Ø Õ%Ðç àØ)¨IÕ5õ6à-°Õ=õ>ð .°Õ=õ>ð .Ð0@ÀÕ0IÕIõ	Jð
 .Ð0@ÀÕ0IÕIõJð ð ¤§¢¨Ó 8Õ8Ðç!àØ*¨YÕ6õ7à.°Õ>õ?ð /°Õ>õ?ð /Ð1AÀ'Õ1JÕJõ	Kð
 /Ð1AÀ'Õ1JÕJõKð ð ¤§¢Ð(8Ó 9Õ9ÐçàØ$ yÕ0õ1à(¨=Õ8õ9ð )¨=Õ8õ9ð )Ð+;¸GÕ+DÕDõ	Eð
 )Ð+;¸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´  FNr;   )<rA  Úosr(  Ú	functoolsr   r   Útorch._dynamo.utilsr   Útorch.utils._tritonr   Ú_triton_ops_metar   r   r‚  Úgetenvr   r   r   r   r*   r2   rC   rH   rU   rZ   rc   re   r}   rŠ   rŸ   r¥   r­   r±   rµ   r  r  r  rc  rl  rv  r†  r
  r4  Útriton.languageÚlanguager—  Újitrî  r  r  r  r  r-  r/  r?  rg  r¾   rƒ  rÀ   r–  r6   r   r   Ú<module>r®     s  ðó Û 	Û Ý ã Ý )Ý *ç 8ñ .1Ø‡I‚IÐ:¸AÓ>ó.Ð *ò
ò
òòò"òò*ò&Uòòòò;ô8&ò2 7òF	ò
/òð"w2¸Tô w2ôtpôf`÷FNñ Nñb Ð=Ô>ñ^
ó ?ð^
ôB!ôHHðV&ð
 
ð&ð ð&ð '+ð&ð (,ð&ð  $ð&ð ð&ð BFð&ð ÷&ðRnð
 
ðnð ðnð '+ðnð (,ðnð  $ðnð ðnð BFðnð ÷nñb ‡<‚<ÛÝ à‡Z�Zô}4ó ð}4ð~ ‡Z�ZôAPó ðAPòF3GðjUð
 ðUð ðUð $(ðUð "ðUð FJ÷Uðn5Jð $(ð	5Jð
 "ð5Jð FJð5Jð !÷5Jðn ‡Z�ZôVó ðVôpE
÷N0ðd ‡Z�ZôDIó ðDIõL4
ðl ‡Z�ZôoIó ðoI÷bf.ðP ‡Z�ZôUPó òUPðp €KØ€LØ€MØ$(Ð!Ø€LØ€LØ $Òr   