§
     ’Þjî  ã                   óv  — d Z ddlZddlZddlmZ ddlmZ ddlm	Z	m
Z
 ddlmZ ddlmZ  edd	d
¦  «        ZdZg d¢Z ej        e¬¦  «        dej        dej        dej        fd„¦   «         Z ej        g d¢¬¦  «        dej        dej        dej        fd„¦   «         Zd„ Zd„ Zd„ ZddlmZ  e¦   «          dS )a  Kernels of the attention kit's cached prefixes (shared.attention_kit): q8 keys / values (int8 per token, head and 32 channels with
FP16 scales, the LM engines' q8 format) dequantized through strides into a 16-bit buffer (a head group of a prefix into the joined keys
or values, a prefix of one batch item broadcast to all), and two attentions of the same queries over separate key sets merged in place by
their base-2 log-sum-exp. The dequantization rounds the FP32 product once, as the PyTorch path; the merge computes in FP32 (precise
division) and rounds once. BITS: BF16 tensors passed as int16 views and converted by the bit helpers of bf16_bits.py, for GPUs whose
Triton has no BF16 conversion instructions (compute capability 7.5); same roundings (to nearest even, as PyTorch's casts).é    N)Ú	libdeviceé   )Úbf16_bits_to_fp32Úfp32_to_bf16_bits)Údevice_supported)ÚLaunchTunerÚwan2gp_kv_prefixé   Ú	Attentioni   )
ÚTOKENSÚDBÚDTÚDHÚSBÚSTÚSHÚOBÚOTÚOH)Údo_not_specializeÚDÚBLOCK_TÚBITSc                 óž  — t          j        d¦  «                             t           j        ¦  «        |z  t          j        d|¦  «        d d …d f         z   }t          j        d¦  «                             t           j        ¦  «        t          j        d¦  «                             t           j        ¦  «        }}t          j        d|¦  «        d d d …f         }||k     }t          j        | ||z  z   ||z  z   ||z  z   |z   |d¦  «                             t           j        ¦  «        }t          j        |||z  z   ||z  z   ||	z  z   |dz  z   |d¦  «                             t           j        ¦  «        }||z  }t          j        |||
z  z   ||z  z   ||z  z   |z   |rt          |¦  «        n|                     |j	        j
        ¦  «        |¦  «         d S )Nr   r   r
   é    ç        )ÚtlÚ
program_idÚtoÚint64ÚarangeÚloadÚfloat32Ústorer   ÚdtypeÚ
element_ty)ÚDATAÚSCALESÚOUTr   r   r   r   r   r   r   r   r   r   r   r   r   ÚtokensÚheadÚbatchÚcÚvalidÚdataÚscaleÚvalues                           úGC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\kv_prefix_triton.pyÚ_dequantize_q8_kernelr3      sÈ  € õ Œ]˜1ÑÔ× Ò ¥¤Ñ*Ô*¨WÑ4µr´yÀÀGÑ7LÔ7LÈQÈQÈQÐPTÈWÔ7UÑU€FÝ”- Ñ"Ô"×%Ò%¥b¤hÑ/Ô/µ´¸qÑ1AÔ1A×1DÒ1DÅRÄXÑ1NÔ1Nˆ%€DÝ
Œ	�!�Q‰Œ˜˜a˜a˜a˜Ô €AØ�VŠO€EÝŒ7�4˜% "™*Ñ$ v°¡{Ñ2°T¸B±YÑ>ÀÑBÀEÈ1ÑMÔM×PÒPÕQSÔQ[Ñ\Ô\€DÝŒG�F˜U R™ZÑ'¨&°2©+Ñ5¸¸r¹	ÑAÀAÈÁGÑKÈUÐTWÑXÔX×[Ò[Õ\^Ô\fÑgÔg€EØ�5‰L€EÝ„HˆS�5˜2‘:Ñ ¨¡Ñ+¨d°R©iÑ7¸!Ñ;ÐY]ð  >BÕ=NÈuÑ=UÔ=UÐ=UÐch×ckÒckÐloÔluô  mAñ  dBô  dBð  DIñ  Jô  Jð  Jð  Jð  Jó    )	r   r   r   r   ÚPBÚPTÚPHÚLBÚLHc                 óž  — t          j        d¦  «                             t           j        ¦  «        |z  t          j        d|¦  «        z   }t          j        d¦  «                             t           j        ¦  «        t          j        d¦  «                             t           j        ¦  «        }}||k     }t          j        |||z  z   ||z  z   |z   |d¦  «        }t          j        |||z  z   ||z  z   |z   |d¦  «        }t          j        ||¦  «        }t          j        ||z
  ¦  «        t          j        ||z
  ¦  «        }}||z   }t          j	        ||¦  «        d d …d f         t          j	        ||¦  «        d d …d f         }}t          j        d|¦  «        d d d …f         |d d …d f         }}| ||z  z   ||z  z   ||z  z   |z   }t          j        ||d d …d f         d¦  «        }t          j        |||z  z   ||	z  z   ||
z  z   |z   |d d …d f         d¦  «        }|rt          |¦  «        t          |¦  «        }}|                     t           j        ¦  «        |z  |                     t           j        ¦  «        |z  z   }t          j        ||rt          |¦  «        n|                     | j        j        ¦  «        |d d …d f         ¦  «         d S )Nr   r   r
   r   )r   r   r   r    r!   r"   Úmaximumr   Úexp2Údiv_rnr   r#   r$   r   r%   r&   ) ÚOÚPÚLOÚLPr   r   r   r   r5   r6   r7   r8   r9   r   r   r   r*   r+   r,   r.   ÚloÚlpÚtopÚwoÚwpÚtotalr-   ÚrowsÚo_ptrsÚoÚpÚmergeds                                    r2   Ú_merge_lse_kernelrM   #   s™  € åŒ]˜1ÑÔ× Ò ¥¤Ñ*Ô*¨WÑ4µr´yÀÀGÑ7LÔ7LÑL€FÝ”- Ñ"Ô"×%Ò%¥b¤hÑ/Ô/µ´¸qÑ1AÔ1A×1DÒ1DÅRÄXÑ1NÔ1Nˆ%€DØ�VŠO€EÝ	Œ��e˜b‘j‘ 4¨"¡9Ñ,¨vÑ5°u¸cÑ	BÔ	B€BÝ	Œ��e˜b‘j‘ 4¨"¡9Ñ,¨vÑ5°u¸cÑ	BÔ	B€BÝ
Œ*�R˜Ñ
Ô
€CÝŒ^˜B ™HÑ%Ô%¥y¤~°b¸3±hÑ'?Ô'?ˆ€BØ�‰G€EÝŒY�r˜5Ñ!Ô! ! ! ! T 'Ô*­B¬I°b¸%Ñ,@Ô,@ÀÀÀÀDÀÔ,Iˆ€BÝŒi˜˜1‰oŒo˜d A A A˜gÔ&¨¨q¨q¨q°$¨w¬€t€AØ�˜‘‰^˜d R™iÑ'¨$°©)Ñ3°aÑ7€FÝ
Œ�˜˜a˜a˜a ˜gœ¨Ñ*Ô*€AÝ
Œ��E˜B‘J‘ ¨¡Ñ*¨T°B©YÑ6¸Ñ:¸EÀ!À!À!ÀTÀ'¼NÈAÑNÔN€AØð :Ý  Ñ#Ô#Õ%6°qÑ%9Ô%9ˆ1ˆØ�TŠT•"”*ÑÔ Ñ" Q§T¢T­"¬*Ñ%5Ô%5¸Ñ%:Ñ:€FÝ„HˆV°$ÐYÕ& vÑ.Ô.Ð.¸F¿IºIÀaÄgÔFXÑ<YÔ<YÐ[`ÐabÐabÐabÐdhÐahÔ[iÑjÔjÐjÐjÐjr4   c                 ól   — | j         t          j        k    ot          | j        t          j        ¦  «         S )z^A BF16 tensor on a GPU below the BF16 floor (Turing): passed to the kernels as its int16 view.)r%   ÚtorchÚbfloat16r   Údevice)Úxs    r2   Ú_bitsrS   8   s(   € àŒ7•e”nÒ$ÐWÕ-=¸a¼hÍÌÑ-WÔ-WÐ)WÐWr4   c           	      óp  ‡ ‡‡‡‡‡‡	‡
‡— ‰ j         \  Š}Š
Š	‰j         d         dk    o‰dk    Št          ‰ ¦  «        Šˆˆˆˆ	ˆ
fd„Šˆˆˆ	ˆ
ˆˆ ˆfd„}t                               d‰rdnd› �‰	‰ j        d||t
          k    ¦  «        } ‰‰ ‰‰|¦  «         t          j        j         	                    ‰ ¦  «         ‰ S )	z¾out (batch, tokens, heads, head_dim) 16-bit = data (batch or 1, tokens, heads, head_dim) int8 times scales (..., head_dim // 32)
    FP16, through their strides (last dimensions contiguous).r   r   c                 ó  •— |\  }}t          t          j        | j        d         |¦  «        ‰
‰f         ||‰r|                      t
          j        ¦  «        n| | j        d         ‰rdn|                     d¦  «        |                     d¦  «        |                     d¦  «        ‰rdn|                     d¦  «        |                     d¦  «        |                     d¦  «        g
|                      ¦   «         d d…         ¢‰	‘|‘‰‘R d|iŽ d S )Nr   r   r
   é   Ú	num_warps)r3   ÚtritonÚcdivÚshapeÚviewrO   Úint16Ústride)ÚtargetÚsourceÚsource_scalesÚconfigÚblockÚwarpsr,   ÚbitsÚ	broadcastÚdimÚheadss         €€€€€r2   Úlaunchz"dequantize_q8_into.<locals>.launchD   s–  ø€ Ø‰ˆˆuÝ�vœ{¨6¬<¸¬?¸EÑBÔBÀEÈ5ÐQÔRØ�M¸tÐ#O 6§;¢;­u¬{Ñ#;Ô#;Ð#;ÈÐQWÔQ]Ð^_ÔQ`Ðgpð  cGÐbcÐbcÐv|÷  wDò  wDð  EFñ  wGô  wGð  IO÷  IVò  IVð  WXñ  IYô  IYØ�MŠM˜!ÑÔ 9ÐI˜a˜a°-×2FÒ2FÀqÑ2IÔ2IÈ=×K_ÒK_Ð`aÑKbÔKbÐdq×dxÒdxÐyzÑd{Ôd{ð	#ð E÷  Lò  Lñ  Nô  Nð  OQð  PQð  OQô  Rð	#ð TWð	#ð Y^ð	#ð ð	#ð 	#ð 	#ð "ð	#ð 	#ð 	#ð 	#ð 	#r4   c                  ój   •‡ — t          j        ‰t          ‰‰f‰j        ‰j        ¬¦  «        Š ˆˆˆˆ fd„S )N)r%   rQ   c                 ód   •—  ‰‰‰d d …d t           …f         ‰d d …d t           …f         | ¦  «        S ©N)Ú_TUNE_TOKENS)ra   r/   rh   Úscalesr^   s    €€€€r2   ú<lambda>z5dequantize_q8_into.<locals>.scratch.<locals>.<lambda>M   s=   ø€ ˜f˜f V¨T°!°!°!°]µl°]Ð2BÔ-CÀVÈAÈAÈAÈ}ÕP\È}ÐL\ÔE]Ð_eÑfÔf€ r4   )rO   Úemptyrl   r%   rQ   )r^   r,   r/   rf   rg   rh   Úoutrm   s   @€€€€€€€r2   Úscratchz#dequantize_q8_into.<locals>.scratchK   sA   øø€ Ý”˜e¥\°5¸#Ð>ÀcÄiÐX[ÔXbÐcÑcÔcˆØfÐfÐfÐfÐfÐfÐfÐfr4   Údequantize_q8rS   Ú ©)r   é   )é   ru   )é@   ru   )rw   é   ©
rZ   rS   Ú_TUNERra   rQ   rl   rO   ÚautogradÚgraphÚincrement_version)rp   r/   rm   r*   rq   ra   r,   rd   re   rf   rg   rh   s   ```   @@@@@@r2   Údequantize_q8_intor~   =   s%  øøøøøøøøø€ ð !$¤	Ñ€Eˆ6�5˜#Ø”
˜1” Ò"Ð0 u¨q¢y€IÝ�‰:Œ:€Dð#ð #ð #ð #ð #ð #ð #ð #ð #ðgð gð gð gð gð gð gð gð gð gð gõ �]Š]ÐD°dÐ+B¨7¨7ÀÐDÐDÀcÈ3Ì:ÐW{ð  ~EØ!¥\Ò1ñ3ô 3€Fà
€Fˆ3��f˜fÑ%Ô%Ð%Ý	„NÔ×*Ò*¨3Ñ/Ô/Ð/Ø€Jr4   c           	      ó>  ‡ ‡‡‡‡‡‡	‡
‡— ‰ j         \  Š}Š
Š	t          ‰ ¦  «        Šˆˆˆ	ˆ
fd„Šˆˆˆˆˆ fd„}t                               d‰rdnd› �‰	‰ j        d||t
          k    ¦  «        } ‰‰ ‰‰‰|¦  «         t          j        j         	                    ‰ ¦  «         ‰ S )zýIn place: out (batch, tokens, heads, head_dim, through its strides) = the attention over both key sets of the queries whose
    attention over one set is out with lse (batch, heads, tokens) base-2 log-sum-exp, and over the other is other with other_lse.c                 óú  •— |\  }}‰r>|                       t          j        ¦  «        |                      t          j        ¦  «        fn| |f\  } }t          t	          j        | j        d         |¦  «        ‰
‰f         | |||| j        d         g|                      ¦   «         d d…         ¢|                     ¦   «         d d…         ¢|                     d¦  «        ‘|                     d¦  «        ‘‰	‘|‘‰‘R |ddœŽ d S )Nr   rV   r   F)rW   Úenable_fp_fusion)r[   rO   r\   rM   rX   rY   rZ   r]   )r^   r_   Ú
target_lseÚ
source_lsera   rb   rc   r,   rd   rf   rg   s          €€€€r2   rh   zmerge_lse_.<locals>.launch\   sK  ø€ Ø‰ˆˆuØQUÐk˜&Ÿ+š+¥e¤kÑ2Ô2°F·K²KÅÄÑ4LÔ4LÐMÐMÐ\bÐdjÐ[k‰ˆ�Ý�6œ; v¤|°A¤¸Ñ>Ô>ÀÀuÐMÔNØ�F˜J¨
°F´LÀ´Oð	]ØFLÇmÂmÁoÄoÐVXÐWXÐVXÔFYð	]Ø\b×\iÒ\iÑ\kÔ\kÐlnÐmnÐlnÔ\oð	]Øq{÷  rCò  rCð  DEñ  rFô  rFð	]à×Ò˜aÑ Ô ð	]à"%ð	]à',ð	]à.2ð	]ð 	]à>CÐV[ð	]ð 	]ð 	]ð 	]ð 	]ð 	]r4   c                  óV  •‡ ‡‡‡— ‰d d …d t           …f                              ¦   «         ‰d d …d t           …f                              ¦   «         cŠŠ ‰d d …d d …d t           …f                              ¦   «         ‰d d …d d …d t           …f                              ¦   «         cŠŠˆˆ ˆˆˆfd„S )Nc                 ó"   •—  ‰‰‰‰‰| ¦  «        S rk   © )ra   rh   r_   rƒ   r^   r‚   s    €€€€€r2   rn   z-merge_lse_.<locals>.scratch.<locals>.<lambda>f   s   ø€ ˜f˜f V¨V°ZÀÈVÑTÔT€ r4   )rl   ÚcloneÚ
contiguous)	r_   rƒ   r^   r‚   rh   ÚlseÚotherÚ	other_lserp   s	   @@@@€€€€€r2   rq   zmerge_lse_.<locals>.scratchc   sÆ   øøøøø€ Ø˜Q˜Q˜Q ¥ Ð-Ô.×4Ò4Ñ6Ô6¸¸a¸a¸aÀÅ,ÀÐ>NÔ8O×8ZÒ8ZÑ8\Ô8\ˆˆ�Ø!$ Q Q Q¨¨¨¨=­L¨=Ð%8Ô!9×!DÒ!DÑ!FÔ!FÈ	ÐRSÐRSÐRSÐUVÐUVÐUVÐXeÕYeÐXeÐReÔHf×HqÒHqÑHsÔHsÐˆ
�JØTÐTÐTÐTÐTÐTÐTÐTÐTr4   Ú	merge_lserS   rs   rt   ry   )rp   r‰   rŠ   r‹   r*   rq   ra   r,   rd   rf   rg   rh   s   ````   @@@@@r2   Ú
merge_lse_r�   V   s  øøøøøøøøø€ ð !$¤	Ñ€Eˆ6�5˜#Ý�‰:Œ:€Dð]ð ]ð ]ð ]ð ]ð ]ð ]ð ]ðUð Uð Uð Uð Uð Uð Uð Uð Uõ
 �]Š]Ð@°$Ð'> w w¸BÐ@Ð@À#ÀsÄzÐSwð  zAð  CIõ  MYò  CYñ  Zô  Z€FØ
€Fˆ3��s˜I vÑ.Ô.Ð.Ý	„NÔ×*Ò*¨3Ñ/Ô/Ð/Ø€Jr4   )Ú!install_triton_compilation_logger)Ú__doc__rO   rX   Útriton.languageÚlanguager   Útriton.language.extrar   Ú	bf16_bitsr   r   Údispatchr   Útriton_tunerr   rz   rl   Ú_RUNTIMEÚjitÚ	constexprr3   rM   rS   r~   r�   Ú%shared.kernels.triton_compilation_logrŽ   r†   r4   r2   ú<module>rš      s¼  ðð~ð ~ð €€€Ø €€€Ø Ð Ð Ð Ð Ð Ø +Ð +Ð +Ð +Ð +Ð +à ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ø &Ð &Ð &Ð &Ð &Ð &Ø %Ð %Ð %Ð %Ð %Ð %à	ˆÐ'¨¨KÑ	8Ô	8€Ø€ØKÐKÐK€ð €„˜hÐ'Ñ'Ô'ð	JÐ\^Ô\hð 	JÐsuÔsð 	JØ "¤ð	Jð 	Jð 	Jñ (Ô'ð	Jð €„ÐXÐXÐXÐYÑYÔYðkÈrÌ|ð kÐfhÔfrð kÐz|ô  {Gð kð kð kñ ZÔYðkð(Xð Xð Xð
ð ð ð2ð ð ð0 TÐ SÐ SÐ SÐ SÐ Sà !Ð !Ñ #Ô #Ð #Ð #Ð #r4   