+
    "Ý³jfX  ã                   ób  a € R" t0 t ^ RIt^ RIHt ^ RIHt ^ RIHtHt ^ RI	t	^ RI
Ht ^ RIHtHtHt ^ RIHt ]P$                  ! RR	4      tR
^R^RRRRRRRR	RR	RR	RR/	R ltR#R ltR
^R^RRRRRRRR	RRRR/R R llt]! RR7       ! R R4      4       tRs] ^ k ]! 4       tR R ltR  R! ltR# )$é    N)Ú	dataclass)ÚLock)ÚCallableÚOptional)Ú	warn_once)Úget_metaÚminimizeÚupdate)Ú
has_tritonÚBSR_AUTOTUNEFÚbetaÚalphaÚ
left_alphaÚright_alphaÚoutÚstoreÚverboseÚforceÚopnamec       	        ó�  aaaa#a$€ ^ RI o$Vf   RpVP                  R,          pVP                  4       pVP                  4       pVP	                  4       ^,
          pVP                  Wÿ^,            w  ppVP                  V^,           V^,            w  pp\        ^^^\        VV,          ^4      R7      p\        ^VP                  4       V,          V,          VV,          ,          ,
          ^4      pVP                  pVf   TpMVP                  pVVJ d   TpMVV3p^ VV3pVVVVVS^ 8H  S^8H  V^8H  3p\        VVVRR7      pVf   Rp\        VV^ VR3RR7      pVf   TpMV
'       g   V# RpWW$V3VVVV$3R llpVVVVV3R	 lp\        VVVV^V	R
7      w  o#pp p!V	'       d   \        RV! RVR RV R R24       V'       d…   V'       d   S#V8X  d   VVJgp   \        P                  P                  4       p"\!        TT"TT\"        ;QJ d     . V#3R l\%        S#4       4       F  NK  	  5M! V#3R l\%        S#4       4       4      4       S## )zžTune bsr_dense_addmm kernel parameters against the given inputs.

When store is True, the tuning results will be stored in the
database of kernel parameters.
NÚbsr_dense_addmm)ÚGROUP_SIZE_ROWÚ
num_stagesÚ	num_warpsÚSPLIT_NT)ÚversionÚexactFç      à?c           	      óh   <a aaaaa€ VVVVVVV VV	3	R  lpS
P                   P                  VR^dR7      # )c                  ó,   <	€ \        SSSSS SSSSR 7	      # ))r   r   r   r   Úmetar   )r   )	r   r   ÚbsrÚdenseÚinputr   r!   r   r   s	   €€€€€€€€€Ú^/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torchao/kernel/bsr_triton_ops.pyÚ	test_funcÚ6tune_bsr_dense_addmm.<locals>.bench.<locals>.test_funcZ   s+   ø€ Ü"ØØØØØØ%Ø'ØØô
ð 
ó    iô  )ÚwarmupÚrep)ÚtestingÚdo_bench)r!   r$   r"   r#   r   r   r&   r   r   r   Útritons   ffffff €€€€r%   ÚbenchÚ#tune_bsr_dense_addmm.<locals>.benchY   s+   þ€ ÷	õ 	ð �~‰~×&Ñ& y¸À#Ð&ÓFÐFr(   c	                 ó²  € V R9   p	\        ^^^^R7      V ,          p
\        \        WW,          ^4      R7      P                  V 4      p\        ^^^^R7      V ,          pV	'       d0   V^ 8”  d   WV,          ,          MW\        V4      ,          ,          pMWV,          ,           pV
e   \        WÚ4      pVe   \	        WÛ4      pV R 8X  d   W],          ^ 8w  d   V# V# )r   )r   r   r   r   )r   >   r   r   )ÚdictÚmaxÚgetÚabsÚmin)ÚnameÚvalueÚ	directionr!   ÚMÚNÚKÚBMÚBKÚis_logÚ	min_valueÚ	max_valueÚ
value_stepÚ
next_values   &&&&&&&&&     r%   Ústep_meta_parameterÚ1tune_bsr_dense_addmm.<locals>.step_meta_parameterj   sÐ   € ð Ð1Ñ1ˆÜ ¨a¸AÈaÔPÐQUÕVˆ	Ü¤ Q¥W¨a£Ô1×5Ñ5°dÓ;ˆ	Ü !¨q¸QÈqÔQÐRVÕWˆ
ßð ˜q”=ð  IÕ-Ö-à¬S°«^Õ;Õ<ñ ð ¨iÕ!7Õ7ˆJØÒ Ü˜ZÓ3ˆJØÒ Ü˜ZÓ3ˆJØ�9Ô ¥°1Ô!4ØˆLØÐr(   )Úmax_stepr   z-> z
, speedup=z.1fz %, timing=z.3fz msc              3   ó6   <"  € T F  pSV,          x € K  	  R # 5i©N© )Ú.0Úkr!   s   & €r%   Ú	<genexpr>Ú'tune_bsr_dense_addmm.<locals>.<genexpr>—   s   øé € Ð0¡<˜a�$�q—'’'£<ùs   ƒéÿÿÿÿ)r-   ÚshapeÚvaluesÚcrow_indicesÚdimr1   r2   ÚroundÚ_nnzÚdtyper   r	   ÚprintÚtorchÚcudaÚget_device_namer
   ÚtupleÚsorted)%r$   r"   r#   r   r   r   r   r   r   r   r   r   r:   rO   rP   Ú
batch_ndimr9   r;   r<   r=   Úreference_metaÚsparsityrT   Ú	out_dtypeÚversion_dtyper   ÚkeyÚinitial_metaÚmay_skip_updater.   rC   ÚspeedupÚtimingÚsensitivity_messageÚdevice_namer!   r-   s%   &&&d$dd$$$$$                       @@r%   Útune_bsr_dense_addmmrg      sR  ü€ ó( à‚~Ø"ˆà�‰�B�€AØ�Z‰Z‹\€FØ×#Ñ#Ó%€LØ×!Ñ!Ó# aÕ'€JØ�9‰9�Z¨q¥.Ð1�D€A€qØ�\‰\˜* q�.¨:¸­>Ð:�F€Bˆô Ø Q°!¼SÀÀbÅÈ!»_ô€Nô
 �Q˜Ÿ™› b�¨2Õ-°°QµÕ7Õ7¸Ó;€HØ�I‰I€EØ
‚{Ø‰	à—I‘Iˆ	Ø�EÓØ‰à 	Ð*ˆØ�- Ð*€GØˆa��B˜˜D A™I t¨q¡y°%¸1±*Ð
=€Cô ˜F C°ÀÔE€LØÒØˆÜ ¨°a¸À°_ÈDÔQˆØÒØ)ˆLøßØÐàˆð  °È÷ Gò Gð" =>ÀÀaÈBÐSUô ô2 2:ØØØØØØô2Ñ.€Dˆ'�6Ð.÷ Ü�Ð'Ð(¨¨7°C¨.¸¸VÀS¸MÈÐMÔNçß˜D LÔ0°\ÈÓ5Wä—j‘j×0Ñ0Ó2ˆÜØØØØßŒEÔ0¤6¨$¤<Ó0�E‰EÔ0¤6¨$¤<Ó0Ó0ô	
ð €Kr(   c                 ó&  € Vf   \         P                  pVf   TpVf   RpWyW¨0R 08X  Ed•   \         P                  P                  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: R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   r   )r   :Né   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=zC. To find optimal triton kernel parameters, run with BSR_AUTOTUNE=1r   r   r   rH   )
rV   Úfloat16rW   rX   r   rZ   r1   r
   r   r2   )r9   r;   r:   ÚMsÚKsr   r   r   r   r   r   r]   rT   r^   Ú_versionÚextrarf   r`   r_   r!   Úmatching_metaÚmkeyÚmeta_ÚnÚsplit_nÚcs   &&&&&&&&&&&&&&&,          r%   Úbsr_dense_addmm_metarw   �   s£  € ð* ‚}Ü—‘ˆØÒØˆ	ØÒØˆØ˜JÐ7¸D¸6ÕAÜ—j‘j×0Ñ0Ó2ˆØ�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É\ð ZTðTôð ×(Ð(œ˜Q�W a›€GØ#×(Ð( q€NØ—�˜q€JØ—�˜Q€IÜñ Øðà%ðð ðð ð	ð
 ñð r(   Úskip_checksÚmax_gridr!   c                óØ  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\        \         P                  ,          R\        \         P                  ,          R\        \         P                  ,          R\        R\        \        \        \
        ,          \        \
        ,          \        \
        ,          3,          ,          R	\        \        ,          /	# )
ri   r$   r"   r#   r   r   r   rx   ry   r!   )rV   ÚTensorr   ÚboolrY   Úintr1   )Úformats   "r%   Ú__annotate__r   þ   sµ   € ÷ Cñ CÜ�<‰<ðCä	�‰ðCô �<‰<ðCô œŸ™Õ&ðCô œ%Ÿ,™,Õ'ðCô 
”%—,‘,Õ	ðCô ðCô ”uœX¤c�]¬H´S­M¼8ÄC½=ÐHÕIÕJðCô ”4�.ñCr(   c               ó˜
  aaa
a$a%a&a'a(a)a*a+a,€ \        4       o(S(P                  '       g   \        R4      h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^ RIH	o, ^ RI
HpHpHpHo*Hp V! 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# S
f§   \+        ^VP#                  4       V^ ,          ,          V^,          ,          VV,          ,          ,
          ^4      p\,        '       d   \/        V VVSSVVVRRRRR7      o
M7\1        VVVV^ ,          V^,          SSVVP2                  VP2                  R7
      o
Ro)Ro+Vf*   Ro)VP!                  R4      P4                  ! . VOVNVN5!  pM,VP6                  ! . VOVN^N5!  P4                  ! . VOVNVN5!  pVf*   Ro+VP!                  R4      P4                  ! . VOVNVN5!  pM,VP6                  ! . VO^NVN5!  P4                  ! . VOVNVN5!  pVP9                  4       R,          ^ 8X  g   Q hVP9                  4       R,          ^ 8X  g   Q hTpV! WW%Wg4      w  pppp ppppVw  o%o$S
P;                  R	\=        VS%,          ^4      4      pVV,          o&TpV! VS%S&34      pV! VS$S&34      pV! V S%S&34      p V! VS%S&34      pV! VS%S&34      p\>        P@                  S,PB                  \>        PD                  S,PB                  \>        PB                  S,PF                  \>        PF                  S,PF                  \>        PH                  S,PJ                  \>        PJ                  S,PJ                  /VP2                  ,          o'VPM                  ^ 4      pVPM                  R4      ^,
          pVPM                  R4      pVVV3p V	eC   \O        V	R
,          RRR1,          4      R^\Q        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^ 8w  g   Q hV$V%V&VVV'V(V)V
V*V+V,3R lp#V! V#V"V V!4       VPS                  4       VPS                  4       8w  d+   VP'                  VP7                  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.
zunsupported without tritonr   N)Úbroadcast_batch_dimsÚlaunch_kernelÚprepare_inputsÚptr_stride_extractorÚtile_to_blocksizeTF)	r   r   r   r   r   r   r   r   r   )r]   rT   r^   r   :Nrk   Nc                 ó¤   <€ SP                   V ,          ! . S! 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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_COLÚ
allow_tf32Ú	acc_dtypeN)ÚkernelÚfloat32)ÚgridÚsliced_tensorsr=   r<   ÚBNr   r   Údot_out_dtypeÚimplsrŠ   r!   r„   r‹   Útls   &*€€€€€€€€€€€€r%   r‘   Úbsr_dense_addmm.<locals>.kernel¨  s¶   ø€ Ø�‰�TÖð 	
Ù! >Ñ2ð	
àð	
ð ò	
ð  ™	ð		
ð
 ! A™Ið	
ð  !™ð	
ð 0ð	
ð  2ð	
ð ð	
ð ð	
ð ð	
ð %¨¯
©
Ñ2ð	
ð $ð	
ð ô	
r(   rM   rH   éþÿÿÿéýÿÿÿrG   )r   NN)r   NrM   )r   r›   éüÿÿÿ)r   r›   N)*Ú_get_triton_implsÚ	availableÚAssertionErrorrO   rP   Úcol_indicesrQ   rN   Útriton.languageÚlanguageÚtorch.sparse._triton_opsr�   r‚   rƒ   r„   r…   Ú	new_emptyrS   Úzero_Úcopy_Úmul_rR   ÚAUTOTUNErg   rw   rT   ÚexpandÚviewÚstrider3   r2   rV   rl   r’   Úbfloat16Úfloat64Úint8Úint32ÚsizerY   ÚlenÚdata_ptr)-r$   r"   r#   r   r   r   r   r   rx   ry   r!   Úf_namerO   rP   r    r[   r9   r;   Ú	blocksizer:   r�   r‚   rƒ   r…   Úoriginal_batch_dims_broadcastedr]   Ú
out_backupr   Úout_untiledÚ	n_batchesÚn_block_rowsÚn_block_colsÚ	full_gridÚgrid_blocksÚtensor_dims_mapr‘   r=   r<   r•   r–   r—   rŠ   r„   r‹   r˜   s-   &&&dd$$$$$d                         @@@@@@@@@r%   r   r   þ   sO  ÿû€ ô, Ó€EØ�?�?ˆ?ÜÐ9Ó:Ð:à€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Ø—‘˜”Øˆ
à‚|Ü˜˜SŸX™X›Z¨)°A­,Õ6¸À1½ÕEÈÈQÍÕOÕOÐQRÓSˆß‹8Ü'ØØØØØØ%Ø'ØØØØØ(ô‰Dô (ØØØØ˜!•Ø˜!•ØØØ!Ø—k‘kØŸ)™)ôˆDð ÐØÐØÒØ ÐØ—_‘_ 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ó
ˆð ×ÑÓ˜rÕ" aÔ'Ð'Ð'Ø×ÑÓ Õ# qÔ(Ð(Ð(à€Jñ 	�s 5°kÓGñ	ØØØØØØØØð �F€BˆØ�h‰h�y¤# a¨2¥g¨q£/Ó2€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ð �AŒ:Ðˆ:÷
÷ 
ð 
ñ$ �&˜/¨9°kÔBà
‡|�|ƒ~˜×,Ñ,Ó.Ô.ð 	×Ñ˜×)Ñ)¨*×*:Ñ*:Ó;Ô<àÐr(   T)Úfrozenc                   ó,   a € ] tR tRt o V 3R ltRtV tR# )Ú_BsrTritonImplsiÄ  c                óB   <€ V ^8„  d   Qh/ S[ ;R&   S[S[,          ;R&   # )ri   rž   r‘   )r|   r   r   )r~   Ú__classdict__s   "€r%   r   Ú_BsrTritonImpls.__annotate__Ä  s!   ø‡ ‚ á�Oñ ñ ‘XÕÑò r(   rH   N)Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__annotate_func__Ú__static_attributes__Ú__classdictcell__)rÂ   s   @r%   rÀ   rÀ   Ä  s   ø‡ ‡ „ r(   rÀ   c                ó$   € V ^8„  d   QhR\         /# ©ri   Úreturn©rÀ   )r~   s   "r%   r   r   Î  s   € ÷ 	ñ 	œ?ñ 	r(   c                  ó    € \         p V f0   \        ;_uu_ 4        \         p V f   \        4       p V s R R R 4       V # V #   + '       g   i     T # ; irG   )Ú_triton_implsÚ_triton_impls_lockÚ_build_triton_impls)r—   s    r%   r�   r�   Î  sI   € ä€EØ‚}ßÖÜ!ˆEØŠ}Ü+Ó-�Ø %�÷	  ð
 €Lˆ5€L÷  Öð
 €Lús	   ™<¼A	c                ó$   € V ^8„  d   QhR\         /# rÌ   rÎ   )r~   s   "r%   r   r   Ú  s   € ÷ _Añ _Aœ_ñ _Ar(   c                  óž   a€ \        4       '       g   \        R R4      # ^ RIp ^ RIHo V P
                  V3R lV3R ll4       p\        RV4      # )FNc                 óš  <€ V ^8„  d   QhRSP                   RSP                   RSP                   RSP                   RSP                   RSP                   RSP                   RSP                   R	SP                   R
SP                   RSP                   RSP                   RSP                   RSP                   RSP                   RSP                   /# )ri   Ú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�   r�   r   r   )Ú	constexpr)r~   r˜   s   "€r%   r   Ú)_build_triton_impls.<locals>.__annotate__â  s  ø€ ÷ U
ñ U
ðL &(§\¡\ðMU
ðP &(§\¡\ðQU
ðZ ')§l¡lð[U
ð^ ')§l¡lð_U
ðx —\‘\ðyU
ðz Ÿ™ð{U
ð| —l‘lð}U
ð~ Ÿ<™<ðU
ð@ ŸL™LðAU
ðB —|‘|ðCU
ðD —|‘|ðEU
ðF Ÿ™ðGU
ðH —<‘<ðIU
ðJ —L‘LðKU
ðL Ÿ™ðMU
ðN —‘ñOU
r(   c7           	      óx	  <€ V^ 8X  g   Q hV^ 8X  g   Q hV^ 8X  g   Q hV!^ 8X  g   Q hSPP                  ^R7      p7SPP                  ^ R7      p8SPP                  ^R7      p9SPP                  ^ R7      p:SPP                  ^R7      p;SPP                  V8V9V:V;V54      w  p8p9VVV7,          ,           VV8,          ,           p<SPP                  V<4      p=SPP                  V<V,           4      p>V>V=,
          p?SPP	                  ^ V04      p@SPP	                  ^ V24      pAV1^8  g   V1^,          ^ 8w  d   ^pBMT1pBSPP	                  ^ XB4      pCV VV7,          ,           VV=,          ,           VX@R,          ,          ,           VXAR,          ,          ,           pDVVV7,          ,           VV9,          ,           VVAR,          ,          ,           VVCR,          ,          ,           pEV#V$V7,          ,           V%V8,          ,           V&V9,          ,           V'V@R,          ,          ,           V(VCR,          ,          ,           pFVV	V7,          ,           V
V=,          ,           pGSPP                  V0VB3V3R7      pH\        V?4       F~  pISPP                  XD4      pJSPP                  XG4      pKSPP                  XEVVK,          ,           XCR,          V18  R7      pLXHSPP                  VJVLV4V3R7      ,          pHVDV,          pDVGV
,          pGK€  	  V-'       g
   XHV*,          pHV.'       go   VVV7,          ,           VV8,          ,           VV9,          ,           VX@R,          ,          ,           VXCR,          ,          ,           pMXHSPP                  VM4      ,          pHV/'       go   VVV7,          ,           VV8,          ,           V V9,          ,           V!X@R,          ,          ,           V"XCR,          ,          ,           pNXHSPP                  VN4      ,          pHV,'       d—   VVV7,          ,           VV8,          ,           VV9,          ,           VX@R,          ,          ,           VXCR,          ,          ,           pOV+'       d   XHSPP                  XO4      ,          pHMXHV)SPP                  XO4      ,          ,          pHSPP                  XFXHP                  V#P                  P                  4      XCR,          V18  R7       R# )r   )ÚaxisN)rT   )Úmask)r�   r^   )ºNNNN)Nrß   )Ú
program_idÚnum_programsÚ	swizzle2dÚloadÚarangeÚzerosÚrangeÚdotr   ÚtorT   Ú
element_ty)QÚ
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Ú	input_ptrÚinput_batch_strideÚinput_tiled_row_strideÚinput_tiled_col_strideÚinput_row_block_strideÚinput_col_block_strideÚ	dense_ptrÚdense_batch_strideÚdense_tiled_row_strideÚdense_tiled_col_strideÚdense_row_block_strideÚdense_col_block_strideÚ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_strideÚ
output_ptrÚoutput_batch_strideÚoutput_tiled_row_strideÚoutput_tiled_col_strideÚoutput_row_block_strideÚoutput_col_block_strider   r   r‡   rˆ   r‰   rŠ   r‹   rŒ   rŽ   r�   r�   r�   r   r   Ú	batch_pidÚrow_block_pidÚcol_block_pidr¹   rº   Úcrow_indices_offset_ptrÚ
nnz_offsetÚnnz_offset_nextÚrow_nnzÚrow_block_arangeÚinner_block_arangeÚPADDED_BLOCKSIZE_COLÚcol_block_arangeÚvalues_block_ptrsÚdense_block_ptrsÚoutput_ptrsÚcol_index_nnz_ptrÚoutput_acc_blockÚ_Úvalues_blockÚdense_row_idxÚdense_blockÚleft_alpha_ptrsÚright_alpha_ptrsÚ
input_ptrsr˜   sQ   &&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&&                         €r%   Ú_bsr_strided_addmm_kernel_implÚ;_build_triton_impls.<locals>._bsr_strided_addmm_kernel_implá  s  ø€ ðV +¨aÔ/Ð/Ð/Ø*¨aÔ/Ð/Ð/Ø+¨qÔ0Ð0Ð0Ø+¨qÔ0Ð0Ð0à—M‘M q�MÓ)ˆ	ØŸ™¨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¨/Ó:Ðà˜2Ô °Õ!3°qÔ!8Ø13Ñ à1>Ð àŸ9™9 QÐ(<Ó=Ðð Ø! 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ØÐ0Ð1¸ð $ó 
Ðô �w–ˆAØŸ7™7Ð#4Ó5ˆLð ŸG™GÐ$5Ó6ˆMØŸ'™'Ø Ð#9¸MÕ#IÕIØ% gÕ.°Ñ>ð "ó ˆ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°:Ó+>Õ$>Õ>Ð ð 	�‰ØØ×Ñ 
× 0Ñ 0× ;Ñ ;Ó<Ø! 'Õ*¨]Ñ:ð 	ö 	
r(   T)r   rÀ   r-   r¡   r¢   Újit)r-   r&  r˜   s     @r%   rÒ   rÒ   Ú  sL   ø€ Ü�<Š<Ü˜u dÓ+Ð+ãÝ à‡Z�Z÷U
ò U
ó ðU
ôn ˜4Ð!?Ó@Ð@r(   c                óV   € V ^8„  d   Qh/ ^ \         9   d   \        \        ,          ;R&   # )ri   rÐ   )Ú__conditional_annotations__r   rÀ   )r~   s   "r%   r   r      s   € ‡	�	÷D 0Ò /ŒxœÕ(Ñ /òE 
r(   )NNNNNNNr   )r*  ÚosÚdataclassesr   Ú	threadingr   Útypingr   r   rV   Útorch._dynamo.utilsr   Útorch.sparse._triton_ops_metar   r	   r
   Útorch.utils._tritonr   Úgetenvr¨   rg   rw   r   rÀ   rÐ   rÑ   r�   rÒ   r   )r*  s   @r%   Ú<module>r3     sC  øð÷ 
€	Ý !Ý ß %ã Ý )ß DÑ DÝ *à�9Š9�^ UÓ+€ðEð
 
ðEð ðEð ðEð ðEð 	ðEð ðEð ðEð ðEð ôEôP^ðBCð
 
ðCð ðCð *.ðCð +/ðCð #'ðCð ðCð NRðCð  ÷CñL �$Ô÷ð ó ðð
 ,0€Ó /Ù“VÐ õ	÷_Ar(   