,`i|=GddlmZmZddlZddlmZddlmZddlm Z ddl m Z ddl m Z ddlmZeeZe jerdZn dd lmZn#e$r dd lmZYnwxYw dEd ejd ejdejdejdededejdejdededejdzdedejdejdedededededdf(d Z dEd ejd!ejd"ejd#ejd ejdejdejdededejdejdededejdzdedejdejdedededededdf.d$Zdejrd%nd&fd ejd!ejd"ejd#ejd ejdejdejdededejdejd'ejdzdededejdzdedejdejd(ejdzd)eddf*d*Zd ejd ejd+ejdedejdejddfd,Z dFd-ejd.ejd/ejd0ejd1ejd2ejdzddfd3Z dGd5ejd6ejd7ejd8ejd9ed:ed;edejd?ejd9ed:ed;ededdoejdej,de dej,dzdejf dZ?ed dHdnejdejdejdedejdejdej,dzdedzdejfdZ@eddfejdejfdZAeddejdejfdZBeddejdejfdZCe3ej4j5drbeddnejdejd}ejdejdzdej9dbej9dej9dej9dej9de!de!dejfdZDe3ej4j5dred dFdejdedej9dej9dej,dzdejf dÄZEedĦdejdejdedej9dejf dDŽZFedȦdejdejdedej9dejf dɄZGedʦdejdejdejdejdejdedej9dej9dej9dejfdЄZHe3ej4j5dѦrCedҦdejdejdejdededej9dej9dejfdԄZIdede!fdքZJdnejdejdejdejdejdej,dejfdۄZKdede!fd܄ZLdede!fd݄ZM dFdnejdejdejdejdej,dejdzdejfdZN dHdnejdejdejdejdej,dejdejdzdejdzdejfdZOdede!fdZPdede!fdZQdnejde"ejejffdZR dFdnejdejdejdejdejdej,dejdzdejfdZS dFdejdejdejdejdejdejdedededejdzfdZTdejdejdejdedede!f dZUdejdejfdZVdejdejdejdejdedededefdZWdejdejdejdejd}ejdejdejdejdejdejde!de!fdZXdejdejdejdejd}ejdejdejdejd ejfd ZY dLdoejdejdeded ed e!dejfd ZZe3ej4j5drCed dLdoejdejdej9dej9d ed e!dejfdZ[ dLdoejdeded ed e!dejf dZ\e3ej4j5dr<ed dLdoejdej9dej9d ed e!dejf dZ] dLdoejdejdeded ed e!dejfdZ^ dLdoejdejdeded ed e!dejfdZ_ dMdeejdejdzde!fdZ` dJdnejdejdzdoejdejdzd}ejdejdzdejdzdejdzdejdzdejdzd~ejde dededede!de!de!de!dejf(dZa dNdej,de dej,dzdej,dzdej,dzdej,dzdej,dzde.efdZb dKdnejdejde dej,dzdejdzdejdzdedzdejdzdejdzdedzdejfdZcdoejdej,de dej,dzdejf d Zd dHdnejdejdejdedejdejdej,dzdedzdejfd!Zedfejdejfd"Zfdejdejfd#Zg dFdejdejdejdejd}ejdejdedejdejdejdejdejd$ejdedzfd%Zhdejde"ejejffd&Zie3ej4j5d'r,ed(dnejdejdejfd)Zjdnejdejdejfd*Zk dOdGejd,ejd-e!d.ede"ejejff d/Zldejd,ejdejdejd0ede"ejejff d1Zmdejd,ejdejdejd0ede"ejejff d2Zn dPdGejdejdzd3edzd_ejdzd4e!d-ejdzd5e"eefdzde"ejejffd6Zo dMdeejdejd7ejdzde!de"ejejejff d8Zpdnejdejd}ejdejdzdedbedededede!de!dejfd9Zq dQdGejdejdzdejdzd:e!de"ejejejdzff d;Zrdejdedededej,dzdejf d<Zsdejdejdededejf d=Ztdejdejdededejf d>Zudejdejdejdejdejdededededejfd?Zvdejdejdejdededej9dej9dejfd@ZwdedefdAZx dRdCejdDejdEejdFejdGejdHejdzdIejdzdJejdzdKe!d'ejdzdLejdzdMejdzdNejdOededPejdzdQejdzdRejdzf$dSZydnejdejde"ejejffdTZzdnejdejdUedejfdVZ{ dFdnejdejdWedejdejf dXZ| dFdnejdejdWedejdejf dYZ} dFdnejdejdej,dejdejdWedejdejfdZZ~dGejd-ejfd[Z dFdejdededejd\ejd]ejd^ejdzddfd_Zd`ededejdaejdejd]ejddfdbZ dFdejdcejdededdedeedfedejd\ejd]ejdgejdhejd^ejdzddfdiZdGejd-ejdejd}ejdejdzdjejdzdejd\ejd]ejdedkedledmeduedejfdnZ dSdjejdejdoejdpejdqe!drejdzddfdsZ dSdjejdejdoejdpejdqe!drejdzddfdtZ dTduejdvedwed0edqe!dxedejdyefdzZ dUdGejd-ejdzdejdejdzd}ejdejdzdejdzdejdzdejdzdejdzd~ejdejdejd|ejdjejd}eded~e!de dededede!de!de!de!dedededejf<dZe3ej4drVe3ej4jdr>eddnejdejdaejdjejdejd}ejdejdejdejd~ejde dej9dej9dej9de!ded0ed}ede!de!dejf*dZeddGejd-ejdzdejdejdzd}ejdejdzdejdzdejdzdejdzdejdzd~ejdejdejd|ejdjejd}eded~e!de dededede!de!de!de!f4dZdBejdejdejdejdejdedejdejddfdZdBejdejdejdejdejdedejdejddfdZdejdejd+ejdejdedejddfdZdAejdejdejdejdDejdEe!dejd+ejdedejddfdZdejdejdedejddf dZ dVd-ejdGejdededdf dZ dFdejdejdejdejdejdededejdejdzddfdZ dFdejdejdejdejdedejdzddfdZdejdejdejdejdejdeddfdZdejd+ejdejdededdf dZd+ejdejdejdejdejddf dZdededefdZdedefdZde.ejdejdede!def dZdedejd ejdededdf dZdeddfdZdefdZdede.eddfdZdede"e.ee.effdZdede.e.ede.e.eddfdÄZdede"eejffdńZdejfdDŽZdeddfdɄZdFdedededzdefd̄Zdeddfd̈́Z dLdedejd ejdede!ddf dЄZdedejfdфZdede.ejddfd҄ZdefdӄZdejdedOede"ejejffdքZ dMdejdejdejdejdedejdejdedzdejdejf d?Zdnejdejdejde"ejejffd@ZƐdGdejde!dejfdAZe3ej4j5dBr+edCdejde!dejfdDZdSdS(X) TYPE_CHECKINGLiteralN) init_logger)current_platform) ScalarType)$flashinfer_quant_nvfp4_8x4_sf_layout)cdivcfdS)NcSN)namefns d/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/vllm/_custom_ops.pyzregister_fake..sBr )rs`r register_fakersr)r) impl_abstract@outquery key_cache value_cache num_kv_headsscale block_tablesseq_lens block_size max_seq_len alibi_slopeskv_cache_dtypek_scalev_scaletp_rankblocksparse_local_blocksblocksparse_vert_strideblocksparse_block_sizeblocksparse_head_sliding_stepreturncrtjj|||||||||| | | | | |||||dSr )torchops_Cpaged_attention_v1)rrrrrrrrrrr r!r"r#r$r%r&r'r(s rr.r.!sc* IL##    %'rexp_sum max_logitstmp_outcxtjj|||||||||| | | | | ||||||||dSr )r+r,r-paged_attention_v2)rr/r0r1rrrrrrrrrr r!r"r#r$r%r&r'r(s rr3r3Msl0 IL##    %-rfp8f16query_start_loc fp8_out_scale mfma_typecttjj|||||||||| | | | | ||||||dSr )r+r,_rocm_Cpaged_attention)rr/r0r1rrrrrrrr6rrr r!r"r#r7r8s rpaged_attention_rocmr<sg, I%%   )rkv_cachecXtjj||||||dSr )r+r,_C_cpumla_decode_kvcache)rrr=rrrs rmla_decode_kvcache_cpurAs8 I'' UHe\8routput prefix_output prefix_lse suffix_output suffix_lse output_lsecXtjj||||||dSr )r+r,r-merge_attn_states)rBrCrDrErFrGs rrIrIs7 IL"" M:}jrT q_seqlens kv_seqlensvertical_indexes slash_indexes context_size block_size_M block_size_NcausalcF|d}|d} |d} |d} ||zdz |z} tj|| | |j|j} tj|| | | |j|j}tj|| | |j|j}tj|| | | |j|j}tjj| ||||||||||| | |||fSNrdtypedevice)sizer+zerosrWrXr,r-convert_vertical_slash_indexes)rJrKrLrMrNrOrPrQ batch_size num_heads nnz_slash nnz_verticalnum_rows block_count block_offset column_count column_indexs rr[r[sc##A&&J""1%%I""1%%I#((++L|+a/L@H+IxyyGWK;o L;IxyyGWL;o L IL//     lL @@rvertical_indices_countslash_indices_countc J|d} |d} |d} |d} ||zdz |z}tj| | ||j|j}tj| | || |j|j}tj| | ||j|j}tj| | || |j|j}tjj|||||||||||||| ||||fSrS)rYr+emptyrWrXr,r-(convert_vertical_slash_indexes_mergehead)rJrKrLrMrerfrNrOrPrQr\r]r^r_r`rarbrcrds rririsi##A&&J""1%%I""1%%I#((++L|+a/L@H+IxyyGWK;o L;IxyyGWL;o L IL99  lL @@r positionskey head_size cos_sin_cacheis_neoxcXtjj||||||dSr )r+r,r-rotary_embedding)rjrrkrlrmrns rrprp>s7 IL!!5#y-rinputweightepsiloncTtjj||||dSr )r+r,r-rms_norm)rrqrrrss rruruLs( IL#ufg66666rresidualcTtjj||||dSr )r+r,r-fused_add_rms_norm)rqrvrrrss rrxrxRs( IL##E8VWEEEEErqkv num_heads_q num_heads_k num_heads_vhead_dimepsq_weightk_weight position_idsc btjj|||||||||| | dSr )r+r,r-fused_qk_norm_rope) ryrzr{r|r}r~rrrmrnrs rrrXsK IL##       rlogits prompt_mask output_maskrepetition_penaltiesc|dd|d}tj||z|d}tj|dkd|z |}||z}dS)NrT)dim?r) unsqueezerepeatrYr+where)rrrr penaltiesscalings r apply_repetition_penalties_torchrts~ 099a9@@GG 6;;q>> K+57KSQQIk&1*cIoyAAG gFFFrcTtjj||||dSr )r+r,r-apply_repetition_penalties_rrrrs rapply_repetition_penalties_cudars4  IL,, [*>rc|jr(|rt||||dSt||||dS)awApply repetition penalties to logits in-place. Args: logits: The logits tensor of shape [num_seqs, vocab_size]. prompt_mask: A boolean tensor indicating which tokens appear in the prompt. output_mask: A boolean tensor indicating which tokens appear in the output. repetition_penalties: The repetition penalties of shape (num_seqs, ). N)is_cuda is_contiguousrrrs rapply_repetition_penaltiesrso~ &..00 ' K.B      ) K.B     r quant_dtypescale_ubc tj||}tj||jdzdf|jtj}tjj |||||||||fS)NrWrTrXrW) r+ empty_likerhnumelshaperXfloat32r,r- rms_norm_dynamic_per_token_quant)rqrrrsrrrvrBscaless rrrs e; 7 7 7F [ %+b/ )1-el%-F IL11vvw( 6>rF group_sizeis_scale_transposedc bt|dksJtj||}|rqtj|jd|dz||jdzf|jtjdd} n\tj||jdz|jd|dzf|jtj} tj j |||| ||||d| || fS)NrUrrrTrr) lenr+rrhrrrXr transposer,r-rms_norm_per_block_quant) rqrrrsrrrrvrrBrs rrrs+ z??a      e; 7 7 7F   [_ 1 -u{{}} B/O P<-    )Aq//   [[]]ek"o -u{2*Q-/O P<-     IL)) 1     6>rqweightrrZ split_k_itersthxthyctjrddlm}||||Stjj||||||S)Nr)awq_dequantize_triton)envsVLLM_USE_TRITON_AWQ2vllm.model_executor.layers.quantization.awq_tritonrr+r,r-awq_dequantize)rrrZrrrrs rrrsg =      %$Wfe<<< 9< & &w}cSV W WWrqzerosctjrddlm}||||||Stjj|||||S)Nr)awq_gemm_triton)rrrrr+r,r-awq_gemm)rqrrrrrs rrrsZ NVVVVVVugvv}MMM 9<  O OOra b_q_weight b_gptq_qzeros b_gptq_scalesb_g_idx use_exllama use_v2_formatbitc Xtjj||||||||Sr )r+r,r- gptq_gemmrrrrrrrrs rrrs8 9< ! !     rrz _C::gptq_gemmctj|d|df|j|jS)NrrTrVr+rhrYrWrXrs r_gptq_gemm_fakersC{ VVAYY ** +1718    rq_permcRtjj|||dSr )r+r,r- gptq_shuffle)rrrs rrr-s$ ILh44444rb_metab_scales workspaceb_q_typesize_msize_nsize_kc dtjj||||||j||| Sr )r+r,r-gptq_marlin_24_gemmid rrrrrrrrrs rrr2s6 9< + + :vxHKQW  rrz_C::gptq_marlin_24_gemmc Htj||f|j|jSNr)r+rhrXrWrs r_gptq_marlin_24_gemm_fakerDs${FF+AHAGLLLLrz_C::marlin_gemmcb_biasa_scales global_scaleb_zerosg_idxperm b_q_type_id is_k_fulluse_atomic_adduse_fp32_reduce is_zp_floatc|j}|tjtjfvr|j}tj| | f|j|Sr)rWr+halfbfloat16rhrX)rrrrrrrrrrrrrrrrrrrrWs r_marlin_gemm_fakerRsD, U^4 4 4NE{FF+AHEJJJJrz_C::awq_dequantizec|d}|d}|dz}tj||f|j|jSNrrTrVrYr+rhrWrX) rrrZrrrin_cqout_cout_cs r_awq_dequantize_fakermsJ||Aa {D%= V]SSSSrz _C::awq_gemmc|d}tj|||ddzf|j|jdSr)rYr+rhrWrXsum)rqrrrr num_in_featss r_awq_gemm_faker{s\zz!}} { L',,q//A*= >+<    #a&&  rz_C::machete_mmb_qb_typeout_typeb_group_scales b_group_zeros b_group_sizeb_channel_scalesa_token_scalesschedulec |d} |d} tj| | f|j|jSNrrTrrYr+rhrXrW) rrrrrrrrrrmns rmachete_mm_fakersA FF1II HHQKK{Aq6!(!'BBBBrz_C::machete_prepack_Ba_typegroup_scales_typecBtj|tjSN) memory_formatr+rcontiguous_formatrrrrs rmachete_prepack_B_faker s %:QRRRRrz_C::cutlass_w4a8_mmmaybe_schedulec|d}|d} ||n tj} tj|| f|j| Sr)rYr+rrhrX) rrrrrrrr rr out_dtypes rcutlass_w4a8_mm_fakersO FF1II HHQKK ( 4HH%. {Aq6!()DDDDrz_C::cutlass_pack_scale_fp8cBtj|tjSrrrs rcutlass_pack_scale_fp8_fakerse6MNNNNrz$_C::cutlass_encode_and_reorder_int4bbcBtj|tjSrrrs r%cutlass_encode_and_reorder_int4b_faker1HIIIIrz,_C::cutlass_encode_and_reorder_int4b_groupedcBtj|tjSrrrs r-cutlass_encode_and_reorder_int4b_grouped_fakerrrallspark_w8a16_gemmz_C::allspark_w8a16_gemm b_qweightb_qzerosrsm_count sm_versionCUBLAS_M_THRESHOLDhas_zpn32k16_reorderc r|d} tj| |f|j|jS)Nrrr) rrrrrrrrrr r!rs r_allspark_w8a16_gemm_faker#s2 FF1II{Aq6!(!'BBBBrggml_dequantizez_C::ggml_dequantizeW quant_typerrWcRtj||ftj|jSNrV)r+rhfloat16rXr%r&rrrWs r_ggml_dequantize_faker+s#{Aq6qxHHHHrz_C::ggml_mul_mat_vec_a8Xrowc^tj|jd|f|j|jSNrrV)r+rhrrWrXr%r,r&r-s r_ggml_mul_mat_vec_a8_faker1s*{AGAJ,AGAHMMMMrz_C::ggml_mul_mat_a8cr|d}tj||f|j|jSr/r)r%r,r&r-batchs r_ggml_mul_mat_a8_faker4s2q {E3r8r&r-r9s r_ggml_moe_a8_vec_fakerAs7{FUNC0QQQQrcuda_device_capabilitycJtjj|Sr )r+r,r-cutlass_scaled_mm_supports_fp4rBs rrDrD 9< 6 67M N NNr block_scale_a block_scale_balpharc|jdkr |jdksJ|jd|jd}}tj||f||j}tjj|||||||S)NrUrrV)ndimrr+rhrXr,r-cutlass_scaled_fp4_mm) rrrGrHrIrrrrs rrLrLsw 6Q;;16Q;;; & 71:qwqzqA +q!fIah ? ? ?C IL&&sAq-PUVVV JrcJtjj|Sr )r+r,r-cutlass_scaled_mm_supports_fp8rEs rrNrN*rFrcJtjj|Sr )r+r,r-$cutlass_scaled_mm_supports_block_fp8rEs rrPrP.s 9< < <=S T TTrscale_ascale_bbiasc|tjus|tjusJ|0||jdkr |j|ksJg|jdd|jdR}|d|jd}|jddzdko|jddzdk}tjs|sddl m }|||||||} n]tj |jd|jdf||j } tj j| |||||| j|S)a `cutlass_scaled_mm` implements a fused version of `output = torch.mm((scale_a * a), (scale_b * b)).to(out_dtype)` where scale_a * a and scale_b * b are implemented using numpy-style broadcasting. In order to support blockwise scaling like found in DeepSeek V3 we also support extended "group" broadcast rules. We extend the numpy-style broadcasting rules with the following rule: "if the extent of a dimension in the source shape is between 1 and corresponding extent in the target shape we repeat each element along that dimension src_shape[dim] // target_shape[dim] times consecutively" example if we have: a = [[1, 2], and target_shape = (2, 4) [3, 4]] then we would expand a to: a = [[1, 1, 2, 2], [3, 3, 4, 4]] currently we only support the case: scale_a.shape * [1, 128] == a.shape scale_b.shape * [128, 128] == b.shape NrTrr)triton_scaled_mmrV)r+rr)rrrWviewris_rocmKvllm.model_executor.layers.quantization.compressed_tensors.triton_scaled_mmrVrhrXr,r-cutlass_scaled_mm) rrrQrRrrS target_shapecutlass_compatible_brVrs rrZrZ2si<  & &)u}*D*D*D D <4::<<171:55$* :Q:Q:Q Q/QWSbS\.171:..L r172;A71:?a/HAGAJOq4H!!J)=J      q!Wgy$GGk171:qwqz2)AHUUU &&sAq'7DIII 38\ ""razp_adjazpc |jddzdkr|jddzdksJ|tjus|tjusJ|0||jdkr |j|ksJg|jdd|jdR}|d|jd}|%||jdksJtj|jd|jdf||j} tj j | |||||||| j|S)z :param azp_adj: In the per-tensor case, this should include the azp. Always per-channel. :param azp: Only set in the per-token case. Per-token if set. rrUrTNrrV) rr+rr)rrWrWrhrXr,r-cutlass_scaled_mm_azp) rrrQrRrr]r^rSr[rs rr`r`es@ 71:?a  AGAJOq$8$8$8 8  & &)u}*D*D*D D <4::<<171:55$* :Q:Q:Q Q/QWSbS\.171:..L r172;A ;#))++333 3 +qwqz171:.i Q Q QC IL&&sAq'7GSRVWWW 38\ ""rcJtjj|Sr )r+r,r-"cutlass_sparse_scaled_mm_supportedrEs rrbrbs 9< : :;Q R RRrcn tjj|S#t$rYdSwxYwNF)r+r,r-cutlass_group_gemm_supportedAttributeErrorrEs rreresBy|889OPPP uus #& 44c|jtjtjtjtjfvsJ|sJd}|jdd|zzdksJtjj |S)a Compresses a sparse matrix for use with Cutlass sparse operations. This function takes a dense tensor and compresses it into two components: non-zero elements and metadata. The compressed representation is compatible with Cutlass sparse kernels. Args: a (torch.Tensor): The input tensor to be compressed. Must have one of the following data types: - `torch.int8` - `torch.float8_e4m3fn` - `torch.bfloat16` - `torch.float16` Returns: tuple[torch.Tensor, torch.Tensor]: A tuple containing: - `a_nzs` (torch.Tensor): A tensor containing non-zero elements of `a`. - `a_meta` (torch.Tensor): A tensor containing metadata for the sparse representation. Raises: ValueError: If the compression operation fails. Notes: - The `a_meta` tensor has a data type of `torch.uint8`. - Each metadata element encodes the sparsity of 4 non-zero elements (i.e., `elemsPerMetaElem = 4`). - The shape of `a_nzs` is `(m, k // 2)`, where `m` and `k` are the dimensions of the input tensor. - The shape of `a_meta` is `(m, k // 2 // elemsPerMetaElem)`. rTrUr) rWr+int8 float8_e4m3fnrr)rrr,r-cutlass_sparse_compress)relemsPerMetaElems rrkrks> 7uz5#6 V V V V V ??    71:-- .! 3 3 3 3 9< / / 2 22rbt_nzsbt_metac |jddzdkr|jddzdksJ|tjus|tjusJ|)|jd|jdkr |j|ksJ|jd}|jd}tj||f||j} tjj | ||||||| S)aG Performs a scaled sparse matrix multiplication using Cutlass. Steps: 1. Create a dense matrix `a` of shape (m, k) on the CUDA device: `a = torch.randn((m, k), device='cuda')`. 2. Create a dense matrix `b` of shape (k, n) on the CUDA device: `b = torch.randn((k, n), device='cuda')`. 3. Prune matrix `b` to 2:4 sparsity along the specified dimension: `b = prune_to_2_4(b, dim=0)`. 4. Compress the transposed sparse matrix `b.t()`: `bt_nzs, bt_meta = cutlass_sparse_compress(b.t())`. 5. Perform sparse matrix multiplication using the compressed matrix, applying scaling factors for `a` and `b`, and the output data type: `out = cutlass_scaled_sparse_mm(a, bt_nzs, bt_meta, scale_a, scale_b, out_dtype)`. Returns: - The result of the scaled sparse matrix multiplication. rrUrTNrV) rr+rr)rWrhrXr,r-cutlass_scaled_sparse_mm) rrmrnrQrRrrSrrrs rrprps@ <?R 1 $ $a2)=)B)B)B B  & &)u}*D*D*D D <4:a=FLO;; i@W@W@W W  A QA +q!fIah ? ? ?C IL)) Q'4 Jrexpert_offsetsproblem_sizes1problem_sizes2input_permutationoutput_permutation num_expertskblockscale_offsetsc \tjj|||||||||| S)a1 Prepare data necessary to perform CUTLASS grouped matrix multiplications used in CUTLASS-based fused MoE. The function takes in topk_ids (token-expert mapping) and uses it to compute: - expert_offsets: Indices that mark at which token index each expert begins its computation after the input is sorted with input_permutation. The number of tokens computed with expert E is expert_offsets[E + 1] - expert_offsets[E] - problem_sizes1, problem_sizes2: MxNxK sizes of each expert's multiplication in two grouped MMs used in the fused MoE operation. - input_permutation: Permutation that must be used to shuffle the input before executing the MMs. - output_permutation: Permutation that must be used to shuffle the output after executing the MMs. - blockscale_offsets: Optional argument passed for fp4 moe. Indices that mark at which block scale index each expert begins its computation. The number of block scale rows computed with expert E is blockscale_offsets[E + 1] - blockscale_offsets[E] )r+r,r-get_cutlass_moe_mm_data) r>rqrrrsrtrurvrrwrxs rrzrzs?F 9< / /     rexpert_first_token_offsetswap_abcTtjj||||||S)zICompute per-expert (M, N, K) problem sizes from expert_first_token_offset)r+r,r-4get_cutlass_moe_mm_problem_sizes_from_expert_offsets)r{rrrsrrwr|s rr~r~s2 9< L L!     r input_tensor dst2src_mapc|jd}tj||jdf|j|j}tjj||||S)z Shuffle and expand the input tensor according to the dst2src_map and store the result in output_tensor. This is used in MoE to permute the input tensor before performing grouped matrix multiplications. rrTr)rr+rhrXrWr,_moe_C shuffle_rows)rrnum_tokens_permuted output_tensors rrr*si &+A.K l034" M  I!!, ]KKK rexpert_num_tokensnum_local_expertspadded_mc Xtjj||||||||S)a} Prepare data necessary to perform CUTLASS grouped matrix multiplications used in CUTLASS-based fused MoE. The function takes in expert_num_tokens (token count per expert) and non_zero_expert_idxs (consecutive indices of experts with non-zero token counts) and uses them to compute: - expert_offsets: Indices that mark at which token index each expert begins its computation. - problem_sizes1, problem_sizes2: MxNxK sizes of each expert's multiplication in two grouped MMs used in the fused MoE operation. )r+r,r-get_cutlass_pplx_moe_mm_data)rqrrrsrrrrrws rrr9s8. 9< 4 4     r out_tensors a_tensors b_tensors problem_sizes a_strides b_strides c_strides per_act_token per_out_chc `tjj|||||||||| | | S)aY A single grouped matrix multiplication used in CUTLASS-based fused MoE. The function executes fp8-quantized OUT = AB matrix multiplication. - expert_offsets: Indices that mark at which token index each expert begins its computation. The number of tokens computed with expert E is expert_offsets[E + 1] - expert_offsets[E] - problem_sizes: MxNxK sizes of each expert's multiplication in two grouped MMs used in the fused MoE operation. - a/b/c_strides: The data strides passed to grouped matrix multiplication. )r+r,r-cutlass_moe_mm) rrrrrrqrrrrrrs rrr\sD2 9< & &   ralphas sf_offsetsc Ztjj||||||||| S)av An FP4 Blockscaled Group Gemm that takes in a_tensors, b_tensors and runs the gemms for each combination based on the specified problem sizes. This is used as the MoE gemm during NVFP4 Quantized FusedMoE forward. - a/b_tensors: the NVFP4 a_ptrs and b_ptrs tensors which are quantized input and expert weights. - a_/b_scales: The blockscales in FP8-E4M3 precision - expert_offsets/sf_offsets: Indices that mark at which token index each expert begins its computation. The number of tokens computed with expert E is expert_offsets[E + 1] - expert_offsets[E] And the sf_size per expert is sf_offset[E+1] - sf_offset[E] - problem_sizes: MxNxK sizes of each expert's multiplication in two grouped MMs used in the fused MoE operation. )r+r,r-cutlass_fp4_group_mm) rrrrrrrrqrs rcutlass_fp4_moe_mmrs;6 9< , ,   rnum_bits is_a_8bitcTtjj||||||Sr )r+r,r-gptq_marlin_repack)rrrrrrs rrrs- 9< * *D&&(I  rrz_C::gptq_marlin_repackchd|z}d}tj||z||z|zf|j|jSN rUrVr+rhrWrX)rrrrrr pack_factormarlin_tile_sizes r_gptq_marlin_repack_fakersQHn { ' '2B)Bk)Q R"$    rcRtjj|||||Sr )r+r,r-awq_marlin_repack)rrrrrs rrrs+ 9< ) )FFHi  rrz_C::awq_marlin_repackchd|z}d}tj||z||z|zf|j|jSrr)rrrrrrrs r_awq_marlin_repack_fakersQHn { ' '2B)Bk)Q R"$    rc $|jd}|dzdksJtj||dz||dzzf|j|j}t |D]:}tjj||||||||||<;|SNrrUrUr) rr+rhrXrWranger,r-r rrrrrrrvrBes rgptq_marlin_moe_repackrs"1%K B;!     [ flFh!m$<= F ;    IL33 qM47FFHi  q  Mrc|jd}|dzdksJtj||dz||dzzf|j|j}t |D]3}tjj||||||||<4|Sr) rr+rhrXrWrr,r-rrs rawq_marlin_moe_repackrs"1%K B;!     [ flFh!m$<= F ;    IL22 qM668Y  q  Mrqzeros_or_noneinplacecNtjj|||Sr )r+r,r-marlin_int4_fp8_preprocess)rrrs rrrs 9< 2 27NG T TTrcxtjj|||||||||| | | j| | |||||Sr )r+r,r- marlin_gemmr)rrrrrrrrrrrrrrrrrrrs rrr#s[* 9< # #     '  rgroup_zeros_typechannel_scales_typetoken_scales_typec `tjj||j|||||Sr )r+r,r-machete_supported_schedulesr)rrrrrrrs rrrPs7 9< 3 3   rc ftjj|||j||||||| Sr )r+r,r- machete_mmr) rrrrrrrrrrs rrrds@ 9< " "      rcZtjj|||j|Sr )r+r,r-machete_prepack_Brr s rrrs, 9< ) )FFI'8  rc Xtjj||||||||Sr )r+r,r-cutlass_w4a8_mm)rrrrrrrr s rrrs8 9< ' '     rcJtjj|Sr )r+r,r-cutlass_pack_scale_fp8rs rrrs 9< . .v 6 66rcJtjj|Sr )r+r,r- cutlass_encode_and_reorder_int4brs rrrs 9< 8 8 ; ;;rgroup_scale_stridescdtjj|||||||||| | | | | S)a. Executes the CUTLASS-based fused-MoE grouped matrix multiplication for the W4A8 quantization scheme. Uses group-wise quantization (INT4 -> FP8) and both per-channel + per-token scaling in the epilogue. Args: out_tensors: Output buffer for all experts (updated in-place). a_tensors: FP8 (E4M3FN) activations for all experts. b_tensors: INT4-packed weight matrix for all experts, packed to INT32 a_scales: Per-token FP8 activation scales, applied in the epilogue. b_scales: Per-channel FP8 weight scales for each expert, applied in the epilogue. b_group_scales: FP8 scale values for group-wise INT4 weight blocks. b_group_size: Number of elements grouped under each entry of b_group_scales. expert_offsets: Cumulative token offsets problem_sizes: Per-expert (M, N, K) GEMM sizes used by the grouped GEMM launcher. a/b/c/group_scale_strides: Strides describing the memory layout of the input tensors. maybe_schedule: Optional override to choose a specific kernel or epilogue schedule. Returns: out_tensors updated in-place with the dequantized INT4xFP8 grouped GEMM result. )r+r,r-cutlass_w4a8_moe_mm)rrrrrrrrqrrrrrr s rrrsK` 9< + +  rcJtjj|Sr )r+r,r-(cutlass_encode_and_reorder_int4b_grouped)rs rrrs 9< @ @ K KKr permute_colsz_C::permute_colsc*tj|Sr r+rrrs r_permute_cols_fakers"""rcLtjj||Sr )r+r,r-rrs rrrs 9< $ $Q - --rnoneinput_global_scaleis_sf_swizzled_layoutbackendctjrJ|jdksJd|jd|jdkrdnd}|||jd}|j\}}d}|j}||zdksJd|d|jtjtj fvsJd|jdd |vr|d krd nd } | rt||\} } ntj ||d zf|tj } |rGd} | |d} ||z}| |d}tj | |dzf|tj } n&tj ||dzf|tj } tjj| || ||| tj} | | fS)ay Quantize input tensor to FP4 and return quantized tensor and scale. This function quantizes the last dimension of the given tensor `input`. For every 16 consecutive elements, a single dynamically computed scaling factor is shared. This scaling factor is quantized using the `input_global_scale` and is stored in a swizzled layout (see https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-scale-factor-b-layout-4x). Args: input: The input tensor to be quantized to FP4 input_global_scale: A scalar scaling factor for the entire tensor. use_8x4_sf_layout: Whether to use the 8x4 or 128x4 layout for the scaling Returns: tuple[torch.Tensor, torch.Tensor]: The output tensor in FP4 but every two values are packed into a uint8 and float8_e4m3 scaling factors in the sizzled layout. rTz%input.ndim needs to be >= 1, but got .rrUrz+last dim has to be multiple of 16, but got z-input.dtype needs to be fp16 or bf16 but got trtllmrTFrUrc||zdz |z|zS)NrTr )xys rrz"scaled_fp4_quant..3sQUQY1$4q$8rrh)rrXrKreshaperrXrWr+r)rrrhuint8int32r,r-scaled_fp4_quantrWrj)rqrrr other_dimsrrrrXuse_8x4_sf_layoutrB output_scaleround_up rounded_mscale_n rounded_ns rrrs2 ')))) ) :???QEJQQQ???jAoo2J MM*ek"o 6 6E ;DAqJ \F z>Q    Ra R R R    ;5=%.9 9 9 9F FFF : 9 9!)G 3 3RU C %   aa[u{KKK  W 98H C((I:oG !,,I ;IN+F%+LL!;17|F%+VVVL %% E<);=R    $$U%899L < rtopkc"tjrJ|jdksJd|jdtj}|j\}}|||zksJd|d|d|dz}|dzd z} t j||dz|jt j } t j||z| t j |j } t j j | | ||||| t j} | | fS) a Quantize input tensor to NVFP4 and return quantized tensor and scale, for packed MoE Inputs. Args: input_tensor: The input tensor to be quantized to NVFP4 input_global_scale: A scalar scaling factor for the entire tensor. expert_offsets: The expert offsets tensor blockscale_offsets: The blockscale offsets tensor Outputs: output: The quantized tensor in NVFP4 output_scales: The blockscale tensor in FP8-E4M3 rU%input.ndim needs to be == 2, but got r2m_numtopk must be less than MAX_TOKENS_PER_EXPERT(,) for cutlass_moe_fp4, observed m_numtopk = ;. Use VLLM_MAX_TOKENS_PER_EXPERT_FP4_MOE to set this value.rUrhrrV)rrXrKr"VLLM_MAX_TOKENS_PER_EXPERT_FP4_MOErr+rhrXrrr,r-scaled_fp4_experts_quantrWrj) rrrqrxrMAX_TOKENS_PER_EXPERT m_numtopkrwscales_kpadded_krB output_scaless rrrEsk& ')))) )   ! ! !D 0ADDD " ! !!C%LIq -4 4 4 4 B  B B6? B B B 5 4 4 BwHE"q(H[16,"5U[FK$k" M  IL)) "&&u':;;M =  rcRtjrJ|jdksJd|jdtj}|j\}}|dzdks Jd|dz}|||zksJd|d|d|d z} | d zd z} t j||dz|jt j } t j||z| t j |j } t j j | | ||||| t j} | | fS)a! Fused SiLU+Mul+NVFP4 quantization for MoE intermediate activations. Args: input_tensor: The input tensor with gate || up layout [m_topk, k*2] input_global_scale: A per-expert scaling factor [n_experts] expert_offsets: The expert offsets tensor [n_experts+1] blockscale_offsets: The blockscale offsets tensor [n_experts+1] topk: Number of top-k experts selected Outputs: output: The quantized tensor in NVFP4 [m_topk, k/2] output_scales: The blockscale tensor in FP8-E4M3 rUrrrz,input width must be even (gate || up layout)rrrrUrrhrrV)rrXrKrrrr+rhrXrrr,r-%silu_and_mul_scaled_fp4_experts_quantrWrj) rrrqrxrrr k_times_2rwrrrBrs rrrs( ')))) )   ! ! !D 0ADDD " ! !!C'-Iy q=A   M   QA -4 4 4 4 B  B B6? B B B 5 4 4 BwHE"q(H[16,"5U[FK$k" M  IL66 "&&u':;;M =  rnum_token_paddinguse_per_token_if_dynamic group_shapec|jdksJ|j}tj}|r#t ||jd|df}|t j||j|}n| Jd|j|ksJ||rVt j|ddf|jt j }t j j ||||ntt jd|jt j }t j j |||n't j j ||||||fS)a_ Quantize input tensor to FP8 and return quantized tensor and scale. This function supports both static and dynamic quantization: If you provide the scale, it will use static scaling and if you omit it, the scale will be determined dynamically. The function also allows optional padding of the output tensors for downstream kernels that will benefit from padding. Args: input: The input tensor to be quantized to FP8 (must be 2D: [M, N]) scale: Optional scaling factor for the FP8 quantization. Supports: - 0D or [1]: per-tensor scaling - 1D: requires explicit group_shape to disambiguate per-channel vs per-token (use (-1, 1) for per-channel, (1, -1) for per-token) - 2D [M/group_m, N/group_n]: group scaling (e.g. [M, N/128] for DeepSeek-style (1,128) groups, or [M/128, N/128] for (128,128)) scale_ub: Optional upper bound for scaling factor in dynamic per token case num_token_padding: If specified, pad the first dimension of the output to at least this value. use_per_token_if_dynamic: Whether to do per_tensor or per_token in the dynamic quantization case. group_shape: Optional tuple (group_m, group_n) specifying the group shape for static quantization. Use -1 for "full extent" (e.g., (-1, -1) for per-tensor, (-1, 1) for per-channel, etc.) Required for 1D scales; optional for 2D scales. Returns: tuple[torch.Tensor, torch.Tensor]: The output tensor in FP8 and scaling factor. rUrrTNrz)padding not supported if output passed in)rKrr fp8_dtypemaxr+rhrXrWrr,r-"dynamic_per_token_scaled_fp8_quantdynamic_scaled_fp8_quantstatic_scaled_fp8_quant) rqrrrrrBrrrs rscaled_fp8_quantrsUT :????*/+E-799IC& A77qB ~U5||duks Jdtjj|||||||fStj||jdzdf|j tj }|rdntj|tj }tjj || |||||fS) Quantize the input tensor to int8 and return the quantized tensor and scale, and maybe azp. Args: input: The input tensor to be quantized to int8. scale: Optional scaling factor for the int8 quantization. When not provided, we invoke dynamic-per-token quantization. azp: Optional zero-point for the int8 quantization. Must be provided for asymmetric quantization if `scale` is provided. symmetric: Whether to use symmetric quantization (scale only, azp ignored). Returns: tuple[torch.Tensor, torch.Tensor, torch.Tensor | None] : Output int8 tensor, scales, and optionally azp. rN6azp must only be provided for asymmetric quantization.rrTr)r+rrir,r-static_scaled_int8_quantrhrrrXrrdynamic_scaled_int8_quant contiguous)rqrr^rrB input_scales input_azps rscaled_int8_quantr_s ( e5: 6 6 6F SD[))) D*))  --feUCHHHuc!!; %+b/ )1-el%-L"Xu'7 EK'X'X'XI IL**  ""L) < **rcRtjj|||||Sr )r+r,r-r$r*s rr$r$s$ 9< ' ':q!U C CCrcPtjj||||Sr )r+r,r-ggml_mul_mat_vec_a8r0s rrrs" 9< + +Aq*c B BBrcPtjj||||Sr )r+r,r-ggml_mul_mat_a8r0s rrrs" 9< ' '1j# > >>rc Ztjj||||||||| Sr )r+r,r- ggml_moe_a8r;s rrrs; 9< # #       rc Vtjj|||||||Sr )r+r,r-r=r@s rr=r=s) 9< ' '1hz3PV W WWrcJtjj|Sr )r+r,r-ggml_moe_get_block_size)r&s rr!r!s 9< / / ; ;;rudeltaABCD_z_ delta_bias_delta_softplus cache_indiceshas_initial_state ssm_states pad_slot_idblock_idx_first_scheduled_tokenblock_idx_last_scheduled_tokeninitial_state_idxcptjj|||||||||| | | | | ||||dSr )r+r,r-selective_scan_fwd)r#r$r%r&r'r(r)r*r+r6r,r-r.r/rr0r1r2s rr4r4s`( IL##       '&%rcZ|}|}||fSr )r)rrs r*rocm_enforce_contiguous_skinny_gemm_inputsr6s' A A a4Krrows_per_blockctt||\}}tjj|||Sr )r6r+r,r:LLMM1)rrr7s rr9r9s2 5a ; ;DAq 9  " "1a 8 88rcu_countcvt||\}}tjj||||Sr )r6r+r,r:wvSplitKrrr:rSs rr<r<s6 6a ; ;DAq 9  % %aD( ; ;;rcvt||\}}tjj||||Sr )r6r+r,r: wvSplitKrcr=s rr?r? s6 6a ; ;DAq 9  ' '1dH = ==rc t||\}}tj|jd|jdf||j}tjj||||||||Sr/)r6r+rhrrXr,r: wvSplitKQ)rrrrQrRr:rSrs rrArAsj 6a ; ;DAq +qwqz171:.i Q Q QC I1dC'8LLL JrcPtjj||dSr )r+r,rmoe_sum)rqrBs rrCrC#s# IUF+++++r experts_idsnum_tokens_post_pad expert_mapc Ztjj|||||||dSr )r+r,rmoe_align_block_size)r>rvrr5rDrErFs rrHrH's@ I))rmax_tokens_per_batch sorted_idscXtjj||||||dSr )r+r,rbatched_moe_align_block_size)rIrrrJr6rEs rrLrL;s= I11 rtoken_lora_mapping max_lorasmax_num_tokens_paddedmax_num_m_blocksadapter_enabledlora_idsc ftjj|||||||||| | | | dSr )r+r,rmoe_lora_align_block_size) r>rMrvrrNrOrPr5rDrErQrRrFs rrTrTMsR I..r topk_weights BLOCK_SIZE_M BLOCK_SIZE_N BLOCK_SIZE_Kctjstdtjj|||||||||| | | | | dS)NzGThe optimized moe_wna16_gemm kernel is only available on CUDA platforms)rrNotImplementedErrorr+r,rmoe_wna16_gemm)rqrBrrrrUr5rDrEr8rVrWrXrs rr[r[ms}  # % % ! U    I##   rtoken_expert_indices gating_output renormalizee_score_correction_biascXtjj||||||dSr )r+r,r topk_softmaxrUr>r\r]r^r_s rrara= I!! rcXtjj||||||dSr )r+r,r topk_sigmoidrbs rrerercrscoresnum_expert_group topk_grouprouted_scaling_factor scoring_funcc tjstdtjj||||||||S)a5 Perform grouped top-k routing for mixture of experts. Args: scores: Raw inputs (logits if scoring_func=1, scores if scoring_func=0) num_expert_group: Number of expert groups topk_group: Number of groups to select topk: Number of experts to select per token renormalize: Whether to renormalize the output weights routed_scaling_factor: Scaling factor for routing weights bias: Bias tensor (e_score_correction_bias). Always fused in kernel. scoring_func: 0=none (no activation), 1=sigmoid zAThe fused grouped_topk kernel is only available on CUDA platforms)rrrZr+r,r grouped_topk)rfrgrhrr^rirSrjs rrlrlsa.  # % % ! O    9  ( (     rrnum_tokens_past_paddedmoe_block_sizemul_topk_weightsthread_kthread_n blocks_per_smctjj|||||||||| | | | | |||||j||||||||||Sr )r+r,rmoe_wna16_marlin_gemmr)rqrBrrrrrrrrrr5r6rmrUrnr8rorrrrrrrrrprqrrs rrtrtsz> 9  1 1     ;  rrmarlin_gemm_moez_moe_C::marlin_gemm_moe b_q_weights b_zero_pointsreplicate_input apply_weightscJtj| || f|j|jSr(r)rrvrJrUr>rrwrrrrrrrrrvrrnrxrys rmarlin_gemm_moe_faker{ s&.{FD&1RRRRrz_moe_C::moe_wna16_marlin_gemmcNtj||z|f|j|jSr(r)rqrBrrrrrrrrrr5r6rmrUrnr8rorrrrrrrrs rmoe_wna16_marlin_gemm_faker}9 s1:{ e^V $EK     rvalue slot_mappingc \tjj||||||||dSr )r+r, _C_cache_opsreshape_and_cacherkr~rrrr!r"r#s rrr[ sC I,,       rc \tjj||||||||dSr )r+r,rreshape_and_cache_flashrs rrrq sC I22       rkv_ck_pecXtjj||||||dSr )r+r,rconcat_and_cache_mla)rrr=rr!rs rrr s8 I// dHlNErq_pekv_cache_scalec `tjj|||||||||| dSr )r+r,rconcat_and_cache_mla_rope_fused) rjrrrrmrnrr=r!rs rrr sI I::        rsrcdstblock_size_in_bytes block_mappingcTtjj||||dS)as Copy specific blocks from one tensor to another. This method assumes each of the two input tensors is composed of consecutive contiguous blocks, of size block_size_in_bytes. i.e. the memory layout for each tensor is: [block0] [block1] ... [block N] block_mapping determines the subset of blocks to copy of the source tensor, and their matching destination block number on the destination tensor. block_mapping is expected to be a tensor of shape (num_blocks_to_copy, 2) where each block_mapping[i] represents a single copy operation, copying block #block_mapping[i][0] from the source tensor to block #block_mapping[i][1] on the destination tensor. block_mapping should have dtype int64. The source and the destination tensors can be either on cpu or gpu, but not both on cpu. the block mapping tensor must on cpu. N)r+r,r swap_blocks)rrrrs rrr s*4 I&&sC1DmTTTTTrrkv_dtypecTtjj||||dSr )r+r,r convert_fp8)rBrqrrs rrr s) I&&vueXFFFFFr src_cache block_table cu_seq_lens token_to_seq num_tokens seq_startsc ^tjj||||||||| dSr )r+r,rgather_and_maybe_dequant_cache) rrrrrrr!rrs rrr sF I99       rr\cXtjj||||||dSr )r+r,rcp_gather_cache)rrrrr\rs rrr s8 I**3 [*jrworkspace_startscXtjj||||||dS)aGather and upconvert FP8 KV cache to BF16 workspace. Args: src_cache: FP8 KV cache [num_blocks, block_size, 656] dst: BF16 output workspace [total_tokens, 576] block_table: Block indices [num_reqs, max_blocks] seq_lens: Sequence lengths [num_reqs] workspace_starts: Workspace start offsets [num_reqs] batch_size: Number of requests N)r+r,r$cp_gather_and_upconvert_fp8_kv_cache)rrrrrr\s rrr s9$ I??3 X/?rquant_block_sizecVtjj|||||dSr )r+r,rindexer_k_quant_and_cache)rwr=rrr!s rrr s7 I44 8\#3^rdst_k dst_scalecVtjj|||||dSr )r+r,rcp_gather_indexer_k_quant_cache)r=rrrrs rrr s6 I::%Kr attributerXcLtjj||Sr )r+r, _C_cuda_utilsget_device_attribute)rrXs rrr% s 9 " 7 7 6 J JJrcJtjj|Sr )r+r,r0get_max_shared_memory_per_block_device_attribute)rXs rrr) s" 9 " S S  r ipc_tensors rank_datarankfully_connectedcPtjj||||Sr )r+r, _C_custom_arinit_custom_ar)rrrrs rrr1 s* 9 ! 0 0Yo  rfainp reg_bufferreg_buffer_sz_bytescVtjj|||||dSr )r+r,r all_reduce)rrrrrs rrr< s, I%%b#sJ@STTTTTrcNtjj|dSr )r+r,rdisposers rrrF s! I""2&&&&&rcHtjjSr )r+r,r meta_sizer rrrrJ s 9 ! + + - --rcLtjj||Sr )r+r,rregister_buffer)rrs rrrN s 9 ! 1 1"k B BBrcJtjj|Sr )r+r,rget_graph_buffer_ipc_metars rrrR s 9 ! ; ;B ? ??rhandlesoffsetscRtjj|||dSr )r+r,rregister_graph_buffers)rrrs rrrV s' I11"gwGGGGGrrYcJtjj|Sr )r+r,r!allocate_shared_buffer_and_handle)rYs rrr\ s 9 ! C CD I IIr mem_handlecJtjj|Sr )r+r,ropen_mem_handle)rs rrr` s 9 ! 1 1* = ==rptrcNtjj|dSr )r+r,rfree_shared_buffer)rs rrrd s! I--c22222r world_size qr_max_sizecNtjj|||Sr )r+r,rinit_custom_qr)rrrs rrri s 9 ! 0 0z; O OOrcNtjj|dSr )r+r,r qr_destroyrs rrrm s! I%%b)))))r quant_level cast_bf2halfcVtjj|||||dSr )r+r,r qr_all_reduce)rrrrrs rrrq s+ I((S#{LQQQQQrcJtjj|Sr )r+r,r qr_get_handlers rrr{ s 9 ! / / 3 33rcLtjj||Sr )r+r,rqr_open_handles)rrs rrr s 9 ! 1 1"g > >>rcHtjjSr )r+r,rrr rrrr s 9 ! - - / //r cache_seqlensnum_heads_per_head_kcNtjj|||S)ac Arguments: cache_seqlens: (batch_size), dtype torch.int32. num_heads_per_head_k: Equals to seq_len_q * num_heads_q // num_heads_k. num_heads_k: num_heads_k. Return: tile_scheduler_metadata: (num_sm_parts, TileSchedulerMetaDataSize), dtype torch.int32. num_splits: (batch_size + 1), dtype torch.int32. )r+r,r-get_flash_mla_metadata)rrr{s rrr s( 9< . .+[  rqk_cache head_dim_vtile_scheduler_metadata num_splits softmax_scalec ||jddz}tjj||d||||||| \} } | | fS)a5 Arguments: q: (batch_size, seq_len_q, num_heads_q, head_dim). k_cache: (num_blocks, page_block_size, num_heads_k, head_dim). block_table: (batch_size, max_num_blocks_per_seq), torch.int32. cache_seqlens: (batch_size), torch.int32. head_dim_v: Head_dim of v. tile_scheduler_metadata: (num_sm_parts, TileSchedulerMetaDataSize), torch.int32, return by get_mla_metadata. num_splits: (batch_size + 1), torch.int32, return by get_mla_metadata. softmax_scale: float. The scaling of QK^T before applying softmax. Default to 1 / sqrt(head_dim). causal: bool. Whether to apply causal attention mask. Return: out: (batch_size, seq_len_q, num_heads_q, head_dim_v). softmax_lse: (batch_size, num_heads_q, seq_len_q), torch.float32. Nrg)rr+r,r-flash_mla_fwd_kvcache) rrrrrrrrrQr softmax_lses rflash_mla_with_kvcacher sc6 - y|99    C  rlseq_nopekv_c_and_k_pe_cache page_table num_kv_splitsc `tjj|||||||||| |Sr )r+r,r-sm100_cutlass_mla_decode) rrrrrrrrrrs rrr sE IL))        Jr num_batchescPtjj||||Sr )r+r,r-$sm100_cutlass_mla_get_workspace_size)rrrrs rrr s) 9< < <[(M  rweight_packed_linearz_C::weight_packed_linearmat1mat2is_vnnictj|d|df|j|jSr/r)rrrSrs rweight_packed_linear_faker sA{ YYq\\499Q<< ( 4;    rfused_experts_cpuz_C::fused_experts_cpu hidden_statesw1w2 use_int8_w8a8 use_fp8_w8a16w1_scalew2_scalea1_scalea2_scalec*tj|Sr r)rrrrUr>rrrrrrrr rs rfused_experts_cpu_faker  s" ...rint8_scaled_mm_with_quantz_C::int8_scaled_mm_with_quantscales2c|d}|d}tj||f|S)Nrr)rYr+rh)rrr rSrrMrs rint8_scaled_mm_with_quant_faker s; IIaLL IIaLL{Aq63333rceZdZddZdZdS)CPUDNNLGEMMHandlerr)Nc0d|_d|_d|_dS)Nr)handlerrrwselfs r__init__zCPUDNNLGEMMHandler.__init__ s#' rcj|j+tjj|jdSdSr )rr+r,r-release_dnnl_matmul_handlerrs r__del__zCPUDNNLGEMMHandler.__del__$ s2 < # IL 4 4T\ B B B B B $ #r)r)N)__name__ __module__ __qualname__rrr rrrr s< CCCCCrrcreate_onednn_mm_handlercHtjjSr )r+r,r-is_onednn_acl_supportedr rrr r , s 9< / / 1 11rrprimitive_cache_sizect}|\|_|_tjj|||_|Sr ) rrYrwrr+r,r-rr)rrr!rs rcreate_onednn_mmr#0 sM!""G!;;==GIwyil;;$GO Nr dnnl_handlerrctjg|jdd|jR|j}tjj||d|j ||j |S)Nrrr) r+rhrrrWr,r- onednn_mmrrwr)r$rrSrBs rr&r&< ss [9171R4=9,.99 I I IF IL "ln--t\5I Mr weight_scales output_type dynamic_quantuse_azpct}|\|_|_tjj|||||||_|Sr ) rrYrwrr+r,r-create_onednn_scaled_mm_handlerr)rrr'r(r)r*r!rs rcreate_onednn_scaled_mmr-I sU!""G!;;==GIwyilBB {M7DXGO Nrc:tj|tj}||jdz}|||jdf}|>||duks Jdtjj|||||||fStj |df|j tj }|rdntj|tj }tjj |||||||fS)rrrNrrTr)r+rrirrrWr,r-rrhrXrrr)rqrr^rrB token_numrrs ronednn_scaled_int8_quantr0Y s ( e5: 6 6 6F R0I JJ 5;r?3 4 4E SD[))) D*))  --feUCHHHuc!!; 1~el%-XXXL!Xu'7 EK'X'X'XI IL**65, RRR < **r input_scaleinput_zp input_zp_adjc dtjj|||||||j|Sr )r+r,r-onednn_scaled_mmr)r$rrBr1r2r3rSs rr5r5 s9 IL!!;,l>R Mrnum_reqsr]sliding_window_sizeisaenable_kv_splitc btjj|||||||||| | } | Sr )r+r,r-get_scheduler_metadata) r6r]rr}rrWr6rQr7r8r9sheduler_metadatas rcpu_attn_get_scheduler_metadatar= sG ;;     rcXtjj||||||dSr )r+r,r-cpu_attn_reshape_and_cache)rkr~rrrr8s rr?r? s< IL++   rsliding_windowsoftcapscheduler_metadatas_auxctjj|||||||||| d| d| | | | dS)NrrT)r+r,r-cpu_attention_with_kv_cache)rrrrBr6rrrQr r@rrArBrCs rrErE sa  IL,,  qq rrisa_hintc tj|d|df|j}tjj||||||||| |S)NrrTr)r+rhrYrWr,r-cpu_gemm_wna16) rqrrrZrrSrrFrBs rrHrH sq[%**Q--Q8 L L LF IL        Mrcztj|}tjj||||Sr )r+rr,r-prepack_moe_weight)rrr8rBs rcpu_prepack_moe_weightrK s5 f % %F IL##FFC888 Mrw13w13_biasw2_biasactc tj|} tjj| ||||||||| | Sr )r+rr,r- cpu_fused_moe) rqrLrrMrNrUr>rOr8rBs rrQrQ sU e $ $F IL         Mrmatmul_mxf4_bf16_tnz_qutlass_C::matmul_mxf4_bf16_tna_sfb_sfcj|jg|jdd|jdRdtjiSNrrrW new_emptyrr+rrrrSrTrIs r_fake_matmul_mxf4_bf16_tnrZ ;q{KAGCRCLK!'!*KKKENKKKrcRtjj|||||Sr )r+r, _qutlass_CrRrYs rrRrR+ s% 9  3 3Aq$e L LLrmatmul_ada_mxf4_bf16_tnz#_qutlass_C::matmul_ada_mxf4_bf16_tncj|jg|jdd|jdRdtjiSrVrWrYs r_fake_matmul_ada_mxf4_bf16_tnr`7 r[rcRtjj|||||Sr )r+r,r]r^rYs rr^r^B s% 9  7 71dD% P PPrfusedQuantizeMxQuestz _qutlass_C::fusedQuantizeMxQuestxh_e2m1xh_e8m0c ||fSr r rrrcrds r_fake_fused_quantize_mx_questrgN rfusedQuantizeMxAbsMaxz!_qutlass_C::fusedQuantizeMxAbsMaxc ||fSr r rfs r_fake_fused_quantize_mx_absmaxrkW rhrquest)methodrm)rlabs_maxc|dkrtd|ddzdkr&td|dd|j|jkrtdt jg|jdd|ddzRtj|jd }||dz|ddz}}t|d }t|d }|d z}|d z} t j|| tj |j } ttj d std |dkr'tj j|||| S|dkr'tj j|||| Std|d)Nrz#`a` must have at least 1 dimension.rrz-last dim of `a` must be divisible by 32, got rz'`a` and `b` must be on the same device.rUrVrrhr]zvThe `_qutlass_C` extension is not loaded. Make sure your custom op library is imported before calling fusedQuantizeMx.rlrnzinvalid method z, must be 'quest' or 'abs_max')r ValueErrorrYrXr+rhrrrr float8_e8m0fnuhasattrr, RuntimeErrorr]rbri) rrrmrcrowscols n_row_blocks n_col_blocks padded_rows padded_colsrds rfusedQuantizeMxrz^ s uuww!||>???vvbzzB!VPRVVVWWWx18BCCCk "vvbzzQ.3k!(GaffRjj(!&&***:$Dc??La==L$K"Kk[( inverse) transform. :param x: value to be transformed inplace :param inplace: modify value in place :return: value after transformation )r+r,r-hadacore_transformrrs rrr s 9< * *1g 6 66rrz_C::hadacore_transformc2|stj|n|Sr rrs r_hadacore_transform_faker s*18u"""q8r)rrrrrr )T)NN)NNF)TFFF)NNNNNNN)Frd)NNNN)Tr)NNNFNN)NNT)r"NNN)FN)r)rrr)rr4)r)typingrrr+ vllm.envsr vllm.loggerrvllm.platformsrvllm.scalar_typervllm.utils.flashinferrvllm.utils.math_utilsr rloggerimport_kernelsr torch.library ImportErrorrTensorintfloatstrr.r3VLLM_ROCM_FP8_MFMA_PAGE_ATTNr<rArIbooltupler[rirprurxrrrrrWrlistrrrrrrr,r-rrrSymIntrrrrrr rrrrr#r+r1r4r<rArDrLrNrPrZr`rbrerkrprzr~rrrrrrrrrrrrrrrrrrrrrrrrrrr rrr$rrrr=r!r4r6r9r<r?rArCrHrLrTr[rarerlrtrr{r}rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr rr_supports_onednnr r#r&r-r0r5r=r?rErHrKrQr]rZrRr`r^rgrkrzr~r{rrr rrrs6J*))))))) ######++++++'''''''&&&&& X  !!!AA/////// AAA@@@@@@@@A($%#$"$)*')) ) <)|) )  )  ),)l))),%))\)\)) "!)"!#)$ %)&$'')( )))))|$%#$"$)*-// / \/ /\ / < / | /// /,/l///,%// \!/"\#/$%/&"'/(!)/* +/,$'-/. /////J*."?JUUU)++ + \+ +\ + < + | +++ +,+l+\D(+++,%+ !+"\#+$\%+&<$&'+()+* +++++\    < l    ,  l      ('+   L <   <     t#      *5A5A|5A 5Al5A< 5A  5A  5A5A 5A 5<u|U\ AB5A5A5A5AF:A:A|:A :Al:A< :A "L :A:A:A:A:A :A 5<u|U\ AB:A:A:A:A| |  <      <        7 7#l749L7KP7 7777 F <F#(<F9>FPUF FFFF       ll<, 8  L    ,        L ,    L    ,       <%)$(  < L  lT!  lT!  5< %&4%)$( %$$ <$ L$$ $ S $ lT! $lT!$$ 5< %&$$$$RX \X LX <X X X X \XXXX" P < P \ P L P L P  P  \ P P P P | <<  \     \, 759<%% ]?##  <  L  |  |               $#  55<55C5D5555  |    L l  |          \      759<.//~J],-- M < ML M  M, M < M  M  M  M  M  M M M.- M]$%%"$ %!'KK <K <$ KLK t# K , K ,% KlT)K$K|d"KlT!K<KK K K K !K"#K$%K&'K( )KKK&%K4]'(( T T  T| T| T  T  T  T T T)( T]>"" |      |     #" ]#$$ (,.2-1#'04.2#CC <C\C C +$ C  t+ C|d*CDjC ,-C t+C*C CCC%$C"]*++SLS SS!;- S  SSS,+S]())(,%)EE <E\E E  E  , E E+$Ed E EEE*)E ]/00OELOU\OOO10O]9::JJ%,JJJ;:J]ABBJJ%,JJJCBJ 759<.//C],--C <C<C,C,% C < C L C,CLC"LCCC CCC.-C" 759<*++,X]()) %) II <II <I < I {T! I  III*)I],--N <N <NN\ N  NNN.-N]())I <I <II\ I  III*)I]$%% X < X < X , XL X !& X  X\ X| X  X  X X X&% X 759<*++ R]()) R < R < R, R R  R \ R  R  R R R*) RO3O4OOOO  |  | < <  <  {  \    O3O4OOOOUUUUUU!% 0#0# |0# |0#\0#\ 0# { 0# ,  0# \0#0#0#0#t $ $## |# |#\#\ # { # \ #  # , # \####:SsStSSSS&3u|&3elEL6P0Q&3&3&3&3`!%,, |, L,\,\ , \ , { , , , \,,,,r/3..l.L.L.L . | .  .. . . t+....b$|LL     & u| %,     L L L |            F&&|&|&l & l & L &<&|&|&|&&&&&&R%%|%|%l % l % L %<%L% %%%%^     ,          \     759<-.. ]+,,   L l             -, .             \     759<,-- ]*++    L                ,+  ,   ,        \8   ,        \0+/UU \UL4'UUUUU0 !')) |) |d) ) L4  ) l ) lT! ),%)\D ) <$ ) , )|)) ) ) ) !)"#)$%)&')( \)))))b,0.2,0#' K {T)kD(  t+  {T) kD  #Y2$(*.)-#,0*. |    kD  L4' <$&*lT)L4'Dj \6  K {T)   \ ($(!% | L    l LkD $J \.75<7EL7777< <<<<<$"&???|?|?l ? l ? L ??L?<?|?|?|??$J????DL|L 5< %&LLLL  759<((#]%&&#el#%,#5<###'&#.EL. .....#' D D <D  D  D  D  5< %& D D D D N;!,;! ;!L;! ;!  ;!  5< %& ;!;!;!;!|>!,>! >!L>! >!  >!  5< %& >!>!>!>!H"&$($(%*"&*.BB <B <$ BTzBlT! B # B L4  BsCx4'B 5< %&BBBBR'+ 4>4> \4> <4> t#4>  4>  5<u| 34 4>4>4>4>n ||llT!       \>"&# %+%+ <%+ <$ %+  %+ %+  5<u|d': :; %+%+%+%+RD |D!$D),D14D=B[4=OD \DDDD C |C |CC C  \ CCCC? |? |?? ?  \ ???? | |l  "L       \0 X | X | Xl X  X  X  X L X \ X X X X<<<<<<(;?:>-1%'' |' <' |' | '  | '  t '  t'$''\D('<$&'|d*' ''' &+\D%8!'"%*L4$7#'$|d*%''''\ |  5< %&9U\9el9C9EL9999 KO<< |< <03<;@<< \<<<<KO>> |> >03>;@<> \>>>>   |  | { \  \    ,  \     ,5<,,,,,'+ll      t# (|      >'+l      l\l t# @# <# L#|#l # lT! # ,% #l### #### # \####V37 ,l ,<    #\D0  .37 ,l ,<    #\D0  4$$ L$$$  $  $ ! $ ,$$$$$D;== <= L4 =|= L4  = l = lT! =,%=lT!= <$ = , =|=l= ="L=,= !=" #=$%=&'=( )=* +=, -=./=01=23=45=67=89=:;=< \=====@ 759h: GGEI,<>O$P$P: ],--S <S\SLSl S , S , S|S|SlS<SS S S SS !S"#S$%S&'S()S* +SSS.-S0]233 |  t# <  t#  ,  ,%  lT) ,% |d" lT! <  , L !&  l ! "# $% &' () *+ ,- ./ 01 23 45   43 B  <|  ,   \\ ,  <|  ,   \\ ,  ,  , l ,    <      | , , ,  <   ,lL 4U U UU< U  UUUU<TYGG LG!&G6;GMPG GGGG'+|   ,    < t# <'+  |         t#      | l  l    .  | l ,          l  < |         KCKKKKKKSSel#|    U U U U U  U  UUUU''''''.3....CC$s)CCCCC@#@%S 490D*E@@@@H H49oH04T#YH HHHH JCJE#u|:K4LJJJJ> >>>>3C3D3333 PPP#PC$JPRUPPPP*3*4**** RR R R R R  R  RRRR4c4el4444??d5<&8?T????0S0000< 5< %& 8#')) |) \))< )  ) #\ ) )4<) ) 5< %&))))X   L ,    l  |  \6#&25FI 759</00  ]-.. l l lT!       /.  759<,--/]*++/|/ L/ L/l / , /  ///,%/,%/I$/,%/,%// ///,+/& 759<455 4]233 4l 4l 4 4lT! 4 ; 4  4  4 4 443 4CCCCCCCC4 .HIIJJ222 !$   L       $  |  ,   \    &!$   L <           $"&# #+#+ <#+ <$ #+  #+ #+#+#+#+L $  |  L $  lT!  ,%  ,   \       l  ; \   \:  <|  ,   $  < |   L  \  l     ,% #s(O      <$       F <l L <$   <$   ,   \2 L  \ <   lT!  \D  , l   \6 759 !677 L]455L <L <LlLl L | LLL65LM |M |M ,M , M < M  \ MMMM 759 !:;; L]899L <L <LlLl L | LLL:9LQ |Q |Q ,Q , Q < Q  \ QQQQ 759 !788 ]566 < !L 38< JO,   76  759 !899 ]677 < !L 38< JO,   87 PW#U#U#U |#U #U189K1L#U 5< %&#U#U#U#UL 759 !233  ]011 < <    l    21 V |V V49LV 5< %&VVVV& 7 7%, 7 7 7 7 7 7  759<-..9]+,,9EL949EL999-,99999sA A'&A'