§
     ’Þj/=  ã                   óÄ  — 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	 d\  Z
ZZZe
dededediZed	iZ e	d
dd¦  «        ZdZdZg d¢Zd„ 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        d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        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d#\  ZZ	 	 d1d&„Zd'„ Zd(„ Zd2d%dd*œd+„Zdd,d-d%dd.œd/„Zdd0l m!Z!  e!¦   «          dS )3a�  Staged SageAttention 2 quantization of q and k fused with their weighted RMSNorm and split-half RoPE: one pass over q from the
projection output to INT8 rows and scales, two passes over k (the partial sums of its normalized, rotated rows for the mean SageAttention
subtracts, then the quantization with that mean, the norm and rotation recomputed from the projection output). The norm reproduces the
in-place kernel of shared.kernels.denoiser_triton (FP32 statistics, the norm output rounded to the activation dtype), the rotation runs in
FP32, and the quantization reads the FP32 rotated values where the unfused path reads them rounded to the activation dtype. The
quantization arithmetic (block scales max|x| / 127, half-away-from-zero rounding, the approximate FP32 division) is that of
SageAttention 2's quantizers (Apache 2.0, https://github.com/thu-ml/SageAttention, sageattention/triton/quant_per_thread.py and
quant_per_block.py), so the quantization-only mode reproduces them bit for bit; the row groups of a scale follow the layouts
LAYOUT_PER_BLOCK (one scale per block of rows), LAYOUT_THREAD_Q (eight scales per 32 rows, rows r, r + 8, r + 16, r + 24),
LAYOUT_THREAD_K (four scales per 64 rows, row pairs) and LAYOUT_WARP_Q (four scales per 128 rows, 32 consecutive rows each).
The rotation pairs channel p with p + PAIRS (split halves) or, interleaved, 2p with 2p + 1; the cos / sin of pair p are read at the
table's pair stride, the sine at its last-dimension stride from the cosine, so a (..., pairs, 2) table and the 2 x 2 matrices
[[cos, -sin], [sin, cos]] of a (..., pairs, 2, 2) table both work as they are.

A program owns one block of rows of one head and walks it by sub-blocks of 8 rows, the live tile of the in-place RoPE kernel: the
maxima of the layout's row groups first, then the rows again from the cache for the quantization (the norm and rotation recomputed,
cheaper than holding or storing the block).é    N)Ú	libdeviceé   )ÚLaunchTuner)r   r   é   é   é   é   é€   Úwan2gp_qk_quantr   Ú	Attention))r   r   )r   r   )r   r   i   )ÚSEQÚXBÚXSÚXHÚRBÚRSÚOBÚOSÚOHÚSBÚSHÚMBÚMHÚPBÚPHÚPKc                 ón   — t                                ||¦  «        }t          j        | |¦  «        ||z  z  S )zzBlocks of rows the launch covers (their scales are all written): the tokens padded as SageAttention's quantizer pads them.)ÚPADDED_ROWSÚgetÚtritonÚcdiv)ÚseqÚlayoutÚblockÚpads       úFC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\qk_quant_triton.pyÚ_blocksr'   #   s0   € å
�/Š/˜& %Ñ
(Ô
(€CÝŒ;�s˜CÑ Ô  C¨5¡LÑ1Ð1ó    ÚRPÚRCÚDÚPAIRSÚNORMÚROTATEÚHAS_MEANÚSUBÚLANESÚINTERLEAVEDÚ	SEGMENTEDc                 óò
  — t          j        | |dd…df         z   |ddd…f         z   |dd…df         d¦  «                             t           j        ¦  «        }|rDt          j        ||z   |d¦  «                             t           j        ¦  «        dd…df         |z  }n!t          j        |dgt           j        ¦  «        }|�r¤|rä|dk    rÞt          j        t          j        t          j        ||ddf¦  «        ddd¦  «        |dddf¦  «        }t          j        |¦  «        \  }}t          j        |¦  «        \  } }!t          j        |¦  «        \  }"}#t          j	        t          j
        |#|#t          j
        |!|!t          j
        |"|"| | z  ¦  «        ¦  «        ¦  «        d¦  «        }$nt          j	        ||z  d¦  «        }$t          j        |$|z  |z   ¦  «        dd…df         }%||%z  t          j        ||z   |ddd…f         z   ¦  «                             t           j        ¦  «        z                       | j        j        ¦  «                             t           j        ¦  «        }&n|}&|�rŸ|d|z  k     }'|r(|dz  dk    }(t          j        |'|dz  |¦  «        })|dz  }*nO||k     }(t          j        |(||z   t          j        |'||z
  |¦  «        ¦  «        })t          j        |(|||z
  ¦  «        }*t          j        | |dd…df         z   |)ddd…f         z   |dd…df         d¦  «                             t           j        ¦  «        }+|r�|+|%z  t          j        ||z   |)ddd…f         z   ¦  «                             t           j        ¦  «        z                       | j        j        ¦  «                             t           j        ¦  «        }+|	|z  |dd…df         |z  z   |*ddd…f         |z  z   },|dd…df         |'ddd…f         z  }-t          j        ||,z   |-d¦  «                             t           j        ¦  «        }.t          j        ||,z   |z   |-d	¦  «                             t           j        ¦  «        }/|(ddd…f         }0t          j
        t          j        |0|.|/¦  «        t          j        |0|&|+¦  «        t          j        |0|/ |.¦  «        t          j        |0|+|&¦  «        z  ¦  «        }1t          j        |'ddd…f         |1|&¦  «        }&|rkt          j        |dd…df         |&t          j        ||	|z  z   |
|z  z   |z   ¦  «                             t           j        ¦  «        ddd…f         z
  d	¦  «        }&|&S )
a5  The rows at base (one head): weighted RMSNorm in FP32 rounded to the activation dtype (NORM; the weight row of each token from
    WROW with SEGMENTED), the rotation of the first 2 * PAIRS channels in FP32 (ROTATE, cos / sin pairs of ROPE), minus the head's MEAN
    row (HAS_MEAN); rows beyond SEQ are zeros.Nr   r   r
   r	   é    r   ç      ð?ç        )ÚtlÚloadÚtoÚfloat32Úint64ÚzerosÚreshapeÚpermuteÚsplitÚsumÚfmar   ÚrsqrtÚdtypeÚ
element_tyÚwhere)2ÚXÚWÚWROWÚROPEÚMEANÚbaser"   ÚokÚcÚbatchÚheadr   r   r   r   ÚEPSr)   r*   r+   r,   r-   r.   r/   r0   r1   r2   r3   ÚxÚwrowÚlanesÚevenÚoddÚx0Úx2Úx1Úx3ÚtotalÚinvÚyÚrotatedÚleadingÚpartnerÚpairÚotherÚangleÚmaskÚcosÚsinÚfirstÚouts2                                                     r&   Ú_rowsri   )   s<  € õ 	Œ��D˜˜˜˜D˜”MÑ! A d¨A¨A¨A g¤JÑ.°°1°1°1°d°7´¸QÑ?Ô?×BÒBÅ2Ä:ÑNÔN€AØð ,ÝŒw�t˜c‘z 2 qÑ)Ô)×,Ò,­R¬XÑ6Ô6°q°q°q¸$°wÔ?À!ÑCˆˆåŒx˜˜a˜¥"¤(Ñ+Ô+ˆØñ Øð 	%�Q˜#’X�XÝ”J�rœz­"¬*°Q¸¸aÀ¸Ñ*EÔ*EÀqÈ!ÈQÑOÔOÐRUÐWYÐ[\Ð^_ÐQ`ÑaÔaˆEÝœ ™œ‰IˆD�#Ý”X˜d‘^”^‰FˆB�Ý”X˜c‘]”]‰FˆB�Ý”F�2œ6 " b­"¬&°°R½¼ÀÀBÈÈRÉÑ9PÔ9PÑ*QÔ*QÑRÔRÐTUÑVÔVˆEˆEå”F˜1˜q™5 !Ñ$Ô$ˆEÝŒo˜e a™i¨#™oÑ.Ô.¨q¨q¨q°$¨wÔ7ˆØ�‰W•r”w˜q 4™x¨!¨D°!°!°!¨G¬*Ñ4Ñ5Ô5×8Ò8½¼ÑDÔDÑD×HÒHÈÌÔI[Ñ\Ô\×_Ò_Õ`bÔ`jÑkÔkˆˆàˆØñ /Ø�a˜%‘i’-ˆØð 	3Ø˜!‘e˜q’jˆGÝ”h˜w¨¨A©¨qÑ1Ô1ˆGØ˜‘6ˆDˆDà˜%’iˆGÝ”h˜w¨¨E©	µ2´8¸GÀQÈÁYÐPQÑ3RÔ3RÑSÔSˆGÝ”8˜G Q¨¨E©	Ñ2Ô2ˆDÝ”˜˜D    D œMÑ)¨G°D¸!¸!¸!°GÔ,<Ñ<¸bÀÀÀÀDÀ¼kÈ1ÑMÔM×PÒPÕQSÔQ[Ñ\Ô\ˆØð 	~Ø˜S‘[¥2¤7¨1¨t©8°g¸dÀAÀAÀA¸gÔ6FÑ+FÑ#GÔ#G×#JÒ#JÍ2Ì:Ñ#VÔ#VÑV×ZÒZÐ[\Ô[bÔ[mÑnÔn×qÒqÕrtÔr|Ñ}Ô}ˆEØ˜‘
˜S    D œ\¨BÑ.Ñ.°°d¸A¸A¸A°g´ÀÑ1CÑCˆØ�!�!�!�T�'Œ{˜W T¨1¨1¨1 WÔ-Ñ-ˆÝŒg�d˜U‘l D¨#Ñ.Ô.×1Ò1µ"´*Ñ=Ô=ˆÝŒg�d˜U‘l RÑ'¨¨sÑ3Ô3×6Ò6µr´zÑBÔBˆØ˜˜a˜a˜a˜Ô ˆÝŒf•R”X˜e S¨#Ñ.Ô.µ´¸ÀÀ5Ñ0IÔ0IÍ2Ì8ÐTYÐ\_Ð[_ÐadÑKeÔKeÕhjÔhpÐqvÐx}ð  @Añ  iBô  iBñ  LBñ  Cô  CˆÝŒH�W˜T 1 1 1˜WÔ% s¨AÑ.Ô.ˆØð oÝŒH�R˜˜˜˜4˜”[ !¥b¤g¨d°U¸R±ZÑ.?À$ÈÁ)Ñ.KÈaÑ.OÑ&PÔ&P×&SÒ&SÕTVÔT^Ñ&_Ô&_Ð`dÐfgÐfgÐfgÐ`gÔ&hÑ"hÐjmÑnÔnˆØ€Hr(   )Údo_not_specializeÚdo_not_specialize_on_alignmentÚSUM_ONLYÚLAYOUTÚBLOCKÚGROUPSc*           !      óà
  — t          j        d¦  «                             t           j        ¦  «        }*t          j        d¦  «                             t           j        ¦  «        }+t          j        d¦  «                             t           j        ¦  «        },t          j        d|¦  «        }-t          j        d|&¦  «        }.t          j        d|%¦  «        }/|!r¿t          j        |gt           j        ¦  «        }0t          d|$|&¦  «        D ]a}1|*|$z  |1z   |.z   }2|2|k     }3|0t          j        t          | |||||,|	z  |2|
z  z   |+|z  z   |2|3|-|,|+||||||||||| |"|&|'|(|)¦  «        d¦  «        z  }0Œbt          j
        ||,|z  z   |+|z  z   |*|z  z   |-z   |0¦  «         dS t          j        |%gt           j        ¦  «        }4t          d|$|&¦  «        D �]Œ}1|*|$z  |1z   |.z   }2|2|k     }3t          | |||||,|	z  |2|
z  z   |+|z  z   |2|3|-|,|+||||||||||| |"|&|'|(|)¦  «        |z  }5t          j        t          j        |5¦  «        d¦  «        }6|#dk    rAt          j        |4t          j        t          j        |6|&dz  |%f¦  «        d¦  «        ¦  «        }4Œ½|#dk    rVt          j        |4t          j        t          j        t          j        |6|&dz  |%df¦  «        d¦  «        d¦  «        ¦  «        }4�Œ|#dk    rEt          j        |/|1dz  k    t          j        |4t          j        |6d¦  «        ¦  «        |4¦  «        }4�Œdt          j        |4t          j        |6d¦  «        ¦  «        }4�ŒŽ|4dz  |z   }7t          j
        ||,|z  z   |+|z  z   |*|%z  z   |/z   |7¦  «         t          d|$|&¦  «        D �]Ö}1|*|$z  |1z   |.z   }2|2|k     }3t          | |||||,|	z  |2|
z  z   |+|z  z   |2|3|-|,|+||||||||||| |"|&|'|(|)¦  «        |z  }5|#dk    rF|5t          j        t          j        |7ddd…f         |&dz  |%f¦  «        |&f¦  «        dd…df         z  }8nµ|#dk    r]|5t          j        t          j        t          j        |7|7¦  «        ddd…dd…f         |&dz  |%df¦  «        |&f¦  «        dd…df         z  }8nR|#dk    r4|5t          j        t          j        |/|1dz  k    |7d	¦  «        d¦  «        z  }8n|5t          j        |7d¦  «        z  }8|8d
t          j        |8dk    dd¦  «        z  z  }8t          j
        ||,|z  z   |2dd…df         |z  z   |+|z  z   |-ddd…f         z   |8                     t           j        ¦  «        |3dd…df         ¦  «         �ŒØdS )a  One block of BLOCK rows (tokens) of one head of one batch item of X (batch, tokens, heads, D), by steps of SUB rows (_rows): the
    sum of the rows into SUMS (SUM_ONLY), or the rows times SM_SCALE as INT8 into OUT with the GROUPS scales of the LAYOUT into SCALE.r   r   r   Nr   r   r5   g     À_@r7   g      à?éÿÿÿÿ)r8   Ú
program_idr:   r<   Úaranger=   r;   ÚrangerA   ri   ÚstoreÚmaxÚabsÚmaximumr>   rF   Úbroadcast_toÚjoinÚint8)9rG   rH   rI   rJ   ÚOUTÚSCALErK   ÚSUMSr   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   rQ   ÚSM_SCALEÚ	SCALE_EPSr)   r*   r+   r,   r-   r.   rl   r/   rm   rn   ro   r0   r1   r2   r3   ÚblkrP   rO   rN   ÚrÚgr[   Ústartr"   rM   ÚamaxÚvÚmÚscaleÚqs9                                                            r&   Ú_norm_rope_quantrŠ   [   s  € õ Œ-˜Ñ
Ô
×
Ò
�bœhÑ
'Ô
'€CÝŒ=˜ÑÔ×Ò�rœxÑ(Ô(€DÝŒM˜!ÑÔ×Ò¥¤Ñ)Ô)€EÝ
Œ	�!�Q‰Œ€AÝ
Œ	�!�SÑÔ€AÝ
Œ	�!�VÑÔ€AØð Ý”˜!˜�bœjÑ)Ô)ˆÝ˜1˜e SÑ)Ô)ð 	tð 	tˆEØ˜‘+ Ñ%¨Ñ)ˆCØ�s’ˆBØ•R”V�E ! Q¨¨d°D¸%À"¹*ÀsÈRÁxÑ:OÐRVÐY[ÑR[Ñ:[Ð]`ÐbdÐfgÐinÐptÐvxÐz|ð  Að  CEð  GJØ"$ b¨!¨U°D¸&À(ÈCÐQVÐXcÐenñpô pØqrñtô tñ tˆEˆEå
Œ�˜ ™
Ñ" T¨B¡YÑ.°°r±Ñ9¸AÑ=¸uÑEÔEÐEØˆÝŒ8�V�H�bœjÑ)Ô)€DÝ�q˜% Ñ%Ô%ð 2ñ 2ˆØ�E‰k˜EÑ! AÑ%ˆØ�3ŠYˆÝ�!�Q˜˜d D¨%°"©*°s¸R±xÑ*?À$ÈÁ)Ñ*KÈSÐRTÐVWÐY^Ð`dÐfhÐjlÐnpÐrtÐvyØ�b˜!˜U D¨&°(¸CÀÈÐU^ñ`ô `ØbjñkˆåŒF•2”6˜!‘9”9˜aÑ Ô ˆØ�QŠ;ˆ;Ý”:˜d¥B¤F­2¬:°a¸#À¹(ÀFÐ9KÑ+LÔ+LÈaÑ$PÔ$PÑQÔQˆDˆDØ�qŠ[ˆ[Ý”:˜d¥B¤F­2¬6µ"´*¸QÀÈÁÈ6ÐSTÐ@UÑ2VÔ2VÐXYÑ+ZÔ+ZÐ\]Ñ$^Ô$^Ñ_Ô_ˆD‰DØ�qŠ[ˆ[Ý”8˜A ¨"¡Ò,­b¬j¸½r¼vÀaÈ¹|¼|Ñ.LÔ.LÈdÑSÔSˆD‰Då”:˜d¥B¤F¨1¨a¡L¤LÑ1Ô1ˆD‰DØ�4‰K˜)Ñ#€EÝ„HˆU�U˜R‘ZÑ $¨¡)Ñ+¨c°F©lÑ:¸QÑ>ÀÑFÔFÐFÝ�q˜% Ñ%Ô%ð lñ lˆØ�E‰k˜EÑ! AÑ%ˆØ�3ŠYˆÝ�!�Q˜˜d D¨%°"©*°s¸R±xÑ*?À$ÈÁ)Ñ*KÈSÐRTÐVWÐY^Ð`dÐfhÐjlÐnpÐrtÐvyØ�b˜!˜U D¨&°(¸CÀÈÐU^ñ`ô `Øbjñkˆà�QŠ;ˆ;Ø•B”J�rœ¨u°T¸1¸1¸1°W¬~ÀÀqÁÈ&Ð?QÑRÔRÐUXÐTZÑ[Ô[Ð\]Ð\]Ð\]Ð_cÐ\cÔdÑdˆAˆAØ�qŠ[ˆ[Ø•B”J�rœ­r¬w°u¸eÑ/DÔ/DÀTÈ1È1È1ÈaÈaÈaÀZÔ/PÐSVÐZ[ÑS[Ð]cÐefÐRgÑhÔhÐknÐjpÑqÔqÐrsÐrsÐrsÐuyÐryÔzÑzˆAˆAØ�qŠ[ˆ[Ø•B”F�2œ8 A¨°"©Ò$4°e¸SÑAÔAÀ1ÑEÔEÑEˆAˆAà•B”F˜5 !Ñ$Ô$Ñ$ˆAØ	ˆS•2”8˜A šF A rÑ*Ô*Ñ*Ñ*ˆÝ
Œ��u˜r‘zÑ! C¨¨¨¨4¨¤L°2Ñ$5Ñ5¸¸r¹	ÑAÀAÀdÈAÈAÈAÀgÄJÑNÐPQ×PTÒPTÕUWÔU\ÑP]ÔP]Ð_aÐbcÐbcÐbcÐeiÐbiÔ_jÑkÔkÐkÑkðlð lr(   )r   Tr   Fc                 ó†  — |€| n|}|j         \  }}}}|st          |||¦  «        nt          j        ||¦  «        }|�|j         d         dk    rdn|                     d¦  «        }|�@|||                     d¦  «        |                     d¦  «        |                     d¦  «        fn|ddddf}|� |g|                     ¦   «         d d…         ¢R n|dddf}|� |g|                     ¦   «         d d…         ¢R n|ddf} |� |g|                     ¦   «         d d…         ¢R n|ddf}!|� |g|                     ¦   «         d d…         ¢R n|dddf}"t          |||f         g |‘|�|n|‘|�|n|‘|d         ‘|d         ‘| d         ‘|!d         ‘|"d         ‘|‘|                     ¦   «         d d…         ¢|d         ‘|d         ‘|dd …         ¢| dd …         ¢|!dd …         ¢|"dd …         ¢|‘|‘|	‘|d         ‘|d         ‘|‘|
‘|‘|‘|‘|d u‘|‘|‘t          |         ‘|€t          n|‘|€t          n|‘|‘|d u‘R ||ddœŽ d S )Nr   r   r   r	   r   F)Ú	num_warpsÚ
num_stagesÚenable_fp_fusion)	Úshaper'   r    r!   ÚstriderŠ   ÚSCALES_PER_BLOCKÚSUB_ROWSÚ	LANE_SUMS)#rR   ÚweightÚroperh   rˆ   ÚmeanÚsumsÚepsÚsm_scaleÚ	scale_epsÚpairsÚnormÚrotateÚsum_onlyr#   r$   ÚwarpsÚscratchÚstagesÚsubrT   ÚinterleavedÚweight_rowsÚxxrO   r"   ÚheadsÚdimÚblocksÚ
rope_batchÚ	rope_argsÚout_argsÚ
scale_argsÚ	mean_argsÚ	sums_argss#                                      r&   Ú_launchr¯   –   sô  € àˆoˆˆ 7€BØœXÑ€Eˆ3��sØ08ÐU�W�S˜& %Ñ(Ô(Ð(½f¼kÈ#ÈuÑ>UÔ>U€FØ�l d¤j°¤m°qÒ&8Ð&8��¸d¿kºkÈ!¹n¼n€JØVZÐVf��z 4§;¢;¨q¡>¤>°4·;²;¸q±>´>À4Ç;Â;ÈqÁ>Ä>ÐRÐRÐmoÐqrÐtuÐwxÐz{Ðl|€IØ+.¨?�Ð'�c—j’j‘l”l 2 A 2Ô&Ð'Ð'Ð'ÀÀQÈÈ1À€HØ16Ð1B�%Ð-˜%Ÿ,š,™.œ.¨¨!¨Ô,Ð-Ð-Ð-ÈÈQÐPQÈ
€JØ.2Ð.>�Ð*˜Ÿš™œ b q bÔ)Ð*Ð*Ð*ÀRÈÈAÀJ€IØ.2Ð.>�Ð*˜Ÿš™œ b q bÔ)Ð*Ð*Ð*ÀRÈÈAÈqÀM€IÝ�f˜e UÐ+Ô,ð DØ
ðDØÐ(ˆFˆF¨bðDØALÐAX°+°+Ð^`ðDØbkÐlmÔbnðDØpxÐyzÔp{ðDð ~Hð  IJô  ~KðDð MVð  WXô  MYðDð [dð  efô  [gðDàðDà�iŠi‰kŒk˜"˜1˜"ŒoðDà(¨œ|ðDà-6°q¬\ðDà<DÀQÀRÀR¼LðDàKUÐVWÐVXÐVXÌ>ðDà\eÐfgÐfhÐfhÔ\iðDàluÐvwÐvxÐvxÔlyðDà{~ðDð AIðDð KTðDð 	�!ŒðDð   ”lðDð %(ðDð */ðDð 15ðDð 7=ðDð ?GðDð IMÐTXÐHXðDð [aðDð chðDõ jzð  {Aô  jBðDð �K�ˆ Sð	Dð 8=°}­)¨)È%ð	Dð R]ð	Dð _jÐquÐ^uð	Dð Dð
  F¸UðDð Dð Dð Dð Dð Dr(   c                 ó$   — | €dn| j         d         S )ziRotated channel pairs of a table (1 | batch, tokens, 1, pairs, 2) or (1 | batch, tokens, 1, pairs, 2, 2).Nr   r   )r�   )r•   s    r&   Ú_pairsr±   ©   s   € à�ˆ1ˆ1 $¤*¨Q¤-Ð/r(   c	                 ób  ‡ ‡‡‡‡‡— ‰ j         \  }	ŠŠŠˆˆˆˆˆˆ fd„}
t                               ||› d‰ j        › d‰› d‰› d‰                      d¦  «        › dt          |¦  «        › dt          |¦  «        › dt          |¦  «        › d	|› d
‰› �‰ j        t          |
|	‰z  ‰z  t          k    ¦  «        S )Nc                  óJ  •‡— t          ‰t          ‰t          ‰z  ¦  «        ¦  «        ‰z  ‰z  } ‰                     dt          ‰| ¦  «        ‰‰                     d¦  «        f¦  «        dd ‰…f         Š‰                     ‰d d…d ‰j        d         …f         ¦  «         ˆˆfd„S )Nr   r   .c                 ó   •—  ‰‰| ¦  «        S )N© )ÚconfigÚlaunchr    s    €€r&   ú<lambda>z2_config.<locals>.prepare_scratch.<locals>.<lambda>µ   s   ø€ ˜f˜f W¨fÑ5Ô5€ r(   )rv   ÚminÚ
_TUNE_ROWSÚ	new_emptyr�   Úcopy_r�   )Útokensr    r$   r§   r¦   r·   r"   rR   s    @€€€€€€r&   Úprepare_scratchz _config.<locals>.prepare_scratch±   s¦   øø€ Ý�U�C ¥Z°5Ñ%8Ñ9Ô9Ñ:Ô:¸eÑCÀeÑKˆØ—+’+˜q¥# c¨6Ñ"2Ô"2°E¸1¿8º8ÀA¹;¼;ÐGÑHÔHÈÈdÈsÈdÈÔSˆØ�Š�a˜˜˜Ð-˜Wœ]¨1Ô-Ð-Ð-Ô.Ñ/Ô/Ð/Ø5Ð5Ð5Ð5Ð5Ð5r(   ú|rR   z|sr   z|NÚRÚSÚLÚB)	r�   Ú_TUNERr¶   rD   r�   ÚintÚdeviceÚ_LAUNCH_CONFIGSrº   )rR   ÚnameÚidentityr$   rœ   r�   rž   r#   r·   rO   r¾   r§   r¦   r"   s   `  `    `  @@@r&   Ú_configrÊ   ®   sB  øøøøøø€ ØœWÑ€Eˆ3��sð6ð 6ð 6ð 6ð 6ð 6ð 6ð 6ð 6ð 6õ �=Š=˜ (ð   Pð   P¨Q¬Wð   Pð   P°uð   Pð   P¸sð   Pð   PÀaÇhÂhÈqÁkÄkð   Pð   PÕUXÐY]ÑU^ÔU^ð   Pð   PÕadÐekÑalÔalð   Pð   PÕorÐs{Ño|Ôo|ð   Pð   Pð  @Fð   Pð   Pð  INð   Pð   PØœ¥?°OÀUÈSÁ[ÐSXÑEXÕ\fÒEfñhô hð hr(   é@   )r£   r¤   c                óÄ  ‡ ‡‡‡‡‡‡‡‡— ‰ j         \  }}}	}
t          j        ||	t          j        |‰¦  «        |
ft          j        ‰ j        ¬¦  «        Št          ‰¦  «        Šˆˆˆˆˆˆˆˆˆ f	d„} |‰ t          ‰ dd‰› dt          ‰¦  «        › dt          ‰du¦  «        › �‰‰du‰dudd	|¦	  «	        ¦  «         ‰ 
                    d
¬¦  «                             |¦  «        S )aµ  Mean over the tokens of the normalized, rotated rows of x (batch, tokens, heads, D) per batch item and head, (batch, heads, D) FP32,
    from FP32 partial sums per block of rows; x is not modified. weight None skips the norm, rope None the rotation; interleaved: pairs
    (2p, 2p + 1) instead of (p, p + pairs); weight_rows (tokens,) int32: the row of weight (rows, D) of each token (segments of tokens
    with their own norm weights).©rD   rÆ   c                 óâ   •	— t          ‰
‰‰d d d | ‰
u r‰n.‰d d…d d …d t          j        | j        d         ‰¦  «        …f         ‰dd‰‰d u‰d udd‰|d         | ‰
u rd n| |d         ‰‰	¬¦  «         d S )Nr   r6   r7   Tr   ©r    r¡   r£   r¤   )r¯   r    r!   r�   )r¥   r¶   r$   r˜   r£   r›   r•   r—   r”   r¤   rR   s     €€€€€€€€€r&   r·   znorm_rope_mean.<locals>.launchÄ   sÓ   ø€ Ý��6˜4  t¨T¸2À¸7¸7°4°4ÈÈRÈaÈRÐQRÐQRÐQRÐTtÕU[ÔU`ÐacÔaiÐjkÔalÐnsÑUtÔUtÐTtÐMtÔHuÐwzÐ|ð  BEð  GLØ˜dÐ" D°Ð$4°d¸A¸uÀfÈQÄiÐacÐghÐahÐahÐY]ÐY]ÐnpÐyð  ABô  zCð  Q\Ø'ð	)ñ 	)ô 	)ð 	)ð 	)ð 	)r(   Únorm_rope_sumsÚpÚIÚGNTr   r   )r§   )r�   ÚtorchÚemptyr    r!   r;   rÆ   r±   rÊ   rÅ   rA   Údiv_)rR   r”   r•   r˜   r$   r£   r¤   rO   r"   r¦   r§   r·   r›   r—   s   ```````     @@r&   Únorm_rope_meanr×   »   sW  øøøøøøøøø€ ð
 œWÑ€Eˆ3��sÝŒ;˜˜u¥f¤k°#°uÑ&=Ô&=¸sÐCÍ5Ì=ÐabÔaiÐjÑjÔj€DÝ�4‰LŒL€Eð)ð )ð )ð )ð )ð )ð )ð )ð )ð )ð )ð )ð )ð
 €Fˆ1�g�aÐ)Ð+h¨uÐ+hÐ+hµs¸;Ñ7GÔ7GÐ+hÐ+hÍ#ÈkÐaeÐNeÑJfÔJfÐ+hÐ+hÐjoÐqwð  @Dð  rDð  FJð  RVð  FVð  X\ð  ^_ð  agñ  hô  hñ  iô  ið  iØ�8Š8˜ˆ8‰?Œ?×Ò Ñ$Ô$Ð$r(   r6   gH¯¼šò×z>)r–   r™   rš   r£   r¤   c                ó  ‡ ‡‡‡‡‡‡‡‡‡	‡
‡‡‡— ‰ j         \  }}}}t          j        ‰ j         t          j        ‰ j        ¬¦  «        Št          j        ||t          |‰‰¦  «        t          ‰         z  ft          j        ‰ j        ¬¦  «        Št          ‰¦  «        Šˆˆˆ	ˆˆˆˆˆˆˆˆˆˆ
ˆ fd„} |‰ t          ‰ dd‰› dt          ‰du¦  «        › dt          ‰	¦  «        › dt          ‰
du¦  «        › �‰‰du‰dud	‰|¦	  «	        ¦  «         ‰‰fS )
a!  INT8 rows and FP32 scales (batch, heads, blocks * scales per block) of the normalized (weight, eps), rotated (rope, interleaved
    as norm_rope_mean) rows of x, minus mean (batch, heads, D) when given and times sm_scale, in the given scale layout and block of
    rows; x is not modified.rÍ   c                 ól  •— | ‰u r5t          ‰‰‰
‰‰‰d ‰‰‰‰	‰d u‰
d ud‰‰|d         |d         ‰‰¬¦  «         d S | j        d         }t          ‰‰‰
‰d d…d |…f         ‰d d…d d …d t          |‰‰¦  «        t          ‰         z  …f         ‰d ‰‰‰‰	‰d u‰
d ud‰‰|d         | |d         ‰‰¬¦  «         d S )NFr   r   )r¡   r£   r¤   rÏ   )r¯   r�   r'   r‘   )r¥   r¶   Úrowsr$   r˜   r£   r#   r–   rh   r›   r•   rˆ   rš   r™   r”   r¤   rR   s      €€€€€€€€€€€€€€r&   r·   z"norm_rope_quantize.<locals>.launchÖ   sa  ø€ Ø�ˆ7ˆ7Ý�A�v˜t S¨%°°t¸SÀ(ÈIÐW\Ð^dÐlpÐ^pÐrvð  Cð  sCð  EJð  LRð  TYØ˜1”I f¨Q¤i¸[ÐVaðcñ cô cð cð cð cð ”8˜A”;ˆDÝ�A�v˜t S¨¨!¨¨U¨d¨U¨¤^°U¸2¸A¸2¸q¸q¸qÐBzÅ7È4ÐQWÐY^ÑC_ÔC_ÕbrÐsyÔbzÑCzÐBzÐ;zÔ5{ð  ~Bð  DHð  JMð  OWð  YbØ˜6¨Ð-¨t¸4Ð/?ÀÈÐPUÐW]Ð^_ÔW`ÐjlÐu{Ð|}Ôu~ð  MXØ +ð-ñ -ô -ð -ð -ð -r(   Únorm_rope_quantrÑ   ÚMNrÒ   rÓ   F)r�   rÔ   rÕ   r{   rÆ   r'   r‘   r;   r±   rÊ   rÅ   )rR   r”   r•   r˜   r#   r$   r–   r™   rš   r£   r¤   rO   r"   r¦   r§   r·   rh   r›   rˆ   s   ```````````     @@@r&   Únorm_rope_quantizerÝ   Í   s²  øøøøøøøøøøøøøø€ ð œWÑ€Eˆ3��sÝ
Œ+�a”g¥U¤Z¸¼Ð
AÑ
AÔ
A€CÝŒK˜ ¥w¨s°F¸EÑ'BÔ'BÕEUÐV\ÔE]Ñ']Ð^ÕfkÔfsÐ|}ô  }Eð  Fñ  Fô  F€EÝ�4‰LŒL€Eð-ð -ð -ð -ð -ð -ð -ð -ð -ð -ð -ð -ð -ð -ð -ð -ð -ð -ð €Fˆ1�g�aÐ*ð  -B°ð  -Bð  -B½¸DÈÐ<LÑ8MÔ8Mð  -Bð  -BÕPSÐT_ÑP`ÔP`ð  -Bð  -BÕcfÐgrÐz~Ðg~ÑcÔcð  -Bð  -Bð  DIð  KQð  Y]ð  K]Ø $Ð&¨¨v°vñ?ô ?ñ @ô @ð @à�ˆ:Ðr(   )Ú!install_triton_compilation_logger)Nr   NNFN)rË   )"Ú__doc__rÔ   r    Útriton.languageÚlanguager8   Útriton.language.extrar   Útriton_tunerr   ÚLAYOUT_PER_BLOCKÚLAYOUT_THREAD_QÚLAYOUT_THREAD_KÚLAYOUT_WARP_Qr‘   r   rÄ   rÇ   rº   Ú_RUNTIMEr'   ÚjitÚ	constexprri   rŠ   r’   r“   r¯   r±   rÊ   r×   rÝ   Ú%shared.kernels.triton_compilation_logrÞ   rµ   r(   r&   ú<module>rì      s%  ðð/ð /ð" €€€Ø €€€Ø Ð Ð Ð Ð Ð Ø +Ð +Ð +Ð +Ð +Ð +à %Ð %Ð %Ð %Ð %Ð %àDNÑ AÐ �/ ?°MØ$ a¨¸!¸_ÈaÐQ^Ð`aÐbÐ Ø Ð$€Ø	ˆÐ&¨¨;Ñ	7Ô	7€Ø*€Ø€
ØlÐlÐl€ð2ð 2ð 2ð „ð.ÐZ\ÔZfð .ÐlnÔlxð .Ð}ô  ~Jð .Ø”ð.Ø%'¤\ð.Ø;=¼<ð.ØSUÔS_ð.ØfhÔfrð.Ø{}ô  |Hð.à”|ð.à02´ð.ð .ð .ñ „ð.ðb €„˜hÀxÐPÑPÔPð4lØœð4lØ+-¬<ð4lØ<>¼Lð4lØQSÔQ]ð4lØegÔeqð4lØ{}ô  |Hð4là!œ|ð4là79´|ð4làMOÌ\ð4làbdÔbnð4làxzô  yEð4lð œ,ð4lð 02¬|ð4lð KMÌ,ð4lð dfÔcoð4lð 4lð 4lñ QÔPð4lðn Ñ €ˆ)ð MQØKOðDð Dð Dð Dð&0ð 0ð 0ð

hð 
hð 
hð%À%ÐUYð %ð %ð %ð %ð %ð$ EIÐSVÐbfÐtyð  HLð ð ð ð ð ð0 TÐ SÐ SÐ SÐ SÐ SØ !Ð !Ñ #Ô #Ð #Ð #Ð #r(   