Ë
    S²j³^ ã            !       óª  — d dl Z d dlZd dlZd dlmZ d dlZd dlmZ d dlm	Z	 ddl
mZ  e ej                  dd«      «      Zd	„ Zd
„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zdqd„Zd„ Zd„ Zd„ Zd„ Zddœd„Z 	 	 	 	 	 	 drd„Z!	 	 	 	 	 	 	 	 dsd„Z" G d„ d«      Z# ee¬ «      d!„ «       Z$	 dtd"„Z%dud#„Z&dddddd$ddd%œd&ejN                  d'ejN                  d(ejN                  d)ejN                  dz  d*ejN                  dz  d+ejN                  dz  d,e(d-e)edz  edz  edz  f   dz  d.e*dz  fd/„Z+dddddd$ddd%œd&ejN                  d'ejN                  d(ejN                  d)ejN                  dz  d*ejN                  dz  d+ejN                  dz  d,e(d-e)edz  edz  edz  f   dz  d.e*dz  fd0„Z, e	«       �rd dl-Z-d dl.m/Z0 e-jb                  d1e0jd                  d2e0jd                  d3e0jd                  d4e0jd                  d5e0jd                  d6e0jd                  fd7„«       Z3e-jb                  d2e0jd                  d3e0jd                  d5e0jd                  d6e0jd                  d8e0jd                  f
d9„«       Z4d:„ Z5d;d;dd$dd<œd&ejN                  d=ejN                  d>ejN                  d+ejN                  dz  d,e(d-e)edz  edz  edz  f   dz  fd?„Z6dd$ddd@œd'ejN                  d(ejN                  d+ejN                  dz  d,e(d-e)edz  edz  edz  f   dz  d.e*dz  fdA„Z7e-jb                  dBe0jd                  dCe0jd                  fdD„«       Z8dqdE„Z9	 	 	 dvdFejN                  dGejN                  dHejN                  dIejN                  dz  dJe:dKe(dLe:dz  fdM„Z;e-jb                  dNe0jd                  dOe0jd                  dPe0jd                  dQe0jd                  dRe0jd                  dSe0jd                  d6e0jd                  fdT„«       Z<dUejN                  dVejN                  dWejN                  dXejN                  dYejN                  f
dZ„Z=e-jb                  d[e0jd                  d\e0jd                  dQe0jd                  d]e0jd                  dRe0jd                  dSe0jd                  d^e0jd                  d6e0jd                  fd_„«       Z>	 dwdUejN                  dVejN                  d`ejN                  daejN                  dbejN                  dcejN                  d.e*dYejN                  dde(fde„Z?e-jb                  dfe0jd                  dge0jd                  dhe0jd                  die0jd                  dje0jd                  dke0jd                  dle0jd                  dme0jd                  dne0jd                  d2e0jd                  d3e0jd                  doe0jd                  d5e0jd                  d6e0jd                  d8e0jd                  d]e0jd                  f dp„«       Z@ydZ9dZ7dZ6dZ;dZ=dZ?dZ@y)xé    N)Ú	lru_cache)Ú	warn_once)Ú
has_tritoné   )Úget_metaÚ*TORCH_SPARSE_BSR_SCATTER_MM_LRU_CACHE_SIZEé   c                 ó   — | st        |«      ‚y ©N)Ú
ValueError)ÚcondÚmsgs     úW/var/www/html/ia-enhanced/venv/lib/python3.12/site-packages/torch/sparse/_triton_ops.pyÚcheckr      s   € ÙÜ˜‹oÐð ó    c                 óX   — t        |j                  t        j                  k(  | › d�«       y )Nz@(): only BSR sparse format is supported for the sparse argument.)r   ÚlayoutÚtorchÚ
sparse_bsr)Úf_nameÚts     r   Úcheck_bsr_layoutr      s'   € Ü	Ø	�‰”E×$Ñ$Ñ$Øˆ(ÐRÐSõr   c                 ór   — t        |j                  |k(  xr |j                  j                  dk(  | › d�«       y )NÚcudaz9(): all inputs are expected to be on the same GPU device.)r   ÚdeviceÚtype)r   r   r   s      r   Úcheck_devicer       s3   € Ü	Ø	�‰�FÑÒ6˜qŸx™xŸ}™}°Ñ6Øˆ(ÐKÐLõr   c           	      ó*  — t        |j                  «       dk\  xr |j                  «       dk\  | › d|j                  «       › d|j                  «       › d�«       |j                  dd  \  }}|j                  dd  \  }}t        ||k(  | › d|› d|› d�«       y )Nr	   zc(): all inputs involved in the matrix product are expected to be at least 2D, but got lhs.dim() == z and rhs.dim() == ú.éþÿÿÿzw(): arguments' sizes involved in the matrix product are not compatible for matrix multiplication, got lhs.shape[-1] == z( which is not equal to rhs.shape[-2] == )r   ÚdimÚshape)r   ÚlhsÚrhsÚ_mÚklÚkrÚ_ns          r   Úcheck_mm_compatible_shapesr)   '   s¥   € Ü	Ø�‰‹	�Q‰Ò)˜3Ÿ7™7›9¨™>Øˆ(ð  Ø #§¡£	˜{Ð*<¸S¿W¹W»Y¸KÀqð	Jôð �Y‰Y�r�sˆ^�F€BˆØ�Y‰Y�r�sˆ^�F€Bˆä	Ø
ˆb‰Øˆ(ð  Ø "˜tÐ#KÈBÈ4Èqð	Rõr   c           	      óè   — t        |j                  |k(  xrD |j                  t        j                  t        j                  t        j
                  ft        |Ž z   v | › d|› d|j                  › d�«       y )Nz\(): all inputs are expected to be of the same dtype and one of (half, bfloat16, float32) or z, but got dtype == r   )r   Údtyper   ÚhalfÚbfloat16ÚfloatÚtuple)r   r   r+   Úadditional_dtypess       r   Úcheck_dtyper1   8   sp   € Ü	Ø	�‰�5Ñò 	SØ�G‰GÜ�Z‰ZœŸ™¬¯©Ð5¼Ð?PÐ8QÑQðSàˆ(ð 3Ø3DÐ2Eð FØŸG™G˜9 Að	'õ	r   c           	      ó¢   ‡— t        |«      dk7  rt        dt        |«      › �«      ‚d„ Šˆfd„}t         ||«      | › d|d   › d|d   › d	�«       y )
Nr	   z"blocksize must have length 2, got c                 ó   — | | dz
  z   S ©Nr   © )Úvs    r   Úis_power_of_twoz(check_blocksize.<locals>.is_power_of_twoG   s   € Ø˜˜Q™‘KÐ Ð r   c                 ó@   •— d}| D ]  }|dk\  xr  ‰|«      xr |}Œ |S )NTé   r5   )ÚbÚresÚ	blocksizer7   s      €r   Úis_compatible_blocksizez0check_blocksize.<locals>.is_compatible_blocksizeJ   s8   ø€ ØˆØò 	KˆIà ‘?ÒA¡°yÓ'AÒJÀs‰Cð	Kð ˆ
r   z(): sparse inputs' blocksize (r   z, r   z;) should be at least 16 and a power of 2 in each dimension.)ÚlenÚAssertionErrorr   )r   r<   r=   r7   s      @r   Úcheck_blocksizer@   C   sj   ø€ Ü
ˆ9ƒ~˜ÒÜÐAÄ#ÀiÃ.ÐAQÐRÓSÐSò!ôô 
Ù 	Ó*Øˆ(Ð0°¸1±°¸bÀÈ1ÁÀð ODð 	Dõr   c                 ó^   — t        | j                  «       «      dkD  r| j                  «       S | S )a¹  Return input as a triton-contiguous tensor.

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

    While triton kernels support triton-non-contiguous tensors (all
    strides being greater than 1) arguments, a considerable slow-down
    occurs because tensor data is copied element-wise rather than
    chunk-wise. Zero strides is assumed to not have this defect.
    r   )ÚminÚstrideÚ
contiguous)r   s    r   Úmake_triton_contiguousrE   X   s)   € ô ˆ1�8‰8‹:ƒ˜Òð �|‰|‹~Ðàˆr   c                 ór   — 	 t        j                  d„ |D «       Ž S # t        $ r t        d| › d�«       Y y w xY w)Nc              3   ó:   K  — | ]  }|j                   d d –— Œ y ­w©Nr    ©r"   ©Ú.0r   s     r   ú	<genexpr>z'broadcast_batch_dims.<locals>.<genexpr>m   s   è ø€ Ò'F¸¨¯©°°¬Ñ'Fùó   ‚Fz3(): inputs' batch dimensions are not broadcastable!)r   Úbroadcast_shapesÚ	Exceptionr   )r   Útensorss     r   Úbroadcast_batch_dimsrQ   k   sB   € ðUÜ×%Ñ%Ñ'F¸gÔ'FÐGÐGøÜò UÜˆe˜�xÐRÐSÖTðUús   ‚ ›6µ6c              '   ól   K  — |D ]+  }t        d «      g|j                  «       z  }||| <   ||   –— Œ- y ­wr   )Úslicer!   )r!   Úslice_rangerP   r   Úslicess        r   ÚslicerrV   r   s=   è ø€ Øò ˆÜ˜“+� §¡£Ñ(ˆØ!ˆˆs‰Ø�‰i‹ñùs   ‚24c              '   ó°   K  — |D ]M  }t        d «      g|j                  «       z  }t        | |d¬«      D ]  \  }}|€Œ	|||<   Œ |t        |«         –— ŒO y ­w)NF©Ústrict)rS   r!   Úzipr/   )ÚdimsrU   rP   r   ÚsÚdÚd_slices          r   Úmultidim_slicerr_   y   se   è ø€ Øò ˆÜ�4‹[ˆM˜AŸE™E›GÑ#ˆÜ˜d F°5Ô9ò 	‰JˆAˆwØ‰}Ø��!’ð	ð ”�a“‰kÓñùs
   ‚9A¼Ac               '   óV   K  — | D ]  }|–— |j                  «       E d {  –—†  Œ  y 7 Œ­wr   )rC   )rP   r   s     r   Úptr_stride_extractorra   ‚   s,   è ø€ Øò ˆØŠØ—8‘8“:×Ññàús   ‚)Ÿ' )c           
   #   ó,  ‡ ‡‡K  — t        ‰ «      dk  st        ‰ «      dkD  rt        dt        ‰ «      › �«      ‚t        ‰«      dk  st        ‰«      dkD  rt        dt        ‰«      › �«      ‚dd l}ˆ ˆfd„}ˆfd„} |j                   |«       Ž D ]t  }t	        ‰ |‰d¬«      D ���	cg c]  \  }}}	t        ||z
  |	«      ‘Œ }
}}}	t	        ||
d¬«      D ��cg c]  \  }}t        |||z   «      ‘Œ }}}|
d d d	…   g ||«      ¢­–— Œv y c c}	}}w c c}}w ­w)
Nr   é   z"full_grid length must be 0-3, got z$grid_blocks length must be 0-3, got c               3   óX   •K  — t        ‰‰d¬«      D ]  \  } }t        d| |«      –— Œ y ­w)NFrX   r   )rZ   Úrange)ÚfgÚmgÚ	full_gridÚgrid_blockss     €€r   Úgenerate_grid_pointsz.grid_partitioner.<locals>.generate_grid_points�   s2   øè ø€ Ü˜) [¸Ô?ò 	#‰FˆB�Ü˜˜2˜rÓ"Ó"ñ	#ùs   ƒ'*c              3   ón   •K  — ‰j                  «       D ]  \  }}t        t        || |«      «      –— Œ y ­wr   )ÚitemsÚnextr_   )rU   r   Út_dimsÚtensor_dims_maps      €r   Úgenerate_sliced_tensorsz1grid_partitioner.<locals>.generate_sliced_tensors”   s7   øè ø€ Ø(×.Ñ.Ó0ò 	;‰IˆAˆvÜ” v¨v°qÓ9Ó:Ó:ñ	;ùs   ƒ25FrX   éÿÿÿÿ)r>   r?   Ú	itertoolsÚproductrZ   rB   rS   )rh   ri   ro   rr   rj   rp   Ú
grid_pointrf   Úgprg   ÚgridÚgrU   s   ```          r   Úgrid_partitionerrx   ˆ   s-  úè ø€ Ü
ˆ9ƒ~˜ÒœS ›^¨aÒ/ÜÐAÄ#ÀiÃ.ÐAQÐRÓSÐSÜ
ˆ;Ó˜!Òœs ;Ó/°!Ò3ÜÐCÄCÈÓDTÐCUÐVÓWÐWãõ#ô;ð (�i×'Ñ'Ñ)=Ó)?Ð@ò 	;ˆ
ô " )¨Z¸ÈUÔS÷
ð 
á��B˜ô ��R‘˜Õð
ˆò 
ô 25°ZÀÈeÔ1T×U©¨¨A”%˜˜B ™FÕ#ÐUˆÑUð ‘4�R�4‰jÐ:Ñ2°6Ó:Ñ:Ó:ñ	;ùô
ùó Vùs   …B DÂ%D
Â?DÃDÃ-'Dc                 ó    ‡— dd d d…   }|€|}n#d„ Št        ˆfd„t        ||d¬«      D «       «      }t        |||«      D ]  ^}} | |g|¢­Ž  Œ y )N)iÿÿÿéÿÿ  rz   rq   c                 ó6   — | €|S t        dt        | |«      «      S r4   )ÚmaxrB   )rw   rg   s     r   Úvalid_grid_dimz%launch_kernel.<locals>.valid_grid_dim«   s!   € ØˆyØ�	ô ˜1œc ! R›jÓ)Ð)r   c              3   ó6   •K  — | ]  \  }} ‰||«      –— Œ y ­wr   r5   )rK   rw   rg   r}   s      €r   rL   z launch_kernel.<locals>.<genexpr>²   s$   øè ø€ ò 
á��2ñ ˜1˜b×!ñ
ùs   ƒFrX   )r/   rZ   rx   )Úkernelro   rh   ri   Úcuda_max_gridrv   Úsliced_tensorsr}   s          @r   Úlaunch_kernelr‚   ¤   st   ø€ à.©t°¨tÑ4€MØÐØ#‰ò	*ô ó 
ä˜[¨-ÀÔFô
ó 
ˆô
 "2Ø�; ó"ò &Ðˆˆ~ñ 	ˆtÐ%�nÔ%ñ&r   c           
      ó8  — | j                  «       j                  d«      }| j                  «       j                  d«      }t        | j	                  «       j                  d«      «      }|D �cg c]  }t        |j                  d«      «      ‘Œ }}t        j                  |j                  d d gd„ |D «       ¢­Ž }d„ } |||d«      } |||d«      } ||||j                  dd  «      }|D �cg c]  } ||||j                  dd  «      ‘Œ }}|||g|¢­S c c}w c c}w )Nr   éýÿÿÿc              3   ó:   K  — | ]  }|j                   d d –— Œ y ­wrH   rI   rJ   s     r   rL   z!prepare_inputs.<locals>.<genexpr>Æ   s   è ø€ Ò;¨a˜QŸW™W S bœ\Ñ;ùrM   c                 ób   — | j                  ||z   «      j                  dt        |«      dz
  «      S )Nr   r   )Úbroadcast_toÚflattenr>   )r   Ú
batch_dimsÚinvariant_dimss      r   Úbatch_broadcast_and_squashz2prepare_inputs.<locals>.batch_broadcast_and_squashË   s1   € Ø�~‰~˜j¨>Ñ9Ó:×BÑBØŒs�:‹ Ñ"ó
ð 	
r   ©rq   r    )Úcrow_indicesÚ	unsqueezeÚcol_indicesrE   Úvaluesr   rN   r"   )	ÚbsrÚdense_tensorsr�   r�   r�   r   rP   Úbatch_dims_broadcastedr‹   s	            r   Úprepare_inputsr”   ½   s6  € à×#Ñ#Ó%×/Ñ/°Ó2€LØ—/‘/Ó#×-Ñ-¨aÓ0€KÜ# C§J¡J£L×$:Ñ$:¸1Ó$=Ó>€FØ?LÖM¸!Ô% a§k¡k°!£nÕ5ÐM€GÐMô #×3Ñ3Ø�‰�S�bÐðÙ;°7Ô;òÐò
ñ
 .ØÐ,¨eó€Lñ -¨[Ð:PÐRWÓX€KÙ'ØÐ&¨¯©°R°SÐ(9ó€Fð
 öàñ 	# 1Ð&<¸a¿g¹gÀbÀc¸lÕKð€Gð ð
 ˜ fÐ6¨wÑ6Ð6ùò7 Nùò,s   Á+!DÃ*Dc                 óŠ  — t        | |g|¢­Ž }|j                  «       j                  |dz   «      }|j                  «       j                  |dz   «      }|j	                  «       j                  ||j	                  «       j
                  dd  z   «      }||j
                  dd  z   }t        j                  |||||j                  ¬«      S )NrŒ   r„   r    ©Úsizer   )	rQ   r�   r‡   r�   r�   r"   r   Úsparse_compressed_tensorr   )r   r‘   rP   Úbatch_shaper�   r�   r�   r—   s           r   Úbroadcast_batch_dims_bsrrš   à   s±   € Ü& v¨sÐ=°WÒ=€Kà×#Ñ#Ó%×2Ñ2°;ÀÑ3FÓG€LØ—/‘/Ó#×0Ñ0°¸uÑ1DÓE€KØ�Z‰Z‹\×&Ñ& {°S·Z±Z³\×5GÑ5GÈÈÐ5LÑ'LÓM€FØ˜Ÿ™ 2 3˜Ñ'€DÜ×)Ñ)Ø�k 6°¸S¿Z¹Zôð r   c                 óš   — | j                   �^ }}}|||d   z  |d   ||d   z  |d   gz   }| j                  |«      j                  dd«      S )Nr   r   r„   r    )r"   ÚviewÚ	transpose)r   r<   ÚrestÚmÚnÚ	new_shapes         r   Útile_to_blocksizer¢   í   sd   € Ø—'‘'�K€Tˆ1ˆaØØ	ˆY�q‰\ÑØ�!‰Ø	ˆY�q‰\ÑØ�!‰ð	ñ €Ið �6‰6�)Ó×&Ñ& r¨2Ó.Ð.r   c                 ó  — | j                   dk  r!| j                  d«      } | j                   dk  rŒ!| j                   dkD  r| j                  d| j                   dz
  «      } | j                   dk7  rt        d| j                  › �«      ‚| S )zÙReturn tensor as 3D tensor by either prepending new dimensions to
    the tensor shape (when ``tensor.ndim < 3``), or by collapsing
    starting dimensions into the first dimension (when ``tensor.ndim >
    3``).
    rc   r   z3tensor should have 3 dimensions after reshape, got )ÚndimrŽ   rˆ   r?   r"   )Útensors    r   Ú	as1Dbatchr¦   ú   s}   € ð �+‰+˜Š/Ø×!Ñ! !Ó$ˆð �+‰+˜‹/à‡{�{�Q‚Ø—‘  6§;¡;°¡?Ó3ˆØ‡{�{�aÒÜØAÀ&Ç,Á,ÀÐPó
ð 	
ð €Mr   ©Úaccumulatorsc                óJ  — |d   }| j                   dk7  rt        d| j                   › d�«      ‚| j                  \  }}}|dk(  �r]|dd \  }}	|j                   dk7  rt        d|j                   › d�«      ‚|j                  \  }
}}||k7  rt        d	|› d
|› d�«      ‚|€B|j                  d   dz
  }t        j                  |||f| j
                  | j                  ¬«      }n>|j                  \  }}}||k7  rt        d|› d|› d�«      ‚||k7  rt        d|› d|› d�«      ‚|dz  s|dz  s|dz  st        €^t        |j                  d   dz
  «      D ]>  }||   }||dz      }t        ||«      D ]   }|	|   \  }}||xx   | |   ||   z  z  cc<   Œ" Œ@ |S t        | |||	|«       |S |dk(  �r|j                  }t        |«      }|j                  \  }}}||z  dk7  rt        d|› d|› d�«      ‚|dd \  }}}}}|d   }|€^||j                  «       j                  «       dz   |z  z   } t        j                  g |dd ¢| ‘|‘­| j
                  | j                  ¬«      }n)|j                  dd \  } }!|!|k7  rt        d|!› d|› d�«      ‚|j                  }"t        |«      }||z  }|dz  s|dz  s|dz  st        €î|j                  «        t        |«      D ]Ï  }#t        |j                  d   «      D ]²  }||   j                  «       }$||   j                  «       }||dz      j                  «       }t        |$|«      \  }%}&||#|%|%|z   …|&|&|z   …f   }'t        ||«      D ]D  }||   ||   }}t        |j                  «       |«      \  }(})|'| |   ||#|(|(|z   …|)|)|z   …f   z  z  }'ŒF Œ´ ŒÑ nt        | |||||||«       |j                  |"«      S |dk(  �rQ|j                  }t        |«      }|j                  \  }}}||z  dk7  rt        d|› d|› d�«      ‚|dd \  }}}}|d   }|€^||j                  «       j                  «       dz   |z  z   } t        j                  g |dd ¢| ‘|‘­| j
                  | j                  ¬«      }n)|j                  dd \  } }!|!|k7  rt        d|!› d|› d�«      ‚|j                  }"t        |«      }||z  }|dz  s|dz  s|dz  st        €øt        |«      D ]é  }#t        t!        |«      «      D ]Ð  }*t        ||*   j                  «       |«      \  }%}&|%|z  }+|&|z  },||+   j                  «       }-||+dz      j                  «       }.||#|%|%|z   …|&|&|z   …f   }'t#        t        |-|.«      «      D ]Q  \  }/}||,|.z  ||,z
  |-z  z   |/z      j                  «       }t        ||«      \  }(})|'| |   ||#|(|(|z   …|)|)|z   …f   z  z  }'ŒS ŒÒ Œë n>t        j$                  d|j
                  |j                  ¬«      }t        | |||||||«       |j                  |"«      S t'        |«      ‚)aw  Scattered matrix multiplication of tensors.

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

    accumulators (Tensor, optional): a tensor of matrix product
      accumulators. If ``indices_data[0]=="scatter_mm"``, the tensor
      is a 1-D batch tensor of output matrices. Otherwise, output
      matrices are slices of the :attr:`accumulators` tensor.
    r   rc   zblocks must be 3D, got ÚDÚ
scatter_mmr   Nzothers must be 3D, got z
blocks K (z) != others K (ú)©r+   r   zaccumulators Ms (z) != blocks Ms (zaccumulators Ns (z) != others Ns (r9   Úbsr_strided_mmzK (z) must be divisible by Ks (ÚSPLIT_Nr    úaccumulators N (ú) != others N (Úbsr_strided_mm_compressed)r   )r¤   r?   r"   r   Úzerosr+   r   Ú_scatter_mm2re   r¦   r|   ÚitemÚ_scatter_mm6Úzero_Údivmodrœ   r>   Ú	enumerateÚemptyÚNotImplementedError)0ÚblocksÚothersÚindices_datar¨   Úindices_formatÚ_PÚMsÚKsÚ	c_offsetsÚpqÚ_QÚKs_ÚNsÚRÚMs_ÚNs_ÚrÚg0Úg1rw   Ú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«     sP  € ðH " !‘_€Nà‡{�{�aÒÜÐ6°v·{±{°mÀ1ÐEÓFÐFØ—‘�J€BˆˆBà˜Ó%Ø$ Q RÐ(‰ˆ	�2à�;‰;˜!ÒÜ Ð#:¸6¿;¹;¸-ÀqÐ!IÓJÐJØ—l‘l‰ˆˆC�Ø�Š9Ü  :¨b¨T°ÀÀÀQÐ!GÓHÐHàÐØ—‘ Ñ" QÑ&ˆAÜ Ÿ;™;Ø�B˜� 6§<¡<¸¿¹ô‰Lð '×,Ñ,‰KˆAˆs�CØ�bŠyÜ$Ð'8¸¸Ð=MÈbÈTÐQRÐ%SÓTÐTØ�bŠyÜ$Ð'8¸¸Ð=MÈbÈTÐQRÐ%SÓTÐTà�Š7�b˜2’g  b¢¬LÐ,@Ü˜9Ÿ?™?¨1Ñ-°Ñ1Ó2ò =�Ø˜q‘\�Ø˜q 1™uÑ%�Ü˜r 2›ò =�AØ˜a™5‘D�A�qà  “O v¨a¡y°6¸!±9Ñ'<Ñ<”Oñ=ð=ð Ðô ˜ ¨°B¸ÔEØÐà	Ð+Ó	+Ø—|‘|ˆÜ˜6Ó"ˆà—,‘,‰ˆˆ1ˆaØˆ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�b˜2’g  b¢¬LÐ,@Ø×ÑÔ Ü˜1“Xò 
Q�Ü˜yŸ™¨qÑ1Ó2ò 	Q�AØ" 1™×*Ñ*Ó,�BØ" 1™×*Ñ*Ó,�BØ" 1 q¡5Ñ)×.Ñ.Ó0�BÜ# B¨›]‘F�B˜Ø& q¨"¨r°B©w¨,¸¸RÀ"¹W¸Ð'DÑE�CÜ" 2 r›]ò Q˜Ø(¨™|¨Y°q©\˜1˜Ü!'¨¯©«°!Ó!4™˜˜BØ˜v a™y¨6°!°R¸"¸r¹'°\À2ÈÈRÉÀ<Ð2OÑ+PÑPÑP™ñQñ	Qñ
Qô ØØØØØØØØô	ð × Ñ Ð!3Ó4Ð4à	Ð6Ó	6Ø—|‘|ˆÜ˜6Ó"ˆà—,‘,‰ˆˆ1ˆaØˆr‰6�QŠ;Ü  3 q cÐ)DÀRÀDÈÐ!JÓKÐKà0<¸Q¸RÐ0@Ñ-ˆ	�9˜i¨Ø�y‘/ˆàÐØ�i—m‘m“o×*Ñ*Ó,¨qÑ0°QÑ6Ñ6ˆAÜ Ÿ;™;Ø*�,˜s Ð#Ð* QÐ*¨Ñ*°&·,±,ÀvÇ}Á}ô‰Lð !×&Ñ& r sÐ+‰EˆAˆrØ�QŠwÜ$Ð'7¸°t¸?È1È#ÈQÐ%OÓPÐPà)×/Ñ/ÐÜ  Ó.ˆà�'‰\ˆà�Š7�b˜2’g  b¢¬LÐ,@Ü˜1“Xò Q�Üœs 9›~Ó.ò 
Q�AÜ# I¨a¡L×$5Ñ$5Ó$7¸Ó;‘F�B˜Ø˜b™�AØ˜b™�AØ" 1™×*Ñ*Ó,�BØ" 1 q¡5Ñ)×.Ñ.Ó0�BØ& q¨"¨r°B©w¨,¸¸RÀ"¹W¸Ð'DÑE�CÜ )¬%°°B«-Ó 8ò Q™˜˜1Ø% a¨"¡f°¸!±¸rÑ/AÑ&AÀAÑ&EÑF×KÑKÓM˜Ü!'¨¨1£™˜˜BØ˜v a™y¨6°!°R¸"¸r¹'°\À2ÈÈRÉÀ<Ð2OÑ+PÑPÑP™ñQñ
QñQô Ÿ™Ø˜IŸO™O°I×4DÑ4DôˆIô ØØØØØØØØô	ð × Ñ Ð!3Ó4Ð4ô " .Ó1Ð1r   c           
      ó\  — ||||	|
|hd hk(  �r.t         j                  j                  «       }t        d| ||||f|dt         j                  df¬«      }|� |j
                  d*i |¤Ž |S | ||fdk(  rU||fdk(  rd}d}d}d	}d}
d	}	�n¸||fd
k(  rd}d}d}d	}d}
d	}	�n£||fdk(  rd}d}d}d	}d}
d	}	�nŽ||fdk(  �r†d}d}d}d}d}
d	}	�nx| ||fdk(  rU||fdk(  rd}d}d}d}d}
d}	�n[||fd
k(  rd}d}d}d	}d}
d}	�nF||fdk(  rd	}d}d}d	}d}
d	}	�n1||fdk(  �r)d}d}d}d	}d}
d	}	�n| ||fdk(  rd||fdk(  rd	}d}d}d}d}
d}	nÿ||fd
k(  rd}d}d}d}d}
d}	në||fdk(  rd}d}d}d	}d}
d}	n×||fdk(  rd}d}d}d	}d}
d	}	nÃ||fdk(  r¼d}d}d}d}d}
d	}	n¯| ||fdk(  rd||fdk(  rd	}d}d}d}d}
d}	n“||fd
k(  rd	}d}d}d	}d}
d}	n||fdk(  rd	}d}d}d	}d}
d	}	nk||fdk(  rd}d}d}d	}d}
d	}	nW||fdk(  rPd	}d}d}d}d}
d	}	nC| ||fdk(  r;||fdk(  rd}d}d}d}d}
d}	n'||fd
k(  rd}d}d}d}d}
d}	n||fdk(  rd}d}d}d}d}
d	}	|€(ddd	dddddddœ	j                  |d«      }|dk\  r|dk\  rd}||z  }|€t        |dk  rdnd|«      }|€t        |dk  rdnd|«      }|
xs d}
|	€†t        | |«      dkD  rddddœj                  |d	«      }	n`t        | |«      dk(  rddddœj                  |d	«      }	n:t        | |«      dk(  rdd	dœj                  |d	«      }	ndddœj                  |d	«      }	|xs d	}||kD  rt        d|› d|› d �«      ‚||kD  rt        d!|› d"|› d �«      ‚|| kD  rt        d#|› d$| › d �«      ‚||kD  rt        d%|› d&|› d �«      ‚||kD  rt        d'|› d(|› d �«      ‚t        d*||||
|	|d)œ|¤ŽS )+Nr«   r   ç      à?©Úversion)é   rê   rê   )r9   r9   r   r9   é   )é    rì   r	   rì   )é@   rí   )é€   rî   )é   rï   rï   é   rí   rî   )é   rñ   rñ   )rê   rê   )é   rò   rò   )é   ró   ró   rê   )	r9   rì   rí   rî   rê   rï   rñ   ró   i    rï   rò   rñ   )r9   rì   rí   )r9   rì   zTILE_M (z) must be <= Ms (r¬   zTILE_N (z) must be <= Ns (zMs (z) must be <= M (zNs (z) must be <= N (zKs (z) must be <= K ()ÚTILE_MÚTILE_NÚ
GROUP_SIZEÚ
num_stagesÚ	num_warpsr¯   r5   )
r   r   Úget_device_namer   Úfloat16ÚupdateÚgetrB   r?   Údict)rÙ   rÒ   rÓ   rÁ   rÂ   rö   rô   rõ   r¯   rø   r÷   ÚextraÚdevice_namerØ   rÇ   s                  r   Úscatter_mm_metar     s1  € ð 	�˜ ¨J¸
ÐCÈÀvÓMÜ—j‘j×0Ñ0Ó2ˆÜØØ��1�b˜"ÐØØœŸ™ sÐ+ô	
ˆð ÐØˆD�K‰KÑ ˜%Ò ØˆKð ˆq�!ˆ9˜
Ò"Ø�Bˆx˜8Ò#Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�b�˜ZÓ'Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø��AˆY˜*Ò$Ø�Bˆx˜8Ò#Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø�b�˜ZÓ'Ø�Ø�Ø�Ø�
Ø�
Ø’	Ø��AˆY˜+Ò%Ø�Bˆx˜8Ò#Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜ZÒ'Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜ZÒ'Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø��AˆY˜+Ò%Ø�Bˆx˜8Ò#Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜ZÒ'Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜ZÒ'Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø��AˆY˜+Ò%Ø�Bˆx˜8Ò#Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø‘	Ø�b�˜XÒ%Ø�Ø�Ø�Ø�
Ø�
Ø�	à€ð ØØØØØØØØñ

÷ ‰#ˆa�‹*ð 	ð �Š9˜˜dšØˆGØ	
ˆg‰€BØ€~Ü˜2 š8‘R¨¨RÓ0ˆØ€~Ü˜2 š8‘R¨¨RÓ0ˆØ’˜q€JØÐÜˆq�!‹9�tÒØ A¨1Ñ-×1Ñ1°"°aÓ8‰IÜ��A‹Y˜$ÒØ A¨1Ñ-×1Ñ1°"°aÓ8‰IÜ��A‹Y˜#ÒØ A™×*Ñ*¨2¨qÓ1‰Ià A™×*Ñ*¨2¨qÓ1ˆIØ’˜q€Jà�‚{Ü˜x¨ xÐ/@ÀÀÀAÐFÓGÐGØ�‚{Ü˜x¨ xÐ/@ÀÀÀAÐFÓGÐGØ	ˆA‚vÜ˜t B 4Ð'7¸°s¸!Ð<Ó=Ð=Ø	ˆA‚vÜ˜t B 4Ð'7¸°s¸!Ð<Ó=Ð=Ø	ˆA‚vÜ˜t B 4Ð'7¸°s¸!Ð<Ó=Ð=äð ØØØØØØñð ñð r   c                 óL  — |€t         j                  }|€|}|€d}||	|
|hd hk(  �rKt         j                  j                  «       }| |||||dk(  |dk(  |dk(  f}||u r|}n||f}t	        d|||||f¬«      }|€|dk7  rt	        d||||df¬«      }|€||urt	        d||||df¬«      }|€Žt	        dg |d d ¢d‘|dd  ¢­|||df¬«      }|€#||urt	        dg |d d ¢d‘|dd  ¢­|||df¬«      }t        |xs i «      D ]8  }||   }|d   }|d	   }||z  }||z  dk(  sŒ ||k  sŒ&t        |«      }||z  |d	<   Œ: |� |j                  di |¤Ž |S t        d
| ›d|›d|›d|›d|›d|›d|›d|›d|›�«       |xs t        ||z  d«      }|xs d}|
xs d}
|	xs d}	t        d|||
|	dœ|¤ŽS )Nrç   r   r   Úbsr_dense_addmmrè   r	   Ú*rc   r¯   z@bsr_dense_addmm uses non-optimal triton kernel parameters for M=z K=z N=z Ms=z, Ks=z beta=z alpha=z dtype=z out_dtype=rë   )r¯   ÚGROUP_SIZE_ROWr÷   rø   r5   )
r   rú   r   rù   r   Úsortedrý   rû   r   r|   )rÙ   rÒ   rÓ   rÁ   rÂ   ÚbetaÚalphar¯   r  rø   r÷   Úsparsityr+   Ú	out_dtypeÚ_versionrþ   rÿ   ÚkeyÚversion_dtyperØ   Úmatching_metaÚmkeyÚmeta_r    Úsplit_nÚcs                             r   Úbsr_dense_addmm_metar  ø  s¥  € ð* €}Ü—‘ˆØÐØˆ	ØÐØˆØ�˜J¨Ð7¸D¸6ÓAÜ—j‘j×0Ñ0Ó2ˆØ�!�Q˜˜B ¨¡	¨4°1©9°e¸q±jÐAˆØ�IÑØ!‰Mà! 9Ð,ˆMÜØØØØ˜}¨hÐ7ô	
ˆð ˆ<˜H¨šOÜØ!ØØØ! =°#Ð6ô	ˆDð ˆ<˜E¨Ñ2ÜØ! 3¨¸hÈÈsÐ=SôˆDð ˆ<ä$Ø!Ø)�#�b�q�'Ð)˜3Ð)  Q R Ñ)ØØ! =°#Ð6ô	ˆMð Ð$¨°iÑ)?Ü (Ø%Ø-�c˜"˜1�gÐ-˜sÐ- S¨¨ WÑ-ØØ% u¨cÐ2ô	!�ô ˜}Ò2°Ó3ò -�Ø% dÑ+�Ø˜‘G�Ø 	Ñ*�Ø˜‘L�Ø�q‘5˜A“: ! q£&Ü ›;�DØ&'¨1¡f�D˜’Oð-ð ÐØˆD�K‰KÑ ˜%Ò ØˆKô
 ðØ�t˜4˜Q˜D  !  U r e¨6¨b¨U°'°D°7¸(¸E¸8À8ÀUÀHÈLÈiÈ\ð[ôð
 Ò(œ˜Q "™W a›€GØ#Ò( q€NØ’˜q€JØ’˜Q€IÜð ØØ%ØØñ	ð
 ñð r   c                   ó2   — e Zd ZdZd„ Zd„ Zd„ Zed„ «       Zy)ÚTensorAsKeyaR  A light-weight wrapper of a tensor that enables storing tensors as
    keys with efficient memory reference based comparison as an
    approximation to data equality based keys.

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

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

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

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