§
     ’ÞjeR  ã                   ó  — d Z ddlZddlZ	 ddlZddlmZ ddlmZ n# e	$ r dZY nw xY wej
        defd„¦   «         Ze��ej        d„ ¦   «         Zej        d„ ¦   «         Zej        d„ ¦   «         Zej        d	„ ¦   «         Zej        d
„ ¦   «         Zej        d„ ¦   «         Zej        d„ ¦   «         Zej        dej        fd„¦   «         Zej        d„ ¦   «         Z ej        d¬¦  «        dej        dej        dej        fd„¦   «         Z ej        d¬¦  «        dej        dej        fd„¦   «         Z ej        dd¬¦  «        dej        dej        fd„¦   «         Zej        dej        fd„¦   «         Zej        d„ ¦   «         Zej        d„ ¦   «         Zej        dej        fd „¦   «         Zej        dej        fd!„¦   «         Zej        dej        fd"„¦   «         Z ej        dej        fd#„¦   «         Z! ej        d$d$¬¦  «        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        fd.„¦   «         Z#d/„ Z$d0„ Z%d1„ Z&d2„ Z'd3Z(d4Z)d5Z*d6„ Z+e�dd7l,m-Z-  e-¦   «          dS dS )8aÈ  Triton kernels for GPUs that run Triton without BF16 conversion instructions (Turing, sm_75): BF16 tensors are read and written as
their bits (shift into FP32 bits; round to nearest even as PyTorch's casts, NaN -> 0x7FC0). For the Qwen3.5/3.8 language models:
- the norms and the activation, which otherwise run as chains of small PyTorch kernels there. The norms round once to BF16 (the PyTorch
  route rounds each FP32 step, then to BF16): RMSNorm keeps sums and products exact as FP32 pairs (error-free transformations; Turing
  runs FP64 at 1/32 of the FP32 rate) with the inverse square root in FP64, the gated norm of 128 values computes in FP64; SiLU*mul
  reproduces the PyTorch route bit for bit;
- the paged Q8 attention of decoding and speculative verification, on FP16 tensor cores (sm_75 has no BF16 ones), each K/V page read
  once for the query heads of a KV head and all the queries of a verification, every row computed alike whatever the query count.é    N)Ú	libdeviceÚreturnc                  óº   — t           duoRt          j                             ¦   «         o4t          j        j        du o!t          j                             ¦   «         dk    S )zGValidated on compute capability 7.5 only; other GPUs keep their routes.N)é   é   )ÚtritonÚtorchÚcudaÚis_availableÚversionÚhipÚget_device_capability© ó    ú@C:\Users\johnn\OneDrive\Documents\AI\shared\kernels\bf16_bits.pyÚeligibler      si   € õ ˜Ðð  K¥%¤*×"9Ò"9Ñ";Ô";ð  KÅÄÔ@QÐUYÐ@Yð  KÕ^cÔ^h×^~Ò^~ñ  _Aô  _Að  EKò  _Kð  Kr   c                 ó„   — |                       t          j        ¦  «        dz                        t          j        d¬¦  «        S )Né   T©Úbitcast)ÚtoÚtlÚint32Úfloat32)Úbitss    r   Úbf16_bits_to_fp32r      s0   € à—’�œÑ!Ô! RÑ'×+Ò+­B¬JÀÐ+ÑEÔEÐEr   c                 ó  — |                       t          j        d¬¦  «        }t          j        | | k    d|dz   |dz	  dz  z   dz	  ¦  «        }|                      t          j        ¦  «                              t          j        d¬¦  «        S )NTr   iÀ  iÿ  r   é   )r   r   Úuint32ÚwhereÚuint16Úint16)Úxr   Úroundeds      r   Úfp32_to_bf16_bitsr%   !   sq   € à�tŠt•B”I tˆtÑ,Ô,ˆÝ”(˜1 š6 6¨D°6©M¸dÀb¹jÈAÑ=MÑ,NÐSUÑ+UÑVÔVˆØ�zŠz�"œ)Ñ$Ô$×'Ò'­¬¸$Ð'Ñ?Ô?Ð?r   c           	      óÔ  — |                       t          j        ¦  «        }|                      t          j        ¦  «        }|                      t          j        d¬¦  «        }t          j        || k    t          j        t          j        |¦  «        t          j        | ¦  «        k    |dz
  |¦  «        dz  |¦  «        }t          |                      t          j        d¬¦  «        ¦  «        S )NTr   r   )r   r   r   Úfloat64r   r    Úabsr%   )r#   ÚfÚbackr   s       r   Úfp64_to_bf16_bitsr+   '   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Ð[_Ñ`Ô`ˆÝ  §¢­¬¸T Ñ!BÔ!BÑCÔCÐCr   c                 ó6   — | |z   }|| z
  }|| ||z
  z
  ||z
  z   fS ©Nr   )ÚaÚbÚsÚbbs       r   Ú_two_sumr2   0   s0   € à�‰EˆØ�‰UˆØ�1˜˜B™‘< A¨¡FÑ+Ð+Ð+r   c                 ó<   — t          | |¦  «        \  }}||||z   z   fS r-   )r2   )Úa_hiÚa_loÚb_hiÚb_lor0   Úes         r   Ú
_add_pairsr9   6   s)   € å˜˜dÑ#Ô#‰ˆˆ1Ø�!�t˜d‘{Ñ#Ð#Ð#r   c                 ó¸  — t          | |¦  «        \  }}t          j        t          j        |¦  «        t	          d¦  «        k     |d¦  «        }|                     t          j        d¬¦  «        }t          j        |dk    t          j        |dk     |dk     k    |dz
  |¦  «        dz  |¦  «        }t          |                     t          j        d¬¦  «        ¦  «        S )NÚinfç        Tr   r   r   )	r2   r   r    r(   Úfloatr   r   r%   r   )ÚhiÚlor0   r8   r   s        r   Úpair_to_bf16_bitsr@   ;   s°   € õ ˜˜BÑÔ‰ˆˆ1ÝŒH•R”V˜A‘Y”Y¥ u¡¤Ò-¨q°#Ñ6Ô6ˆØ�tŠt•B”I tˆtÑ,Ô,ˆÝŒx˜˜Qš¥¤¨!¨aª%°Q¸²UÒ);¸TÀA¹XÀtÑ LÔ LÈqÑ PÐRVÑWÔWˆÝ  §¢­¬¸T Ñ!BÔ!BÑCÔCÐCr   c                 ó‚   — t          t          j        | |d¦  «        ¦  «                             t          j        ¦  «        S ©Nr   )r   r   Úloadr   r'   )ÚPTRÚmasks     r   Ú_load_bf16_fp64rF   D   s.   € å ¥¤¨¨d°AÑ!6Ô!6Ñ7Ô7×:Ò:½2¼:ÑFÔFÐFr   ÚW_BITSc                 ó¾   — |r&t          t          j        | |z   |d¦  «        ¦  «        S t          j        | |z   |d¦  «                             t          j        ¦  «        S rB   )r   r   rC   r   r   )ÚWÚoffsrE   rG   s       r   Ú_weightrK   H   sS   € àð 	AÝ$¥R¤W¨Q°©X°t¸QÑ%?Ô%?Ñ@Ô@Ð@ÝŒw�q˜4‘x  qÑ)Ô)×,Ò,­R¬ZÑ8Ô8Ð8r   c                 ó$  — t          j        t          j        | | z  d¬¦  «        |                     t          j        ¦  «        ¦  «        }t          j        t          j        ddt          j        ¦  «        t          j        ||z   ¦  «        ¦  «        S )Nr   ©Úaxisr   ç      ð?)r   Údiv_rnr   Úsumr   r'   ÚfullÚsqrt)r#   Ún_colsÚepsÚmeans       r   Ú_inverse_rmsrW   N   sg   € åÔ¥¤ q¨1¡u°1Ð 5Ñ 5Ô 5°v·y²yÅÄÑ7LÔ7LÑMÔMˆÝÔ¥¤¨¨Cµ´Ñ <Ô <½i¼nÈTÐTWÉZÑ>XÔ>XÑYÔYÐYr   )Úx_strideÚr_strideÚy_strideÚ	ro_stride)Údo_not_specializeÚHAS_RESIDUALÚBLOCKc                 ó¬  — t          j        d¦  «                             t           j        ¦  «        }t          j        d|¦  «        }||	k     }t          t          j        | ||z  z   |z   |d¦  «        ¦  «        }t          j        |¦  «        }|r£t          |t          t          j        |||z  z   |z   |d¦  «        ¦  «        ¦  «        \  }}t          j	        t          j
        |¦  «        t          d¦  «        k     |d¦  «        }t          j        |||z  z   |z   t          ||¦  «        |¦  «         t          j        t          j
        |¦  «        d¬¦  «                             t           j        d¬¦  «        dz	                       t           j        ¦  «        }t          j        t          j        |d¦  «        d	¦  «        }d
|z
  dz                       t           j        d¬¦  «        }||z  ||z  }}||z  }t          j        |t          j        ||| ¦  «        d|z  |z  z   fdt,          ¦  «        \  }}t          j	        t          j
        |¦  «        t          d¦  «        k     |d¦  «        }|dz                        t           j        ¦  «        dz                       t           j        d¬¦  «        }t1          j        |                     t           j        ¦  «        |                     t           j        ¦  «        z   |z  |z  |	                     t           j        ¦  «        ¦  «        }t1          j        t          j        ddt           j        ¦  «        t1          j        ||
z   ¦  «        ¦  «        }|                     t           j        ¦  «        }||                     t           j        ¦  «        z
                       t           j        ¦  «        }||z  }t          j        ||| ¦  «        ||z  ||z  z   z   } t9          ||||¦  «        }!||!z  }"t          j        |||z  z   |z   t          |"t          j        ||!|" ¦  «        | |!z  z   ¦  «        |¦  «         d S )Nr   r;   r<   rM   Tr   é   r   éý   éþ   é   i€  é4   r   rO   )r   Ú
program_idr   Úint64Úaranger   rC   Ú
zeros_liker2   r    r(   r=   Ústorer@   Úmaxr   r   ÚminimumÚmaximumr   ÚreduceÚfmar9   r'   r   rP   rR   rS   rK   )#ÚXÚRESrI   ÚYÚRES_OUTrX   rY   rZ   r[   rT   rU   r]   rG   r^   ÚrowrJ   rE   r#   r?   ÚtopÚscaleÚxsÚlsÚsquareÚtotalÚtotal_loÚunscalerV   ÚinverseÚ
inverse_hiÚ
inverse_lor.   r5   ÚwÚys#                                      r   Ú_rms_norm_kernelr�   S   sœ  € õ
 Œm˜AÑÔ×!Ò!¥"¤(Ñ+Ô+ˆÝŒy˜˜EÑ"Ô"ˆØ�fŠ}ˆÝ�bœg a¨#°©.Ñ&8¸4Ñ&?ÀÀqÑIÔIÑJÔJˆÝŒ]˜1ÑÔˆØð 	WÝ˜QÕ 1µ"´'¸#ÀÀhÁÑ:NÐQUÑ:UÐW[Ð]^Ñ2_Ô2_Ñ `Ô `ÑaÔa‰EˆAˆrÝ”�"œ& ™)œ)¥e¨E¡l¤lÒ2°B¸Ñ<Ô<ˆBÝŒH�W˜s Y™Ñ.°Ñ5Õ7HÈÈBÑ7OÔ7OÐQUÑVÔVÐVåŒv•b”f˜Q‘i”i aÐ(Ñ(Ô(×+Ò+­B¬I¸tÐ+ÑDÔDÈÑJ×NÒNÍrÌxÑXÔXˆÝŒj�œ C¨Ñ+Ô+¨SÑ1Ô1ˆØ˜‘) Ñ"×&Ò&¥r¤z¸4Ð&Ñ@Ô@ˆØ�U‘˜B ™JˆBˆØ�b‘ˆÝœ) V­R¬V°B¸¸V¸GÑ-DÔ-DÀqÈ2ÁvÐPRÁ{Ñ-RÐ$SÐUVÕXbÑcÔc‰ˆˆxÝ”8�BœF 5™MœM­E°%©L¬LÒ8¸(ÀCÑHÔHˆØ˜#‘I—>’>¥"¤(Ñ+Ô+¨rÑ1×5Ò5µb´jÈ$Ð5ÑOÔOˆÝÔ §¢­"¬*Ñ!5Ô!5¸¿ºÅBÄJÑ8OÔ8OÑ!OÐSZÑ ZÐ]dÑ dÐfl×foÒfoÕprÔpzÑf{Ôf{Ñ|Ô|ˆÝÔ"¥2¤7¨2¨sµB´JÑ#?Ô#?ÅÄÐPTÐWZÑPZÑA[ÔA[Ñ\Ô\ˆØ—Z’Z¥¤
Ñ+Ô+ˆ
Ø 
§¢­b¬jÑ 9Ô 9Ñ9×=Ò=½b¼jÑIÔIˆ
Ø�
‰NˆÝŒv�a˜ a RÑ(Ô(¨A°
©N¸RÀ*¹_Ñ,LÑMˆÝ�A�t˜T 6Ñ*Ô*ˆØ�‰EˆÝ
Œ��S˜8‘^Ñ# dÑ*Õ,=¸aÅÄÈÈ1ÈqÈbÑAQÔAQÐTXÐ[\ÑT\ÑA\Ñ,]Ô,]Ð_cÑdÔdÐdÐdÐdr   )rX   Úg_stridec
                 óž  — t          j        d¦  «                             t           j        ¦  «        }
t          j        d|	¦  «        }||k     }t          | |
|z  z   |z   |¦  «        }t          t          j        ||
|z  z   |z   |d¦  «        ¦  «        }t          j        |dt          j
        | ¦  «        z   ¦  «                             t           j        ¦  «        }|t          ||||¦  «                             t           j        ¦  «        z  t          |||¦  «        z  |z  }t          j        ||
|z  z   |z   t          |¦  «        |¦  «         d S )Nr   rO   )r   re   r   rf   rg   rF   r   rC   rP   r   Úexpr'   rK   rW   ri   r+   )ro   ÚGrI   rq   rX   r‚   rT   rU   rG   r^   rs   rJ   rE   r#   ÚgateÚsilur€   s                    r   Ú_gated_rms_norm_kernelrˆ   t   s'  € õ Œm˜AÑÔ×!Ò!¥"¤(Ñ+Ô+ˆÝŒy˜˜EÑ"Ô"ˆØ�fŠ}ˆÝ˜A  h¡Ñ.°Ñ5°tÑ<Ô<ˆÝ ¥¤¨¨S°8©^Ñ);¸dÑ)BÀDÈ!Ñ!LÔ!LÑMÔMˆÝŒy˜˜s¥Y¤]°D°5Ñ%9Ô%9Ñ9Ñ:Ô:×=Ò=½b¼jÑIÔIˆØ•˜˜4  vÑ.Ô.×1Ò1µ"´*Ñ=Ô=Ñ=ÅÈQÐPVÐX[Ñ@\Ô@\Ñ\Ð_cÑcˆÝ
Œ��S˜6‘\Ñ! DÑ(Õ*;¸AÑ*>Ô*>ÀÑEÔEÐEÐEÐEr   )ÚTOTAL)r\   Údo_not_specialize_on_alignmentÚCOLUMNSc                 ó  — t          j        d¦  «        |z  t          j        d|¦  «        z   }||z  d|z  z  ||z  z   }t          t          j        | |z   ||k     d¦  «        ¦  «        }t          t          j        | |z   |z   ||k     d¦  «        ¦  «        }t          t          t          j        |dt          j        | ¦  «        z   ¦  «        ¦  «        ¦  «        }	t          j	        ||z   t          |	|z  ¦  «        ||k     ¦  «         d S )Nr   rc   rO   )
r   re   rg   r   rC   r%   rP   r   r„   ri   )
ro   ÚOUTr‹   r‰   r^   ÚindexÚsourcer†   ÚvalueÚ	activateds
             r   Ú_silu_mul_kernelr’   €   sõ   € õ ”˜aÑ Ô  5Ñ(­2¬9°Q¸Ñ+>Ô+>Ñ>ˆØ˜'Ñ! Q¨¡[Ñ1°E¸G±OÑCˆÝ ¥¤¨¨V©°U¸U²]ÀAÑ!FÔ!FÑGÔGˆÝ!¥"¤'¨!¨f©*°wÑ*>ÀÈÂÈqÑ"QÔ"QÑRÔRˆÝ%Õ&7½¼	À$ÈÍiÌmÐ]aÐ\aÑNbÔNbÑHbÑ8cÔ8cÑ&dÔ&dÑeÔeˆ	Ý
Œ��u‘Õ/°	¸EÑ0AÑBÔBÀEÈEÂMÑRÔRÐRÐRÐRr   Úoffsetc                 óV   — |dz   | z
  dz                        t          j        d¬¦  «        S )Nrb   r`   Tr   )r   r   r   )rt   r“   s     r   Ú_power_of_twor•   Š   s-   € ð ˜#‘ Ñ#¨Ñ*×.Ò.­r¬zÀ4Ð.ÑHÔHÐHr   c                 óÂ   — |                       t          j        ¦  «        }|| |                      t          j        ¦  «        z
                        t          j        ¦  «        fS r-   )r   r   Úfloat16r   )r#   r>   s     r   Ú
_fp16_pairr˜   �   sC   € à�TŠT•"”*ÑÔˆØ�A˜Ÿš�bœjÑ)Ô)Ñ)×-Ò-­b¬jÑ9Ô9Ð9Ð9r   c                 ó„   — |                       t          j        d¬¦  «        dz	                        t          j        ¦  «        S )NTr   r`   )r   r   r   r   )r#   s    r   Ú	_exponentrš   ”   s/   € à—’•R”Y¨�Ñ-Ô-°Ñ3×7Ò7½¼ÑAÔAÐAr   r/   c           	      ó  — |dz  t          j        dd¦  «        z   }t          t          j        | |d d …d f         z   |d d d …f         z   |d d …d f         d¬¦  «        ¦  «        |d d …d f         z                       t           j        ¦  «        S ©Né    r   ©rE   Úother)r   rg   r   rC   r   r—   )ÚQÚrow_offsetsÚ
valid_rowsru   r/   Údimss         r   Ú_query_blockr¤   ˜   sÑ   € à�2‰v�œ	 ! RÑ(Ô(Ñ(ˆÝ!¥"¤'¨!¨k¸!¸!¸!¸T¸'Ô.BÑ*BÀTÈ$ÐPQÐPQÐPQÈ'Ä]Ñ*RÐYcÐdeÐdeÐdeÐgkÐdkÔYlÐtuÐ"vÑ"vÔ"vÑwÔwÐzð  ABð  ABð  ABð  DHð  AHô  {Iñ  I÷  Mò  Mõ  NPô  NXñ  Yô  Yð  	Yr   c                 ó˜  — t          j        ||d d d …f         z   |dz  z   t          j        dd¦  «        d d …d f         z   |d d d …f         d¬¦  «                             t           j        ¦  «        }t          j        | |¦  «        t          j        ||z   |z   |d¬¦  «                             t           j        ¦  «        d d d …f         z  S rœ   )r   rC   rg   r   r—   Údotr   )ÚqÚKÚKSÚoffsetsÚscale_offsetsÚvalid_columnsr/   Úks           r   Ú_scores_blockr®   �   sØ   € õ ŒG�A˜  a a a Ô(Ñ(¨1¨r©6Ñ1µB´I¸aÀÑ4DÔ4DÀQÀQÀQÈÀWÔ4MÑMÐTaÐbfÐhiÐhiÐhiÐbiÔTjÐrsÐtÑtÔt×wÒwÕxzô  yCñ  Dô  DˆÝŒv�a˜‰|Œ|�bœg b¨=Ñ&8¸1Ñ&<À=ÐXYÐZÑZÔZ×]Ò]Õ^`Ô^hÑiÔiÐjnÐpqÐpqÐpqÐjqÔrÑrÐrr   c	                 óì  — |t          j        ||z   |z   |d¬¦  «                             t           j        ¦  «        d d d …f         z  }	t          j        t          t          j        |	d¦  «        ¦  «        d¦  «        }
t          |	t          |
d¦  «        d d …d f         z  ¦  «        \  }}t          j        ||d d …d f         z   |dz  z   t          j	        dd¦  «        d d d …f         z   |d d …d f         d¬¦  «                             t           j
        ¦  «        }t          j        ||¦  «        }t          j        |||¦  «        }| |d d …d f         z  |t          d|
z
  d¦  «        d d …d f         z  z   S )Nr   rž   r   r�   rb   )r   rC   r   r   rl   rš   rj   r˜   r•   rg   r—   r¦   )ÚaccÚpÚ
correctionÚVÚVSrª   r«   r¬   r/   Úweightsrt   r>   r?   ÚvÚweighteds                  r   Ú_values_blockr¸   £   sƒ  € ð •b”g˜b =Ñ0°1Ñ4¸=ÐPQÐRÑRÔR×UÒUÕVXÔV`ÑaÔaÐbfÐhiÐhiÐhiÐbiÔjÑjˆÝŒj�¥2¤6¨'°1Ñ#5Ô#5Ñ6Ô6¸Ñ:Ô:ˆÝ˜G¥m°C¸Ñ&;Ô&;¸A¸A¸A¸t¸GÔ&DÑDÑEÔE‰ˆˆBÝŒG�A˜    4 Ô(Ñ(¨1¨r©6Ñ1µB´I¸aÀÑ4DÔ4DÀTÈ1È1È1ÀWÔ4MÑMÐTaÐbcÐbcÐbcÐeiÐbiÔTjÐrsÐtÑtÔt×wÒwÕxzô  yCñ  Dô  DˆÝ”6˜"˜a‘=”=ˆÝ”6˜"˜a Ñ*Ô*ˆØ�Z    4 Ô(Ñ(¨8µmÀCÈ#ÁIÈqÑ6QÔ6QÐRSÐRSÐRSÐUYÐRYÔ6ZÑ+ZÑZÐZr   c                 ó´   — t          j        | |d d …d f         dz  z   |dz  z   t          j        dd¦  «        d d d …f         z   ||d d …d f         ¬¦  «         d S )Né   r�   r   ©rE   )r   ri   rg   )ÚPARTIALrŽ   r¢   r°   r/   s        r   Ú_store_blockr½   ®   ss   € å
Œ�˜5    D œ>¨CÑ/Ñ/°!°b±&Ñ8½2¼9ÀQÈÑ;KÔ;KÈDÐRSÐRSÐRSÈGÔ;TÑTÐVYÐ`jÐklÐklÐklÐnrÐkrÔ`sÐtÑtÔtÐtÐtÐtr   )Ú	Q_PER_SEQÚTABLE_WIDTHÚH_QÚH_KVÚPAGEÚSPLITSÚMÚNÚSCALEc                 ó,  — t          j        d¦  «        t          j        d¦  «        t          j        d¦  «        }}}||z  ||z  }}||z  }||z  t          j        d|¦  «        z   }||
z  ||z  z   }||z  ||z  z   }||
|z  k     }||z  |z   dz  }t          t          j        | |d d …d f         z   t          j        dd¦  «        d d d …f         z   |d d …d f         d¬¦  «        ¦  «        }t          j        t          j        t          t          j        t          j	        |¦  «        d¦  «        ¦  «        d¦  «        d¦  «        }t          |d¦  «        } |t          d	|z
  d
¦  «        z  }!t          | ||| d¦  «        }"t          | ||| d¦  «        }#t          | ||| d¦  «        }$t          | ||| d¦  «        }%t          | ||| d¦  «        }&t          | ||| d¦  «        }'t          | ||| d¦  «        }(t          | ||| d¦  «        })t          j        ||z   ¦  «        }*|*|
z
  }+t          j        t          j        |*|¦  «        |¦  «        },||,z  |z  }-t          j        |*|dz   |,z  |z  ¦  «        }.t          j        |fdt           j        ¦  «        }/t          j        |ft           j        ¦  «        }0t          j        |dft           j        ¦  «        }1t          j        |dft           j        ¦  «        }2t          j        |dft           j        ¦  «        }3t          j        |dft           j        ¦  «        }4t          j        |dft           j        ¦  «        }5t          j        |dft           j        ¦  «        }6t          j        |dft           j        ¦  «        }7t          j        |dft           j        ¦  «        }8t          j        |-|.|¦  «        D �]ë}9|9t          j        d|¦  «        z   }:|:|.k     };t          j        |||z  z   |:|z  z   |;d¬¦  «        }<|<|z  |:|z  z                        t           j        ¦  «        }=|=|z  |z   dz  }>|=|z  |z   dz  }?t'          |"|||>|?|;d¦  «        }@|@t'          |#|||>|?|;d¦  «        z  }@|@t'          |$|||>|?|;d¦  «        z  }@|@t'          |%|||>|?|;d¦  «        z  }@|@t'          |&|||>|?|;d¦  «        z  }@|@t'          |'|||>|?|;d¦  «        z  }@|@t'          |(|||>|?|;d¦  «        z  }@|@t'          |)|||>|?|;d¦  «        z  }@|d d …d f         |;d d d …f         z  |:d d d …f         |+|d d …d f         |z  z   k    z  }At          j        |A|@|!d d …d f         z  d¦  «        }@t          j        |/t          j        |@d¦  «        ¦  «        }Bt          j        |/|Bz
  dz  ¦  «        }Ct          j        |At          j        |@|Bd d …d f         z
  dz  ¦  «        d¦  «        }D|0|Cz  t          j        |Dd¦  «        z   }0t/          |1|D|C|||>|?|;d¦	  «	        }1t/          |2|D|C|||>|?|;d¦	  «	        }2t/          |3|D|C|||>|?|;d¦	  «	        }3t/          |4|D|C|||>|?|;d¦	  «	        }4t/          |5|D|C|||>|?|;d¦	  «	        }5t/          |6|D|C|||>|?|;d¦	  «	        }6t/          |7|D|C|||>|?|;d¦	  «	        }7t/          |8|D|C|||>|?|;d¦	  «	        }8|B}/�Œí||z  |z   |z  |z   }Et1          ||E||1d¦  «         t1          ||E||2d¦  «         t1          ||E||3d¦  «         t1          ||E||4d¦  «         t1          ||E||5d¦  «         t1          ||E||6d¦  «         t1          ||E||7d¦  «         t1          ||E||8d¦  «         t          j        ||Ez   |/|¬¦  «         t          j        |	|Ez   |0|¬¦  «         d S )Nr   r   rc   rº   rž   é   ra   é   rb   iòÿÿÿé   é   r   é   r   g@Œµx¯Är�   é   çþ‚+eG÷?r<   r»   )r   re   rg   r   rC   rk   rl   rš   rj   r(   r•   r¤   ÚcdivrR   r   ÚzerosÚranger   rf   r®   r    Úexp2rQ   r¸   r½   ri   )Fr    r¨   r³   r©   r´   ÚTABLESÚLENGTHSr¼   ÚMAXIMUMÚSUMr¾   r¿   rÀ   rÁ   rÂ   rÃ   rÄ   rÅ   rÆ   ÚsequenceÚkv_headÚtileÚsplitÚquery_blockÚgroupÚrowsÚquery_indexÚ
query_headr¢   r¡   r§   rt   Úq_scaleÚ	row_scaleÚq0Úq1Úq2Úq3Úq4Úq5Úq6Úq7Úcontext_lenÚ
prefix_lenÚtiles_per_splitÚbeginÚendrl   ÚdenominatorÚacc0Úacc1Úacc2Úacc3Úacc4Úacc5Úacc6Úacc7ÚstartÚcolumnsr¬   ÚpageÚslotrª   r«   ÚscoresÚvalidÚnew_maximumr²   r±   rŽ   sF                                                                         r   Ú_q8_attention_partials_kernelrÿ   ²   s  € õ #%¤-°Ñ"2Ô"2µB´MÀ!Ñ4DÔ4DÅbÄmÐTUÑFVÔFV˜4�'ˆØ! F™]¨D°F©Nˆ{ˆØ! T™kˆØ˜Q‰¥¤¨1¨a¡¤Ñ0ˆØ Ñ*¨T°U©]Ñ:ˆØ˜u‘_ t¨e¡|Ñ3ˆ
Ø˜I¨Ñ-Ò-ˆ
Ø" SÑ(¨:Ñ5¸Ñ<ˆÝ�bœg a¨+°a°a°a¸°gÔ*>Ñ&>ÅÄÈ1ÈcÑARÔARÐSWÐYZÐYZÐYZÐSZÔA[Ñ&[ÐblÐmnÐmnÐmnÐptÐmtÔbuÐ}~ÐÑÔñ  Aô  AˆÝŒj�œ¥I­b¬fµR´V¸A±Y´YÀÑ.BÔ.BÑ$CÔ$CÀRÑHÔHÈ#ÑNÔNˆÝ  RÑ(Ô(ˆØ�M¨#°©)°SÑ9Ô9Ñ9ˆ	Ý˜!˜[¨*°g¸qÑAÔAˆÝ˜!˜[¨*°g¸qÑAÔAˆÝ˜!˜[¨*°g¸qÑAÔAˆÝ˜!˜[¨*°g¸qÑAÔAˆÝ˜!˜[¨*°g¸qÑAÔAˆÝ˜!˜[¨*°g¸qÑAÔAˆÝ˜!˜[¨*°g¸qÑAÔAˆÝ˜!˜[¨*°g¸qÑAÔAˆÝ”g˜g¨Ñ0Ñ1Ô1ˆØ  9Ñ,ˆ
Ýœ'¥"¤'¨+°qÑ"9Ô"9¸6ÑBÔBˆØ˜Ñ'¨!Ñ+ˆÝŒj˜ u¨q¡y°OÑ&CÀaÑ&GÑHÔHˆÝ”'˜1˜$ ­¬Ñ4Ô4ˆÝ”h ˜t¥R¤ZÑ0Ô0ˆÝŒx˜˜B˜¥¤Ñ,Ô,ˆÝŒx˜˜B˜¥¤Ñ,Ô,ˆÝŒx˜˜B˜¥¤Ñ,Ô,ˆÝŒx˜˜B˜¥¤Ñ,Ô,ˆÝŒx˜˜B˜¥¤Ñ,Ô,ˆÝŒx˜˜B˜¥¤Ñ,Ô,ˆÝŒx˜˜B˜¥¤Ñ,Ô,ˆÝŒx˜˜B˜¥¤Ñ,Ô,ˆÝ”X˜e S¨!Ñ,Ô,ð 	"ñ 	"ˆEØ�bœi¨¨1™oœoÑ-ˆGØ# cšMˆMÝ”7˜6 H¨{Ñ$:Ñ:¸WÈ¹_ÑLÐS`ÐhiÐjÑjÔjˆDØ˜4‘K '¨D¡.Ñ0×4Ò4µR´XÑ>Ô>ˆDØ˜d‘{ WÑ,°Ñ3ˆGØ! D™[¨7Ñ2°aÑ7ˆMÝ" 2 q¨"¨g°}ÀmÐUVÑWÔWˆFØ•m B¨¨2¨w¸À}ÐVWÑXÔXÑXˆFØ•m B¨¨2¨w¸À}ÐVWÑXÔXÑXˆFØ•m B¨¨2¨w¸À}ÐVWÑXÔXÑXˆFØ•m B¨¨2¨w¸À}ÐVWÑXÔXÑXˆFØ•m B¨¨2¨w¸À}ÐVWÑXÔXÑXˆFØ•m B¨¨2¨w¸À}ÐVWÑXÔXÑXˆFØ•m B¨¨2¨w¸À}ÐVWÑXÔXÑXˆFØ˜q˜q˜q $˜wÔ'¨-¸¸a¸a¸a¸Ô*@Ñ@ÀGÈDÐRSÐRSÐRSÈGÔDTÐXbÐeiÐjkÐjkÐjkÐmqÐjqÔerÐv{Ñe{ÑX{ÒD{Ñ|ˆEÝ”X˜e V¨i¸¸¸¸4¸Ô.@Ñ%@À'ÑJÔJˆFÝœ* W­b¬f°V¸QÑ.?Ô.?Ñ@Ô@ˆKÝœ '¨KÑ"7Ð;MÑ!MÑNÔNˆJÝ”˜¥¤¨°+¸a¸a¸aÀ¸gÔ2FÑ)FÐJ\Ñ(\Ñ ]Ô ]Ð_bÑcÔcˆAØ%¨
Ñ2µR´V¸A¸q±\´\ÑAˆKÝ   q¨*°a¸¸WÀmÐUbÐdeÑfÔfˆDÝ   q¨*°a¸¸WÀmÐUbÐdeÑfÔfˆDÝ   q¨*°a¸¸WÀmÐUbÐdeÑfÔfˆDÝ   q¨*°a¸¸WÀmÐUbÐdeÑfÔfˆDÝ   q¨*°a¸¸WÀmÐUbÐdeÑfÔfˆDÝ   q¨*°a¸¸WÀmÐUbÐdeÑfÔfˆDÝ   q¨*°a¸¸WÀmÐUbÐdeÑfÔfˆDÝ   q¨*°a¸¸WÀmÐUbÐdeÑfÔfˆDØ!ˆG‰GØ˜sÑ" ZÑ/°6Ñ9¸EÑAˆÝ�W˜e Z°°qÑ9Ô9Ð9Ý�W˜e Z°°qÑ9Ô9Ð9Ý�W˜e Z°°qÑ9Ô9Ð9Ý�W˜e Z°°qÑ9Ô9Ð9Ý�W˜e Z°°qÑ9Ô9Ð9Ý�W˜e Z°°qÑ9Ô9Ð9Ý�W˜e Z°°qÑ9Ô9Ð9Ý�W˜e Z°°qÑ9Ô9Ð9Ý
Œ�˜5‘ '°
Ð;Ñ;Ô;Ð;Ý
Œ��u‘˜k°
Ð;Ñ;Ô;Ð;Ð;Ð;r   ÚDc                 óÌ  — t          j        d¦  «        }t          j        d|¦  «        }t          j        d|¦  «        }t          j        |||z  z   |z   ¦  «        }	t          j        |	t          j        |	d¦  «        z
  dz  ¦  «        }
t          j        t          j        |||z  z   |z   ¦  «        |
z  d¦  «        }t          j        | ||z  |d d …d f         z   |z  z   |d d d …f         z   ¦  «        }t          j        |||z  z   |z   t          t          j	        t          j        ||
d d …d f         z  d¦  «        t          j
        |d¦  «        ¦  «        ¦  «        ¦  «         d S )Nr   rÎ   g#B’¡œÇ;)r   re   rg   rC   rÒ   rj   rQ   ri   r%   rP   rl   )r¼   rÕ   rÖ   r�   rÃ   r   rs   Úsplitsr£   rl   rµ   rï   Úpartials                r   Ú_q8_attention_reduce_kernelr    s€  € åŒm˜AÑÔˆÝ”˜1˜fÑ%Ô%ˆÝŒy˜˜A‰ŒˆÝ”'˜' C¨&¡LÑ0°6Ñ9Ñ:Ô:ˆÝ”'˜7¥R¤V¨G°QÑ%7Ô%7Ñ7Ð;MÑMÑNÔNˆÝ”f�RœW S¨3°©<Ñ%7¸&Ñ%@ÑAÔAÀGÑKÈQÑOÔOˆÝ”'˜' S¨6¡\°F¸1¸1¸1¸d¸7´OÑ%CÀqÑ$HÑHÈ4ÐPTÐVWÐVWÐVWÐPWÌ=ÑXÑYÔYˆÝ
Œ��s˜Q‘w‘ Ñ%Õ'8½¼Å2Ä6È'ÐT[Ð\]Ð\]Ð\]Ð_cÐ\cÔTdÑJdÐfgÑChÔChÕjlÔjtð  vAð  CJñ  kKô  kKñ  :Lô  :Lñ  (Mô  (Mñ  	Nô  	Nð  	Nð  	Nð  	Nr   c                 ón   — | j         t          j        k    r|                      t          j        ¦  «        n| S r-   )Údtyper	   Úbfloat16Úviewr"   )Útensors    r   Ú_bitsr
    s)   € Ø'-¤|µu´~Ò'EÐ'Eˆ6�;Š;•u”{Ñ#Ô#Ð#È6ÐQr   c                 ó¨  — |                       d| j        d         ¦  «        }|€|n |                      d|j        d         ¦  «        }t          j        |j        | j        | j        ¬¦  «        }|�t          j        |¦  «        n|}t          |j        d         f         t          |¦  «        t          |¦  «        t          |¦  «        t          |¦  «        t          |¦  «        | 	                    d¦  «        | 	                    d¦  «        | 	                    d¦  «        | 	                    d¦  «        |j        d         ||du|j        t          j
        k    t          j        |j        d         ¦  «        |j        d         dk    rdndd¬	¦  «         |€|                     | ¦  «        S |                     | ¦  «        |                     | ¦  «        fS )
zXBF16 x (and residual) with unit last stride: y, or (y, x + residual) as RMSNorm.forward.éÿÿÿÿN©Údevicer  r   r   i   rÍ   F)Ú	num_warpsÚenable_fp_fusion)ÚreshapeÚshaper	   Úemptyr  r  Ú
empty_liker�   r
  Ústrider  r   Únext_power_of_2Úview_as)r#   ÚresidualÚweightrU   Úx_2dÚresidual_2dr€   Úresidual_outs           r   Úrms_normr    s¶  € à�9Š9�R˜œ œÑ%Ô%€DØ"Ð*�$�$°×0@Ò0@ÀÀXÄ^ÐTVÔEWÑ0XÔ0X€KÝŒ�D”J q¤x°q´wÐ?Ñ?Ô?€AØ*2Ð*>•5Ô# AÑ&Ô&Ð&ÀA€LÝ�d”j ”mÐ%Ô&¥u¨T¡{¤{µE¸+Ñ4FÔ4FÍÈfÉÌÕW\Ð]^ÑW_ÔW_ÕafÐgsÑatÔatØ'+§{¢{°1¡~¤~°{×7IÒ7IÈ!Ñ7LÔ7LÈaÏhÊhÐWXÉkÌkÐ[g×[nÒ[nÐopÑ[qÔ[qÐswÔs}Ð~ô  tAð  CFØ'/°tÐ';¸V¼\ÍUÌ^Ò=[Õ]cÔ]sÐtxÔt~ð  @Aô  uBñ  ^Cô  ^CØ6:´jÀ´mÀdÒ6JÐ6J°°ÐPQÐdiðkñ kô kð kð ÐØ�yŠy˜‰|Œ|ÐØ�9Š9�Q‰<Œ<˜×-Ò-¨aÑ0Ô0Ð0Ð0r   c                 óÔ  — t          j        | j        | j        | j        ¬¦  «        }t          | j        d         f         t          | ¦  «        t          |¦  «        t          |¦  «        t          |¦  «        |                      d¦  «        |                     d¦  «        | j        d         ||j        t           j        k    t          j
        | j        d         ¦  «        d¬¦  «         |S )z`BF16 x and gate of shape (rows, dim) with unit last stride, as Qwen3_5RMSNormGated.forward_list.r  r   r   ©r  )r	   r  r  r  r  rˆ   r
  r  r  r   r  )r#   r†   r  rU   r€   s        r   Úgated_rms_normr   '  sÔ   € åŒ�A”G A¤H°A´GÐ<Ñ<Ô<€AÝ˜AœG AœJ˜=Ô)­%°©(¬(µE¸$±K´KÅÀvÁÄÕPUÐVWÑPXÔPXÐZ[×ZbÒZbÐcdÑZeÔZeÐgk×grÒgrÐstÑguÔguÐwxÔw~ð  @Aô  xBð  DGØ*0¬,½%¼.Ò*HÍ&ÔJ`ÐabÔahÐijÔakÑJlÔJlÐxyð{ñ {ô {ð {à€Hr   c                 óœ  — |                       ¦   «         } t          j        g | j        dd…         ¢| j        d         dz  ‘R | j        | j        ¬¦  «        }t          t          j        | 	                    ¦   «         d¦  «        f         t          | ¦  «        t          |¦  «        |j        d         | 	                    ¦   «         d¦  «         |S )zABF16 x = [gate, value] on the last dimension: silu(gate) * value.Nr  rc   r  rº   )Ú
contiguousr	   r  r  r  r  r’   r   rÏ   Únumelr
  )r#   Úoutputs     r   Úsilu_mulr%  /  s©   € à	�Š‰Œ€AÝŒ[Ð:˜1œ7 3 B 3œ<Ð:¨¬°¬¸Ñ)9Ð:Ð:À1Ä8ÐSTÔSZÐ[Ñ[Ô[€FÝ•f”k &§,¢,¡.¤.°#Ñ6Ô6Ð8Ô9½%À¹(¼(ÅEÈ&ÁMÄMÐSYÔS_Ð`bÔScÐek×eqÒeqÑesÔesÐuxÑyÔyÐyØ€Mr   r�   r   é@   c                 óî  — | j         d         |                     ¦   «         z  }t          j        || j         d         |j         d         z  z  t          ¦  «        }	t          j        g | j         dd…         ¢t          ‘| j         d         ‘R | j        t
          j	        ¬¦  «        }
t          j        g | j         dd…         ¢t          ‘R | j        t
          j	        ¬¦  «        }t          j
        |¦  «        }t          j
        | ¦  «        }t          |                     ¦   «         |j         d         |	t          z  f         t          | ¦  «        |||||||
||||j         d         | j         d         |j         d         |j         d         t          t          t          |dd¬¦  «         t          | j         d         | j         d         z  f         |
||t          |¦  «        t          | j         d         d¬¦  «         |S )	zöBF16 q (sequences * queries, heads, 256) against the paged Q8 cache (int8 values, FP16 scales per 32 values): BF16 output of q's
    shape. context_lens counts each sequence's keys including its queries, which attend causally (decode: one query).r   r   rc   Nr  )r  Ú
num_stagesrË   r  )r  r#  r   rÏ   ÚQ8_ATTENTION_ROWSr	   r  ÚQ8_ATTENTION_SPLITSr  r   r  rÿ   r
  ÚQ8_ATTENTION_KEYSr  )r§   Úk_cacheÚv_cacheÚk_scaleÚv_scaleÚblock_tablesÚcontext_lensÚsoftmax_scaleÚqueries_per_sequenceÚquery_blocksr  rl   rï   r$  s                 r   Úq8_paged_attentionr5  <  sð  € ð œ7 1œ:¨×);Ò);Ñ)=Ô)=Ñ=ÐÝ”;Ð3°q´w¸q´zÀWÄ]ÐSTÔEUÑ7UÑVÕXiÑjÔj€LÝŒkÐI˜AœG B Q BœKÐIÕ)<ÐI¸a¼gÀa¼jÐIÐIÐRSÔRZÕbgÔboÐpÑpÔp€GÝŒkÐ=˜AœG B Q BœKÐ=Õ)<Ð=Ð=ÀaÄhÕV[ÔVcÐdÑdÔd€GÝÔ" 7Ñ+Ô+€KÝÔ˜aÑ Ô €FÝ! <×#5Ò#5Ñ#7Ô#7¸¼ÀqÔ9IÈ<ÕZmÑKmÐ"nÔoÝˆa‰Œ�'˜7 G¨W°lÀLÐRYÐ[bÐdoð  rFØÔ˜1Ô˜qœw qœz¨7¬=¸Ô+;¸W¼]È1Ô=MÕObÕduõ  xIØ ¨qð2ñ 2ô 2ð 2õ   ¤¨¤¨a¬g°a¬jÑ!8Ð :Ô;¸GÀWÈkÕ[`ÐagÑ[hÔ[hÕj}ð  @Aô  @Gð  HIô  @Jð  VWð  Xñ  Xô  Xð  XØ€Mr   )Ú!install_triton_compilation_logger).Ú__doc__Ú	functoolsr	   r   Útriton.languageÚlanguager   Útriton.language.extrar   ÚImportErrorÚcacheÚboolr   Újitr   r%   r+   r2   r9   r@   rF   Ú	constexprrK   rW   r�   rˆ   r’   r•   r˜   rš   r¤   r®   r¸   r½   rÿ   r  r
  r  r   r%  r*  r)  r+  r5  Ú%shared.kernels.triton_compilation_logr6  r   r   r   ú<module>rB     sX  ððEð Eð Ð Ð Ð à €€€ðØ€M€M€MØ Ð Ð Ð Ð Ð Ø/Ð/Ð/Ð/Ð/Ð/Ð/øØð ð ð Ø€F€F€Fðøøøð „ðK�$ð Kð Kð Kñ „ðKð
 
ÑØ„ZðFð Fñ „ZðFð „Zð@ð @ñ „Zð@ð
 „ZðDð Dñ „ZðDð „Zð,ð ,ñ „Zð,ð
 „Zð$ð $ñ „Zð$ð „ZðDð Dñ „ZðDð „ZðGð Gñ „ZðGð „Zð9 r¤|ð 9ð 9ð 9ñ „Zð9ð
 „ZðZð Zñ „ZðZð €V„ZÐ"SÐTÑTÔTðeØ')¤|ðeØ=?¼\ðeØRTÔR^ðeð eð eñ UÔTðeð@ €V„ZÐ":Ð;Ñ;Ô;ð	FÐTVÔT`ð 	FÐikÔiuð 	Fð 	Fð 	Fñ <Ô;ð	Fð €V„Z *ÈZÐXÑXÔXðS¨"¬,ð SÀbÄlð Sð Sð Sñ YÔXðSð „ZðI 2¤<ð Ið Ið Iñ „ZðIð „Zð:ð :ñ „Zð:ð „ZðBð Bñ „ZðBð „ZðY¸2¼<ð Yð Yð Yñ „ZðYð „ZðsÈ"Ì,ð sð sð sñ „Zðsð
 „Zð[Ð[]Ô[gð [ð [ð [ñ „Zð[ð „Zðu¸¼ð uð uð uñ „Zðuð €V„ZÐ">Ð_{Ð|Ñ|Ô|ðS<Ø+-¬<ðS<Ø?A¼|ðS<ØSUÔS_ðS<ØikÔiuðS<à)+¬ðS<à:<¼,ðS<àOQÌ|ðS<ð S<ð S<ñ }Ô|ðS<ðj „ZðNÈÌð NÐY[ÔYeð Nð Nð Nñ „ZðNðRð Rð Rð1ð 1ð 1ðð ð ðð ð ð Ð ØÐ ØÐ ðð ð ð$ 
ÐØWÐWÐWÐWÐWÐWØ%Ð%Ñ'Ô'Ð'Ð'Ð'ð Ðs   Œ �'¦'