§
     ’Þ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	m
Z
 ddlmZmZmZmZmZmZ  ej        d¦  «        Zej        d„ ¦   «         Z ej        dgdg¬	¦  «        d
ej        dej        dej        f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        f
d„¦   «         Zd„ Z ej        ddgddg¬	¦  «        d
ej        dej        dej        f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        dej        fd#„¦   «         Zd&d$„Zdd%lmZ  e¦   «          dS )'a‘  H3 denoiser fusions of denoiser_triton.py for GPUs whose Triton has no BF16 conversion instructions (compute capability 7.5): BF16
tensors are passed as int16 views and converted by the bit helpers of bf16_bits.py (shift to FP32 bits; round to nearest even back,
as PyTorch's casts). Same operations and rounding points as the original kernels: the modulation is bit-identical to the eager path,
the Q/K RMSNorm + RoPE rounds once after the rotation (error against FP64 below the eager path's: 2.22e-3 vs 2.44e-3 at the H3 block
shape on an RTX 2080 Ti, specs/TURING_H3_FL2VA_AUDIT_20261021.md). ops.py selects them where the device is below the BF16 floor.é    N)Ú	libdevice©Úbf16_bits_to_fp32Úfp32_to_bf16_bits)Ú_ROPE_CONFIGSÚ
_TUNE_ROWSÚ_TUNERÚ	_axis_idsÚ	_axis_rowÚmodulation_stridei€?  c                 ó:   — t          t          | ¦  «        ¦  «        S ©Nr   )Úvs    úKC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\denoiser_bits_triton.pyÚ_round_bf16r      s   € åÕ.¨qÑ1Ô1Ñ2Ô2Ð2ó    ÚN)Údo_not_specializeÚdo_not_specialize_on_alignmentÚDÚ	ROWS_BITSÚBLOCKc                 óF  — 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          t          j        | |z   |d¦  «        ¦  «        }
|rMt          t          j        ||	z   |d¦  «        ¦  «        }t          t          j        ||	z   |d¦  «        ¦  «        }nLt          t          j        ||	z   |d¦  «        ¦  «        }t          t          j        ||	z   |d¦  «        ¦  «        }t          d|z   ¦  «        }t          |
|z  ¦  «        }t          j        | |z   t          ||z   ¦  «        |¦  «         d S )Nr   ç      ð?©
ÚtlÚ
program_idÚtoÚint64Úaranger   Úloadr   Ústorer   )ÚXÚSCALEÚSHIFTr   r   r   r   ÚiÚvalidÚcolÚxÚscaleÚshiftÚfactorÚscaleds                  r   Ú_modulate_bitsr.      sx  € å
Œ�aÑÔ×Ò�BœHÑ%Ô%¨Ñ-µ´	¸!¸UÑ0CÔ0CÑC€AØ�ŠE€EØˆ5�y�A‚~€~ÝŒ}˜QÑÔ 1¨¡:Ñ.°%Ñ7½"¼)ÀAÀuÑ:MÔ:MÑMˆˆà�!‰eˆÝ�"œ' ! a¡%¨°Ñ2Ô2Ñ3Ô3€AØð <Ý!¥"¤'¨%°#©+°u¸aÑ"@Ô"@ÑAÔAˆÝ!¥"¤'¨%°#©+°u¸aÑ"@Ô"@ÑAÔAˆˆå�BœG E¨C¡K°¸Ñ:Ô:Ñ;Ô;ˆÝ�BœG E¨C¡K°¸Ñ:Ô:Ñ;Ô;ˆÝ˜˜u™Ñ%Ô%€FÝ˜˜V™Ñ$Ô$€FÝ„HˆQ�‰UÕ% f¨u¡nÑ5Ô5°uÑ=Ô=Ð=Ð=Ð=r   c                 ó,  —  || ¦  «        }|                      t          j        ¦  «        }| j        d         }|j        t          j        k    }|r>|                      t          j        ¦  «        |                      t          j        ¦  «        fn||f\  }}|D ]Y\  }	}
}|
|	k    rM|
|	z
  |z  }t          t          j        |d¦  «        f         ||	|
…         ||         ||         |||dd¬¦  «         ŒZt          j	        j
                             |¦  «         |S )zÞnorm(x) (BF16), then per segment x * (1 + scale[row]) + shift[row] with the eager path's BF16 roundings, in place on the norm
    output; the modulation rows FP32 (rounded to BF16 as the eager cast) or BF16 (read as bits).éÿÿÿÿé   F©Úenable_fp_fusion)ÚviewÚtorchÚint16ÚshapeÚdtypeÚbfloat16r.   ÚtritonÚcdivÚautogradÚgraphÚincrement_version)r)   Únormr+   r*   ÚsegmentsÚoutputÚbitsÚwidthÚ	rows_bitsÚstartÚstopÚrowÚsizes                r   Únorm_modulaterI   +   s"  € ð ˆT�!‰WŒW€FØ�;Š;•u”{Ñ#Ô#€DØŒG�BŒK€EØ”�uœ~Ò-€IØIRÐf�E—J’J�uœ{Ñ+Ô+¨U¯ZªZ½¼Ñ-DÔ-DÐEÐEÐY^Ð`eÐXf�L€Eˆ5Ø$ð Wð WÑˆˆt�SØ�%Š<ˆ<Ø˜5‘L EÑ)ˆDÝ�FœK¨¨dÑ3Ô3Ð5Ô6°t¸EÀ$¸JÔ7GÈÈsÌÐUZÐ[^ÔU_ÐaeÐglÐnwÐy}ð  QVð  Wñ  Wô  Wð  WøÝ	„NÔ×*Ò*¨6Ñ2Ô2Ð2Ø€Mr   )ÚROWSÚHEADSÚSEQÚXBÚXSÚXHÚRBÚRSÚRPÚRCÚPAIRSc                 óø  — 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          t          j        ||d d …d f         d¦  «        ¦  «        }|d|z  k     }t          j        ||k     ||z   t          j        |||z
  |¦  «        ¦  «        }t          t          j        | |d d …d f         z   |d d d …f         z   |d d …d f         d¦  «        ¦  «        }|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         }&t          t          j        ||z   ¦  «        ¦  «        d d d …f         }'t          t          j        ||z   ¦  «        ¦  «        d d d …f         }(t          ||&z  |'z  ¦  «        })t          ||&z  |(z  ¦  «        }|d d …d f         |	z  |d d …d f         |
z  z   t          j        ||k     |||z
  ¦  «        d d d …f         |z  z   }*|d d …d f         |d d d …f         z  }+t          t          j        ||*z   |+t           ¦  «        ¦  «        },t          t          j        ||*z   |z   |+d¦  «        ¦  «        }-|d d d …f         |k     }.t          j        t          j        |.|,|-¦  «        t          j        |.|)|¦  «        t          j        |.|- |,¦  «        t          j        |.||)¦  «        z  ¦  «        }/t          j        |t%          t          j        |d d d …f         |/|)¦  «        ¦  «        |d d …d f         ¦  «         d S )Nr   é   é€   é   é    é   )r   r   r   r   r    r   r!   ÚwhereÚreshapeÚpermuteÚsplitÚsumÚfmar   Úrsqrtr   ÚBF16_ONEr"   r   )0r#   ÚWÚROPErJ   rK   rL   rM   rN   rO   rP   rQ   ÚEPSrR   rS   r   rT   r   ÚrowsÚokÚheadÚtokenÚseqÚbatchÚcÚbaseÚptrr)   ÚrotatedÚpartnerÚotherÚlanesÚevenÚoddÚx0Úx2Úx1Úx3ÚtotalÚinvÚwÚw_otherÚyÚangleÚmaskÚcosÚsinÚfirstÚouts0                                                   r   Ú_rms_rope_bitsr„   ;   sU  € õ Œ=˜ÑÔ×Ò�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¨!¨!¨!¨T¨'¤{°AÑ6Ô6Ñ7Ô7€AØ�!�e‘)Šm€GÝŒh�q˜5’y ! e¡)­R¬X°g¸qÀ5¹yÈ!Ñ-LÔ-LÑMÔM€GÝ�bœg a¨$¨q¨q¨q°$¨w¬-Ñ&7¸'À$ÈÈÈÀ'Ô:JÑ&JÈBÈqÈqÈqÐRVÈwÌKÐYZÑ[Ô[Ñ\Ô\€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Ý�"œ' ! a¡%™.œ.Ñ)Ô)¨$°°°¨'Ô2€AÝ¥¤¨¨G©Ñ 4Ô 4Ñ5Ô5°d¸A¸A¸A°gÔ>€GÝ�A˜‘G˜a‘KÑ Ô €AÝ˜ ™ gÑ-Ñ.Ô.€EØ�!�!�!�T�'ŒN˜RÑ # a a a¨ g¤,°Ñ"3Ñ3µb´h¸qÀ5ºyÈ!ÈQÐQVÉYÑ6WÔ6WÐX\Ð^_Ð^_Ð^_ÐX_Ô6`ÐceÑ6eÑe€EØˆaˆaˆa�ˆgŒ;˜  q q q Ô)Ñ)€DÝ
�BœG D¨5¡L°$½ÑAÔAÑ
BÔ
B€CÝ
�BœG D¨5¡L°2Ñ$5°t¸QÑ?Ô?Ñ
@Ô
@€CØˆd�A�A�AˆgŒJ˜Ò€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Õ#¥B¤H¨W°T¸1¸1¸1°WÔ-=¸sÀAÑ$FÔ$FÑGÔGÈÈAÈAÈAÈtÈGÌÑUÔUÐUÐUÐUr   c                 ó�  ‡ ‡‡‡‡‡	‡
‡‡‡‡— ‰ j         \  }ŠŠŠ‰j         d         dk    rdn‰                     d¦  «        Š
‰                      t          j        ¦  «        |                     t          j        ¦  «        ‰                     t          j        ¦  «        cŠŠŠˆˆˆˆˆ
ˆˆfd„Š	ˆˆˆ	ˆˆ ˆfd„}t          j        d‰ j        › d‰› d‰› d‰j         d	         › d
‰                      d¦  «        › �	‰ j        t          ||‰z  ‰z  t          k    ¦  «        } ‰	‰|¦  «         t          j        j                             ‰ ¦  «         ‰ S )a  In place on Q or K (batch, tokens, heads, head_dim) in BF16, any strides with contiguous channels; weight the BF16 RMSNorm
    weight; rope the BF16 (1 | batch, tokens, 1, pairs, 2) cos / sin table of the split-half rotation of the first 2 * pairs channels.r   rZ   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   rZ   é   rX   éþÿÿÿF)Ú	num_warpsr3   )r7   r„   r:   r;   Ústride)ÚxxÚconfigÚblockÚwarpsrf   ÚdimÚepsÚheadsÚropeÚ
rope_batchÚ	rope_bitsÚw_bitss        €€€€€€€r   Úlaunchzrms_rope_.<locals>.launchh   sE  ø€ Ø‰ˆˆuØŒx˜Œ{˜RœX aœ[Ñ(¨5Ñ0ˆÝ�œ D¨%Ñ0Ô0Ð2Ô3°B¸À	È4ÐQVÐXZÔX`ÐabÔXcð 	YÐfh×foÒfoÑfqÔfqÐrtÐstÐrtÔfuð 	Yð  xBð 	Yð  DH÷  DOò  DOð  PQñ  DRô  DRð 	Yð  TWð 	YØ48·K²KÀ±N´Nð	YØDHÇKÂKÐPQÁNÄNð	YØTWð	YØY]ÔYcÐdfÔYgð	YØinð	Yð 	YØzð  SXð	Yð 	Yð 	Yð 	Yð 	Yð 	Yr   c                  ó  •‡— t          dt          ‰t          ‰z  ¦  «        ¦  «        } ‰                     d| ‰‰                     d¦  «        f¦  «        dd ‰…f         Š‰                     ‰d d…d | …f         ¦  «         ˆˆfd„S )NrZ   rV   .c                 ó   •—  ‰‰| ¦  «        S r   © )rŒ   r–   Úscratchs    €€r   ú<lambda>z4rms_rope_.<locals>.prepare_scratch.<locals>.<lambda>r   s   ø€ ˜f˜f W¨fÑ5Ô5€ r   )ÚmaxÚminr   Ú	new_emptyrŠ   Úcopy_)Útokensrš   r�   r‘   r–   rj   r)   Úx_bitss    @€€€€€€r   Úprepare_scratchz"rms_rope_.<locals>.prepare_scratchn   sŠ   øø€ Ý�Q�˜C¥¨uÑ!4Ñ5Ô5Ñ6Ô6ˆØ×"Ò" A v¨u°a·h²h¸q±k´kÐ#BÑCÔCÀCÈÈ#ÈÀIÔNˆØ�Š�f˜R˜a˜R  & ˜[Ô)Ñ*Ô*Ð*Ø5Ð5Ð5Ð5Ð5Ð5r   Úrms_rope_bitsú|r)   z|prˆ   z|srV   )r7   rŠ   r4   r5   r6   r	   rŒ   r8   Údevicer   r   r<   r=   r>   )r)   Úweightr’   r�   rk   r¢   rŒ   r�   r‘   r–   r“   r”   rj   r•   r¡   s   ` ``   @@@@@@@@r   Ú	rms_rope_r§   a   s�  øøøøøøøøøøø€ ð œWÑ€Eˆ3��sØ”j ”m qÒ(Ð(��¨d¯kªk¸!©n¬n€JØ !§¢¥u¤{Ñ 3Ô 3°V·[²[ÅÄÑ5MÔ5MÈtÏyÊyÕY^ÔYdÑOeÔOeÐ€FˆF�IðYð Yð Yð Yð Yð Yð Yð Yð Yð Yð Yð6ð 6ð 6ð 6ð 6ð 6ð 6ð 6ð 6ð 6õ Œ]˜?¨q¬wÐ,gÐ,g¸Ð,gÐ,gÀÐ,gÐ,gÈÌ
ÐSUÌÐ,gÐ,gÐZ[×ZbÒZbÐcdÑZeÔZeÐ,gÐ,gÐijÔiqõ  tAð  CRØ  3™;¨Ñ.µ*Ò<ñ>ô >€Fà
€Fˆ6�6ÑÔÐÝ	„NÔ×*Ò*¨1Ñ-Ô-Ð-Ø€Hr   Ú
MOD_STRIDEÚ	PER_TOKENc	                 óP  — 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          t          j        | |z   |	|k     d¦  «        ¦  «        }t          t          j        ||z   |	|k     d¦  «        ¦  «        }t          t          j        ||z   |	|k     d¦  «        ¦  «        }t          d|z   ¦  «        }t          ||z  ¦  «        }t          j        ||z   t          ||z   ¦  «        |	|k     ¦  «         d S )Nr   rZ   r   r   )r#   r$   r%   ÚYr   r¨   r   r©   r   r&   ÚgroupÚoffsetÚsr)   r*   r+   r,   Úvalues                     r   Ú_scale_shift_bitsr°   {   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°QÑ7Ô7Ñ8Ô8€AÝ�bœg e¨a¡i°°Q²¸Ñ:Ô:Ñ;Ô;€EÝ�bœg e¨a¡i°°Q²¸Ñ:Ô:Ñ;Ô;€EÝ˜˜u™Ñ%Ô%€FÝ˜˜F™
Ñ#Ô#€EÝ„HˆQ�‰ZÕ*¨5°5©=Ñ9Ô9¸1¸qº5ÑAÔAÐAÐAÐAr   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	        ¦  «        |                     t           j	        ¦  «        |                     t           j	        ¦  «        |                     t           j	        ¦  «        |t          |¦  «        | j        d         |dd¬¦
  «
         |r$t           j        j                             |¦  «         |S )z�denoiser_triton.scale_shift for BF16 x, scale and shift (same shapes; the modulation rows may be strided views): in place or into a new tensor.rZ   r   r1   rV   Fr2   )r5   Ú
empty_liker7   Úis_contiguousÚnumelr°   r:   r;   r4   r6   r   r<   r=   r>   )r)   r*   r+   Úin_placerƒ   Ú	per_tokenÚgroupsrH   s           r   Úscale_shiftr¸   �   sJ  € àÐ
0ˆ!ˆ!�UÔ-¨aÑ0Ô0€CØ”˜A” !¤'¨!¤*Ò,ÐF°×1DÒ1DÑ1FÔ1F€IØÐ@ˆQˆQ ¤¨Q¤°%´+¸a´.Ñ!@€FØ�7Š7‰9Œ9˜Ñ€DÝ•v”{ 4¨Ñ.Ô.°Ð7Ô8Ø	�Š�uŒ{ÑÔ˜UŸZšZ­¬Ñ4Ô4°e·j²jÅÄÑ6MÔ6MÈsÏxÊxÕX]ÔXcÑOdÔOdÐfjÕl}ð  Dñ  mEô  mEð  GHô  GNð  OPô  GQØ�4Øð ñ  ô  ð  ð ð 4ÝŒÔ×.Ò.¨sÑ3Ô3Ð3Ø€Jr   )r   ÚTOKENSÚWIDTHÚG0ÚG1ÚG2ÚBATCH_STRIDEÚ
HEAD_STARTÚHALFÚPADÚAXESÚA0ÚA1ÚA2Ú
FREQ_WIDTHÚFP32Ú	FREQ_BITSc                 óê  — 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    rRt          |||||¦  «        }&t          j	        ||&|z  z   |!z   |%d¦  «        }'t          j	        ||&|z  z   |!z   |%d¦  «        }(n©|$dk    rRt          |||||¦  «        }&t          j	        ||&|z  z   |!z   |%d¦  «        }'t          j	        ||&|z  z   |!z   |%d¦  «        }(nQt          |||||¦  «        }&t          j	        ||&|z  z   |!z   |%d¦  «        }'t          j	        ||&|z  z   |!z   |%d¦  «        }(|rt          |'¦  «        t          |(¦  «        }*})n>|'                     t           j        ¦  «        |(                     t           j        ¦  «        }*})t          j        |%|)|"¦  «        }"t          j        |%|*|#¦  «        }#�Œ¥t          t          j	        | |z   ||k     d¦  «        ¦  «        }+t          t          j	        | |z   |
z   ||k     d¦  «        ¦  «        },t          j        |+|"¦  «        t          j        |,|"¦  «        }-}$|st          |$¦  «        }$t          |-¦  «        }-||k     ||z   |z  dk    ||k    z  z  }%t          j        | |z   t!          t          j        |$t          j        |,|#¦  «        ¦  «        ¦  «        |%¦  «         t          j        | |z   |
z   t!          t          j        |-t          j        |+|#¦  «        ¦  «        ¦  «        |%¦  «         d S )Nr   rV   r   g        rZ   )r   r   r   r   r    ÚfullÚfloat32Ústatic_ranger   r!   r   r[   r   Ú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   r&   Úfri   rk   rh   Úlocalr­   Úaxisr(   rl   r®   Úar'   rG   ÚcrÚsrÚcvÚsvru   rw   Úbs.                                                 r   Ú_split_rope_bitsrß   ž   sû  € õ 	Œ�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Ý”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?ˆBˆBØ�!ŠVˆVÝ˜E 2 r¨2¨rÑ2Ô2ˆCÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?ˆBÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?ˆBˆBå˜E 2 r¨2¨rÑ2Ô2ˆCÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?ˆBÝ”˜˜c JÑ.Ñ.°Ñ4°e¸QÑ?Ô?ˆBØð 	:Ý& rÑ*Ô*Õ,=¸bÑ,AÔ,A�ˆBˆBà—U’U�2œ:Ñ&Ô&¨¯ª­b¬jÑ(9Ô(9�ˆBÝŒH�U˜B Ñ"Ô"ˆÝŒH�U˜B Ñ"Ô"ˆ‰Ý	�2œ7 1 v¡:¨q°1ªu°aÑ8Ô8Ñ	9Ô	9€BÝ	�2œ7 1 v¡:°Ñ#4°a¸!²e¸QÑ?Ô?Ñ	@Ô	@€BÝÔ˜B Ñ"Ô"¥IÔ$4°R¸Ñ$;Ô$;€q€AØð Ý˜‰NŒNˆÝ˜‰NŒNˆØ�ŠU˜ 
Ñ*¨dÑ2°aÒ7¸EÀSºLÑIÑJ€EÝ„HˆQ�‰ZÕ*­9Ô+;¸A½yÔ?OÐPRÐTUÑ?VÔ?VÑ+WÔ+WÑXÔXÐZ_Ñ`Ô`Ð`Ý„HˆQ�‰Z˜$ÑÕ 1µ)Ô2BÀ1ÅiÔFVÐWYÐ[\ÑF]ÔF]Ñ2^Ô2^Ñ _Ô _ÐafÑgÔgÐgÐgÐgr   c           	      ó´  ‡‡‡	— |j         |j        cŠŠ	|j         Št          ˆˆˆ	fd„t	          d¦  «        D ¦   «         ¦  «        }|                      ¦   «         dz  }t          t          j        |d¦  «        f         |  	                    t          j        ¦  «        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 )
z{denoiser_triton.split_rope for a BF16 x (in place); the cos / sin tables FP32 (cache.use_fp32_freqs) or BF16, read as bits.c           	   3   óü   •K  — | ]v}‰t          |t          ‰¦  «        d z
  ¦  «                 ‰t          |t          ‰¦  «        d z
  ¦  «                 fD ]'}‰r|                     t          j        ¦  «        n|V — Œ(ŒwdS )rZ   N)r�   Úlenr4   r5   r6   )Ú.0rÙ   ÚtrB   r€   r�   s      €€€r   ú	<genexpr>zsplit_rope.<locals>.<genexpr>Ó   sÒ   øè è € ð  Ið  I¸!ÐRUÕVYÐZ[Õ]`ÐadÑ]eÔ]eÐhiÑ]iÑVjÔVjÔRkÐmpÕqtÐuvÕx{Ð|ñ  yAô  yAð  DEñ  yEñ  rFô  rFô  nGð  RHð  Ið  IÈA¨$Ð5�1—6’6�%œ+Ñ&Ô&Ð&°Að  Ið  Ið  Ið  Ið  Ið  Ið  Ir   r‡   rV   é   rZ   )rZ   r   r3   F)Úaxis_cosÚaxis_sinÚuse_fp32_freqsÚtupleÚranger´   rß   r:   r;   r4   r5   r6   r7   Únum_attention_headsÚpad_sizerâ   Ú
grid_sizesr
   Ú	rope_axesrŠ   )
r)   ÚcacheÚlayoutÚ
head_startÚdim_headÚtablesrH   rB   r€   r�   s
          @@@r   Ú
split_roperõ   Ï   sï  øøø€ àŒ ¤€H€CˆØÔ#Ð#€DÝð  Ið  Ið  Ið  Ið  Ið  I½uÀQ¹x¼xð  Iñ  Iô  Iñ  Iô  I€FØ�7Š7‰9Œ9˜‰>€DÝ�Vœ[¨¨sÑ3Ô3Ð5Ô6Ø	�Š�uŒ{ÑÔðZØ$ðZØ&*ðZØ,-¬G°A¬JðZØ89¼À¼
ðZØEMÐExÐQRÔQXÐYZÔQ[Ð_dÔ_xÑQxÐ}~ÑD~ðZàŒðZå˜C™œðZà$*Ô$5¸ÀÅCÈÔHYÑDZÔDZÑ@ZÑ8[Ñ$[ðZå_hÐinÔixÑ_yÔ_yðZà{~ð  @Aô  |Bô  |Hð  IJô  |KðZð 	
�Š�‰ŒðZð  ðZð "'Ô!5ðZð 8<ðZð >AðZð Zð Zð TYðZð Zð Zr   )Ú!install_triton_compilation_logger)r   N) Ú__doc__r5   r:   Útriton.languageÚlanguager   Útriton.language.extrar   Úshared.kernels.bf16_bitsr   r   Úshared.kernels.denoiser_tritonr   r   r	   r
   r   r   Ú	constexprrb   Újitr   r.   rI   r„   r§   r°   r¸   rß   rõ   Ú%shared.kernels.triton_compilation_logrö   r™   r   r   ú<module>r      sT  ððDð Dð
 €€€Ø €€€Ø Ð Ð Ð Ð Ð Ø +Ð +Ð +Ð +Ð +Ð +à IÐ IÐ IÐ IÐ IÐ IÐ IÐ IØ uÐ uÐ uÐ uÐ uÐ uÐ uÐ uÐ uÐ uÐ uÐ uÐ uÐ uÐ uÐ uàˆ2Œ<˜ÑÔ€ð „ð3ð 3ñ „ð3ð €„˜s˜eÀSÀEÐJÑJÔJð>¨"¬,ð >À2Ä<ð >ÐXZÔXdð >ð >ð >ñ KÔJð>ð&ð ð ð  €„ÐTÐTÐTð  vlð  vlð  vlð  mñ  mô  mð"VÈbÌlð "VÐ`bÔ`lð "VÐqsÔq}ð "VØœ,ð"VØ/1¬|ð"Vð "Vð "Vñ mô  mð"VðJð ð ð4 €„˜s LÐ1ÐSVÐXdÐReÐfÑfÔfðB¸B¼Lð BÐUWÔUað BÐjlÔjvð Bð Bð Bñ gÔfðBð&ð ð ð €„ÐfÐfÐfð  HPð  HPð  HPð  Qñ  Qô  Qð-hØ"$¤,ð-hØ57´\ð-hàœ<ð-hà9;¼ð-hàKMÌ<ð-hà]_Ô]ið-hð "$¤ð-hð œ<ð	-hð 57´Lð	-hð JLÌð	-hð -hð -hñ Qô  Qð-hð`	Zð 	Zð 	Zð 	Zð TÐ SÐ SÐ SÐ SÐ Sà !Ð !Ñ #Ô #Ð #Ð #Ð #r   