o )i'6@sddlZddlmZmZmZddlZddlmZddlm Z ddl m Z ddl m Z e eZe sUe sUzddlZWneyTZz edeWYdZ[ndZ[wwdZeeddlZdZWdn1smwYeryd d Znzdd lmZWneydd lmZYnw d*dejdejdejdejdededejdejdededeejde dejdejdedededed ed!df(d"d#Z! d*dejd$ejd%ejd&ejdejdejdejdededejdejdededeejde dejdejdedededed ed!df.d'd(Z" d+dejd$ejd%ejd&ejdejdejdejdededejdejd)eejdededeejde dejdejd*eejd!df(d+d,Z#dejdejd-ejdedejdejd!dfd.d/Z$ d+d0ejd1ejd2ejd3ejd4ejd5eejd!dfd6d7Z% d,d8ejd9ejd:ejd;ejded?e&d!e'ejejejejffd@dAZ( d,d8ejd9ejd:ejd;ejdBejdCejded?e&d!e'ejejejejffdDdEZ)dFejdejdGeejdHedIejdJe&d!dfdKdLZ*dFejdejdGeejdHedIejdJe&dMedNejd!dfdOdPZ+dejdQejdRejdSed!df dTdUZ,dQejdVejdRejdSed!df dWdXZ-dYejdZejd[ejd\ejd!df d]d^Z.dYejdZejd[ejd\ejd!df d_d`Z/dYejdZejd[ejd\ejd!df dadbZ0  d-dQejdRejdSedcej1ddeejdVeejd!e'ejejffdedfZ2dgejdhejdiejdjedkedled!ejfdmdnZ3dQejdgejdoejdhejdjed!ejf dpdqZ4drejdsejdtejduejdvejdwe&dxed!ejfdydzZ5e6ej7j8dzred{drejdsejdtejduejdvejdwe&dxed!ejfd|d}Z9d~ejdejdxed!dfddZ:drejdsejdejdejdededed!ejfddZ;drejdsejdejdejdejde dededed!ejfddZed    d.drejdeejdsejdeejdejdeejdeejdeejdeejdejdedej=dej=dej=de&de&de&de&d!ejf&ddZ?eddrejdsejdejdejdejdejdej=dej=dej=d!ejfddZ@eddrejdsejdejdejdej=dej=dej=d!ejfddZAeddgejdhejdiejdjej=dkedled!ejfddZBeddQejdgejdoejdhejdjej=d!ejf ddZCed       d/drejdejde deej1deejdeejdeedeejdeejdee d!ejfddZDeddsejdej1de deej1d!ejf ddZEe6ej7j8dreddrejdejdejdeejdej=dej=dej=dej=dej=de&de&d!ejfddʄZFe6ej7j8d˃r\ed̃ d+dejdedej=dej=deej1d!ejf dd҄ZGedӃdejdejdedej=d!ejf ddׄZHed؃dejdejdedej=d!ejf ddڄZIedۃdejdejdejdejdejdedej=dej=dej=d!ejfddZJe6ej7j8dreddejdejdejdededej=dej=d!ejfddZKded!e&fddZLd0ejdrejdejdejdejdejdejfddZMdrejdejdejdejdejdej1d!ejfddZNded!e&fddZOded!e&fddZP d+drejdejdejdejdej1deejd!ejfddZQ  d-drejdejdejdejdej1dejdeejdeejd!ejfddZRded!e&fddZSded!e&fddZTdrejd!e'ejejffd d ZU d+drejd ejd ejdejdejdej1deejd!ejfd dZV d+dejdejdejdejdejdejdedededeejfddZWdejdejfddZXdejdejdejdejdedededefdd ZYd!ejd"ejd#ejd$ejdejdejdejd%ejd&ejd'ejd(e&d)e&fd*d+ZZd!ejd"ejd#ejd$ejdejd,ejdejdejd-ejfd.d/Z[dsejdejdeded0ed!ejf d1d2Z\dsejdeded0ed!ejf d3d4Z]dsejdejdeded0ed!ejf d5d6Z^dsejdejdeded0ed!ejf d7d8Z_    d.drejdeejdsejdeejdejdeejdeejdeejdeejdejde dededede&de&de&de&d!ejf&d9d:Z`    d0dej1de deej1d;eej1dd?Zb       d/drejdejde deej1deejdeejdeedeejdeejdee d!ejfd@dAZcdsejdej1de deej1d!ejf dBdCZde6ej7j8dDrjedEdrejdejd!ejfdFdGZedrejdejd!ejfdHdDZfdQejdIejd!e'ejejffdJdKZgdejdIejdejdejdLed!e'ejejff dMdNZh     d1dQejdeejdOeeddeejdPe&d0eejd!e'ejejffdQdRZi  d2dgejdejdSeejde&d!e'ejejejff dTdUZjdrejdejdejdeejdededededede&de&d!ejfdVdZk   d3dQejdeejdeejdWe&d!e'ejejeejff dXdYZldrejdsejdejdejdejdejdededed!ejfdZd[Zmdejdedededeej1d!ejf d\d˄Zndejdejdeded!ejf d]d^Zodejdejdeded!ejf d_d`Zpdejdejdejdejdejdedededed!ejfdadbZqdejdejdejdededej=dej=d!ejfdcdZrded!efdddeZsdfejdgejdhejdiejdjejdkeejdleejdmeejdne&d)eejdoeejdpeejdqejdrefdsdtZtdrejdejdued!ejfdvdwZudrejdejdxed!ejfdydzZvdrejdejdej1dejdejdxed!ejfd{d|ZwdQejd0ejfd}d~Zxdejdededejdejdejd!dfddZydQejd0ejdejdejdeejdeejdejdejdejdededededxed!ejfddZzdejdejdejdejd!df ddZ{dQejd0eejdejdeejdejdeejdeejdeejdeejdejdejdejdejdejdedede&de&de dededede&de&de&de&d!ejf6ddZ|e r e6ej7j}d r eddrejdejdejdejdejdejdejdejdejdejde dej=dej=dej=de&dedLedede&de&d!ejf*ddZ~eddQejd0eejdejdejdeejdeejdeejdejdejdejdejdejdedede&de&de dededede&de&de&de&d!ejf2ddZdGejdejdejdejdejde dejdejd!dfddZdGejdejdejdejdejde dejdejd!dfddZdejdejd-ejdejde dejd!dfddZdeaejdeaejdejd!dfddZdeaejdejd!dfddZdejdejdejd!dfddZ  d4d0ejdQejdede d!df ddZ d+dejdejdejdejdedeejd!dfddZdeded!efdÐdĄZded!efdŐdƄZdeaejdejdede&d!ef dːd̄Zdedejdejdeded!df dѐd҄Zded!dfdӐdԄZd!efdՐdքZdedeaed!dfdאd؄Zded!e'eaeeaeffdِdڄZdedeaeaedeaeaed!dfdݐdބZded!e'eejffddZdejfddZded!dfddZ d+dededeed!efddZded!dfddZ d5dedejdejdede&d!df ddZded!ejfddZdedeaejd!dfddZd!efddZdejdeded!e'ejejffddZ  d2dejdejdejdejdedejdejdeed?e&d!e'ejejffddZdejdejdejdejdejdejded!ejfdd Zdejdejdejdejdejdejdejded ed!ejfd d Zded eded ed!ef ddZe6ej7j8dreddejdejdeejde&d!ejf ddZe6ej7j8drFeddejdejdejdejdejde&de&de&deejd eejdeeaed!eejd"eejde&d!ejfd#d$Ze6ej7j8d%rwed&dejdejd'ejdeejdej1de&d!ejfd(d)ZdSdS(6N) TYPE_CHECKINGOptionalUnion) init_logger)current_platform) ScalarTypez%Failed to import from vllm._C with %rFTcs fddS)NcsSN)namefnr \/home/app/PaddleOCR-VL-test/.venv_paddleocr/lib/python3.10/site-packages/vllm/_custom_ops.pyszregister_fake..r r r r r register_fake r)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_stepreturncC6tjj|||||||||| | | | | |||||dSr)torchops_Cpaged_attention_v1)rrrrrrrrrrrrrr r!r"r#r$r%r r r r+'s r+exp_sum max_logitstmp_outcCs<tjj|||||||||| | | | | ||||||||dSr)r(r)r*paged_attention_v2)rr,r-r.rrrrrrrrrrrrr r!r"r#r$r%r r r r/Ds  r/query_start_loc fp8_out_scalecCr'r)r(r)_rocm_CZpaged_attention)rr,r-r.rrrrrrrr0rrrrrr r1r r r paged_attention_rocmdsr3kv_cachecCtjj||||||dSr)r(r)Z_C_cpuZmla_decode_kvcache)rrr4rrrr r r mla_decode_kvcache_cpur6output prefix_output prefix_lse suffix_output suffix_lse output_lsecCstjj||||||dSr)r(r)r*merge_attn_states)r8r9r:r;r<r=r r r r>r> q_seqlens kv_seqlensvertical_indexes slash_indexes context_size block_size_M block_size_NcausalcCs|d}|d} |d} |d} ||d|} tj|| | |j|jd} tj|| | | |j|jd}tj|| | |j|jd}tj|| | | |j|jd}tjj| ||||||||||| | |||fSNrdtypedevice)sizer(zerosrLrMr)r*convert_vertical_slash_indexes)r@rArBrCrDrErFrG batch_size num_heads nnz_slash nnz_verticalnum_rows block_count block_offset column_count column_indexr r r rPsJ     rPvertical_indices_countslash_indices_countc Cs|d} |d} |d} |d} ||d|}tj| | ||j|jd}tj| | || |j|jd}tj| | ||j|jd}tj| | || |j|jd}tjj|||||||||||||| ||||fSrH)rNr(emptyrLrMr)r*(convert_vertical_slash_indexes_mergehead)r@rArBrCrZr[rDrErFrGrQrRrSrTrUrVrWrXrYr r r r]sJ      r] positionskey head_size cos_sin_cacheis_neoxcCr5r)r(r)r*rotary_embedding)r^rr_r`rarbr r r rcr7rcrot_dimcos_sin_cache_offsetsc C tjj||||||||dSr)r(r)r*batched_rotary_embedding)r^rr_r`rarbrdrer r r rgsrginputweightepsiloncCs |}tjj||||dSr) contiguousr(r)r*rms_norm)rrhrirjZinput_contiguousr r r rlsrlresidualcCtjj||||dSr)r(r)r*fused_add_rms_norm)rhrmrirjr r r rosrologits prompt_mask output_maskrepetition_penaltiescCsN|jddd|d}t||B|d}t|dkd||}||9}dS)NrI)dim?r)Z unsqueezerepeatrNr(where)rprqrrrsZ penaltiesZscalingr r r apply_repetition_penalties_torchs    rxcCrnr)r(r)r*Zapply_repetition_penalties_rprqrrrsr r r apply_repetition_penalties_cuda(rzcCs2|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_contiguousrzrxryr r r apply_repetition_penalties/s r~ quant_dtypescale_ubc CsVtj||d}tj||jddf|jtjd}tjj |||||||||fS)NrLrIrMrL) r( empty_liker\numelshaperMfloat32r)r* rms_norm_dynamic_per_token_quant)rhrirjrrrmr8scalesr r r rCsrqweightrrO split_k_itersthxthycCs6tjrddlm}||||Stjj||||||S)Nr)awq_dequantize_triton)envsVLLM_USE_TRITON_AWQ2vllm.model_executor.layers.quantization.awq_tritonrr(r)r*awq_dequantize)rrrOrrrrr r r rXs   rqzeroscCs8tjrddlm}||||||Stjj|||||S)Nr)awq_gemm_triton)rrrrr(r)r*awq_gemm)rhrrrrrr r r rcs ra b_q_weight b_gptq_qzeros b_gptq_scalesb_g_idx use_exllamabitc Ctjj|||||||Sr)r(r)r* gptq_gemmrrrrrrrr r r rmsrz _C::gptq_gemmcCs$tj|d|df|j|jdS)NrrIrKr(r\rNrLrMrr r r _gptq_gemm_fakewsrq_weightq_permcCtjj|||dSr)r(r)r* gptq_shuffle)rrrr r r rrb_scales workspacesize_msize_nsize_kc Crr)r(r)r* marlin_gemmrrrrrrrr r r rsrb_metab_q_typec Cs tjj||||||j||| Sr)r(r)r*gptq_marlin_24_gemmid rrrrrrrrrr r r rsrz_C::gptq_marlin_24_gemmc Cstj||f|j|jdSNrr(r\rMrLrr r r _gptq_marlin_24_gemm_fakerz_C::gptq_marlin_gemmcb_bias global_scaleb_zerosg_idxperm b_q_type_id is_k_fulluse_atomic_adduse_fp32_reduce is_zp_floatcCstj| | f|j|jdSrr)rrrrrrrrrrrrrrrrrrr r r _gptq_marlin_gemm_fakesrz_C::marlin_qqq_gemms_toks_chs_groupc Cstj||ftj|jdSNrKr(r\float16rM rrrrrrrrrr r r _marlin_qqq_gemm_fakes rz_C::marlin_gemmcCstj||ftj|jdSrrrr r r _marlin_gemm_fakes rz_C::awq_dequantizec Cs4|d}|d}|d}tj||f|j|jdSNrrIrKrNr(r\rLrM) rrrOrrrZin_cZqout_cZout_cr r r _awq_dequantize_fakes   rz _C::awq_gemmcCs4|d}tj|||ddf|j|jddSr)rNr(r\rLrMsum)rhrrrrZ num_in_featsr r r _awq_gemm_fakes rz_C::machete_mmb_qb_typeout_typeb_group_scales b_group_zeros b_group_sizeb_channel_scalesa_token_scalesschedulec Cs,|d} |d} tj| | f|j|jdS)NrrIrrNr(r\rMrL) rrrrrrrrrrmnr r r machete_mm_fakes  rz_C::machete_prepack_Ba_typegroup_scales_typecCstj|tjdS)N)Z memory_format)r(rZcontiguous_formatrrrrr r r machete_prepack_B_fakesrallspark_w8a16_gemmz_C::allspark_w8a16_gemm b_qweightb_qzerosr group_sizesm_count sm_versionCUBLAS_M_THRESHOLDhas_zpn32k16_reorderc Cs"|d} tj| |f|j|jdS)Nrrr) rrrrrrrrrrrrr r r _allspark_w8a16_gemm_fakes rggml_dequantizez_C::ggml_dequantizeW quant_typerrLcCstj||ftj|jdSrrrrrrrLr r r _ggml_dequantize_fake rrz_C::ggml_mul_mat_vec_a8XrowcCstj|jd|f|j|jdSNrrK)r(r\rrLrMrrrrr r r _ggml_mul_mat_vec_a8_fakesrz_C::ggml_mul_mat_a8cCs"|d}tj||f|j|jdSrr)rrrrbatchr r r _ggml_mul_mat_a8_fakes rz_C::ggml_moe_a8sorted_token_ids expert_idsnum_tokens_post_paddedtop_ktokensc Cs&|d}tj|||ftj|jdSr)rNr(r\rrM rrrrrrrrrr r r _ggml_moe_a8_fake's rggml_moe_a8_vecz_C::ggml_moe_a8_vectopk_idscCs&|d}tj|||f|j|jdSrrrrrrrrrr r r _ggml_moe_a8_vec_fake;s rcuda_device_capabilitycCtjj|Sr)r(r)r*cutlass_scaled_mm_supports_fp4rr r r rLrbscales_ascales_b problem_sizesexpert_offsetsc Cstjj|||||||dSr)r(r)r*#cutlass_blockwise_scaled_grouped_mm)r8rr r r r r r r r rPs r block_scale_a block_scale_balpha out_dtypec Cs`|jdkr |jdks J|jd|jd}}tj||f||jd}tjj|||||||S)NrJrrK)ndimrr(r\rMr)r*cutlass_scaled_fp4_mm) rr rrrrrrrr r r r^srcCrr)r(r)r*cutlass_scaled_mm_supports_fp8rr r r rjrrcCrr)r(r)r*$cutlass_scaled_mm_supports_block_fp8rr r r rnrscale_ascale_bbiasc Cs|tjus |tjus J|dus ||jdkr|j|ks Jg|jdd|jdR}|d|jd}|jdddkoI|jdddk}tsP|s`ddl m }|||||||} ntj |jd|jdf||j d} 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 NrIrr)triton_scaled_mmrK)r(bfloat16rrrrLviewris_rocmZKvllm.model_executor.layers.quantization.compressed_tensors.triton_scaled_mmrr\rMr)r*cutlass_scaled_mm) rr rrrr target_shapeZcutlass_compatible_brrr r r r ss"$   r azp_adjazpc Cs|jdddkr|jdddksJ|tjus |tjus J|dus4||jdkr2|j|ks4Jg|jdd|jdR}|d|jd}|dus[||jdks[Jtj|jd|jdf||jd} 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. rrrINrrK) rr(rrrrLrr\rMr)r*cutlass_scaled_mm_azp) rr rrrr"r#rr!rr r r r$s"(  r$cCrr)r(r)r*"cutlass_sparse_scaled_mm_supportedrr r r r%rr%cCrr)r(r)r*cutlass_group_gemm_supportedrr r r r&rr&cCsV|jtjtjtjtjfvsJ|sJd}|jdd|dks$Jtjj |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)`. rIrJr) rLr(int8 float8_e4m3fnrrr}rr)r*cutlass_sparse_compress)rZelemsPerMetaElemr r r r*s  r*bt_nzsbt_metac Cs|jdddkr|jdddksJ|tjus |tjus J|dus5|jd|jdkr3|j|ks5J|jd}|jd}tj||f||jd} 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. rrrINrK) rr(rrrLr\rMr)r*cutlass_scaled_sparse_mm) rr+r,rrrrrrrr r r r-s(  r-problem_sizes1problem_sizes2input_permutationoutput_permutation num_expertskblockscale_offsetsc Cs 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) rr r.r/r0r1r2rr3r4r r r r5&s !r5 input_tensor dst2src_mapcCs>|jd}tj||jdf|j|jd}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. rrIr)rr(r\rMrLr)_moe_C shuffle_rows)r6r7Znum_tokens_permutedZ output_tensorr r r r9Os r9expert_num_tokensnum_local_expertspadded_mc Cstjj||||||||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)r r.r/r:r;r<rr3r r r r=\sr= out_tensors a_tensors b_tensorsa_scales a_strides b_strides c_strides per_act_token per_out_chc Cs$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) r>r?r@rArr r rBrCrDrErFr r r rGts rGalphas sf_offsetsc Ctjj||||||||| 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*Zcutlass_fp4_group_mm) r>r?r@rArrHr r rIr r r cutlass_fp4_moe_mms rKnum_bitscCtjj|||||Sr)r(r)r*gptq_marlin_repack)rrrrrLr r r rNrNcCtjj||||Sr)r(r)r*awq_marlin_repack)rrrrLr r r rQrQcCst|jd}|ddks Jtj||d||df|j|jd}t|D]}tjj|||||||||<q$|SNrrrJr) rr(r\rMrLranger)r*rNrrrrrLr2r8er r r gptq_marlin_moe_repacks   rWcCsn|jd}|ddks Jtj||d||df|j|jd}t|D]}tjj|||||||<q$|SrS) rr(r\rMrLrTr)r*rQrUr r r awq_marlin_moe_repacks   rXcCs2tjj|||||||||| | j| | | ||||Sr)r(r)r*gptq_marlin_gemmr)rrrrrrrrrrrrrrrrrrr r r rYsrYgroup_zeros_typechannel_scales_typetoken_scales_typec Cstjj||j|||||Sr)r(r)r*machete_supported_schedulesr)rrrrZr[r\rr r r r]s r]c Cs"tjj|||j||||||| Sr)r(r)r* machete_mmr) rrrrrrrrrrr r r r^s r^cCstjj|||j|Sr)r(r)r*machete_prepack_Brrr r r r_ rOr_ permute_colsz_C::permute_colscC t|Srr(rrrr r r _permute_cols_fakes rdcCtjj||Sr)r(r)r*r`rcr r r r`input_global_scalec Cs trJ|jdksJd|jd|jdkrdnd}|||jd}|j\}}d}|j}||dks>Jd|d|jtjtj fvsPJd|jdtj ||d f|tj d }d d }||d } ||} || d} tj | | df|tj d } tj j||| || tj} || fS)a' 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. 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. rIz%input.ndim needs to be >= 1, but got .rrrz+last dim has to be multiple of 16, but got z-input.dtype needs to be fp16 or bf16 but got rJrcSs||d||S)NrIr )xyr r r rHsz"scaled_fp4_quant..r')rrrZreshaperrMrLr(rrr\uint8int32r)r*scaled_fp4_quantrr)) rhrgZ other_dimsrrrrMr8Zround_upZ rounded_mZscale_nZ rounded_nZ output_scaler r r rns:         rntopkc CstrJ|jdksJd|jdtj}|j\}}|||ks-Jd|d|d|d}|dd } tj||d|jtj d } tj||| tj |jd } tj j | | ||||| tj} | | fS) a Quantize input tensor to FP4 and return quantized tensor and scale, for packed MoE Inputs. Args: input_tensor: The input tensor to be quantized to FP4 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 FP4 output_scales: The blockscale tensor in FP8-E4M3 rJz%input.ndim needs to be == 2, but got rhz2m_numtopk must be less than MAX_TOKENS_PER_EXPERT(z,) for cutlass_moe_fp4, observed m_numtopk = z;. Use VLLM_MAX_TOKENS_PER_EXPERT_FP4_MOE to set this value.rr'rrK)rrrrZ"VLLM_MAX_TOKENS_PER_EXPERT_FP4_MOErr(r\rMrlrmr)r*scaled_fp4_experts_quantrr)) r6rgr r4roZMAX_TOKENS_PER_EXPERTZ m_numtopkr3Zscales_kZpadded_kr8Z output_scalesr r r rqVs>       rqnum_token_paddinguse_per_token_if_dynamiccCs|jdksJ|j}t}|rt||jd|df}|dur*tj||j|d}n|dus2Jd|j|ks9J|durr|r[tj|ddf|jtj d}tj j ||||||fStjd|jtj d}tj j |||||fS|dks~J|jtj 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 scale: Optional scaling factor for the FP8 quantization 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. Returns: tuple[torch.Tensor, torch.Tensor]: The output tensor in FP8 and scaling factor. rJrrINrz)padding not supported if output passed in)rrrZ fp8_dtypemaxr(r\rMrLrr)r*Z"dynamic_per_token_scaled_fp8_quantZdynamic_scaled_fp8_quantrZstatic_scaled_fp8_quant)rhrrrrrsr8rrr r r scaled_fp8_quants6   ru zero_pointc Cs|jd}|jd}|dddd}tj||f|j|jd}tjd|f|j|jd}d} |rD|dus8Jdtjd|f|j|jd} tjj||||||| ||| ||| fS)a Rearrange qweight, scale, and zero_point(if asymmetric) to n32k16 format for Ampere W8A16 Fused Gemm kernel Args: qweight: uint8 weight tensor, original k x n format. scale: fp16/bf16 weight scale tensor, 1 x n format. zero_point: fp16/bf16 weight zero_point tensor, 1 x n format. Must be provided for asymmetric quantization. has_zp: if use symmetric quantization, has_zp = False. if use asymmetric quantization, has_zp = True. Returns: tuple[torch.Tensor, torch.Tensor, Optional[torch.Tensor]] : rearranged weight, scale, and optionally zero_point. rrI rNz8zero_point must be provided for asymmetric quantization.)rr(r\rMrLr)r*Z#rearrange_kn_weight_as_n32k16_order) rrrvrKNZ N_32alignZqweight_reorderZ scale_reorderZzero_point_reorderr r r allspark_repack_weights2        rzc Cs"tjj|||||||||| | Sr)r(r)r*r) rrrrrrrrrrrr r r rs  symmetriccCstj|tjd}|dur%||duksJdtjj|||||||fStj||jddf|j tj d}|rtj|jd|jdf||jd}tjj|||||||Sr)r(r\rrMr)r2 wvSplitKQ)rr rrrrrr r r rs rcCtjj||dSr)r(r)r8moe_sum)rhr8r r r rsr experts_idsnum_tokens_post_padcCr5r)r(r)r8moe_align_block_size)rr2rrrrr r r rsr topk_weights BLOCK_SIZE_M BLOCK_SIZE_N BLOCK_SIZE_KcCs<tstdtjj|||||||||| | | | | dS)NzGThe optimized moe_wna16_gemm kernel is only available on CUDA platforms)rr|NotImplementedErrorr(r)r8moe_wna16_gemm)rhr8rrrrrrrrrrrrr r r rsrtoken_expert_indices gating_outputcCrnr)r(r)r8 topk_softmax)rrrrr r r rr{rnum_tokens_past_paddedmoe_block_sizemul_topk_weightsis_epcCsBtjj|||||||||| | | | | |||||j|||||||Sr)r(r)r8moe_wna16_marlin_gemmr)rhr8rrrrrrrrrrrrrrrrrrrrrrrrr r r rs rZmarlin_gemm_moez_moe_C::marlin_gemm_moe b_q_weights sorted_ids b_zero_pointsreplicate_input apply_weightscCstj| || f|j|jdSrr(r\rLrM)rrrrrrrrrrrrrrrr2rorrrr r r marlin_gemm_moe_fakes rz_moe_C::moe_wna16_marlin_gemmcCstj|| |f|j|jdSrr)rhr8rrrrrrrrrrrrrrrrrrrrrrr r r moe_wna16_marlin_gemm_fakesrvalue slot_mappingc Crfr)r(r) _C_cache_opsreshape_and_cacher_rrrrrrr r r r rs rc Crfr)r(r)rreshape_and_cache_flashrr r r rs  rkv_ck_pecCr5r)r(r)rconcat_and_cache_mla)rrr4rrrr r r r"sr key_caches value_caches block_mappingcCrr)r(r)r copy_blocks)rrrr r r r/sr kv_cachescCrr)r(r)rcopy_blocks_mla)rrr r r r5rRrsrcdstcCrr)r(r)r swap_blocks)rrrr r r r:rrrufp8kv_dtypecCrnr)r(r)r convert_fp8)r8rhrrr r r r?sr src_cache block_table cu_seq_lensrQ seq_startscCr5r)r(r)r gather_cache)rrrrrQrr r r rFr?r attributerMcCrer)r(r) _C_cuda_utilsget_device_attribute)rrMr r r rPrfrcCrr)r(r)r0get_max_shared_memory_per_block_device_attribute)rMr r r rTsr ipc_tensors rank_datarankfully_connectedcCrPr)r(r) _C_custom_arinit_custom_ar)rrrrr r r r[srfainp reg_bufferreg_buffer_sz_bytescCtjj|||||dSr)r(r)r all_reduce)rrrrrr r r rasrcCtjj|dSr)r(r)rdisposerr r r rgrrcC tjjSr)r(r)r meta_sizer r r r rkrrcCrer)r(r)rregister_buffer)rrr r r rorfrcCrr)r(r)rget_graph_buffer_ipc_metarr r r rsrrhandlesoffsetscCrr)r(r)rregister_graph_buffers)rrrr r r rwrrrNcCrr)r(r)r!allocate_shared_buffer_and_handle)rNr r r r|rr mem_handlecCrr)r(r)ropen_mem_handle)rr r r rrrptrcCrr)r(r)rfree_shared_buffer)rr r r rrr world_size qr_max_sizecCrr)r(r)rinit_custom_qr)rrrr r r rsrcCrr)r(r)r qr_destroyrr r r rrr quant_level cast_bf2halfcCrr)r(r)r qr_all_reduce)rrrrrr r r rsrcCrr)r(r)r qr_get_handlerr r r rrrcCrer)r(r)rqr_open_handles)rrr r r rrfrcCrr)r(r)rrr r r r rr cache_seqlensnum_heads_per_head_k num_heads_kcCr)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)rrrr r r rs rqk_cache head_dim_vtile_scheduler_metadata num_splits softmax_scalec CsB|dur |jdd}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*Zflash_mla_fwd_kvcache) rrrrrrrrrGrZ softmax_lser r r flash_mla_with_kvcaches rq_nopeq_pekv_c_and_k_pe_cache page_tablec Cstjj||||||||Sr)r(r)r*cutlass_mla_decode)rrrrrrrr r r rsr num_kv_splitsc Cs"tjj||||||||| |Sr)r(r)r*sm100_cutlass_mla_decode) rrrrrrrrrr r r rs r num_batchescCrPr)r(r)r*$sm100_cutlass_mla_get_workspace_size)rrrrr r r rsrZweight_packed_linearz_C::weight_packed_linearmat1mat2is_vnnicCs$tj|d|df|j|jdSrr)r r rr r r r weight_packed_linear_fakesr Zfused_experts_cpuz_C::fused_experts_cpu hidden_statesw1w2inplace use_int8_w8a8 use_fp8_w8a16w1_scalew2_scalea1_scalea2_scalecCrarrb)r rrrrrrrrrrrrr r r r fused_experts_cpu_fakes rZint8_scaled_mm_with_quantz_C::int8_scaled_mm_with_quantscales2cCs&|d}|d}tj||f|dS)Nrr)rNr(r\)r r rrrr Mryr r r int8_scaled_mm_with_quant_fake(s r)rrrrrr)T)NN)TFFF)NNNNNNN)NNNN)NNNFN)NF)NNT)rur)F) contextlibtypingrrrr(Z vllm.envsrZ vllm.loggerrZvllm.platformsrZvllm.scalar_typer__name__loggerZis_tpuZis_xpuZvllm._CZvllm ImportErrorrVwarningZsupports_moe_opssuppressZ vllm._moe_CrZ torch.libraryrZTensorintfloatstrr+r/r3r6r>booltuplerPr]rcrgrlrorxrzr~rLrrrrhasattrr)r*rrrrZSymIntrrrrrrrrrrrrrrrrrrrr r$r%r&r*r-r5r9r=rGrKrNrQrWrXrYlistr]r^r_rdr`rnrqrurzrr|r}rr~rrrrrrrrrrrrrr8rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr rrr r r r s               /      3          9    2                                                       8 37  )            99B/ )        &                           "( " $  ,