§
     ’Þj4  ã            	       óÀ  — 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	  e	ddd¦  «        Z
 ej        g d¢g d¢¬	¦  «        d
ej        dej        dej        fd„¦   «         Z ej        dgdg¬	¦  «        d
ej        dej        fd„¦   «         Z ej        dgdg¬	¦  «        d
ej        dej        fd„¦   «         Z eed¦  «        Zd„ Zd„ Zd„ Z ej        dgdg¬	¦  «        d
ej        dej        fd„¦   «         Z ej        dgdg¬	¦  «        d
ej        dej        fd„¦   «         Zd„ Zej        d
ej        dej        dej        fd„¦   «         Z ej        g d¢g d¢¬	¦  «        d
ej        dej        dej        fd„¦   «         ZdZdZd„ Zd„ Zd „ Z ej        g d!¢g d!¢¬	¦  «        d"ej        d#ej        d$ej        dej        fd%„¦   «         Zd&„ Z dd'l!m"Z"  e"¦   «          dS )(zIH3 decoder inference fusions. Preserve the eager path's dtype boundaries.é    N)Ú	libdevice)ÚLaunchTunerÚwan2gp_h3_vae_kernelsé   zH3 VAE)ÚROWSÚTOKENSÚHEADSÚSBÚSTÚSHÚCBÚCT)Údo_not_specializeÚdo_not_specialize_on_alignmentÚDÚRÚBLOCKc           	      ó>  — t          j        d¦  «        |z  t          j        d|¦  «        z   }t          j        d|¦  «        }||z  |z  }|||z  z  }||z  ||z  z   ||z  |	z  z   }| }|}t          j        ||d d …d f         z   |d d d …f         z   |d d …d f         |k     d¦  «                             t           j        ¦  «        }|                     |j        j        ¦  «                             t           j        ¦  «        }|�rWt          j        ||dz  k     ||dz  z   ||dz  z
  ¦  «        }t          j        ||k     ||¦  «        }t          j	        |t          j
        |d d d …f         ||f¦  «        d¦  «        }|d d …d f         |
z  |d d …d f         |z  z   |d d d …f         z   }t          j        ||z   |d d …d f         |k     |d d d …f         |k     z  d¦  «                             |j        j        ¦  «                             t           j        ¦  «        }t          j        ||z   |d d …d f         |k     |d d d …f         |k     z  d¦  «                             |j        j        ¦  «                             t           j        ¦  «        }||z                       |j        j        ¦  «                             t           j        ¦  «        }||z                       |j        j        ¦  «                             t           j        ¦  «        }t          j        |d d d …f         |dz  k     ||z
  ||z   ¦  «        }t          j        |d d d …f         |k     ||¦  «        }t          j        ||d d …d f         |z  z   |d d d …f         z   ||d d …d f         |k     ¦  «         d S )Nr   é   é   )ÚtlÚ
program_idÚarangeÚloadÚtoÚfloat32ÚdtypeÚ
element_tyÚwhereÚgatherÚbroadcast_toÚstore)ÚXÚOUTÚCÚSr   r   r	   r
   r   r   r   r   r   r   r   ÚrowsÚcolsÚtokenÚbatchÚoffsetsÚptrÚoutÚxÚpartnerÚotherÚcoÚcosÚsinÚaÚbÚrotateds                                  úCC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\video_triton.pyÚ
_cast_roper8      s|  € õ Œ=˜ÑÔ˜eÑ#¥b¤i°°5Ñ&9Ô&9Ñ9€DÝŒ9�Q˜‰?Œ?€DØ�E‰M˜FÑ"€EØ�V˜e‘^Ñ$€EØ�b‰j˜5 2™:Ñ%¨°©¸Ñ(;Ñ;€GØ
€CØ
€CÝ
Œ��g˜a˜a˜a ˜gÔ&Ñ&¨¨d°A°A°A¨g¬Ñ6¸¸Q¸Q¸QÀ¸W¼ÈÒ8LÈaÑPÔP×SÒSÕTVÔT^Ñ_Ô_€Að 	
�ŠˆSŒYÔ!Ñ"Ô"×%Ò%¥b¤jÑ1Ô1€AØñ 
4Ý”(˜4 ! q¡&š=¨$°°a±©-¸ÀÀQÁ¹ÑGÔGˆÝ”(˜4 !š8 W¨dÑ3Ô3ˆÝ”	˜!�Rœ_¨W°T¸1¸1¸1°WÔ-=ÀÀq¸zÑJÔJÈAÑNÔNˆØ�1�1�1�d�7Œ^˜bÑ  5¨¨¨¨D¨¤>°BÑ#6Ñ6¸¸dÀAÀAÀA¸g¼ÑFˆÝŒg�a˜"‘f˜t A A A t Gœ}¨tÒ3¸¸TÀ1À1À1¸W¼ÈÒ8IÑJÈAÑNÔN×QÒQÐRUÔR[ÔRfÑgÔg×jÒjÕkmÔkuÑvÔvˆÝŒg�a˜"‘f˜t A A A t Gœ}¨tÒ3¸¸TÀ1À1À1¸W¼ÈÒ8IÑJÈAÑNÔN×QÒQÐRUÔR[ÔRfÑgÔg×jÒjÕkmÔkuÑvÔvˆØ�‰W�LŠL˜œÔ-Ñ.Ô.×1Ò1µ"´*Ñ=Ô=ˆØ�S‰[×Ò˜SœYÔ1Ñ2Ô2×5Ò5µb´jÑAÔAˆÝ”(˜4  a a a œ=¨1°©6Ò1°1°q±5¸!¸a¹%Ñ@Ô@ˆÝŒH�T˜$   ˜'”] QÒ&¨°Ñ3Ô3ˆÝ„HˆS�4˜˜˜˜4˜”= 1Ñ$Ñ$ t¨D°!°!°!¨G¤}Ñ4°a¸¸a¸a¸aÀ¸g¼ÈÒ9MÑNÔNÐNÐNÐNó    ÚNc                 óZ  — t          j        d¦  «        |z  t          j        d|¦  «        z   }||z  d|z  z  ||z  z   }t          j        | |z   ||k     d¦  «                             t           j        ¦  «        }t          j        | |z   |z   ||k     d¦  «                             t           j        ¦  «        }|dt          j        | ¦  «        z   z                       | j        j        ¦  «                             t           j        ¦  «        }t          j	        | |z   ||z  ||k     ¦  «         d S ©Nr   r   ç      ð?©
r   r   r   r   r   r   Úexpr   r   r"   )	r#   r:   r   r   ÚiÚoffsetÚgateÚvalueÚ	activateds	            r7   Ú_swiglurE   *   sÿ   € å
Œ�aÑÔ˜5Ñ ¥2¤9¨Q°Ñ#6Ô#6Ñ6€AØ�1‰f˜˜Q™Ñ ! a¡%Ñ'€FÝŒ7�1�v‘:˜q 1šu aÑ(Ô(×+Ò+­B¬JÑ7Ô7€DÝŒG�A˜‘J ‘N A¨¢E¨1Ñ-Ô-×0Ò0µ´Ñ<Ô<€EØ˜�rœv t e™}œ}Ñ,Ñ-×1Ò1°!´'Ô2DÑEÔE×HÒHÍÌÑTÔT€IÝ„HˆQ�‰Z˜ UÑ*¨A°ªEÑ2Ô2Ð2Ð2Ð2r9   c                 ó®  — t          j        d¦  «                             t           j        ¦  «        |z  t          j        d|¦  «        z   }t          j        | |z   ||k     d¦  «                             t           j        ¦  «        }t          j        |||z  z   ¦  «                             t           j        ¦  «        }t          j        ||z   ||k     d¦  «                             t           j        ¦  «        }	||z                       | j        j        ¦  «                             t           j        ¦  «        }
t          j	        | |z   |
|	z   ||k     ¦  «         dS )zebranch * scale + residual with the two roundings of branch.mul_(scale).add_(residual): bit-identical.r   N)
r   r   r   Úint64r   r   r   r   r   r"   )r#   ÚSCALEÚRESIDUALr:   r   r   r@   r.   ÚscaleÚresidualÚscaleds              r7   Ú	_residualrM   4   s  € õ 	Œ�aÑÔ×Ò�BœHÑ%Ô%¨Ñ-µ´	¸!¸UÑ0CÔ0CÑC€AÝ
Œ��A‘�q˜1’u˜aÑ Ô ×#Ò#¥B¤JÑ/Ô/€AÝŒG�E˜A ™E‘MÑ"Ô"×%Ò%¥b¤jÑ1Ô1€EÝŒw�x !‘| Q¨¢U¨AÑ.Ô.×1Ò1µ"´*Ñ=Ô=€HØ�%‰i�^Š^˜AœGÔ.Ñ/Ô/×2Ò2µ2´:Ñ>Ô>€FÝ„HˆQ�‰U�F˜XÑ% q¨1¢uÑ-Ô-Ð-Ð-Ð-r9   r    c           	      ó¸  — | j         \  }}}}|€| | fn|\  }}|€dn|j         d         }t          j        | j         | j        t          j        ¬¦  «        }	t          t          j        ||z  |z  d¦  «        f         | |	||||z  |z  ||g|                      ¦   «         d d…         ¢|                     d¦  «        ‘|                     d¦  «        ‘|‘|‘d‘R ddiŽ |	S )	Nr   éÿÿÿÿ©Údevicer   é   é   r   Úenable_fp_fusionF)	ÚshapeÚtorchÚemptyrQ   Úbfloat16r8   ÚtritonÚcdivÚstride)
ÚqÚrotaryr5   ÚtÚheadsÚdimr2   r3   ÚrÚoqs
             r7   Ú	cast_roperc   B   sI  € Ø”wÑ€A€qˆ%�Ø˜��1ˆvˆv¨V�H€CˆØˆ^ˆˆ ¤¨2¤€AÝ	Œ�Q”W Q¤XµU´^Ð	DÑ	DÔ	D€BÝ•”˜A ™E E™M¨2Ñ.Ô.Ð0Ô1°!°R¸¸cÀ1ÀqÁ5È5Á=ÐRSÐUZð  oÐ]^×]eÒ]eÑ]gÔ]gÐhjÐijÐhjÔ]kð  oÐmp×mwÒmwÐxyÑmzÔmzð  oÐ|÷  }Gò  }Gð  HIñ  }Jô  }Jð  oð  LOð  oð  QRð  oð  TVð  oð  oð  oð  inð  oð  oð  oØ€Ir9   c                 ó  — | j         d         dz  }|                      ¦   «         dz  }t          t          j        |d¦  «        f         | ||dd¬¦  «         t
          j        j                             | ¦  «         | dd |…f         S )NrO   r   é   F©rT   .)	rU   ÚnumelrE   rY   rZ   rV   ÚautogradÚgraphÚincrement_version)Úexpandedr`   Úns      r7   Úswiglu_rm   K   s€   € Ø
Œ.˜Ô
 Ñ
!€CØ�ŠÑÔ˜AÑ€AÝ�VŒ[˜˜CÑ Ô Ð"Ô# H¨a°°cÈEÐRÑRÔRÐRÝ	„NÔ×*Ò*¨8Ñ4Ô4Ð4Ø�C˜˜#˜�IÔÐr9   c           	      óì   — |                       ¦   «         }t          t          j        |d¦  «        f         | |||| j        d         dd¬¦  «         t
          j        j                             | ¦  «         | S )Nre   rO   Frf   )	rg   rM   rY   rZ   rU   rV   rh   ri   rj   )ÚbranchrJ   rK   rl   s       r7   Ú	residual_rp   S   sg   € Ø�Š‰Œ€AÝ�vŒ{˜1˜cÑ"Ô"Ð$Ô% f¨e°X¸qÀ&Ä,ÈrÔBRÐTWÐjoÐpÑpÔpÐpÝ	„NÔ×*Ò*¨6Ñ2Ô2Ð2Ø€Mr9   c                 óZ  — t          j        d¦  «        |z  t          j        d|¦  «        z   }||z  d|z  z  ||z  z   }t          j        | |z   ||k     d¦  «                             t           j        ¦  «        }t          j        | |z   |z   ||k     d¦  «                             t           j        ¦  «        }|dt          j        | ¦  «        z   z                       | j        j        ¦  «                             t           j        ¦  «        }	t          j	        ||z   |	|z  ||k     ¦  «         d S r<   r>   ©
r#   r$   r:   r   r   r@   rA   rB   rC   rD   s
             r7   Ú_swiglu_contiguousrs   Z   sÿ   € å
Œ�aÑÔ˜5Ñ ¥2¤9¨Q°Ñ#6Ô#6Ñ6€AØ�1‰f˜˜Q™Ñ ! a¡%Ñ'€FÝŒ7�1�v‘:˜q 1šu aÑ(Ô(×+Ò+­B¬JÑ7Ô7€DÝŒG�A˜‘J ‘N A¨¢E¨1Ñ-Ô-×0Ò0µ´Ñ<Ô<€EØ˜�rœv t e™}œ}Ñ,Ñ-×1Ò1°!´'Ô2DÑEÔE×HÒHÍÌÑTÔT€IÝ„HˆS�1‰W�i %Ñ'¨¨QªÑ/Ô/Ð/Ð/Ð/r9   c                 ó´  — t          j        d¦  «                             t           j        ¦  «        |z  t          j        d|¦  «        z   }||z  d|z  z  ||z  z   }t          j        | |z   ||k     d¦  «                             t           j        ¦  «        }t          j        | |z   |z   ||k     d¦  «                             t           j        ¦  «        }t          j        |dt          j	        | ¦  «        z   ¦  «                             | j
        j        ¦  «                             t           j        ¦  «        }	t          j        ||z   |	|z  ||k     ¦  «         dS )z˜silu(gate) rounded to the activation dtype, then times the value: F.silu(gate).mul_(value) bit for bit (IEEE division and
    expf as PyTorch's kernel).r   r   r=   N)r   r   r   rG   r   r   r   Údiv_rnr   r?   r   r   r"   rr   s
             r7   Ú
_swiglu_rnrv   d   s  € õ 	Œ�aÑÔ×Ò�BœHÑ%Ô%¨Ñ-µ´	¸!¸UÑ0CÔ0CÑC€AØ�1‰f˜˜Q™Ñ ! a¡%Ñ'€FÝŒ7�1�v‘:˜q 1šu aÑ(Ô(×+Ò+­B¬JÑ7Ô7€DÝŒG�A˜‘J ‘N A¨¢E¨1Ñ-Ô-×0Ò0µ´Ñ<Ô<€EÝ”	˜$ ¥i¤m°T°EÑ&:Ô&:Ñ :Ñ;Ô;×>Ò>¸q¼wÔ?QÑRÔR×UÒUÕVXÔV`ÑaÔa€IÝ„HˆS�1‰W�i %Ñ'¨¨QªÑ/Ô/Ð/Ð/Ð/r9   c                 ó  — | j         d         dz  }t          j        g | j         dd…         ¢|‘R | j        | j        ¬¦  «        }|                     ¦   «         }t          t          j        |d¦  «        f         | |||dd¬¦  «         |S )ztRows of [gate | value] to silu(gate) * value in a new contiguous tensor (the twice-wide projection can be released).rO   r   NrP   i   Frf   )	rU   rV   rW   rQ   r   rg   rv   rY   rZ   ©rk   r`   r-   rl   s       r7   Ú	swiglu_rnry   p   sŠ   € à
Œ.˜Ô
 Ñ
!€CÝ
Œ+Ð1˜œ s¨ sÔ+Ð1¨SÐ1Ð1¸(¼/ÐQYÔQ_Ð
`Ñ
`Ô
`€CØ�	Š	‰Œ€AÝ•”˜A˜tÑ$Ô$Ð&Ô'¨°#°q¸#¸tÐV[Ð\Ñ\Ô\Ð\Ø€Jr9   ÚHALFc           
      ób  — t          j        d|¦  «        }| |dd…df         z   |ddd…f         z   }t          j        ||dd…df         d¦  «                             t           j        ¦  «        }t          j        | |dd…df         z   |ddd…f         z   |dd…df         d¦  «                             t           j        ¦  «        }|dk    rqt          j        t          j        t          j        ||
ddf¦  «        ddd¦  «        ¦  «        \  }}t          j        ||||z  ¦  «        }t          j	        |d¦  «        }nt          j	        ||z  d¦  «        }t          j        ||z  |z   ¦  «        dd…df         }||z                       | j        j        ¦  «                             t           j        ¦  «        }||z                       | j        j        ¦  «                             t           j        ¦  «        }|ddd…f         |	k     }t          j        t          j        |||¦  «        t          j        |||¦  «        t          j        || |¦  «        t          j        |||¦  «        z  ¦  «        }t          j        |t          j        |ddd…f         ||¦  «        |dd…df         ¦  «         dS )a=  Rows of one head, loaded whole: RMSNorm without weight in FP32 rounded to the activation dtype (the norm output), then the
    split-half rotation of the rotary channels in FP32 with one final rounding; the other channels keep the norm output. The partner
    channels are a second load of the same rows (cache hits).r   Né@   r   é    r   )r   r   r   r   r   ÚsplitÚtransÚreshapeÚfmaÚsumr   Úrsqrtr   r   r   r"   )r#   ÚbaseÚokr2   r3   r/   ÚkeepÚEPSr   rz   r   Úcr,   r.   r0   ÚfirstÚsecondÚsquaresÚtotalÚinvÚyÚ
first_halfr6   s                          r7   Ú_norm_rope_rowsr�   y   sf  € õ
 	Œ	�!�Q‰Œ€AØ
ˆd�1�1�1�d�7ŒmÑ
˜a  a a a œjÑ
(€CÝ
Œ��R˜˜˜˜4˜”[ !Ñ$Ô$×'Ò'­¬
Ñ3Ô3€AÝŒG�A˜˜Q˜Q˜Q ˜WœÑ%¨°°a°a°a°Ô(8Ñ8¸"¸Q¸Q¸QÀ¸W¼+ÀqÑIÔI×LÒLÍRÌZÑXÔX€EØˆB‚w€wõ œ¥¤­"¬*°Q¸ÀÀ2¸Ñ*GÔ*GÈÈAÈqÑ!QÔ!QÑRÔR‰ˆˆvÝ”&˜ ¨°©Ñ7Ô7ˆÝ”�w Ñ"Ô"ˆˆå”�q˜1‘u˜aÑ Ô ˆÝ
Œ/˜% !™) c™/Ñ
*Ô
*¨1¨1¨1¨d¨7Ô
3€CØ	
ˆS‰�Š�Q”WÔ'Ñ(Ô(×+Ò+­B¬JÑ7Ô7€AØ�S‰[×Ò˜QœWÔ/Ñ0Ô0×3Ò3µB´JÑ?Ô?€Eð �4˜˜˜�7”˜dÒ"€JÝŒf•R”X˜j¨#¨sÑ3Ô3µR´X¸jÈ!ÈUÑ5SÔ5SÝ”X˜j¨3¨$°Ñ4Ô4µr´xÀ
ÈEÐSTÑ7UÔ7UÑUñWô W€Gå„HˆS•"”(˜4  a a a œ=¨!¨WÑ5Ô5°r¸!¸!¸!¸T¸'´{ÑCÔCÐCÐCÐCr9   )r   r	   ÚSEQÚQBÚQSÚQHÚKBÚKSÚKHr   ÚCSc                 ó  — t          j        d¦  «                             t           j        ¦  «        |z  t          j        d|¦  «        z   }||k     }||z  }||z  }||z  }||z  }t          j        d|¦  «        }|d|z  k    }t          j        ||k     ||z   t          j        ||||z
  ¦  «        ¦  «        }|d d …d f         |d d d …f         d|z  k     z  }|d d …d f         |z  |d d …d f         |z  z   |d d d …f         z   }t          j        ||z   |d¦  «        }t          j        ||z   |d¦  «        }t          | ||z  ||z  z   ||	z  z   |||||||||¦  «         t          |||
z  ||z  z   ||z  z   |||||||||¦  «         d S )Nr   r   )r   r   r   rG   r   r   r   r�   ) ÚQÚKÚCOSÚSINr   r	   r‘   r’   r“   r”   r•   r–   r—   r   r˜   r‡   r   rz   r   r'   r…   Úheadr)   Úseqr*   rˆ   r†   r/   r6   Úangler2   r3   s                                    r7   Ú_qk_rms_roper¡   •   sÅ  € õ Œ=˜ÑÔ×Ò�rœxÑ(Ô(¨5Ñ0µ2´9¸QÀÑ3FÔ3FÑF€DØ	�Š€BØ�%‰<€DØ�E‰M€EØ
�#‰+€CØ�S‰L€EÝ
Œ	�!�Q‰Œ€AØ��D‘Š=€DÝŒh�q˜4’x  T¡­2¬8°D¸!¸QÀ¹XÑ+FÔ+FÑGÔG€GØ����D�Œk˜Q˜t Q Q Q˜wœZ¨!¨d©(Ò2Ñ3€GØ�!�!�!�T�'ŒN˜RÑ # a a a¨ g¤,°Ñ"3Ñ3°a¸¸a¸a¸a¸´jÑ@€EÝ
Œ'�#˜‘+˜w¨Ñ
*Ô
*€CÝ
Œ'�#˜‘+˜w¨Ñ
*Ô
*€CÝ�A�u˜r‘z C¨"¡HÑ,¨t°b©yÑ8¸"¸cÀ3ÈÐQUÐWZÐ\]Ð_cÐejÑkÔkÐkÝ�A�u˜r‘z C¨"¡HÑ,¨t°b©yÑ8¸"¸cÀ3ÈÐQUÐWZÐ\]Ð_cÐejÑkÔkÐkÐkÐkr9   ))é   r   )rR   r   )rR   é   )r}   r£   )r|   r¢   i   c                 ó8   — | | dz
  z  dk    o|dz  dk    o|| k    S )zPHead widths the kernel takes: a power of two, an even number of rotary channels.r   r   r   © )r`   Úrots     r7   Úcan_rms_roper§   ­   s+   € à�#˜‘'‰?˜aÒÐ? C¨!¡G¨q¢LÐ?°S¸C²ZÐ?r9   c                 ó   ‡ ‡‡‡‡‡	‡
‡‡‡— |d         |d         cŠŠ‰ j         \  ŠŠŠ
Š	‰‰z  ‰
z  }ˆˆ	ˆˆ
ˆfd„Šˆˆ	ˆ
ˆˆˆ ˆfd„}t                               d‰ j        › d‰
› d‰	› d‰j         d	         › d
‰                      d¦  «        › �	‰ j        t          ||t          k    ¦  «        } ‰‰ ‰|¦  «         t          j	        j
                             ‰ ¦  «         t          j	        j
                             ‰¦  «         ‰ ‰fS )aW  In place on q and k (batch, tokens, heads, head_dim), any strides with contiguous channels (views of a fused projection): per head
    RMSNorm without weight, then the split-half RoPE of the first rotary channels with the angles of rotary (cos, sin), FP32 (batch,
    tokens, 1, rotary channels) whose two halves are equal. H3's 64-channel heads preserve Comfy Kitchen's reduction and rotation
    rounding, including its FP32 fused multiply-add order. Not bit-identical to the PyTorch 16-bit rotation but more precise (one rounding
    after the rotation instead of rounded cos / sin and products).r   r   c           
      ó²  •— |\  }}| j         d         | j         d         z  ‰	z  }t          t          j        ||¦  «        f         | |‰‰
|‰	| j         d         g|                      ¦   «         d d…         ¢|                     ¦   «         d d…         ¢‰                     d¦  «        ‘‰                     d¦  «        ‘‰‘‰‘‰j         d         dz  ‘|‘R |ddœŽ d S )Nr   r   rS   rO   r   F)Ú	num_warpsrT   )rU   r¡   rY   rZ   r[   )ÚqqÚkkÚconfigÚblockÚwarpsrl   r2   r`   Úepsr_   r3   s         €€€€€r7   Úlaunchzrms_rope_.<locals>.launch¼   s?  ø€ Ø‰ˆˆuØŒH�QŒK˜"œ( 1œ+Ñ%¨Ñ-ˆÝ•f”k ! UÑ+Ô+Ð-Ô.¨r°2°s¸CÀÀEÈ2Ì8ÐTUÌ;ð 	MÐY[×YbÒYbÑYdÔYdÐegÐfgÐegÔYhð 	MÐkm×ktÒktÑkvÔkvÐwyÐxyÐwyÔkzð 	MÐ|÷  }Gò  }Gð  HIñ  }Jô  }Jð 	MØ03·
²
¸1±´ð	MØ?Bð	MØDGð	MØILÌÐSUÌÐZ[ÑI[ð	MØ]bð	Mð 	MØnsð  GLð	Mð 	Mð 	Mð 	Mð 	Mð 	Mr9   c            	      óh  •‡ ‡‡‡— t          dt          ‰
t          ‰z  ¦  «        ¦  «        Št          dt          ‰t          ‰‰z  z  ¦  «        ¦  «        Š ˆˆˆ ˆfd„‰	‰fD ¦   «         \  ŠŠ‰                     ‰	d ‰ …d ‰…f         ¦  «         ‰                     ‰d ‰ …d ‰…f         ¦  «         ˆˆˆfd„S )Nr   c           	   3   ó„   •K  — | ]:}|                      ‰‰‰|                     d ¦  «        f¦  «        dd‰…f         V — Œ;dS )r   .N)Ú	new_emptyr[   )Ú.0r.   r`   r_   ÚitemsÚtokenss     €€€€r7   ú	<genexpr>z5rms_rope_.<locals>.prepare_scratch.<locals>.<genexpr>Å   sU   øè è € ÐlÐlÐ`a §¢¨U°F¸EÀ1Ç8Â8ÈAÁ;Ä;Ð,OÑ PÔ PÐQTÐVZÐWZÐVZÐQZÔ [ÐlÐlÐlÐlÐlÐlr9   c                 ó   •—  ‰‰‰| ¦  «        S )Nr¥   )r­   r±   Ú	scratch_kÚ	scratch_qs    €€€r7   ú<lambda>z4rms_rope_.<locals>.prepare_scratch.<locals>.<lambda>È   s   ø€ ˜f˜f Y°	¸6ÑBÔB€ r9   )ÚmaxÚminÚ
_TUNE_ROWSÚcopy_)r¶   rº   r»   r·   r*   r`   r_   Úkr±   r\   rŸ   s   @@@@€€€€€€€r7   Úprepare_scratchz"rms_rope_.<locals>.prepare_scratchÂ   sÐ   øøøøø€ Ý�Q�˜C¥¨uÑ!4Ñ5Ô5Ñ6Ô6ˆÝ�A•s˜5¥*°¸%±Ñ"@ÑAÔAÑBÔBˆØlÐlÐlÐlÐlÐlÐlÐfgÐijÐekÐlÑlÔlÑˆ	�9Ø�Š˜˜&˜5˜& ' 6 '˜/Ô*Ñ+Ô+Ð+Ø�Š˜˜&˜5˜& ' 6 '˜/Ô*Ñ+Ô+Ð+ØBÐBÐBÐBÐBÐBÐBr9   Úqk_rms_ropeú|r.   z|rrO   z|sr   )rU   Ú_TUNERr­   r   r[   rQ   Ú_ROPE_CONFIGSr¿   rV   rh   ri   rj   )r\   rÁ   r]   r°   r'   rÂ   r­   r*   r2   r`   r_   r±   rŸ   r3   s   `` `   @@@@@@@r7   Ú	rms_rope_rÇ   ²   ss  øøøøøøøøøø€ ð �aŒy˜& œ)€H€CˆØœWÑ€Eˆ3��sØ�3‰;˜Ñ€DðMð Mð Mð Mð Mð Mð Mð Mð MðCð Cð Cð Cð Cð Cð Cð Cð Cð Cð Cõ �]Š]˜=¨Q¬WÐ*dÐ*d°uÐ*dÐ*d¸sÐ*dÐ*dÀcÄiÐPRÄmÐ*dÐ*dÐWX×W_ÒW_Ð`aÑWbÔWbÐ*dÐ*dÐfgÔfnÕp}Ø*¨DµJÒ,>ñ@ô @€Fà
€Fˆ1ˆa�ÑÔÐÝ	„NÔ×*Ò*¨1Ñ-Ô-Ð-Ý	„NÔ×*Ò*¨1Ñ-Ô-Ð-Øˆaˆ4€Kr9   c                 ó  — | j         d         dz  }t          j        g | j         dd…         ¢|‘R | j        | j        ¬¦  «        }|                     ¦   «         }t          t          j        |d¦  «        f         | |||dd¬¦  «         |S )zHRelease the twice-wide projection before the following linear operation.rO   r   NrP   re   Frf   )	rU   rV   rW   rQ   r   rg   rs   rY   rZ   rx   s       r7   ÚswiglurÉ   Ò   sŠ   € à
Œ.˜Ô
 Ñ
!€CÝ
Œ+Ð1˜œ s¨ sÔ+Ð1¨SÐ1Ð1¸(¼/ÐQYÔQ_Ð
`Ñ
`Ô
`€CØ�	Š	‰Œ€AÝ�œ A sÑ+Ô+Ð-Ô.¨x¸¸aÀÀcÐ\aÐbÑbÔbÐbØ€Jr9   )r:   ÚTÚHÚWr%   ÚPÚFRONTc
           	      óú  — t          j        d¦  «        |	z  t          j        d|	¦  «        z   }
|
|d|z  z   z  |z
  }|
|d|z  z   z  |d|z  z   z  |z
  }|
|d|z  z   |d|z  z   z  z  ||z   z  |z
  }|
|d|z  z   |d|z  z   z  ||z   z  z  |z  }|
||d|z  z   z  |d|z  z   z  ||z   z  z  }t          j        |dk     | t          j        ||k    d|z  dz
  |z
  |¦  «        ¦  «        }t          j        |dk     | t          j        ||k    d|z  dz
  |z
  |¦  «        ¦  «        }||z  |z   |z  |z   |z  |z   |z  |z   }t          j        | |z   |
|k     |dk    z  d¦  «                             t           j        ¦  «        }| j        j        t           j	        k    r,t          j
        |dt          j        | ¦  «        z   ¦  «        }n|dt          j        | ¦  «        z   z  }t          j        ||
z   ||
|k     ¦  «         d S r<   )r   r   r   r   r   r   r   r   r   Úfloat16ru   r   r?   r"   )r#   r$   r:   r%   rÊ   rË   rÌ   rÍ   rÎ   r   r@   ÚcolÚrowr^   rˆ   r5   rA   r.   rC   s                      r7   Ú_silu_pad_framesrÓ   Û   s&  € õ 	Œ�aÑÔ˜5Ñ ¥2¤9¨Q°Ñ#6Ô#6Ñ6€AØ
ˆq�1�q‘5‰y‰/˜AÑ
€CØ
��A˜‘E‘	Ñ
˜a ! a¡%™iÑ
(¨1Ñ
,€CØ	ˆq�1�q‘5‰y˜Q  Q¡™YÑ'Ñ(¨A°©IÑ6¸Ñ>€AØ	ˆq�1�q‘5‰y˜Q  Q¡™YÑ'¨1¨u©9Ñ5Ñ6¸Ñ:€AØ	ˆa�1�q˜1‘u‘9‰o  Q¨¡U¡Ñ+¨q°5©yÑ9Ñ:€AÝ
Œ(�3˜’7˜S˜D¥"¤(¨3°!ª8°Q¸±U¸Q±YÀ±_ÀcÑ"JÔ"JÑ
KÔ
K€CÝ
Œ(�3˜’7˜S˜D¥"¤(¨3°!ª8°Q¸±U¸Q±YÀ±_ÀcÑ"JÔ"JÑ
KÔ
K€CØ�A‘˜‘	˜Q‰ Ñ" aÑ'¨#Ñ-°Ñ2°SÑ8€FÝ
Œ��F‘
˜Q šU q¨A¢vÑ.°Ñ2Ô2×5Ò5µb´jÑAÔA€AØ„wÔ�RœZÒ'Ð'Ý”	˜!˜S¥9¤=°!°Ñ#4Ô#4Ñ4Ñ5Ô5ˆˆà�S�2œ6 1 "™:œ:Ñ%Ñ&ˆÝ„HˆS�1‰W�e˜Q šUÑ#Ô#Ð#Ð#Ð#r9   c                 ó  — | j         \  }}}}}t          j        ||||z   |d|z  z   |d|z  z   f| j        | j        ¬¦  «        }	|	                     ¦   «         }
t          t          j        |
d¦  «        f         | |	|
||||||dd¬¦  «         |	S )Nr   rP   é   Frf   )	rU   rV   rW   rQ   r   rg   rÓ   rY   rZ   )Ú
normalizedr*   ÚframesÚpaddingÚfrontÚ_ÚchannelsÚheightÚwidthr-   rl   s              r7   Úsilu_pad_framesrÞ   ð   s¹   € Ø$.Ô$4Ñ!€A€x��F˜EÝ
Œ+�u˜h¨°©¸ÀÀWÁÑ8LÈeÐVWÐZaÑVaÑNaÐbØ'Ô.°jÔ6FðHñ Hô H€Cà�	Š	‰Œ€AÝ•f”k ! SÑ)Ô)Ð+Ô,¨Z¸¸aÀÈ6ÐSYÐ[`Ø(/°¸ÈeðUñ Uô Uð Uà€Jr9   )Ú!install_triton_compilation_logger)#Ú__doc__rV   rY   Útriton.languageÚlanguager   Útriton.language.extrar   Úshared.kernels.triton_tunerr   rÅ   ÚjitÚ	constexprr8   rE   rM   ÚhasattrÚCAST_ROPE_SUPPORTEDrc   rm   rp   rs   rv   ry   r�   r¡   rÆ   r¿   r§   rÇ   rÉ   rÓ   rÞ   Ú%shared.kernels.triton_compilation_logrß   r¥   r9   r7   ú<module>rê      sÿ  ðà OÐ OØ €€€Ø €€€Ø Ð Ð Ð Ð Ð Ø +Ð +Ð +Ð +Ð +Ð +à 3Ð 3Ð 3Ð 3Ð 3Ð 3à	ˆÐ,¨a°Ñ	:Ô	:€ð €„ÐWÐWÐWð  yrð  yrð  yrð  sñ  sô  sðOà”ðOà!#¤ðOà68´lðOð Oð Oñ sô  sðOð8 €„˜s˜eÀSÀEÐJÑJÔJð3�R”\ð 3¨"¬,ð 3ð 3ð 3ñ KÔJð3ð €„˜s˜eÀSÀEÐJÑJÔJð.¨¬ð .¸R¼\ð .ð .ð .ñ KÔJð.ð �g˜b (Ñ+Ô+Ð ðð ð ðð ð ðð ð ð €„˜s˜eÀSÀEÐJÑJÔJð0 R¤\ð 0¸"¼,ð 0ð 0ð 0ñ KÔJð0ð €„˜s˜eÀSÀEÐJÑJÔJð0˜Rœ\ð 0°"´,ð 0ð 0ð 0ñ KÔJð0ðð ð ð „ðDÀ"Ä,ð DÐVXÔVbð DÐkmÔkwð Dð Dð Dñ „ðDð6 €„ÐfÐfÐfð  HPð  HPð  HPð  Qñ  Qô  QðlÐ[]Ô[gð lÐoqÔo{ð lØœðlð lð lñ Qô  Qðlð& =€Ø€
ð@ð @ð @ð
ð ð ð@ð ð ð €„Ð2Ð2Ð2ÐSgÐSgÐSgÐhÑhÔhð$ 2¤<ð $Ø œlð$àœLð$à13´ð$ð $ð $ñ iÔhð$ð(ð ð ð TÐ SÐ SÐ SÐ SÐ SØ !Ð !Ñ #Ô #Ð #Ð #Ð #r9   