§
     ’Þj$  ã                   ól  — d Z ddlZddlmZ ddlmZ ddgZej        dej	        fd„¦   «         Z
 ej        ee¬¦  «        d	ej	        d
ej	        dej	        dej	        dej	        dej	        dej	        fd„¦   «         ZddlZddlmZ  eddd¦  «        Zg d¢ZdZdZd„ Zd„ Zd„ ZddlmZ  e¦   «          dS )a  Decoding attention over a fixed-capacity BF16/FP16 cache (the MTP draft cache) whose length is read on the GPU, for CUDA graphs on
GPUs without FlashAttention (Turing): each query row attends the keys up to its own position only, so the slots past the written ones
are never read and no copy of the cache is made. The 16-bit values are decoded from their bits and every product, sum and exponential
is FP32 (IEEE), without BF16 instructions. Split over the keys (flash decoding); the partial results are merged with PyTorch.é    N)Ú	libdeviceÚCOUNTÚSPLITSÚBF16c                 ó  — |rVt          j        | |d¬¦  «                             t           j        ¦  «        dz                       t           j        d¬¦  «        S t          j        | |d¬¦  «                             t           j        ¦  «        S )Nr   )ÚmaskÚotheré   T)Úbitcastç        )ÚtlÚloadÚtoÚuint32Úfloat32)Úptrsr   r   s      úMC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\static_cache_attention.pyÚ_load16r      st   € àð dÝ”˜ 4¨qÐ1Ñ1Ô1×4Ò4µR´YÑ?Ô?À2ÑE×IÒIÍ"Ì*Ð^bÐIÑcÔcÐcÝŒ7�4˜d¨#Ð.Ñ.Ô.×1Ò1µ"´*Ñ=Ô=Ð=ó    )Údo_not_specializeÚdo_not_specialize_on_alignmentÚGROUPSÚKV_HEADSÚHEAD_DIMÚROWS_PADÚBLOCK_RÚBLOCK_Nc                 óÎ  — t          j        d¦  «        t          j        d¦  «        t          j        d¦  «        }}}t          j        |¦  «        }||z   }t          j        t          j        ||¦  «        |¦  «        |z  }||z  }t          j        ||z   |¦  «        }||z  t          j        d|¦  «        z   }|||z  k     }|||z  z   }||z  ||z  z   }t          j        d|¦  «        }t          | ||z  d d …d f         |
z  z   |d d …d f         |z  z   |d d d …f         z   |d d …d f         |¦  «        } t          j        |ft          d¦  «        t           j	        ¦  «        }!t          j
        |ft           j	        ¦  «        }"t          j
        ||ft           j	        ¦  «        }#t          |||¦  «        D �]¾}$|$t          j        d|¦  «        z   }%|%|k     }&|%d d …d f         |z  ||z  z   |d d d …f         z   }'t          ||'z   |&d d …d f         |¦  «        }(t          j        | t          j        |(¦  «        d¬¦  «        |	z  })t          j        |&d d d …f         |%d d d …f         |d d …d f         k    z  |)t          d¦  «        ¦  «        })t          j        |!t          j        |)d¦  «        ¦  «        }*t          j        |*t          d¦  «        k    d|*¦  «        }+t#          j        |!|+z
  ¦  «        },t#          j        |)|+d d …d f         z
  ¦  «        }-t          ||'z   |&d d …d f         |¦  «        }.|"|,z  t          j        |-d¦  «        z   }"|#|,d d …d f         z  t          j        |-|.d¬¦  «        z   }#|*}!�ŒÀ||z  |z   |z  |z   }/t          j        ||/z   |!¦  «         t          j        ||/z   |"¦  «         t          j        ||/d d …d f         |z  z   |d d d …f         z   |#¦  «         d S )Nr   é   é   z-infÚieee)Úinput_precisionr   )r   Ú
program_idr   ÚcdivÚminimumÚaranger   ÚfullÚfloatr   ÚzerosÚrangeÚdotÚtransÚwhereÚmaximumÚmaxr   ÚexpÚsumÚstore)0ÚQÚKÚVÚSEQLENSÚPMÚPLÚPACCr   r   ÚscaleÚ	STRIDE_QCÚ	STRIDE_KNr   r   r   r   r   r   r   ÚkvÚsplitÚ	row_blockÚseqlenÚtotalÚspanÚstartÚstopÚrowsÚrow_okÚpositionÚheadÚdÚqÚmÚlÚaccÚfirstÚnÚkey_okÚoffsetsÚkÚsÚm_newÚm_safeÚalphaÚpÚvÚouts0                                                   r   Ú_static_decoderZ      sì  € õ œ=¨Ñ+Ô+­R¬]¸1Ñ-=Ô-=½r¼}ÈQÑ?OÔ?Oˆyˆ€BÝŒW�WÑÔ€FØ�U‰N€EÝŒ7•2”7˜5 &Ñ)Ô)¨7Ñ3Ô3°gÑ=€DØ�D‰L€EÝŒ:�e˜d‘l EÑ*Ô*€DØ�wÑ¥¤¨1¨gÑ!6Ô!6Ñ6€DØ�F˜U‘NÒ"€FØ˜˜u™Ñ$€HØ�‰;˜ ™Ñ&€DÝ
Œ	�!�XÑÔ€AÝ��T˜E‘\ 1 1 1 d 7Ô+¨iÑ7Ñ7¸$¸q¸q¸qÀ$¸w¼-È(Ñ:RÑRÐUVÐW[Ð]^Ð]^Ð]^ÐW^ÔU_Ñ_ÐagÐhiÐhiÐhiÐkoÐhoÔapÐrvÑwÔw€AÝ
Œ��
�E &™MœM­2¬:Ñ6Ô6€AÝ
Œ�'��RœZÑ(Ô(€AÝ
Œ(�G˜XÐ&­¬
Ñ
3Ô
3€CÝ�u˜d GÑ,Ô,ð ñ ˆØ•B”I˜a Ñ)Ô)Ñ)ˆØ�T’ˆØ�A�A�A�t�G”*˜yÑ(¨2°©=Ñ8¸1¸TÀ1À1À1¸W¼:ÑEˆÝ�A˜‘K ¨¨¨¨4¨¤°$Ñ7Ô7ˆÝŒF�1•b”h˜q‘k”k°6Ð:Ñ:Ô:¸UÑBˆÝŒH�V˜D ! ! !˜G”_¨¨$°°°¨'¬
°h¸q¸q¸qÀ$¸wÔ6GÒ(GÑHÈ!ÍUÐSYÉ]Ì]Ñ[Ô[ˆÝ”
˜1�bœf Q¨™lœlÑ+Ô+ˆÝ”˜%¥5¨¡=¤=Ò0°#°uÑ=Ô=ˆÝ”˜a &™jÑ)Ô)ˆÝŒM˜!˜f Q Q Q¨ WœoÑ-Ñ.Ô.ˆÝ�A˜‘K ¨¨¨¨4¨¤°$Ñ7Ô7ˆØ�‰I�œ˜q !™œÑ$ˆØ�E˜!˜!˜!˜T˜'”NÑ"¥R¤V¨A¨qÀ&Ð%IÑ%IÔ%IÑIˆØˆ‰Ø�8Ñ˜bÑ  HÑ
,¨tÑ
3€CÝ„HˆR�#‰X�qÑÔÐÝ„HˆR�#‰X�qÑÔÐÝ„HˆT�C˜˜˜˜4˜”L 8Ñ+Ñ+¨a°°a°a°a°¬jÑ8¸#Ñ>Ô>Ð>Ð>Ð>r   )ÚLaunchTunerÚwan2gp_static_cache_attentionr   zStatic Cache Attention))r
   é   r   )é    r]   r   )r
   é   r   )r^   r_   r   )r
   r]   r    )é@   r_   r   r
   r    c                 ó–   — t           | d         z   |z  dz  d| d         z  z   }|t          j                             |¦  «        j        k    S )Nr   r]   é€   )Ú_BLOCK_RÚtorchÚcudaÚget_device_propertiesÚshared_memory_per_block_optin)ÚconfigÚhead_dimÚdeviceÚshareds       r   Ú_fitsrl   F   sG   € õ ˜ œÑ" hÑ.°Ñ2°S¸6À!¼9±_ÑD€FØ•U”Z×5Ò5°fÑ=Ô=Ô[Ò[Ð[r   c                 óÜ   — | j         oe| j        t          j        t          j        fv oF|j        | j        k    o6|                     ¦   «         o"|                      ¦   «         o| j        d         dv S )Néÿÿÿÿ)r`   rb   é   )Úis_cudaÚdtyperd   Úbfloat16Úfloat16Úis_contiguousÚshape)ÚqueryÚcaches     r   Úusablerx   M   sˆ   € ØŒ=ð H˜Uœ[­U¬^½U¼]Ð,KÐKð HÐPUÔP[Ð_dÔ_jÒPjð HÐot÷  pCò  pCñ  pEô  pEð HØ×ÒÑ!Ô!ðHØ&+¤k°"¤o¸Ð&GðHr   c                 ó  ‡ ‡‡‡‡‡‡‡‡‡‡‡‡‡‡‡‡— ‰ j         \  }Š}Š‰j         d         Š|‰z  Š‰‰z  }t          j        |t          ¦  «        Š‰t          z  Št	          dt
          j                             ‰ j        ¦  «        j	        t          z  ‰‰z  z  ¦  «        Š‰ j        t
          j        k    Š‰rd„ nd„ } |‰ ¦  «         |‰¦  «         ||¦  «        cŠŠŠˆˆ fd„t          D ¦   «         }	ˆˆˆˆˆˆˆˆˆ ˆˆˆˆˆˆfd„Šˆˆˆ ˆˆfd„Šˆˆfd„}
t                               d	‰ j        › d
‰› d‰› d‰› �‰ j        |	|
d¦  «        } ‰¦   «         \  }}} ‰||||¦  «         |dd…dd…d|…f         |dd…dd…d|…f         |dd…dd…d|…f         }}}t          j        ||                     d¦  «        z
  ¦  «        }|                     d¦  «        |z                       d¦  «                             ||z                       d¦  «                             d¦  «        ¦  «        }|                     ‰‰‰‰¦  «                             dddd¦  «                             d‰|‰¦  «                             ‰ j        ¦  «        S )a  query (1, count, heads, head_dim) at positions seqlens + 0..count-1, whose keys and values are already written; key_cache and
    value_cache (1, capacity, kv_heads, head_dim), contiguous; seqlens: int32 (1,) on the GPU. Returns (1, count, heads, head_dim).r    r   c                 ó@   — |                       t          j        ¦  «        S ©N)Úviewrd   Úuint16©Úts    r   ú<lambda>zattention.<locals>.<lambda>]   s   € �a—f’f�Uœ\Ñ*Ô*€ r   c                 ó   — | S r{   © r~   s    r   r€   zattention.<locals>.<lambda>]   s   € ÀA€ r   c                 ó@   •— g | ]}t          |‰‰j        ¦  «        ¯|‘ŒS r‚   )rl   rj   )Ú.0rh   ri   rv   s     €€r   ú
<listcomp>zattention.<locals>.<listcomp>_   s,   ø€ ÐYÐYÐY˜Vµ5¸ÀÈ5Ì<Ñ3XÔ3XÐY�&ÐYÐYÐYr   c                 óÄ   •— |\  }}}t          ‰‰‰f         ‰‰‰‰| ||‰‰‰‰                     d¦  «        ‰                     d¦  «        ‰	‰‰
‰‰t          |||¬¦  «         d S )Nr   )	r   r   r   r   r   r   r   Ú	num_warpsÚ
num_stages)rZ   Ústriderc   )ÚpmÚplÚpaccrh   Úblock_nÚwarpsÚstagesÚbf16ÚcountÚgroupsri   rR   Ú	key_cacheÚkv_headsrJ   rv   Ú
row_blocksÚrows_padr:   ÚseqlensÚsplitsrX   s          €€€€€€€€€€€€€€€r   Úlaunchzattention.<locals>.launcha   s´   ø€ Ø!'Ñˆ�˜Ý˜ &¨*Ð5Ô6°q¸!¸QÀÈÈRÐQUÐW\Ð^dÐfkÐmr×myÒmyÐz{Ñm|Ôm|ð  H÷  Oò  Oð  PQñ  Rô  RØ>DÈxÐbjÐu}ð  EIÝ?GÐQXÐdiÐv|ð	~ñ 	~ô 	~ð 	~ð 	~ð 	~r   c                  óú   •— t          j        ‰‰‰ft           j        ‰j        ¬¦  «        t          j        ‰‰‰ft           j        ‰j        ¬¦  «        t          j        ‰‰‰‰ ft           j        ‰j        ¬¦  «        fS )N)rq   rj   )rd   Úemptyr   rj   )ri   r”   rv   r–   r˜   s   €€€€€r   Úpartialszattention.<locals>.partialsg   st   ø€ Ý”˜V X¨xÐ8ÅÄÐV[ÔVbÐcÑcÔcÝ”˜V X¨xÐ8ÅÄÐV[ÔVbÐcÑcÔcÝ”˜V X¨x¸ÐBÍ%Ì-Ð`eÔ`lÐmÑmÔmðoð 	or   c                  ó&   •‡ —  ‰¦   «         Š ˆˆ fd„S )Nc                 ó   •—  ‰g ‰¢| ‘R Ž S r{   r‚   )rh   r™   Úscratchs    €€r   r€   z4attention.<locals>.prepare_scratch.<locals>.<lambda>n   s   ø€ ˜f˜fÐ6 gÐ6¨vÐ6Ð6Ð6€ r   r‚   )rŸ   r™   rœ   s   @€€r   Úprepare_scratchz"attention.<locals>.prepare_scratchl   s!   øø€ Ø�(‘*”*ˆØ6Ð6Ð6Ð6Ð6Ð6r   Údecodez|Dz|Gz|KVTNr   rn   é   )ru   Útritonr$   rc   r/   rd   re   rf   rj   Úmulti_processor_countÚ_SPLITS_PER_SMrq   rr   Ú_CONFIGSÚ_TUNERrh   r0   ÚamaxÚ	unsqueezer1   Údiv_r|   ÚpermuteÚreshaper   )rv   r“   Úvalue_cacher—   r:   Ú_ÚheadsrE   ÚbitsÚ
candidatesr    rh   rŠ   r‹   rŒ   ÚweightsÚoutputr�   r‘   r’   ri   rR   r”   r™   rœ   rJ   r•   r–   r˜   rX   s   `` ``            @@@@@@@@@@@@@r   Ú	attentionr´   R   sn  øøøøøøøøøøøøøøøøø€ ð !&¤Ñ€A€uˆe�XØŒ˜qÔ!€HØ�hÑ€FØ�E‰>€DÝ”˜T¥8Ñ,Ô,€JØ�HÑ$€HÝ�•E”J×4Ò4°U´\ÑBÔBÔXÕ[iÑiÐnvð  zDñ  oDñ  Eñ  Fô  F€FØŒ;�%œ.Ò(€DØ/3ÐFÐ*Ð*Ð*¸+¸+€DØˆd�5‰kŒk˜4˜4 	™?œ?¨D¨D°Ñ,=Ô,=€G€A€qˆ!ØYÐYÐYÐYÐY¥xÐYÑYÔY€Jð~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ð ~ðoð oð oð oð oð oð oð oð oð
7ð 7ð 7ð 7ð 7ð 7õ �]Š]˜8¨¬Ð%XÐ%X°xÐ%XÐ%XÀ6Ð%XÐ%XÈhÐ%XÐ%XÐZ_ÔZfÐhrð  uDð  FJñ  Kô  K€FØ�8‘:”:�L€BˆˆDØ
€Fˆ2ˆr�4˜Ñ Ô Ð Ø�a�a�a˜˜˜˜E˜T˜E�k”? B q q q¨!¨!¨!¨U¨d¨U {¤O°T¸!¸!¸!¸Q¸Q¸QÀÀÀ¸+Ô5FˆDˆ€BÝŒi˜˜RŸWšW Q™ZœZ™Ñ(Ô(€GØ×Ò Ñ#Ô# dÑ*×/Ò/°Ñ2Ô2×7Ò7¸À2¹×8JÒ8JÈ1Ñ8MÔ8M×8WÒ8WÐXZÑ8[Ô8[Ñ\Ô\€FØ�;Š;�x ¨°Ñ9Ô9×AÒAÀ!ÀQÈÈ1ÑMÔM×UÒUÐVWÐY^Ð`eÐgoÑpÔp×sÒsÐtyÔtñ  Aô  Að  Ar   )Ú!install_triton_compilation_logger)Ú__doc__r£   Útriton.languageÚlanguager   Útriton.language.extrar   Ú_RUNTIMEÚjitÚ	constexprr   rZ   rd   Úshared.kernels.triton_tunerr[   r§   r¦   rc   r¥   rl   rx   r´   Ú%shared.kernels.triton_compilation_logrµ   r‚   r   r   ú<module>r¿      s§  ððAð Að €€€Ø Ð Ð Ð Ð Ð Ø +Ð +Ð +Ð +Ð +Ð +à�XÐ€ð „ð>˜bœlð >ð >ð >ñ „ð>ð €„˜hÀxÐPÑPÔPð#?ÐgiÔgsð #?ð  @Bô  @Lð #?Øœ\ð#?Ø57´\ð#?ØIKÌð#?Ø`bÔ`lð#?Øwyô  xDð#?ð #?ð #?ñ QÔPð#?ðN €€€Ø 3Ð 3Ð 3Ð 3Ð 3Ð 3à	ˆÐ4°aÐ9QÑ	RÔ	R€àSÐSÐS€Ø€Ø€ð\ð \ð \ðHð Hð Hð
$Að $Að $AðN TÐ SÐ SÐ SÐ SÐ SØ !Ð !Ñ #Ô #Ð #Ð #Ð #r   