§
     ’Þj£/  ã                   óò  — d Z ddlZddlmZ  ej        dgdg¬¦  «        dej        dej        fd„¦   «         Zd-d„Zddl	Z	dd	l
mZ dd
lmZ  eddd¦  «        ZdZdZ ej        g d¢g d¢¬¦  «        dej        dej        dej        dej        dej        dej        fd„¦   «         Zddœd„ZddlmZ  e¦   «           ej        ddgddg¬¦  «        dej        dej        dej        fd„¦   «         Zd„ Zej        dej        fd„¦   «         Zd „ Zd!„ Z ej        g d"¢g 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d.d,„ZdS )/zjH3 denoiser fusions: modulation (native RMSNorm and BF16 rounding boundaries kept) and Q/K RMSNorm + RoPE.é    NÚN)Údo_not_specializeÚdo_not_specialize_on_alignmentÚDÚBLOCKc                 ó  — t          j        d¦  «                             t           j        ¦  «        |z  t          j        d|¦  «        z   }||k     }||z  dk    r4t          j        d¦  «        ||z  z  |z  t          j        d|¦  «        z   }n||z  }t          j        | |z   |d¦  «                             t           j        ¦  «        }	t          j        ||z   |d¦  «                             | j        j        ¦  «                             t           j        ¦  «        }
t          j        ||z   |d¦  «                             | j        j        ¦  «                             t           j        ¦  «        }d|
z                        | j        j        ¦  «                             t           j        ¦  «        }|	|z                       | j        j        ¦  «                             t           j        ¦  «        }t          j	        | |z   ||z   |¦  «         d S )Nr   ç      ð?©
ÚtlÚ
program_idÚtoÚint64ÚarangeÚloadÚfloat32ÚdtypeÚ
element_tyÚstore)ÚXÚSCALEÚSHIFTr   r   r   ÚiÚvalidÚcolÚxÚscaleÚshiftÚfactorÚscaleds                 úFC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\denoiser_triton.pyÚ	_modulater!      s™  € å
Œ�aÑÔ×Ò�BœHÑ%Ô%¨Ñ-µ´	¸!¸UÑ0CÔ0CÑC€AØ�ŠE€EØˆ5�y�A‚~€~ÝŒ}˜QÑÔ 1¨¡:Ñ.°%Ñ7½"¼)ÀAÀuÑ:MÔ:MÑMˆˆà�!‰eˆÝ
Œ��A‘�u˜aÑ Ô ×#Ò#¥B¤JÑ/Ô/€AÝŒG�E˜C‘K ¨Ñ*Ô*×-Ò-¨a¬gÔ.@ÑAÔA×DÒDÅRÄZÑPÔP€EÝŒG�E˜C‘K ¨Ñ*Ô*×-Ò-¨a¬gÔ.@ÑAÔA×DÒDÅRÄZÑPÔP€EØ�E‰k×Ò˜aœgÔ0Ñ1Ô1×4Ò4µR´ZÑ@Ô@€FØ�&‰j�_Š_˜QœWÔ/Ñ0Ô0×3Ò3µB´JÑ?Ô?€FÝ„HˆQ�‰U�F˜U‘N EÑ*Ô*Ð*Ð*Ð*ó    c           
      ób  —  || ¦  «        }|�|                      |¦  «        }| j        d         }|D ]X\  }}	}
|	|k    rL|	|z
  |z  }t          t          j        |d¦  «        f         |||	…         ||
         ||
         ||dd¬¦  «         ŒYt
          j        j                             |¦  «         |S )Néÿÿÿÿé   F©Úenable_fp_fusion)	r   Úshaper!   ÚtritonÚcdivÚtorchÚautogradÚgraphÚincrement_version)r   Únormr   r   Úsegmentsr   ÚoutputÚwidthÚstartÚstopÚrowÚsizes               r    Únorm_modulater7      sË   € ØˆT�!‰WŒW€FØÐØ—’˜5Ñ!Ô!ˆØŒG�BŒK€EØ$ð (ð (Ñˆˆt�SØ�%Š<ˆ<Ø˜‘J %Ñ'ˆDÝ•v”{ 4¨Ñ.Ô.Ð0Ô1Ø�u˜T�zÔ" E¨#¤J°°c´
¸DÀ%ÈØ!&ð(ñ (ô (ð (øõ 
„NÔ×*Ò*¨6Ñ2Ô2Ð2Ø€Mr"   )Ú	libdevice)ÚLaunchTunerÚwan2gp_h3_denoiser_kernelsé   ÚH3))é   r;   )é   é   )r>   é   )é    r@   )é@   r=   i   )ÚROWSÚHEADSÚSEQÚXBÚXSÚXHÚRBÚRSÚRPÚRCÚPAIRSÚINTERLEAVEDc                 óŒ
  — t          j        d¦  «                             t           j        ¦  «        |z  t          j        d|¦  «        z   }||k     }||z  }||z  }||z  }||z  }t          j        d|¦  «        }||z  ||z  z   ||z  z   }| |dd…df         z   |ddd…f         z   }t          j        ||dd…df         d¦  «                             t           j        ¦  «        }|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        ¦  «        } |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   ¦  «                             t           j        ¦  «        ddd…f         z                       | j        j        ¦  «                             t           j        ¦  «        }*| |)z  t          j        ||z   ¦  «                             t           j        ¦  «        ddd…f         z                       | j        j        ¦  «                             t           j        ¦  «        } |dd…df         |	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         }/t          j        t          j        |/|-|.¦  «        t          j        |/|*| ¦  «        t          j        |/|. |-¦  «        t          j        |/| |*¦  «        z  ¦  «        }0t          j        |t          j        |ddd…f         |0|*¦  «        |dd…df         ¦  «         dS )
a  Rows (batch, token, head) of Q or K, loaded whole: weighted RMSNorm in FP32 rounded to the activation dtype (the norm output),
    then the rotation of the first 2 * PAIRS channels in FP32 with one final rounding (cos / sin pairs of ROPE): channel p with
    p + PAIRS (split halves) or, INTERLEAVED, 2p with 2p + 1 (a complex table viewed as real); the other channels keep the norm
    output. 128-channel heads sum the squares as Comfy Kitchen does (lane sums of c, c + 32, c + 64, c + 96, then the lanes): its
    results bit for bit.r   Nr?   r;   é€   r@   rA   r	   ç        )r   r   r   r   r   r   r   ÚwhereÚreshapeÚpermuteÚsplitÚsumÚfmar8   Úrsqrtr   r   r   )1r   ÚWÚROPErC   rD   rE   rF   rG   rH   rI   rJ   ÚEPSrK   rL   r   rM   r   rN   ÚrowsÚokÚheadÚtokenÚseqÚbatchÚcÚbaseÚptrr   ÚrotatedÚleadingÚpartnerÚpairÚotherÚlanesÚevenÚoddÚx0Úx2Úx1Úx3ÚtotalÚinvÚyÚangleÚmaskÚcosÚsinÚfirstÚouts1                                                    r    Ú	_rms_roperz   /   sÜ  € õ Œ=˜ÑÔ×Ò�rœxÑ(Ô(¨5Ñ0µ2´9¸QÀÑ3FÔ3FÑF€DØ	�Š€BØ�%‰<€DØ�E‰M€EØ
�#‰+€CØ�S‰L€EÝ
Œ	�!�Q‰Œ€AØ�2‰:˜˜b™Ñ  4¨"¡9Ñ,€DØ
ˆd�1�1�1�d�7ŒmÑ
˜a  a a a œjÑ
(€CÝ
Œ��R˜˜˜˜4˜”[ !Ñ$Ô$×'Ò'­¬
Ñ3Ô3€AØ�!�e‘)Šm€GØð /Ø�a‘%˜1’*ˆÝ”(˜7 A¨¡E¨1Ñ-Ô-ˆØ�A‰vˆˆà�e’)ˆÝ”(˜7 A¨¡I­r¬x¸ÀÀUÁÈAÑ/NÔ/NÑOÔOˆÝŒx˜  A¨¡IÑ.Ô.ˆÝŒ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ØˆC‚x€xÝ”
�2œ:¥b¤j°°U¸A¸r°NÑ&CÔ&CÀQÈÈ1ÑMÔMÐPUÐWYÐ[\Ð^_ÐO`ÑaÔaˆÝ”H˜U‘O”O‰	ˆˆcÝ”˜$‘”‰ˆˆBÝ”˜#‘”‰ˆˆBÝ”•r”v˜b "¥b¤f¨R°µR´V¸BÀÀBÈÁGÑ5LÔ5LÑ&MÔ&MÑNÔNÐPQÑRÔRˆˆå”�q˜1‘u˜aÑ Ô ˆÝ
Œ/˜% !™) c™/Ñ
*Ô
*¨1¨1¨1¨d¨7Ô
3€CØ	
ˆS‰•2”7˜1˜q™5‘>”>×$Ò$¥R¤ZÑ0Ô0°°q°q°q°Ô9Ñ	9×=Ò=¸a¼gÔ>PÑQÔQ×TÒTÕUWÔU_Ñ`Ô`€AØ�S‰[�2œ7 1 w¡;Ñ/Ô/×2Ò2µ2´:Ñ>Ô>¸tÀQÀQÀQ¸wÔGÑG×KÒKÈAÌGÔL^Ñ_Ô_×bÒbÕceÔcmÑnÔn€EØ�!�!�!�T�'ŒN˜RÑ # a a a¨ g¤,°Ñ"3Ñ3°d¸4ÀÀÀ¸7´mÀbÑ6HÑH€EØˆaˆaˆa�ˆgŒ;˜  q q q Ô)Ñ)€DÝ
Œ'�$˜‘,  cÑ
*Ô
*×
-Ò
-­b¬jÑ
9Ô
9€CÝ
Œ'�$˜‘, Ñ# T¨3Ñ
/Ô
/×
2Ò
2µ2´:Ñ
>Ô
>€CØ�D˜!˜!˜!�GÔ€EÝ
Œ&•”˜%  cÑ*Ô*­B¬H°U¸A¸uÑ,EÔ,EÅrÄxÐPUÐX[ÐW[Ð]`ÑGaÔGaÕdfÔdlÐmrÐtyÐ{|Ñd}Ôd}ÑG}Ñ
~Ô
~€CÝ„HˆS•"”(˜7 4¨¨¨ 7Ô+¨S°!Ñ4Ô4°b¸¸¸¸D¸´kÑBÔBÐBÐBÐBr"   F)Úinterleavedc                óü  ‡ ‡‡‡‡‡‡	‡
‡‡— ‰ j         \  }ŠŠ	Š‰j         d         dk    rdn‰                     d¦  «        Šˆˆˆ	ˆˆˆˆfd„Š
ˆˆ	ˆ
ˆˆ fd„}t                               d‰ j        › d‰	› d‰› d‰j         d	         › d
‰                      d¦  «        › dt          ‰¦  «        › �‰ j        t          ||‰z  ‰	z  t          k    ¦  «        } ‰
‰ |¦  «         t          j
        j                             ‰ ¦  «         ‰ S )aN  In place on Q or K (batch, tokens, heads, head_dim), any strides with contiguous channels (views of a fused projection); weight:
    the RMSNorm weight; rope: (1 | batch, tokens, 1, pairs, 2) cos / sin of the rotation of the first 2 * pairs channels, split halves
    or interleaved pairs (2p, 2p + 1: a complex table viewed as real).r   r;   c           	      ó¤  •— |\  }}| j         d         | j         d         z  ‰z  }t          t          j        ||¦  «        f         | ‰‰	|‰| j         d         g|                      ¦   «         d d…         ¢‰
‘‰	                     d¦  «        ‘‰‘‰	                     d¦  «        ‘‰	                     d¦  «        ‘‰‘‰	j         d         ‘|‘‰‘R |ddœŽ d S )Nr   r;   é   r@   éþÿÿÿF)Ú	num_warpsr'   )r(   rz   r)   r*   Ústride)ÚxxÚconfigÚblockÚwarpsr\   ÚdimÚepsÚheadsr{   ÚropeÚ
rope_batchÚweights        €€€€€€€r    Úlaunchzrms_rope_.<locals>.launchf   sE  ø€ Ø‰ˆˆuØŒx˜Œ{˜RœX aœ[Ñ(¨5Ñ0ˆÝ•6”;˜t UÑ+Ô+Ð-Ô.¨r°6¸4ÀÀuÈbÌhÐWXÌkð 	aÐ\^×\eÒ\eÑ\gÔ\gÐhjÐijÐhjÔ\kð 	aÐmwð 	aÐy}÷  zEò  zEð  FGñ  zHô  zHð 	að  JMð 	aØ/3¯{ª{¸1©~¬~ð	aØ?C¿{º{È1¹~¼~ð	aØORð	aØTXÔT^Ð_aÔTbð	aØdið	aØkvð	að 	að CHð  [`ð	að 	að 	að 	að 	að 	ar"   c                  ó  •‡— t          dt          ‰t          ‰z  ¦  «        ¦  «        } ‰                     d| ‰‰                     d¦  «        f¦  «        dd ‰…f         Š‰                     ‰d d…d | …f         ¦  «         ˆˆfd„S )Nr;   r?   .c                 ó   •—  ‰‰| ¦  «        S ©N© )rƒ   rŒ   Úscratchs    €€r    ú<lambda>z4rms_rope_.<locals>.prepare_scratch.<locals>.<lambda>p   s   ø€ ˜f˜f W¨fÑ5Ô5€ r"   )ÚmaxÚminÚ
_TUNE_ROWSÚ	new_emptyr�   Úcopy_)Útokensr‘   r†   rˆ   rŒ   r`   r   s    @€€€€€r    Úprepare_scratchz"rms_rope_.<locals>.prepare_scratchl   s‡   øø€ Ý�Q�˜C¥¨uÑ!4Ñ5Ô5Ñ6Ô6ˆØ—+’+˜q &¨%°·²¸!±´Ð=Ñ>Ô>¸sÀDÀSÀD¸yÔIˆØ�Š�a˜˜˜˜G˜V˜G˜”nÑ%Ô%Ð%Ø5Ð5Ð5Ð5Ð5Ð5r"   Úrms_ropeú|r   z|pr   z|sr?   z|I)r(   r�   Ú_TUNERrƒ   r   ÚintÚdeviceÚ_ROPE_CONFIGSr•   r+   r,   r-   r.   )r   r‹   r‰   r‡   r{   ra   r™   rƒ   r†   rˆ   rŒ   rŠ   r`   s   `````   @@@@@r    Ú	rms_rope_r    _   sr  øøøøøøøøøø€ ð œWÑ€Eˆ3��sØ”j ”m qÒ(Ð(��¨d¯kªk¸!©n¬n€Jðað að að að að að að að að að að6ð 6ð 6ð 6ð 6ð 6ð 6ð 6ð 6õ �]Š]˜:¨!¬'Ð'vÐ'v°EÐ'vÐ'v¸CÐ'vÐ'vÀ4Ä:ÈbÄ>Ð'vÐ'vÐUV×U]ÒU]Ð^_ÑU`ÔU`Ð'vÐ'vÕdgÐhsÑdtÔdtÐ'vÐ'vÐxyô  yAõ  CPð  RaØ  3™;¨Ñ.µ*Ò<ñ>ô >€Fà
€Fˆ1ˆfÑÔÐÝ	„NÔ×*Ò*¨1Ñ-Ô-Ð-Ø€Hr"   )Ú!install_triton_compilation_loggerÚ
MOD_STRIDEÚ	PER_TOKENc	                 óJ  — t          j        d¦  «                             t           j        ¦  «        |z  t          j        d|¦  «        z   }	t          j        d¦  «                             t           j        ¦  «        }
|
|z  |	z   }|r|	}nN||z  dk    r:|
|z  t          j        d¦  «        ||z  z  |z  z   t          j        d|¦  «        z   }n|
|z  |	|z  z   }t          j        | |z   |	|k     d¦  «                             t           j        ¦  «        }t          j        ||z   |	|k     d¦  «                             t           j        ¦  «        }t          j        ||z   |	|k     d¦  «                             t           j        ¦  «        }d|z                        | j        j        ¦  «                             t           j        ¦  «        }||z                       | j        j        ¦  «                             t           j        ¦  «        }t          j	        ||z   ||z   |	|k     ¦  «         d S )Nr   r;   r	   r
   )r   r   r   ÚYr   r¢   r   r£   r   r   ÚgroupÚoffsetÚsr   r   r   r   Úvalues                     r    Ú_scale_shiftrª   }   sÍ  € å
Œ�aÑÔ×Ò�BœHÑ%Ô%¨Ñ-µ´	¸!¸UÑ0CÔ0CÑC€Aõ ŒM˜!ÑÔ×Ò¥¤Ñ)Ô)€EØ�Q‰Y˜‰]€FØð 'ØˆˆØ	
ˆU‰�aŠˆØ�JÑ¥"¤-°Ñ"2Ô"2°a¸5±jÑ"AÀUÑ!JÑJÍRÌYÐWXÐZ_ÑM`ÔM`Ñ`ˆˆà�JÑ  Q¡Ñ&ˆÝ
Œ��F‘
˜A šE 1Ñ%Ô%×(Ò(­¬Ñ4Ô4€AÝŒG�E˜A‘I˜q 1šu aÑ(Ô(×+Ò+­B¬JÑ7Ô7€EÝŒG�E˜A‘I˜q 1šu aÑ(Ô(×+Ò+­B¬JÑ7Ô7€EØ�E‰k×Ò˜aœgÔ0Ñ1Ô1×4Ò4µR´ZÑ@Ô@€FØ�‰Z�OŠO˜AœGÔ.Ñ/Ô/×2Ò2µ2´:Ñ>Ô>€EÝ„HˆQ�‰Z˜ ™¨¨AªÑ.Ô.Ð.Ð.Ð.r"   c                 óæ  — |r| nt          j        | ¦  «        }|j        d         | j        d         k    o|                     ¦   «         }|rdn|j        d         |j        d         z  }|                      ¦   «         |z  }t          t          j        |d¦  «        |f         | ||||t          |¦  «        | j        d         |dd¬¦
  «
         |r$t           j	        j
                             |¦  «         |S )Nr;   r   r%   r?   Fr&   )r+   Ú
empty_liker(   Úis_contiguousÚnumelrª   r)   r*   Úmodulation_strider,   r-   r.   )r   r   r   Úin_placery   Ú	per_tokenÚgroupsr6   s           r    Úscale_shiftr³   ’   sî   € ØÐ
0ˆ!ˆ!�UÔ-¨aÑ0Ô0€CØ”˜A” !¤'¨!¤*Ò,ÐF°×1DÒ1DÑ1FÔ1F€IØÐ@ˆQˆQ ¤¨Q¤°%´+¸a´.Ñ!@€FØ�7Š7‰9Œ9˜Ñ€DÝ•&”+˜d DÑ)Ô)¨6Ð2Ô3Ø	ˆ5�%˜˜dÕ$5°eÑ$<Ô$<¸a¼gÀa¼jÈ)ÐUYØð ñ  ô  ð  ð ð 4ÝŒÔ×.Ò.¨sÑ3Ô3Ð3Ø€Jr"   ÚAXISc                 óR   — |dk    r| ||z  z  |z  }n|dk    r	| |z  |z  }n| |z  }|S )Nr   r;   r�   )r_   r´   ÚG0ÚG1ÚG2r5   s         r    Ú	_axis_rowr¹   Ÿ   sI   € ð ˆq‚y€yØ˜˜R™Ñ  2Ñ%ˆˆØ	�ŠˆØ�r‰k˜BÑˆˆà�b‰jˆØ€Jr"   c                 óL   — t          | ¦  «        ddt          | ¦  «        z
  z  z   S )zuThe grid axis of each frequency table as three constexpr ints (a constexpr tuple is not subscriptable on Triton 3.2).)r   r~   )ÚtupleÚlen)Ú	rope_axess    r    Ú	_axis_idsr¾   «   s%   € å�ÑÔ˜d a­#¨i©.¬.Ñ&8Ñ9Ñ9Ð9r"   c                 óx   — | j         d         dk    r|                      d¦  «        n|                      d¦  «        S )zpElements between consecutive modulation rows: LTX hands unbind views of a (B, T, K, D) projection, stride K * D.r;   r   )r(   r�   )r   s    r    r¯   r¯   °   s0   € à#œk¨!œn¨qÒ0Ð0ˆ5�<Š<˜‰?Œ?ˆ?°e·l²lÀ1±o´oÐEr"   )r   ÚTOKENSÚWIDTHr¶   r·   r¸   ÚBATCH_STRIDEÚ
HEAD_STARTÚHALFÚPADÚAXESÚA0ÚA1ÚA2Ú
FREQ_WIDTHÚFP32c           
      óJ	  — t          j        d¦  «                             t           j        ¦  «        |z  t          j        d|¦  «        z   }||	dz  z  }||	dz  z  |z  }||	dz  z  |z  }||
z  ||
z  }}||z  ||	z  z   ||
z  dz  z   |z   }|||
z  z   }||z
  |z  }||z
  |z  } t          j        |fdt           j        ¦  «        }!t          j        |fdt           j        ¦  «        }"t          j        |¦  «        D �]ó}#||k     ||k    z  ||#k    z  }$|#dk    r�t          |||||¦  «        }%t          j	        ||%|z  z   | z   |$d¦  «                             t           j        ¦  «        }&t          j	        ||%|z  z   | z   |$d¦  «                             t           j        ¦  «        }'�n|#dk    rŒt          |||||¦  «        }%t          j	        ||%|z  z   | z   |$d¦  «                             t           j        ¦  «        }&t          j	        ||%|z  z   | z   |$d¦  «                             t           j        ¦  «        }'n‹t          |||||¦  «        }%t          j	        ||%|z  z   | z   |$d¦  «                             t           j        ¦  «        }&t          j	        ||%|z  z   | z   |$d¦  «                             t           j        ¦  «        }'t          j
        |$|&|!¦  «        }!t          j
        |$|'|"¦  «        }"�Œõt          j	        | |z   ||k     d¦  «                             t           j        ¦  «        }(t          j	        | |z   |
z   ||k     d¦  «                             t           j        ¦  «        })t          j        |(|!¦  «        t          j        |)|!¦  «        }*}#|sx|#                     | j        j        ¦  «                             t           j        ¦  «        }#|*                     | j        j        ¦  «                             t           j        ¦  «        }*||k     ||z   |z  dk    ||k    z  z  }$t          j        | |z   t          j        |#t          j        |)|"¦  «        ¦  «        |$¦  «         t          j        | |z   |
z   t          j        |*t          j        |(|"¦  «        ¦  «        |$¦  «         d S )Nr   r?   r	   rQ   r;   )r   r   r   r   r   Úfullr   Ústatic_ranger¹   r   rR   r8   Úmul_rnr   r   r   Úsub_rnÚadd_rn)+r   ÚC0ÚS0ÚC1ÚS1ÚC2ÚS2r   rÀ   rÁ   rÄ   rÅ   rÆ   r¶   r·   r¸   rÇ   rÈ   rÉ   rÊ   rÂ   rÃ   rË   r   r   Úfr_   ra   r^   Úlocalr§   Úaxisr   rb   r¨   Úar   r5   ÚcvÚsvrm   ro   Úbs+                                              r    Ú_split_roperß   µ   sB  € õ 	Œ�aÑÔ×Ò�BœHÑ%Ô%¨Ñ-µ´	¸!¸UÑ0CÔ0CÑC€AØ	ˆU�a‰ZÑ€AØ�%˜1‘*Ñ Ñ&€EØ�%˜1‘*Ñ Ñ'€EØ�t‘)˜Q ™Xˆ%€DØ�\Ñ! E¨E¡MÑ1°D¸4±KÀ!±OÑCÀeÑK€FØ	ˆJ˜ÑÑ€AØ�‰G�tÑ€DØˆs‰7�tÑ
€CÝ
Œ��˜#�rœzÑ*Ô*€AÝ
Œ��˜#�rœzÑ*Ô*€AÝŒ_˜TÑ"Ô"ð #ñ #ˆØ�Q’˜1 š8Ñ$¨°ª	Ñ2ˆØ�Š6ˆ6Ý˜E 2 r¨2¨rÑ2Ô2ˆCÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?×BÒBÅ2Ä:ÑNÔNˆBÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?×BÒBÅ2Ä:ÑNÔNˆB‰BØ�!ŠVˆVÝ˜E 2 r¨2¨rÑ2Ô2ˆCÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?×BÒBÅ2Ä:ÑNÔNˆBÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?×BÒBÅ2Ä:ÑNÔNˆBˆBå˜E 2 r¨2¨rÑ2Ô2ˆCÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?×BÒBÅ2Ä:ÑNÔNˆBÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?×BÒBÅ2Ä:ÑNÔNˆBÝŒH�U˜B Ñ"Ô"ˆÝŒH�U˜B Ñ"Ô"ˆ‰Ý	Œ��V‘˜Q šU AÑ	&Ô	&×	)Ò	)­"¬*Ñ	5Ô	5€BÝ	Œ��V‘˜dÑ" A¨¢E¨1Ñ	-Ô	-×	0Ò	0µ´Ñ	<Ô	<€Bõ Ô˜B Ñ"Ô"¥IÔ$4°R¸Ñ$;Ô$;€q€AØð 4Ø�DŠD�”Ô#Ñ$Ô$×'Ò'­¬
Ñ3Ô3ˆØ�DŠD�”Ô#Ñ$Ô$×'Ò'­¬
Ñ3Ô3ˆà�ŠU˜ 
Ñ*¨dÑ2°aÒ7¸EÀSºLÑIÑJ€EÝ„HˆQ�‰Z�Ô)¨!­YÔ-=¸bÀ!Ñ-DÔ-DÑEÔEÀuÑMÔMÐMÝ„HˆQ�‰Z˜$Ñ¥	Ô 0°µIÔ4DÀRÈÑ4KÔ4KÑ LÔ LÈeÑTÔTÐTÐTÐTr"   c           	      ób  ‡‡— |j         |j        cŠŠt          ˆˆfd„t          d¦  «        D ¦   «         ¦  «        }|                      ¦   «         dz  }t          t          j        |d¦  «        f         | g|¢|‘| j        d         ‘| j        d         ‘|p| j        d         |j	        z  dz  ‘|j
        ‘t          ‰¦  «        ‘|j        ddt          |j        ¦  «        z
  z  z   ¢t          |j        ¦  «        ¢‰d         j        d         ‘|                      d¦  «        ‘|‘|j        ‘d‘R dd	iŽS )
Nc           	   3   óº   •K  — | ]U}‰t          |t          ‰¦  «        d z
  ¦  «                 ‰t          |t          ‰¦  «        d z
  ¦  «                 fD ]}|V — ŒŒVdS )r;   N)r”   r¼   )Ú.0rÛ   r©   rv   rw   s      €€r    ú	<genexpr>zsplit_rope.<locals>.<genexpr>ç   sp   øè è € ÐmÐm˜q¸3½sÀ1ÅcÈ#ÁhÄhÈqÁjÑ?QÔ?QÔ;RÐTWÕX[Ð\]Õ_bÐcfÑ_gÔ_gÐhiÑ_iÑXjÔXjÔTkÐ:lÐmÐm°�UÐmÐmÐmÐmÐmÐmÐmr"   r~   r?   é   r;   )r;   r   r'   F)Úaxis_cosÚaxis_sinr»   Úranger®   rß   r)   r*   r(   Únum_attention_headsÚpad_sizer¼   Ú
grid_sizesr¾   r½   r�   Úuse_fp32_freqs)	r   ÚcacheÚlayoutÚ
head_startÚdim_headÚpointersr6   rv   rw   s	          @@r    Ú
split_roperñ   å   s   øø€ ØŒ ¤€H€CˆÝÐmÐmÐmÐmÐm¥E¨!¡H¤HÐmÑmÔmÑmÔm€HØ�7Š7‰9Œ9˜‰>€DÝ�œ D¨#Ñ.Ô.Ð0Ô1Ø	ðTØðTØðTØœG AœJðTØ()¬°¬
ðTØ5=Ð5hÀÄÈÄÈuÔOhÑAhÐmnÑ4nðTàŒðTå˜C™œðTà$*Ô$5¸ÀÅCÈÔHYÑDZÔDZÑ@ZÑ8[Ñ$[ðTå_hÐinÔixÑ_yÔ_yðTà{~ð  @Aô  |Bô  |Hð  IJô  |KðTð 	
�Š�‰ŒðTð  ðTð "'Ô!5ðTð 8;ðTð Tð Tð NSðTð Tð Tr"   r�   )r   N)Ú__doc__r)   Útriton.languageÚlanguager   ÚjitÚ	constexprr!   r7   r+   Útriton.language.extrar8   Úshared.kernels.triton_tunerr9   rœ   rŸ   r•   rz   r    Ú%shared.kernels.triton_compilation_logr¡   rª   r³   r¹   r¾   r¯   rß   rñ   r�   r"   r    ú<module>rú      sX  ðØ pÐ pØ €€€Ø Ð Ð Ð Ð Ð ð €„˜s˜eÀSÀEÐJÑJÔJð+ R¤\ð +¸"¼,ð +ð +ð +ñ KÔJð+ðð ð ð ð  €€€Ø +Ð +Ð +Ð +Ð +Ð +Ø 3Ð 3Ð 3Ð 3Ð 3Ð 3à	ˆÐ1°1°dÑ	;Ô	;€Ø<€Ø€
ð €„ÐTÐTÐTð  vlð  vlð  vlð  mñ  mô  mð,CÈÌð ,CÐ[]Ô[gð ,CÐlnÔlxð ,CØ”\ð,CØ*,¬,ð,CØEGÄ\ð,Cð ,Cð ,Cñ mô  mð,Cð^ 49ð ð ð ð ð ð4 TÐ SÐ SÐ SÐ SÐ SØ !Ð !Ñ #Ô #Ð #ð €„˜s LÐ1ÐSVÐXdÐReÐfÑfÔfð/°r´|ð /ÐPRÔP\ð /ÐegÔeqð /ð /ð /ñ gÔfð/ð(
ð 
ð 
ð „ð˜2œ<ð ð ð ñ „ðð:ð :ð :ð
Fð Fð Fð
 €„ÐfÐfÐfð  HPð  HPð  HPð  Qñ  Qô  Qð,UØœ\ð,UØ02´ð,Uà”lð,Uà46´Lð,UàFHÄlð,UàXZÔXdð,Uð œLð,Uð ”lð	,Uð ,.¬<ð	,Uð ,Uð ,Uñ Qô  Qð,Uð^Tð Tð Tð Tð Tð Tr"   