§
     ’ÞjG)  ã                   ó.  — d 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
mZ ddlmZ n# e$ r dZY nw xY wdZdZd	\  ZZdZe�ˆ ed
dd¦  «        Z ej        g d¢¬¦  «        dej        dej        dej        dej        dej        f
d„¦   «         Zej        d„ ¦   «         Z ej        ddg¬¦  «        dej        dej        fd„¦   «         Zej        d„ ¦   «         Z ej        d¬¦  «        d„ ¦   «         Zd„ Zd„ Zd#d!„Z e�dd"l!m"Z"  e"¦   «          dS dS )$aî  Means over the token axis of attention keys (the smoothing mean SageAttention subtracts before quantizing them), as
x.mean(dim, keepdim=True) gives them but without the partial-sum buffer of PyTorch's split reduction: its FP32 partial sums follow the
shape (twice the keys for an image's 8k tokens, a quarter of them for a 720p video's 75k), allocated before the INT8 keys.

Triton (compute capability 7.5 and up; BF16 read as bits on GPUs without BF16 conversions), for keys of MIN_BYTES and more:
- each program sums a span of tokens for HB adjacent heads (one contiguous run of channels per token), BLOCK_T rows per load, in
  one FP32 accumulator per element of the tile (elementwise operations only: Triton 3.2 miscompiled a reshape fold of vectorized tiles),
  compensated (two-sum: each accumulator carries the rounding error of its additions), so that the sums stay exact to about one FP32
  rounding whatever the span (the precision does not depend on the GPU), then reduced over the rows in FP64 into its partial sums;
- a second pass adds the partial sums in a fixed order in FP64, divides once and rounds once, correctly, to the output dtype (BF16
  written as bits on every GPU, so Turing and the others give the same values). The mean is as close to the exact one as its dtype
  allows: never further than PyTorch's FP32 sums. Deterministic for a given GPU and configuration.
The launch configuration (rows per load, warps, pipelining stages, programs per SM) is tuned once per GPU, head width, heads per program
and size class (keys held by the L2 cache or not), on a slice of a real call (the kernel only reads it), and cached. The span is a
runtime value from the programs per SM: a new resolution or video length compiles nothing (the token, head and batch counts and the
strides are not compiled in either). The partial sums take (batch, heads, spans, head_dim) FP64, a few MB. Smaller keys, where
PyTorch's single launch is faster and its buffer small, and other devices or dtypes, take the PyTorch expression.é    Né   )Údispatch)Úbf16_bits_to_fp32Úfp64_to_bf16_bits)ÚLaunchTuneri   ))é   é   é   r   )r   r	   r
   é   )é   r   r
   r   )r   r   r
   r   )r   r	   r   r   )r   r   r   r   )r   r   r
   r   )r   r	   r
   é    )r	   i   Úwan2gp_token_meané   Ú	Attention)ÚTOKENSÚSPANÚSPANSÚHEADS)Údo_not_specializeÚDÚHBÚBLOCK_TÚSTAGESÚBITSc           	      ó  — t          j        d¦  «        t          j        d¦  «                             t           j        ¦  «        t          j        d¦  «                             t           j        ¦  «        }}}t          j        d|
|	z  ¦  «        }||
z  ||	z  z   }|
dk    r||z  |z   n
||
|	z  z  |z   }|                     t           j        ¦  «        |z  }t          j        ||z   |¦  «        }| ||z  z   }t          j        ||
|	z  ft           j        ¬¦  «        }t          j        ||
|	z  ft           j        ¬¦  «        }t          j        d|||¬¦  «        D ]«}||z   t          j        d|¦  «        z   }t          j	        ||d d …d f         |z  z   |d d d …f         z   ||k     d d …d f         d¦  «        }|rt          |¦  «        n|                     t           j        ¦  «        }||z   }||z
  }||||z
  z
  ||z
  z   z  }|}Œ¬t          j        |||z  |z   |z  |z   |	z  z   ||	z  z   t          j        |                     t           j        ¦  «        |                     t           j        ¦  «        z   d¬¦  «        ¦  «         d S )Nr   r   r   ©Údtype)Ú
num_stages)Úaxis)ÚtlÚ
program_idÚtoÚint64ÚarangeÚminimumÚzerosÚfloat32ÚrangeÚloadr   ÚstoreÚsumÚfloat64)ÚXÚPARTIALr   r   r   r   ÚSA0ÚSA1ÚSTr   r   r   r   r   ÚspanÚgroupÚa0ÚcÚheadÚcolumnÚstartÚendÚbaseÚhiÚloÚoffsetÚrowsÚxÚtotalÚbacks                                 úAC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\token_mean.pyÚ_partial_sums_kernelrC   -   s˜  € õ œ-¨Ñ*Ô*­B¬M¸!Ñ,<Ô,<×,?Ò,?ÅÄÑ,IÔ,IÍ2Ì=ÐYZÑK[ÔK[×K^ÒK^Õ_aÔ_gÑKhÔKh�RˆeˆÝŒI�a˜˜a™Ñ Ô ˆØ�r‰z˜A ™FÑ"ˆà$&¨!¢G G�˜‘˜q‘�°¸"¸q¹&Ñ1AÀAÑ1EˆØ—’�œÑ!Ô! DÑ(ˆÝŒj˜ ™ vÑ.Ô.ˆØ�2˜‘8‰|ˆåŒX�w  Q¡Ð'­r¬zÐ:Ñ:Ô:ˆÝŒX�w  Q¡Ð'­r¬zÐ:Ñ:Ô:ˆÝ”h˜q $¨¸FÐCÑCÔCð 	ð 	ˆFØ˜6‘>¥B¤I¨a°Ñ$9Ô$9Ñ9ˆDÝ”˜˜t A A A t Gœ}¨rÑ1Ñ1°F¸4ÀÀÀ¸7´OÑCÀdÈSÂjÐRSÐRSÐRSÐUYÐRYÔEZÐ\]Ñ^Ô^ˆAØ(,ÐBÕ! !Ñ$Ô$Ð$°!·$²$µr´zÑ2BÔ2BˆAØ˜‘FˆEØ˜2‘:ˆDØ�2˜ ™Ñ&¨1¨t©8Ñ4Ñ4ˆBØˆBˆBÝ
Œ�˜R %™Z¨$Ñ.°%Ñ7¸$Ñ>À!ÑCÑCÀaÈ!ÁeÑKÍRÌVÐTV×TYÒTYÕZ\ÔZdÑTeÔTeÐhj×hmÒhmÕnpÔnxÑhyÔhyÑTyð  ABð  NCñ  NCô  NCñ  	Dô  	Dð  	Dð  	Dð  	Dó    c           	      óº  — |                       t          j        ¦  «        }|                      t          j        ¦  «        }|                      t          j        d¬¦  «        }t          j        || k    t          j        t          j        |¦  «        t          j        | ¦  «        k    |dz
  |¦  «        dz  |¦  «        }|                      t          j        d¬¦  «        S )NT)Úbitcastr   )r"   r    r'   r,   Úuint32ÚwhereÚabs)r?   ÚfrA   Úbitss       rB   Ú_round_to_odd_fp32rL   E   sŸ   € ð �DŠD•”ÑÔˆØ�tŠt•B”JÑÔˆØ�tŠt•B”I tˆtÑ,Ô,ˆÝŒx˜ š	¥2¤8­B¬F°4©L¬L½2¼6À!¹9¼9Ò,DÀdÈQÁhÐPTÑ#UÔ#UÐXYÑ#YÐ[_Ñ`Ô`ˆØ�wŠw•r”z¨4ˆwÑ0Ô0Ð0rD   r   r   ÚKINDc                 ó6  — t          j        d¦  «                             t           j        ¦  «        t          j        d¦  «                             t           j        ¦  «        }}t          j        d|¦  «        }|t          j        d¦  «        z  |z   }	| |	|z  |z  z   |z   }
t          j        |ft           j        ¬¦  «        }t          d|¦  «        D ]}|t          j	        |
||z  z   ¦  «        z  }Œ ||                     t           j        ¦  «        z  }|dk    r |                     t           j
        ¦  «        }nB|dk    r-t          |¦  «                             t           j        ¦  «        }nt          |¦  «        }t          j        ||	|z  z   |z   |¦  «         d S )Nr   r   r   )r    r!   r"   r#   r$   Únum_programsr&   r,   r(   r)   r'   rL   Úfloat16r   r*   )r.   ÚOUTr   r   r   rM   r6   r4   ÚdÚrowÚpartialr@   r2   ÚmeanÚvalues                  rB   Ú_finish_kernelrW   N   s\  € å”= Ñ#Ô#×&Ò&¥r¤xÑ0Ô0µ"´-ÀÑ2BÔ2B×2EÒ2EÅbÄhÑ2OÔ2OˆbˆÝŒI�a˜‰OŒOˆØ•2”? 1Ñ%Ô%Ñ%¨Ñ,ˆØ˜C %™K¨!™OÑ+¨aÑ/ˆÝ”˜!˜¥R¤ZÐ0Ñ0Ô0ˆÝ˜!˜U‘O”Oð 	1ð 	1ˆDØ•R”W˜W t¨a¡xÑ/Ñ0Ô0Ñ0ˆEˆEØ�v—y’y¥¤Ñ,Ô,Ñ,ˆØ�1Š9ˆ9Ø—G’G�BœJÑ'Ô'ˆEˆEØ�QŠYˆYÝ& tÑ,Ô,×/Ò/µ´
Ñ;Ô;ˆEˆEå% dÑ+Ô+ˆEÝ
Œ��s˜Q‘w‘ Ñ" EÑ*Ô*Ð*Ð*Ð*rD   c                 ó\   — t           j                             | ¦  «        }|j        |j        fS ©N)ÚtorchÚcudaÚget_device_propertiesÚmulti_processor_countÚL2_cache_size)ÚindexÚ
propertiess     rB   Ú_devicera   a   s(   € å”×1Ò1°%Ñ8Ô8€JØÔ+¨ZÔ-EÐEÐErD   é   )Úmaxsizec           	      ó”  ‡— ˆfd„t          t          | ¦  «        dz
  ¦  «        D ¦   «         }t          |¦  «        dk    r|                     dd¦  «         |d         €dn| |d                  ||d                  f\  }}| |d                  ||d                  | d         }}}||| ‰         ||||‰         ||k    r|dz  dk    rdndfS )a  The (a0, heads, tokens) view of a mean over dim of a tensor of this shape and these strides: its sizes and strides, and the heads
    per program (4 when the heads are adjacent in memory and a multiple of 4: one contiguous run of channels per token, no head mask).c                 ó    •— g | ]
}|‰k    ¯|‘ŒS © rf   )Ú.0rR   Údims     €rB   ú
<listcomp>z_plan.<locals>.<listcomp>k   s   ø€ Ð;Ð;Ð;�A°!°s²(°(ˆa°(°(°(rD   r   r   N)r   r   éÿÿÿÿr	   )r(   ÚlenÚinsert)	ÚshapeÚstridesrh   Úothersr4   Úsa0ÚheadsÚsa1Úwidths	     `      rB   Ú_planrt   g   sä   ø€ ð <Ð;Ð;Ð;��s 5™zœz¨A™~Ñ.Ô.Ð;Ñ;Ô;€FÝ
ˆ6�{„{�aÒÐØ�Š�a˜ÑÔÐØ˜qœ	Ð)ˆfˆf°°f¸Q´iÔ0@À'È&ÐQRÌ)ÔBTÐ/U�G€BˆØ˜f QœiÔ(¨'°&¸´)Ô*<¸eÀB¼i�ˆ3€EØˆu�e˜C”j %¨¨c°7¸3´<ÀcÈUÂlÀlÐW\Ð_`ÑW`ÐdeÒWeÐWeÀÀÐklÐlÐlrD   c                 óº  — |\  }}}}t          j        ||	¦  «        }t          dt          t          j        ||¦  «        ||z  ||z  z  ¦  «        ¦  «        }t          j        ||¦  «        }t	          j        ||||ft          j        | j        ¬¦  «        }t          |||f         | ||||||||||	|||
|¬¦  «         t          ||f         ||||||d¬¦  «         d S )Nr   ©r   Údevice)Ú	num_warps)
ÚtritonÚcdivÚmaxÚminrZ   Úemptyr,   rw   rC   rW   )ÚsourceÚoutr4   rq   Útokensrs   rp   rr   ÚstÚhbrK   ÚkindÚconfigÚsmsÚblock_tÚwarpsÚstagesÚper_smÚgroupsÚspansr2   rT   s                         rB   Ú_launchrŒ   s   s	  € Ø%+Ñ"€GˆU�F˜FÝŒ[˜ Ñ#Ô#€FÝ�•3•v”{ 6¨7Ñ3Ô3°V¸c±\ÀfÈrÁkÑ5RÑSÔSÑTÔT€EÝŒ;�v˜uÑ%Ô%€DÝŒk˜2˜u e¨UÐ3½5¼=ÐQWÔQ^Ð_Ñ_Ô_€GÝ˜% ¨Ð,Ô-¨f°g¸vÀtÈUÐTYÐ[^Ð`cÐegÐinÐprÐt{ð  ~Dð  FJØ8=ð?ñ ?ô ?ð ?å�E˜2�;Ô ¨¨f°e¸UÀDÐTUÐVÑVÔVÐVÐVÐVrD   c                 ó‚  — | j         d         }t          d uo©|                      ¦   «         dv o“|                      ¦   «         |                      ¦   «         z  t
          k    oa|                      d¦  «        dk    oHd|cxk    odk    nc o8||dz
  z  dk    o,t          j        | t          j
        t          j        fd¬¦  «        S )	Nrj   )r
   r	   r   r   rb   r   )é   r   )ÚdtypesÚcapability_floor)rm   ry   rh   ÚnumelÚelement_sizeÚ	MIN_BYTESÚstrider   Úcuda_inferencerZ   Úbfloat16rP   )r?   rs   s     rB   Ú	_eligibler—   ~   s  € ØŒG�BŒK€EÝ˜$Ðð I 1§5¢5¡7¤7¨fÐ#4ð I¸¿º¹¼ÀQÇ^Â^ÑEUÔEUÑ9UÕYbÒ9bð IÐgh×goÒgoÐprÑgsÔgsÐwxÒgxð IÐ}ð  DIð  ~Pð  ~Pò  ~Pð  ~Pð  MPò  ~Pð  ~Pð  ~Pð  ~Pð IØ˜ ™Ñ# qÒ(ðIÝ-5Ô-DÀQÕPUÔP^Õ`eÔ`mÐOnð  BHð  .Iñ  .Iô  .IðJrD   Fc                 óÈ  ‡ ‡‡‡‡	‡
‡‡‡‡‡‡‡‡‡‡‡— t          ‰ ¦  «        s;|r"t          j        ‰ ‰dt          j        ¬¦  «        n‰                      ‰d¬¦  «        S t	          ‰ j        ‰                      ¦   «         ‰¦  «        \  ŠŠ
ŠŠŠŠŠŠ	t          ‰ j        ¦  «        }d|‰<   t          j        ||rt          j        n‰ j	        ‰ j
        ¬¦  «        }‰ j	        t          j        k    o$t          j        ‰ j
        t          j        ¦  «         Š‰r‰                      t          j        ¦  «        n‰ Š|rdn‰ j	        t          j        k    rdndŠ‰dk    r|                     t          j        ¦  «        n|Št#          ‰ j
        j        ¦  «        \  ŠŠ‰                      ¦   «         ‰                      ¦   «         z  ‰k    Šˆˆˆˆ	ˆ
ˆˆˆˆˆˆˆˆˆˆˆˆ fd„}t*                               d	‰› d
‰	› d
‰rdnd› �‰ j
        t.          |‰                      ¦   «         ‰                      ¦   «         z  t0          k    ¦  «        }t3          ‰‰‰‰
‰‰‰‰‰‰	‰‰t5          |¦  «        ‰¦  «         |S )zåx.mean(dim, keepdim=True) of 16-bit keys or values (3 or 4 dimensions, the head dimension last and contiguous) rounded to x's
    dtype, or unrounded in FP32 with fp32=True (torch.mean(x, dim, keepdim=True, dtype=torch.float32)).T)rh   Úkeepdimr   )rh   r™   r   rv   r   r   c                  óZ  •‡ ‡‡— ‰
r‰ngt          ‰t          j        t          t          ‰	z  t
          ¦  «        ‰                     ¦   «         ‰                     ¦   «         z  ‰z  ¦  «        ¦  «        Š‰                     ‰d‰¦  «        Š t          j
        ‰¦  «        Šˆˆˆˆˆˆ ˆˆˆˆˆˆˆfd„S )Nr   c                 ó<   •— t          ‰‰
‰‰‰‰‰‰	‰‰‰‰| ‰¦  «        S rY   )rŒ   )r„   r4   rK   r‚   rq   rƒ   Úpartr>   rp   rr   Úsampler…   r�   rs   s    €€€€€€€€€€€€€rB   ú<lambda>z-token_mean.<locals>.scratch.<locals>.<lambda>˜   s0   ø€ �g d¨F°B¸¸tÀUÈCÐQTÐVXÐZ\Ð^bÐdhÐjpÐruÑvÔv€ rD   )r|   ry   rz   r{   Ú_TUNE_L2_MULTIPLEÚ_TUNE_MIN_SLICEr‘   r’   ÚnarrowrZ   Ú
empty_like)rœ   r>   r�   r4   rK   rh   r‚   rq   rƒ   Úl2Úresidentrp   rr   r…   r~   r�   Útargetr€   rs   r?   s   @@@€€€€€€€€€€€€€€€€€rB   Úscratchztoken_mean.<locals>.scratch”   sä   øøøø€ Ø!ð  Vˆvˆv¥s¨6µ6´;½sÕCTÐWYÑCYÕ[jÑ?kÔ?kÐmn×mtÒmtÑmvÔmvÐyz÷  zHò  zHñ  zJô  zJñ  nJð  NTñ  nTñ  4Uô  4Uñ  (Vô  (VˆØ�}Š}˜S ! TÑ*Ô*ˆÝÔ! &Ñ)Ô)ˆØvÐvÐvÐvÐvÐvÐvÐvÐvÐvÐvÐvÐvÐvÐvÐvÐvrD   Úpartial_sumsú|r£   Údram)r—   rZ   rU   r'   rt   rm   r”   Úlistr}   r   rw   r–   r   Útriton_dtype_supportedÚviewÚint16rP   ra   r_   r‘   r’   Ú_TUNERr„   Ú_CANDIDATESÚ_TUNE_MIN_BYTESrŒ   Útuple)r?   rh   Úfp32rm   r   r¦   r„   r4   rK   r‚   rq   rƒ   r£   r¤   rp   rr   r…   r~   r�   r¥   r€   rs   s   ``     @@@@@@@@@@@@@@@rB   Ú
token_meanr³   „   sž  øøøøøøøøøøøøøøøøø€ õ �Q‰<Œ<ð tØLPÐs�uŒz˜! ¨d½%¼-ÐHÑHÔHÐHÐVW×V\ÒV\ÐadÐnrÐV\ÑVsÔVsÐsÝ16°q´wÀÇÂÁ
Ä
ÈCÑ1PÔ1PÑ.€Bˆˆv�u˜c 3¨¨BÝ�”‰MŒM€EØ€Eˆ#�JÝ
Œ+�e°DÐ#E¥5¤= =¸a¼gÈaÌhÐ
WÑ
WÔ
W€CØŒ7•e”nÒ$Ðf­XÔ-LÈQÌXÕW\ÔWeÑ-fÔ-fÐ)f€DØ$(Ð/ˆQ�VŠV•E”KÑ Ô Ð ¨a€FØÐ>ˆ1ˆ1˜qœw­%¬-Ò7Ð7˜!˜!¸Q€DØ&*¨a¢i iˆS�XŠX•e”kÑ"Ô"Ð"°S€FÝ�a”h”nÑ%Ô%�G€CˆØ�wŠw‰yŒy˜1Ÿ>š>Ñ+Ô+Ñ+¨rÒ1€Hðwð wð wð wð wð wð wð wð wð wð wð wð wð wð wð wð wð wð wð wð wõ �]Š]˜>¨eÐ+YÐ+Y°bÐ+YÐ+YÀ8Ð;W¸4¸4ÐQWÐ+YÐ+YÐ[\Ô[cÕepÐryØŸ7š7™9œ9 q§~¢~Ñ'7Ô'7Ñ7½?ÒJñLô L€FåˆF�F˜B  v¨u°c¸3ÀÀBÈÈdÕTYÐZ`ÑTaÔTaÐcfÑgÔgÐgØ€JrD   )Ú!install_triton_compilation_logger)F)#Ú__doc__Ú	functoolsrZ   Ú r   ry   Útriton.languageÚlanguager    Ú	bf16_bitsr   r   Útriton_tunerr   ÚImportErrorr“   r¯   rŸ   r    r°   r®   ÚjitÚ	constexprrC   rL   rW   Úcachera   Ú	lru_cachert   rŒ   r—   r³   Ú%shared.kernels.triton_compilation_logr´   rf   rD   rB   ú<module>rÂ      s„  ððuð uð" Ð Ð Ð à €€€à Ð Ð Ð Ð Ð ðØ€M€M€MØ Ð Ð Ð Ð Ð Ø?Ð?Ð?Ð?Ð?Ð?Ð?Ð?Ø)Ð)Ð)Ð)Ð)Ð)Ð)øØð ð ð Ø€F€F€Fðøøøð €	ð E€Ø%1Ñ "Ð �?à€à	ÐØˆ[Ð,¨a°Ñ=Ô=€Fð €V„ZÐ"FÐ"FÐ"FÐGÑGÔGðDÐVXÔVbð DÐhjÔhtð Dð  @Bô  @Lð DØ%'¤\ðDØ9;¼ðDð Dð Dñ HÔGðDð. „Zð1ð 1ñ „Zð1ð €V„Z 8¨WÐ"5Ð6Ñ6Ô6ð+°r´|ð +È2Ì<ð +ð +ð +ñ 7Ô6ð+ð$ „ðFð Fñ „ðFð
 €Ô˜SÐ!Ñ!Ô!ðmð mñ "Ô!ðmðWð Wð WðJð Jð Jðð ð ð ð8 
ÐØWÐWÐWÐWÐWÐWà%Ð%Ñ'Ô'Ð'Ð'Ð'ð Ðs   ’+ «5´5