§
     ’Þj�'  ã                   óÌ  — d Z ddlZddlmZ ddlmZmZ ej        dej	        dej	        dej	        dej	        dej	        d	ej	        d
ej	        dej	        dej	        fd„¦   «         Z
ej        	 ddej	        dej	        dej	        dej	        dej	        dej	        dej	        dej	        dej	        fd„¦   «         Zej        dej	        dej	        dej	        dej	        fd„¦   «         ZdS )aM  
Copyright (c) 2024 by SageAttention team.

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.

The SageAttention (1.0.6) attention kernel of attn_qk_int8_block_varlen.py for GPUs whose Triton has no BF16 conversion instructions
(compute capability 7.5): the output is written as the bits of a BF16 tensor, rounded once from the FP32 accumulator as the kernel's
own BF16 store does on sm_80+, and with V_BITS v is read as the bits of a BF16 tensor and converted to FP16 for the P.V product. INT8
q/k inputs, block scales, tiling and arithmetic are those of the original kernel. shared.sage1_core launches it on such GPUs instead
of writing an FP16 output and copying it back to BF16, and for short sequences instead of copying v to FP16. MASK_MODE adds an
attention mask in SageAttention's (batch, heads, q, k) layout (strides, zero for broadcast dimensions): 1 a boolean mask (uint8
view) whose zero entries get no weight, 2 an additive FP32 mask as SDPA adds it to the scaled scores. RETURN_LSE also writes the base-2
log-sum-exp of each query's scores, to merge attentions over separate key sets (a cached prefix, shared.attention_kit).
é    N)Úbf16_bits_to_fp32Úfp32_to_bf16_bitsÚHÚBLOCK_MÚHEAD_DIMÚBLOCK_NÚSTAGEÚoffs_mÚoffs_nÚV_BITSÚ	MASK_MODEc           	      ój  — d|}}t          |||¦  «        D �]˜}t          j        ||¦  «        }|d d d …f         ||z
  k     }t          j        ||¬¦  «        }t          j        |¦  «        }t          j        ||¦  «                             t          j        ¦  «        |z  |z  }|dk    rPt          j        |
|d d …d f         |k     |z  d¬¦  «        }t          j        |dk    |t          d¦  «        ¦  «        }n4|dk    r.|t          j        |
|d d …d f         |k     |z  d¬¦  «        dz  z   }t          j	        |t          j
        |d¦  «        ¦  «        } || d d …d f         z
  }t          j                             |¦  «        }!t          j        |!d¦  «        }"t          j                             || z
  ¦  «        }#||#z  |"z   }| |#d d …d f         z  } |rRt          t          j        |	|d d …d f         ||z
  k     ¬¦  «        ¦  «                             t          j        ¦  «        }$n't          j        |	|d d …d f         ||z
  k     ¬¦  «        }$|!                     t          j        ¦  «        }!| t          j        |!|$t          j        ¬	¦  «        z  } | }|||z  z  }||z  }|	||z  z  }	|
||z  z  }
�Œš| ||fS )
Nr   ©Úmaské   ©r   Úotherz-infé   g        gþ‚+eG÷?)Ú	out_dtype)ÚrangeÚtlÚmultiple_ofÚloadÚdotÚtoÚfloat32ÚwhereÚfloatÚmaximumÚmaxÚmathÚexp2Úsumr   Úfloat16)%ÚaccÚl_iÚm_iÚqÚq_scaleÚkv_lenÚqo_lenÚK_ptrsÚK_scale_ptrÚV_ptrsÚM_ptrsÚ	stride_knÚ	stride_vnÚ	stride_mkÚstart_mr   r   r   r   r	   r
   r   r   r   ÚloÚhiÚstart_nÚk_maskÚkÚk_scaleÚqkÚkeepÚm_ijÚpÚl_ijÚalphaÚvs%                                        ú;C:\Users\johnn\OneDrive\Documents\AI\shared\sage1_turing.pyÚ_attn_fwd_innerrB       sº  € ð �ˆ€BÝ˜˜R Ñ)Ô)ð %&ñ %&ˆÝ”. ¨'Ñ2Ô2ˆØ˜˜a˜a˜a˜” F¨WÑ$4Ò5ˆÝŒG�F Ð(Ñ(Ô(ˆÝ”'˜+Ñ&Ô&ˆÝŒV�A�q‰\Œ\�_Š_�RœZÑ(Ô(¨7Ñ2°WÑ<ˆØ˜Š>ˆ>õ ”7˜6¨°°°°4°¬¸6Ò)AÀVÑ(KÐSTÐUÑUÔUˆDÝ”˜$ !š) R­¨v©¬Ñ7Ô7ˆBˆBØ˜!Š^ˆ^à•b”g˜f¨F°1°1°1°d°7¬O¸fÒ,DÈÑ+NÐVYÐZÑZÔZÐ]oÑoÑoˆBÝŒz˜#�rœv b¨!™}œ}Ñ-Ô-ˆØ�$�q�q�q˜$�w”-ÑˆÝŒG�LŠL˜ÑÔˆÝŒv�a˜‰|Œ|ˆå”—’˜S 4™ZÑ(Ô(ˆØ�E‰k˜DÑ ˆà�E˜!˜!˜!˜T˜'”NÑ"ˆàð 	Kõ "¥"¤'¨&°v¸a¸a¸aÀ¸g´È&ÐSZÑJZÒ7[Ð"\Ñ"\Ô"\Ñ]Ô]×`Ò`ÕacÔakÑlÔlˆAˆAå”˜ V¨A¨A¨A¨t¨G¤_¸ÀÑ8HÒ%IÐJÑJÔJˆAØ�DŠD•”ÑÔˆà�rŒv�a˜¥b¤jÐ1Ñ1Ô1Ñ1ˆØˆØ�'˜IÑ%Ñ%ˆØ�qÑˆØ�'˜IÑ%Ñ%ˆØ�'˜IÑ%Ñ%ˆ‰Ø��Sˆ=Ðó    FÚnum_kv_groupsÚ
RETURN_LSEc!                 ó®  — t          j        d¦  «        }!t          j        d¦  «                             t           j        ¦  «        }"t          j        d¦  «                             t           j        ¦  «        }#t          j        ||"z   ¦  «        }$t          j        ||"z   dz   ¦  «        }%|%|$z
  }&|!|z  |&k    rd S t          j        ||"z   ¦  «        }'t          j        ||"z   ¦  «        }(|'|z  |#z   |!|z  z   })|(||z  z  |#|z  z   }*t          j        ||"z   ¦  «        }+t          j        ||"z   dz   ¦  «        },|,|+z
  }-|!|z  t          j        d|¦  «        z   }.t          j        d|¦  «        }/t          j        d|¦  «        }0| |$|z  |#|z  z   z   |.d d …d f         |z  z   |0d d d …f         z   }1||)z   }2||+|z  |#|z  |z  z   z   |/d d d …f         |z  z   |0d d …d f         z   }3||*z   }4||+|z  |#|z  |z  z   z   |/d d …d f         |z  z   |0d d d …f         z   }5|	|$|z  |#|z  z   z   |.d d …d f         |z  z   |0d d d …f         z   }6|
|"|z  z   |#|z  z   |.d d …d f         |z  z   |/d d d …f         |z  z   }7t          j        |gt           j        ¬¦  «        t          d¦  «        z
  }8t          j        |gt           j        ¬¦  «        dz   }9t          j        ||gt           j        ¬¦  «        }:t          j        |1|.d d …d f         |&k     ¬¦  «        };t          j        |2¦  «        }<t          |:|9|8|;|<|-|&|3|4|5|7||||!||z  |||d|z
  |.|/||¦  «        \  }:}9}8|:|9d d …d f         z  }:t          j
        |6t          |:¦  «        |.d d …d f         |&k     ¬¦  «         | rLt          j
        ||"|z  |#z   |&z  z   |.z   |8t           j                             |9¦  «        z   |.|&k     ¬¦  «         d S d S )	Nr   r   r   )ÚdtypeÚinfg      ð?r   é   )r   Ú
program_idr   Úint64r   ÚarangeÚzerosr   r   rB   Ústorer   r!   Úlog2)=ÚQÚKÚVÚcu_seqlens_qÚcu_seqlens_kÚQ_scaleÚK_scaleÚcu_seqlens_q_scaleÚcu_seqlens_k_scaleÚOutÚMaskÚLseÚ	stride_qhÚ	stride_qnÚ	stride_khr0   Ú	stride_vhr1   Ú	stride_ohÚ	stride_onÚ	stride_mzÚ	stride_mhÚ	stride_mqr2   r   rD   r   r   r   r	   r   r   rE   r3   Úoff_zÚoff_hÚcu_seqlens_q_startÚcu_seqlens_q_endr+   Úcu_seq_lens_q_scale_startÚcu_seq_lens_k_scale_startÚq_scale_offsetÚk_scale_offsetÚcu_seqlens_k_startÚcu_seqlens_k_endr*   r
   r   Úoffs_kÚQ_ptrsÚQ_scale_ptrr,   r-   r.   ÚO_block_ptrr/   r'   r&   r%   r(   r)   s=                                                                rA   Ú	_attn_fwdrs   R   së  € õ& Œm˜AÑÔ€GåŒM˜!ÑÔ×Ò¥¤Ñ)Ô)€EÝŒM˜!ÑÔ×Ò¥¤Ñ)Ô)€Eåœ °Ñ!5Ñ6Ô6ÐÝ”w˜|¨eÑ3°aÑ7Ñ8Ô8ÐàÐ 2Ñ2€Fà�'Ñ˜fÒ$Ð$Øˆå "¤Ð(:¸UÑ(BÑ CÔ CÐÝ "¤Ð(:¸UÑ(BÑ CÔ CÐà.°Ñ2°UÑ:¸WÀq¹[ÑH€NØ.°!°}Ñ2DÑEÈÐQ^ÑH^Ñ^€Nåœ °Ñ!5Ñ6Ô6ÐÝ”w˜|¨eÑ3°aÑ7Ñ8Ô8ÐàÐ 2Ñ2€Fà�wÑ¥¤¨1¨gÑ!6Ô!6Ñ6€FÝŒY�q˜'Ñ"Ô"€FÝŒY�q˜(Ñ#Ô#€FØÐ$ yÑ0°5¸9Ñ3DÑDÑEÈÈqÈqÈqÐRVÈwÌÐZcÑHcÑcÐflÐmqÐstÐstÐstÐmtÔfuÑu€FØ˜NÑ*€KØÐ$ yÑ0°E¸]Ñ4JÈiÑ3WÑWÑXÐ[aÐbfÐhiÐhiÐhiÐbiÔ[jÐmvÑ[vÑvÐyð  ABð  ABð  ABð  DHð  AHô  zIñ  I€FØ˜NÑ*€KØÐ$ yÑ0°E¸]Ñ4JÈiÑ3WÑWÑXÐ[aÐbcÐbcÐbcÐeiÐbiÔ[jÐmvÑ[vÑvÐyð  AEð  GHð  GHð  GHð  AHô  zIñ  I€FØÐ+¨iÑ7¸%À)Ñ:KÑKÑLÈvÐVWÐVWÐVWÐY]ÐV]ÌÐajÑOjÑjÐmsÐtxÐz{Ðz{Ðz{Ðt{Ôm|Ñ|€Kà�E˜IÑ%Ñ%¨°	Ñ(9Ñ9¸FÀ1À1À1ÀdÀ7¼OÈiÑ<WÑWÐZ`ÐaeÐghÐghÐghÐahÔZiÐluÑZuÑu€Få
Œ(�G�9¥B¤JÐ
/Ñ
/Ô
/µ%¸±,´,Ñ
>€CÝ
Œ(�G�9¥B¤JÐ
/Ñ
/Ô
/°#Ñ
5€CÝ
Œ(�G˜XÐ&­b¬jÐ
9Ñ
9Ô
9€Cå
Œ�˜V A A A t Gœ_¨vÒ5Ð6Ñ6Ô6€AÝŒg�kÑ"Ô"€GÝ# C¨¨c°1°g¸vÀvÈvÐWbÐdjÐlrÐt}ð  @Ið  KTØ$+Ø$%¨Ñ$6Ø$+¨X°wØ$%¨¡I¨v°v¸vÀyñ	&ô &�M€Cˆˆcð ��A�A�A�t�G”Ñ
€Cå„Hˆ[Õ+¨CÑ0Ô0¸¸q¸q¸qÀ$¸w¼È&Ò8PÐRÑRÔRÐRØð mÝ
Œ�˜ ™	 EÑ)¨VÑ3Ñ3°fÑ<¸cÅBÄGÇLÂLÐQTÑDUÔDUÑ>UÐ\bÐekÒ\kÐlÑlÔlÐlÐlÐlðmð mrC   ÚCÚBLKÚHAS_MEANc                 óp  — t          j        d¦  «        }t          j        d¦  «        }t          j        d¦  «        }t          j        ||z   ¦  «        }t          j        ||z   dz   ¦  «        }||z
  }||z  |k    rdS t          j        ||z   ¦  «        }||z  t          j        d|¦  «        z   }t          j        d|¦  «        }| ||z  z   ||z  z   |dd…df         |z  z   |ddd…f         z   }|||	z  z   ||z  z   |dd…df         |	z  z   |ddd…f         z   }|||z  z   |z   ||z  z   }t	          t          j        ||dd…df         |k     d¬¦  «        ¦  «        }|r*|t          j        |||z  z   |z   ¦  «        ddd…f         z
  }||
z  }t          j        t          j        |¦  «        ¦  «        dz  }||z  }|dt          j        |dk    dd¦  «        z  z  }|                     t           j	        ¦  «        }t          j
        |||dd…df         |k     ¬	¦  «         t          j
        ||¦  «         dS )
a%  quant_per_block_int8_kernel of quant_per_block_varlen.py reading a BF16 input as bits (int16 pointer): no FP16 copy on GPUs without
    BF16 conversions. With HAS_MEAN, Mean (heads, C) in FP32 is subtracted first (SageAttention's key smoothing), so k itself is neither
    copied nor modified.r   r   r   Nr   g     À_@g      à?éÿÿÿÿr   )r   rJ   r   rL   r   r    Úabsr   r   Úint8rN   )ÚInputÚOutputÚScaleÚMeanÚcu_seqlens_inputÚcu_seqlens_scaleÚ	stride_ihÚ	stride_inr`   ra   Úsm_scaler   rt   ru   rv   Úoff_blkrf   Úoff_bÚcu_seqlens_input_startÚcu_seqlens_input_endÚLÚcu_seqlens_scale_startr   ro   Ú
input_ptrsÚoutput_ptrsÚ
scale_ptrsÚxÚscaleÚx_int8s                                 rA   Ú quant_per_block_int8_bits_kernelr�   œ   s�  € õ Œm˜AÑÔ€GÝŒM˜!ÑÔ€EÝŒM˜!ÑÔ€EåœWÐ%5¸Ñ%=Ñ>Ô>ÐÝœ7Ð#3°eÑ#;¸aÑ#?Ñ@Ô@ÐàÐ5Ñ5€Aà�#‰˜!ÒÐØˆåœWÐ%5¸Ñ%=Ñ>Ô>Ðà�s‰]�RœY q¨#Ñ.Ô.Ñ.€FÝŒY�q˜!‰_Œ_€FàÐ/°)Ñ;Ñ;¸eÀiÑ>OÑOÐRXÐYZÐYZÐYZÐ\`ÐY`ÔRaÐdmÑRmÑmÐpvÐw{Ð}~Ð}~Ð}~Ðw~ÔpÑ€JØÐ1°IÑ=Ñ=ÀÈ	Ñ@QÑQÐTZÐ[\Ð[\Ð[\Ð^bÐ[bÔTcÐfoÑToÑoÐrxÐy}ð  @Að  @Að  @Að  zAô  sBñ  B€KØÐ/°!Ñ3Ñ3°eÑ;¸gÈ¹kÑI€Jå�"œ' *°6¸!¸!¸!¸T¸'´?ÀQÒ3FÈaÐPÑPÔPÑQÔQ€AØð <Ø•”˜˜u q™yÑ(¨6Ñ1Ñ2Ô2°4¸¸¸°7Ô;Ñ;ˆØˆ�M€AÝŒF•2”6˜!‘9”9ÑÔ Ñ$€EØ�‰Y€FØ
ˆc•B”H˜V qš[¨!¨RÑ0Ô0Ñ0Ñ0€FØ�YŠY•r”wÑÔ€FÝ„Hˆ[˜& v¨a¨a¨a°¨g¤¸Ò':Ð;Ñ;Ô;Ð;Ý„HˆZ˜ÑÔÐÐÐrC   )F)Ú__doc__ÚtritonÚtriton.languageÚlanguager   Úshared.kernels.bf16_bitsr   r   ÚjitÚ	constexprrB   rs   r�   © rC   rA   ú<module>r™      sõ  ððð ð0 €€€Ø Ð Ð Ð Ð Ð à IÐ IÐ IÐ IÐ IÐ IÐ IÐ Ið „ð.ð ”|ð.ð  œ\ð	.ð 68´\ð	.ð MOÌLð	.ð
 œ<ð.ð
 24´ð.ð
 HJÄ|ð.ð
 ^`Ô]ið.ð
 wyô  wCð.ð .ð .ñ „ð.ðb „ð" */ð!Fmð Fmð ”ðFmð /1¬lðFmð œðFmð ”|ðFmð ”|ðFmð ”\ðFmð ”lðFmð œðFmð  œ,ð!Fmð Fmð Fmñ „ðFmðR „ð( ð
 )+¬ð( ð )+¬ð( ð <>¼<ð( ð TVÔS_ð( ð ( ð ( ñ „ð( ð ( ð ( rC   