o پi<@sdddlZddlmZmZmZddlZddlZddlZddlm Z ddl m Z ej de jfddZdr@ is_neox_stylereturncCs|d|j}|d|j}|rtj|ddd\}}n|ddddf}|ddddf}||||}||||}|rLtj||fddStj||fdddS)z Args: x: [num_tokens, num_heads, head_size] cos: [num_tokens, head_size // 2] sin: [num_tokens, head_size // 2] is_neox_style: Whether to use the Neox-style or GPT-J-style rotary positional embeddings. r'dim.Nr) unsqueezer?r r!chunkcatstackflatten)rr>r@rMx1x2o1o2rrr_apply_rotary_embVsr\cseZdZdedededededejddffd d Zdeee fdej fd d Z dej fd dZ  ddej dej de ej de ej deej ej ff ddZZS)RotaryEmbeddingr.rArBrCrMr rNNcsNt||_||_||_||_||_||_|}||j d|dddS)NrFF) persistent) super__init__r.rArBrCrMr _compute_cos_sin_cacheregister_buffer)selfr.rArBrCrMr cache __class__rrr`us zRotaryEmbedding.__init__cCs(d|tjd|jdtjd|j}|S)Nr8rr'r )r!r;rAfloat)rcrCrGrrr_compute_inv_freqs z!RotaryEmbedding._compute_inv_freqcCsR||j}tj|jtjd}td||}|}|}tj ||fdd}|S)zCompute the cos and sin cache.r z i,j -> ijrPrQ) rhrCr!r;rBrgeinsumr>r@rU)rcrGrHrIr>r@rdrrrras z&RotaryEmbedding._compute_cos_sin_cache positionsr4r5offsetscCs.|dur||}|}|jd}|jd|}|jddd\}}|j} ||d|j}|dd|jf} |d|jdf} t| |||j } t j | | fdd | }| |j}|dur|j} ||d|j}|dd|jf} |d|jdf}t| |||j } t j | |fdd | }| |j}||fS)z-A PyTorch-native implementation of forward().Nrr'rPrQ.)rWshaperF index_selectrTr-r.rAr\rMr!rUreshaper?r )rcrjr4r5rk num_tokenscos_sinr>r@ query_shape query_rot query_pass key_shapekey_rotkey_passrrrforward_natives,    zRotaryEmbedding.forward_native)NN)__name__ __module__ __qualname__r boolr!r r`rrgTensorrhrarrrw __classcell__rrrerr]ss> r]cCst||||||d|S)zEInitialize Torch Native RotaryEmbedding based on vLLM implementation.r.rArBrCrMr )r]r?)r.rArBrCrMr rrrrget_torch_rotary_embeddingsrcCsHzddlm}WntytdYnw|||||||d|S)z$Initialize SglKernelRotaryEmbedding.r)SglKernelRotaryEmbeddingzDSglKernelRotaryEmbedding is not available. Test case can be removed.r~)#sgl_kernel.testing.rotary_embeddingr ImportErrorpytestskipr?)r.rArBrCrMr rrrrrget_sgl_rotary_embeddings" rcCs|dur |dus JdS|dusJt|rJdt|r(Jd|tjkr/dnd}|tjkr8dnd}tjj||||ddS)z4Compare results between JIT and SGL implementations.NzNaN in JIT resultszNaN in SGL resultsg{Gz?gh㈵>)atolrtol)r!isnananyr<testing assert_close)jit_outsgl_outr rrrrrcompare_resultss  rzhead_size, rotary_dim, max_position_embeddings, base, is_neox_style, dtype, device, batch_size, seq_len, num_q_heads, num_kv_headsc Cs,g|]\}}dddddtjd||ddf qS)@@Tr)r!bfloat16).0bsslrrr sr))rr) r)r)r)r'r)rrrrTrrrrri'r'rri7'iFr key_is_nonec  Cst|||||| | \} } }t|||||}t|||||||}||_t| }| | }}| | }}td}td}| rMd}d}t ddd||}}tj tj |t ddd||}t ||||||d\}}Wdn1swYtj |t ddd||}|j|||d\}}Wdn1swYtj t|||t|||dS)z8Test correctness of JIT rotary embedding implementation.rNgY@)r)r#rjr4r5r.rFis_neoxrjr4r5)r7rLrrFr! randn_likecloneget_device_moduleStreamr%r synchronizestreamrrwr)r.rArBrCrMr rr/r0r1r2rr4r5r6rFtorch_rotary_embr query_jitkey_jit query_torch key_torch stream_jit stream_kernelr_jitr_torch query_jit_out key_jit_outquery_torch_out key_torch_outrrrtest_correctnesssb/           rc ( Cs\t|||||| | \} } } t|||||}t|||||||}||_d}t|D]&}| | }}t| |||||d| | }}|j| ||dq(d}tj t }t|D]}| | }}t| |||||dq^tj t ||}tj t }t|D]}| | }}|j| ||dqtj t ||}| | }}| | } }!t| |||||d\}"}#|j| | |!d\}$}%t |"|$|t |#|%|||}&t d|d|d|&t d|d d d |d d d |d kr |d kr||ntd}'t d|'dd|d kr*|d ks,JdS)zPPerformance test comparing JIT and SGL implementations with accuracy validation.rrrr(z Performance Test - Batch=z , SeqLen=z , Tokens=zJIT: iz.9fz ms, SGL: rrinfzSpeedup (SGL/JIT): z.2frN)r7rLrrFrrr forward_cudar!rrtimerprintrg)(r.rArBrCrMr rr/r0r1r2r4r5r6rFsgl_rotary_embwarmupr query_warmkey_warmquery_sgl_warm key_sgl_warm iteration start_timerrjit_time query_sglkey_sglsgl_timequery_jit_final key_jit_finalquery_sgl_final key_sgl_finalrr query_sgl_out key_sgl_outr3speeduprrrtest_performancegs               " r)r)$rtypingrrrrr!tritontriton.languagelanguager sglang.jit_kernel.pos_encrjit constexprrrgr%r7rLr|r{r\nnModuler]rrrmark parametrizerr<float16rrrrrrs   T      K