o پi@s ddlZddlZddlZddlZddlmZddlmZmZmZm Z m Z ddl Z ddl Z ddl mZzddlmZWnYddlmZddlmZmZmZmZmZmZmZmZddlmZeZeZ eZ!edomeZ"e rdd l#m$Z$dd l%m&Z'z dd l#m(Z(d Z)Wne*ydd l#m+Z+dZ)YnwerdZ,e"rz ddl-m.Z.m/Z/m0Z0Wne*ye*dwzddl1Z2d Z,Wn e*ydZ,Ynwe3e4Z5ede6fddZ7e7re j8Z9dZ:n e j;Z9e edgdde j?de j?de j?de j?de j?ddf ddZ@e jAdejBfddZCe jAdejBd ejBfd!d"ZDd#e9dddfd$e j?d%eEd&eFd'e jGd(e6d)e6d*e6de e j?e j?ffd+d,ZHeHZId$e j?d%eEd-e jGd(e6d)e6d*e6d.e e j?de e j?e j?ffd/d0ZJ #     dd$e j?d%eEd-e jGd&eFd(e6d)e6d*e6d1e6d.e e j?de e j?e j?ffd2d3ZKd(e6d)e6d*e6fd4d5ZL #      dd$e j?d%eEd&eFd(e6d)e6d*e6d1e6d.e e j?d6e e6fd7d8ZM #      dd$e j?d%eEd-e jGd&eFd(e6d)e6d*e6d1e6d.e e j?d6e e6fd9d:ZNe9fd$e j?d'e jGfd;d<ZOe jAdejBd=ejBfd>d?ZP dd$e j?d@e j?dAe6de e j?e j?ffdBdCZQe jAdDejBdEejBdFejBdGejBdHejBf dIdJZRe jAdDejBdEejBdFejBdGejBdHejBf dKdLZSejdMeEdNeEdOeEdPeEde eeEeff dQdRZTdSdTZUerrdUdVZVdWdTZUe jWfde j?de j?de j?de j?dXeeEdYe jGde eEeEeEffdZd[ZXde j?de j?de j?de j?dXeeEdYe jGde j?fd\d]ZYe jWfde j?de j?de j?de j?dXeeEdYe jGde j?fd^d_ZZe jWfde j?de j?de j?de j?dXeeEdYe jGde j?fd`daZ[e jAdbejBdMejBdNejBdcejBddejBdeejBdfejBdgejBdhejBdiejBdjejBfdkdlZ\dmdndmdodpdqe j?dre j?dse j?dte j?dYe jGdueEdOeEdPeEdveEde j?fdwdxZ]e jAdyejBfdzd{Z^e jAdyejBfd|d}Z_ ~dd$e j?de j?d&eFde e j?e j?ffddZ`e jAdejBdejBfddZadmd~e9fd$e j?d%eEd&eFd'e jGde e j?e j?ff ddZb erddZcddZdddZe   dde j?de e j?de eEde6defe j?e j?ff ddZgn   dde j?de e j?de eEde6defe j?e j?ff ddZge jhdddDgddZie jAdNejBdfejBdejBddejBfddZje!seiejZj dde j?de j?de j?d%eEdeEde e j?e j?ff ddZke jAde j?de j?de j?deEdejBdDejBdFejBfddZl ddqe j?de j?deEde j?fddZmd$e j?fddZne jAdejBdDejBdEejBdFejBdejBdejBf ddZo  dde j?de j?de j?de j?depe jGde e j?deEdeEdeEde j?fddZqe re)re jrsdddZtn e jrsdddZte jrsdddZtdSdS)N) lru_cache)AnyDictListOptionalTuple)TensorDescriptor)deep_gemm_wrapper) ceil_alignget_bool_env_varget_device_core_countget_device_nameis_cpuis_cudais_hiplog_info_on_rank0)register_custom_opSGLANG_USE_AITER)sgl_per_token_quant_fp8)per_tensor_quant_fp8)sgl_per_token_group_quant_8bitT)sgl_per_token_group_quant_fp8F)dynamic_per_tensor_quantdynamic_per_token_scaled_quantstatic_per_tensor_quantz6aiter is required when SGLANG_USE_AITER is set to TruereturncCstr dtjdjvSdS)Ngfx94rF)_is_hiptorchcudaget_device_properties gcnArchNamer"r"]/home/ubuntu/.local/lib/python3.10/site-packages/sglang/srt/layers/quantization/fp8_kernel.py is_fp8_fnuzYsr$l@C) mutates_argsAAsBBscCst||f||f|dSN)r gemm_nt_f8f8bf16)r(r)r*r+r&r"r"r#deep_gemm_fp8_fp8_bf16_ntjsr.BLOCKc Cstd} || |7}|| |7}|| 7}td|} | |k} tj|| | ddtj} ttt| |} | |}d|}t | ||||j j }tj || || dt ||dS)zA Triton-accelerated function to perform per-token-group quantization on a tensor. This function converts the tensor values into float8 values. rmaskother?r2N) tl program_idarangeloadtofloat32maximummaxabsclampdtype element_tystore)y_ptry_q_ptry_s_ptry_strideNepsbit8_minbit8_maxr/g_idcolsr2y_absmaxy_sy_s_invy_qr"r"r#_per_token_group_quant_8bitus    rR SCALE_UE8M0c  Cs td} || tj|7}|| tj|7}||} | | } | | }|| ||7}td| }||k}tj|||ddtj}ttt ||}||}| ret t t t |}t |||||jj}tj||||dt||dS)zA Triton-accelerated function to perform per-token-group quantization on a tensor. This function converts the tensor values into float8 values. rr0r1r5N)r6r7r:int64r8r9r;r<r=r>exp2ceillog2r?r@rArB)rCrDrE group_size y_num_columnsy_s_col_striderHrIrJr/rSrKblocks_per_row scale_col scale_rowrLr2rMrNrOrQr"r"r#$_per_token_group_quant_8bit_colmajors"  r^绽|=xrXrHr@column_major_scalesscale_tma_aligned scale_ue8m0cCs|jd|dks Jd|sJdtr%|tjkrd}nd}| }n|tjkr0t|} nt|} | j}| j}tj ||j |d} t |j|j |||dd } | |} |} t | }tt|d d d }d }|rt| f|| | ||jd | d |||||||d n|rJt| f|| | || ||||||d |rddlm}|dksJ|| d| jd| jd d ||fdd} | | fS)a'Function to perform per-token-group quantization on an input tensor `x`. It converts the tensor values into signed float8 values and returns the quantized tensor along with the scaling factor used for quantization. Args: x: The input tensor with ndim >= 2. group_size: The group size used for quantization. eps: The minimum to avoid dividing zero. dtype: The dype of output tensor. Returns: Tuple[torch.Tensor, torch.Tensor]: The quantized tensor and the scaling factor for quantization. r=the last dimension of `x` cannot be divisible by `group_size``x` is not contiguousg_@r%devicer@Fx_shaperhrXrarbrc)rIrJr/ num_warps num_stagesrS)rIrJr/rnro!transform_sf_into_required_layoutNT num_groupsmnkrecipeis_sfa)shape is_contiguousrrint8iinfofinfor=min empty_likerh-create_per_token_group_quant_fp8_output_scalenumeltritonnext_power_of_2r^striderR deep_gemmrq)r`rXrHr@rarbrcrJrIinfox_qx_sMrGr/rnrorqr"r"r#_per_token_group_quant_8bit_raws         r dst_dtypemasked_mcCsHddlm}ddlm}|sJ|sJ|sJ|dk} | r;|j\} } |d}|dus0Jtj| g|j tj d}tj g|jdd|jddR|j |d} tj g|jdd|jdd|R|j tj d} ||| | |||d|dks~J|| | jd| jd | jdd ||fd d }| r| d} | d}| |fS) Nrrp)"silu_and_mul_masked_post_quant_fwdrgrd)inputoutput output_scalequant_group_sizerrcrrrlTrs)rrq$sglang.srt.layers.moe.ep_moe.kernelsrdimry unsqueezertensorrhint32zerosr;squeeze)r`rXrrarbrcrrqrneeds_unsqueeze num_tokens_routput_scale_for_kernelrr"r"r#-_per_token_group_quant_8bit_fuse_silu_and_mulBsT       $   rfuse_silu_and_mulc Cs0|r t|||||||dSt|||||||dS)N)r`rXrrarbrcr)r`rXrHrarbrcr@)rr) r`rXrrHrarbrcrrr"r"r#per_token_group_quant_8bits&  rcCs6|r>|r|sJ|^}}}||d} } t| d} t| d} tjg|| d| R|tjddddd| ddfS|r|ro|dddd} tj|dd|d|| f|tjdddd|dddfStj|d|f|dd|tjdddStj|dd|d|f|tjdS)Nrrrgrdr.)r remptyint transposer;permute)rjrhrXrarbrcx_batchx_q_mnx_q_kx_s_mnx_s_k aligned_mn aligned_k aligned_sizer"r"r#rsF    r enable_v2c Cs|jd|dks Jd|sJdg|jdd|jd|r%dndR} tj| |jtd} t| |j||||d} |jddkritrZt|| | ||t t ||||d | | fS|r^Jt || | ||t t || | fS) Nrdrrerfrrlrgri)r) ryrzrrrh fp8_dtyper%enable_sgl_per_token_group_quant_8bitrfp8_minfp8_maxr) r`rXrHrarbrcrrr out_shaperrr"r"r# sglang_per_token_group_quant_fp8sF * rc Csfddlm} |tjkr&|rJ|rJ|rJ|dusJ| ||||| dSt||||||||| d S)Nr)!sglang_per_token_group_quant_int8)r`rXrHr@r) r`rXrHrarbrcrrr)*sglang.srt.layers.quantization.int8_kernelrrr{r) r`rXrrHrarbrcrrrrr"r"r#!sglang_per_token_group_quant_8bits0  rcCsR|sJdtj||j|d}tj|jdd|jtjd}t|||||fS)Nrfrgrrl)rzrrrhrryr;r)r`r@rrr"r"r#sglang_per_token_quant_fp89s r REPEAT_SCALEc Cstd} || |7}|| |7}| r|| 7}td|} | |k} tj|| | ddtj} t|tj}d|}t| ||||jj}tj || || d| r^t ||dSdS)zA Triton-accelerated function to perform quantization using the given scale on a tensor This function converts the tensor values into float8 values. rr0r1r4r5N r6r7r8r9r:r;r?r@rArB)rCrDrEy_s_repeat_ptrrFrGrrr/rrKrLr2rMrOrPrQr"r"r#_static_quant_fp8Ls    rr repeat_scalec Cs|sJd|dksJdtj||jtd}||jd}|jd}|r8tj|df|jtjd}nd}t |}t t |ddd}d} t |f||||||tt|||| d |rb|n|}||fS) a@Function to perform static quantization using the given scale on an input tensor `x`. It converts the tensor values into signed float8 values and returns the quantized tensor along with the scaling factor used for quantization. Args: x: The input tensor with ndim >= 2. x_s: The quantization scale. repeat_scale: Whether to broadcast per-tensor scale to per-channel scale. dtype: The dype of output tensor. Returns: Tuple[torch.Tensor, torch.Tensor]: The quantized tensor and the scaling factor for quantization. rfrlzonly supports per-tensor scalergrdNrkrm)rrr/rrnro)rzrrrrhrryrr;rrr~r=rrr) r`rrrrrG x_s_repeatr/rnror"r"r#static_quant_fp8ws>   r BLOCK_SIZE_M BLOCK_SIZE_N BLOCK_SIZE_K GROUP_SIZE_M needs_maskingc6Cstjdd}t||}t||}||}||}||}t|||}|||} |||}!| |td||}"|!|td||}#td|}$||"dddf| |$dddf| }%||$dddf| |#dddf| }&||"|}'|#|}(||(|})|| }*tj||ftjd}+tdt||D]q},|rtj|%|$dddf||,|kdd}-tj|&|$dddf||,|kdd}.n t|%}-t|&}.t|'}/t|)}0|+t |-|.|/dddf|0dddf7}+|%|| 7}%|&|| 7}&|'|*|7}'|)|*|7})q|j j tj kr#|+ tj }1n|j j tjkr2|+ tj}1n|+ tj}1| |td|}2|!|td|}3|||2dddf||3dddf}4|2dddf|k|3dddf|k@}5tj|4|1|5ddS)Triton-accelerated function used to perform linear operations (dot product) on input tensors `A` and `B` with block-wise quantization, and store the result in output tensor `C`. raxisNr@r0r1r5r6r7cdivr~r8rr;ranger9dotr@rAbfloat16r:float16rB)6r(r*r&r)r+rrGKgroup_ngroup_k stride_am stride_ak stride_bk stride_bn stride_cm stride_cn stride_As_m stride_As_k stride_Bs_k stride_Bs_nrrrrrpid num_pid_m num_pid_nnum_pid_in_groupgroup_id first_pid_m group_size_mpid_mpid_noffs_amoffs_bnoffs_ka_ptrsb_ptrsAs_ptrsoffs_bsnBs_ptrs scale_step_k accumulatorrvaba_sb_scoffs_cmoffs_cnc_ptrsc_maskr"r"r#_w8a8_block_fp8_matmulsT &     ,,  (*    0    ,(rc7Cstjdd}t||}t||}||}||}||}t|||}|||} |||}!| |td||}"|!|td||}#td|}$||"dddf| |$dddf| }%||$dddf| |#dddf| }&||"|}'|#|}(||(|})|| }*tj||ftjd}+d},tdt|||,D]}-|rtj|%|$dddf||-|,|kdd}.tj|&|$dddf||-|,|kdd}/n t|%}.t|&}/t|'}0t|)}1|+t |.|/|0dddf|1dddf7}+|%|| 7}%|&|| 7}&|'|*|7}'|)|*|7})|rPtj|%|$dddf||-|,d|kdd}.tj|&|$dddf||-|,d|kdd}/n t|%}.t|&}/t|'}0t|)}1|+t |.|/|0dddf|1dddf7}+|%|| 7}%|&|| 7}&|'|*|7}'|)|*|7})|rtj|%|$dddf||-|,d |kdd}.tj|&|$dddf||-|,d |kdd}/n t|%}.t|&}/t|'}0t|)}1|+t |.|/|0dddf|1dddf7}+|%|| 7}%|&|| 7}&|'|*|7}'|)|*|7})|r@tj|%|$dddf||-|,d |kdd}.tj|&|$dddf||-|,d |kdd}/n t|%}.t|&}/t|'}0t|)}1|+t |.|/|0dddf|1dddf7}+|%|| 7}%|&|| 7}&|'|*|7}'|)|*|7})q|j j tj kr|+ tj }2n|j j tjkr|+ tj}2n|+ tj}2| |td|}3|!|td|}4|||3dddf||4dddf}5|3dddf|k|4dddf|k@}6tj|5|2|6d dS) rrrNrrr0r1rlrrr5r)7r(r*r&r)r+rrGrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr UNROLL_FACTORrvrrrrrrrrrr"r"r#!_w8a8_block_fp8_matmul_unrolledx4 s &     ,,      0    ""    0    ""    0    ""    0    ,(rrGrblock_nblock_kc CstjrdStdd}d|d|d|d|d|d }tjtjtj t d |}tj |r_t |}t td |d d dt|DWdS1sZwYtd|dS)a} Return optimized configurations for the w8a8 block fp8 kernel. The return value will be a dictionary that maps an irregular grid of batch sizes to configurations of the w8a8 block fp8 kernel. To evaluate the kernel on a given batch size bs, the closest batch size in the grid should be picked and the associated configuration chosen to invoke the kernel. N rzN=z,K=z ,device_name=z,dtype=fp8_w8a8,block_shape=[z, z].jsonconfigszUsing configuration from z for W8A8 Block FP8 kernel.cSsi|] \}}t||qSr")r).0keyvalr"r"r# sz.get_w8a8_block_fp8_configs..ziUsing default W8A8 Block FP8 kernel config. Performance might be sub-optimal! Config file not found at %s)r_dynamo is_compilingr replaceospathjoindirnamerealpath__file__existsopenrloggerjsonr9itemswarning)rGrrr device_namejson_file_nameconfig_file_pathfr"r"r#get_w8a8_block_fp8_configss( $     rcCstSr,)rrrGMETAr"r"r##select_w8a8_block_fp8_matmul_kernelsrcCs.t||dt||d}|tkdSNrr)rrr )rrGrnum_workgroupsr"r"r#$use_w8a8_block_fp8_matmul_unrolledx4srcCst|||rtStSr,)rrrrr"r"r#rs  block_size output_dtypec Cst|dksJ|d|d}}|jd|jdksJ|jdd|jddks-J|s3J|jtjkrJt|jd||jdksIJn*|jtjkrrtt|jd|d|jdksqJd|jd|jd|nt | |jd}|j dksJ|sJ|j dksJ|j\} } |jtjkrt| ||jdksJt| ||jdksJn=|jtjkr| |jdksJd |jd |jd|tt| |d|jdksJd |jd |jd|nt |jdd| f} |j | |d } || | | fS) NrrrlrdrzA.shape=z As.shape=z block_size=zB.shape=z Bs.shape=r) lenryrzr@rfloatrrrNotImplementedErrorrndim new_empty) r(r*r)r+rr rrrrGrC_shaper&r"r"r#prepare_block_fp8_matmul_inputss:   " &    ,  r'c CsDt||||||\}}}} | jtjkrtjsJt||||| | Sr,)r'r@rrr ENABLE_JIT_DEEPGEMMr.) r(r*r)r+rr rrGrr&r"r"r#w8a8_block_fp8_matmul_deepgemm0sr)cs$t||||||\}}|\}} t||d|d} | r-| t| fddd} n d|d|dddd d } t|| d dk} fd d } t| }|| |||||||| |d|d|d|d|d|d|d|d|d|dfi| d| i|S)aThis function performs matrix multiplication with block-wise quantization. It takes two input tensors `A` and `B` with scales `As` and `Bs`. The output is returned in the specified `output_dtype`. Args: A: The input tensor, e.g., activation. B: The input tensor, e.g., weight. As: The per-token-group quantization scale for `A`. Bs: The per-block quantization scale for `B`. block_size: The block size for per-block quantization. It should be 2-dim, e.g., [128, 128]. output_dytpe: The dtype of the returned tensor. Returns: torch.Tensor: The result of matmul. rrlcs t|Sr,)r>r`)rr"r#cs z.w8a8_block_fp8_matmul_triton..)r@ rr)rrrrrnrorc"t|dt|dfSrrrrrrGr"r#gridrsz*w8a8_block_fp8_matmul_triton..gridrrdr)r'rr~keysboolrr)r(r*r)r+rr rr&rrrconfigrr2kernelr"r1r#w8a8_block_fp8_matmul_tritonBsT  r7cCs8|tjkrtjrt||||||dSt||||||dS)N)r )rrr r(r)r7)r(r*r)r+rr r"r"r#w8a8_block_fp8_matmuls  r8r output_typeBLOCK_MBLOCK_NBLOCK_Krep_mrep_nrep_k NUM_STAGESc# Cs|dkrtj}n|dkrtj}n|dkrtj}tjdd}t|| }||}||}|| }|| }d}d}|| }|| }d}d}tj| | ftjd}tjdt|| |dD]h}|||g}|||g} |d||ddg}!|d||ddg}"|! | |ddd dd ddd | | |}!|" | |ddd dd ddd | | |}"t ||!d | j |"d |}|| 7}|| 7}||7}qY| ||g||dS) Nrrlrrr-r)rorre4m3)r6r;rrr7rrrr9reshapetrans dot_scaledTrBr:)#a_desc a_scale_descb_desc b_scale_descc_descrrGrr9r:r;r<r=r>r?r@r rrrrrroffs_k_aoffs_k_b offs_scale_m offs_scale_n offs_scale_kVEC_SIZErrvrrscale_ascale_br"r"r#!_mxfp8_block_scaled_matmul_kernelsN     rSrrrkr)block_mrrrora_scalerb_scalerTrocCsB|j\} } |j\} } | | ksJ|tjkrd} n|tjkr d} n|tjkr(d} ntd||d}|d}|dd}t|||g}t|||g}d||ddg}tj||d }d||ddg}tj||d }tj| | f||j d }t|||g}t | |t | |df}t ||||||| | | | ||||||||S) z6Block-scaled matmul for MXFP8 using Triton dot_scaled.rrlrzUnsupported output dtype: rrr-rrk) block_shaper@rh) ryrr;rr ValueErrorr from_tensorrrhrrrS)rrUrrVr rTrrrorrrGK_br9r=r>r?rFrHscale_block_shaperGrIrrJr2r"r"r# mxfp8_block_scaled_matmul_tritonsR      r] BLOCK_SIZEcCstd}td} td|} | |k} || |||7}tj|| | ddtj} ttt| |} t || |dS)Nrrlr0r1) r6r7r8r9r:r;r<r=r> atomic_max)x_ptrx_s_ptr head_size x_stride_h x_stride_srHrr^seq_idhead_idoffsetr2r`rNr"r"r# _per_tensor_quant_mla_fp8_stage1.s  rhc Cstd} td} td| } | |k} t|}d|}|| || |7}|| ||| |7}tj|| | ddtj}t|||||jj}tj || || ddS)Nrrlr4r0r1r5r)r`rax_q_ptrnum_seqrbrcrdrrr^rerfrgr2rx_s_invr`rr"r"r# _per_tensor_quant_mla_fp8_stage2Es   rl-q=x_s_outc Cs|dks Jd|jdkr|jtjkr|j|jksJ|j|td}|j\}}}t |}||f}t ||||| d| d|t |t||||||| d| dtt | ||fS)z This function quantizes input values to float8 values with tensor-wise quantization and specialized for mla absorbed case. r`x` is not a 3d-tensorrlrrrl)rryr@rr;rhr%sizerrrrhrrrlr) r`rnrHrnum_headrjrbr^r2r"r"r#per_tensor_quant_mla_fp8bs>      rs NUM_GROUPcCs td}td}|||||7}|||||7}||| |7}|dkr5t||tdtd|}||k}t|D]J}tj|||||ddtj}t t t || }|| }t ||| | |j j}tj||||||dt||| |qCdS)zA Triton-accelerated function to perform per-token-group quantization on a tensor for deep_gemm grouped_gemm_masked. This function converts the tensor values into float8 values. y and y_q: (b, t, k) y_s: (b, k//group_size, t) rrlr0r1r5N)r6r7rB num_programsr8rr9r:r;r<r=r>r?r@rA)rCrDrE masked_m_ptrrX y_stride_b y_stride_t y_q_stride_b y_q_stride_t y_s_stride_b y_s_stride_grHrrrtr/t_idb_idrLr2gidrMrNrOrQr"r"r#/_per_token_group_quant_mla_deep_gemm_masked_fp8s&    rcCs|dks Jd|j\}}}|ddd}||}|||ks*Jd|d|j|||f|d} |j|||ftjd} |j|ftjd} t|} ||f} t| || | | || d| d | d| d | d| d |t t || | | d d | ||fS) z This function quantizes input values to float8 values with per-token-group-quantization for deep_gemm grouped_gemm_masked and specialized for mla absorbed case. rrorkzk % z must be zerorrrlr) rryr%rr;rrrrrrr)r`rXrHr@rmrv aligned_m num_tiles_krrrr^r2r"r"r#.per_token_group_quant_mla_deep_gemm_masked_fp8s:  rc Csh|j\}}d}|jdddj}tj||d}|t}||t||tt t }||dS)zXNative PyTorch fallback for dynamic per-token FP8 quantization when vLLM is unavailable.rmrlT)rkeepdimr~N) ryr>r=valuesrr?rcopy_rr:r) rrscalerrGrHabsmax scale_val output_datar"r"r##_native_dynamic_per_token_quant_fp8s  rcCsbd}|}tj||d}|t}|d|dt||ttt }||dS)zYNative PyTorch fallback for dynamic per-tensor FP8 quantization when vLLM is unavailable.rmrrdN) r>r=rr?rviewrrr:r)rrrrHrrrr"r"r#$_native_dynamic_per_tensor_quant_fp8s rcCs&t||ttt}||dS)zMNative PyTorch fallback for static FP8 quantization when vLLM is unavailable.N)rr?rrr:rr)rrrrr"r"r#_native_static_quant_fp8!srrrnum_token_paddinguse_per_token_if_dynamiccCs|jdksJd|jd|j}|rt||jd|df}tj||jtd}|dur|rdtj|ddf|jtjd}trHt |||||fSt rZtj j |||d||fSt|||||fStjd|jtjd}trzt|||||fSt rtj j |||||fSt|||||fS|dksJd|trt|||||fSt rtj j |||||fSt|||||fS)NrExpected 2D input tensor, got Drrlrg!Expected scalar scale, got numel=)r$ryr=rrrhrr; _use_aiterr _has_vllmops_C"dynamic_per_token_scaled_fp8_quant contiguousrrrdynamic_scaled_fp8_quantrrrstatic_scaled_fp8_quantrrrrrryrr"r"r#scaled_fp8_quant'sP         rcCs|jdksJd|jd|j}|rt||jd|df}tj||jtd}|dur\|rFtj|ddf|jtjd}t|||||fStj d|jtjd}t |||dd||fS| dkskJd | t |||d d||fS) NrrrrrlrgF) is_staticrT) r$ryr=rrrhrr;rrsgl_per_tensor_quant_fp8rrr"r"r#rYs0   cCs(g|]}dD] }tjd|i|dqqS))rrrmr:)rn)rConfig)rrTrnr"r"r# sr)r-r,rr)rr< M_ALIGNMENT)rrrc Cstd} td} t|| d} t|| tj} t| |t| || |td|} tt| |D]}||td|}|| ||dddf|| dddf}|| kdddf| |kdddf@}tj||dtj }tj t |dd}tj |dt dd}|d |dddftj}|| ||dddf|| dddf}tj|||dt||}|| || | |}tj||d || kdq=dS) Nrrlrr5rg-C6?inf)r~r=g|@)r6r7r9r:rT multiple_ofr8rrr;r=r>r?r" float8e4nvrB)rexpert_offsets problem_sizesa_fp8sfarr<rr:k_offset expert_idrcurrent_expert_offsetcoord_kicoord_mra_maskinpinp_amaxinp_fp8 a_fp8_ptrsrvsfa_ptrsr"r"r#._per_token_group_quant_fp8_hopper_moe_mn_majors.   0( . rrlrrexpert_tokens_alignmentc Cs|dksJ|sJd|jd|dksJdtj||jtd}|jd|jd}}||}tj||f|jtjd} |jd} || f} t | ||||| ||||| fS)Nrz`A` is not contiguousrdrz=the last dimension of `A` cannot be divisible by `group_size`rgrl) rrzryrrrhrrr;r) r(rrrXra_qrrrvr num_expertsr2r"r"r#-per_token_group_quant_fp8_hopper_moe_mn_majors,  rdata_ptrtrans_data_ptrrvcCsbtd}td}td} t||} t||d} | | } t| |t| ||| |} || |}| |td|}||k}td| |tdD]Y}|||td|}|| k}|dddf||dddf}|dddf|dddf| }|dddf|dddf@}tj| ||d}tj||||dqUdS)Nrrlrr5)r6r7r9rr8rrurB)rrrrvrrrrm_idk_idcurr_expert_offsetnext_expert_offsetnum_tokens_of_expertdata_start_ptrtrans_data_start_ptrk_coordk_maskstart_mm_coordm_maskoff trans_offr2datar"r"r#_per_group_transposes*      $$ rc sr|dksJ|sJd|\t|}|ddfdd}t|||||ddd |S) Nrz`a` is not contiguousrrlcs.td|dt|dfS)Nrlrrr/r0rvrrr"r#r+ sz%per_group_transpose..rrm)rr)rrzrqrrr)rrrtrans_ar2r"rr#per_group_transposes  rcCsZ|}|j}|ddko|dtd|dk}|ddko(|dtd|dk}|p,|SNrrl)rryr=)r`stridessizesis_not_transpose is_transposer"r"r#is_weak_contiguouss ""rACCUMULATOR_DTYPEBLOCK_SIZE_SCALE_ABLOCK_SIZE_SCALE_Bc=Cstjdd}t||}||}||}|}tj||f|d}||td|tj}||k}||td|tj}||k}td|tj}| |dddf| |dddf} | |dddf| |dddf}!td||dk||}"|"|k}#td||dk||}$|$|k}%|| }&||!}'||"}(||$})tdt||D]M}*||k}+|dddf|+dddf@},tj|&|,d}-|+dddf|dddf@}.tj|'|.d}/tj |-|/||d}||7}|&|| 7}&|'|| 7}'q|#dddftdddkdddf@}0t|(dddf|0}1|1 |df}1|1|tj }|%dddftdddkdddf@}2t|)dddf|2}3|3 |df}3|3j |tj }||j j}4|r|}5||5}6|5|k}7t|6|7}8|4|87}4||td|tj}9||td|tj}:|9tj}9|:tj}:|| |9dddf||:dddf};|9dddf|k|:dddf|k@}.rXcSs|jddko |jddkSr)ryr*r"r"r#r+si r-)r,r,rk)r,rrrkr,rr)r,rrrr)rrrrrr)rrrrr)ryr@rrBis_floating_pointrrrrhr=rrr6r;rrrr:)rrrQrRrrrrr use_heuristicrr2result has_scalar is_small_Nnext_power_of_2_M tile_shape block_size_sa block_size_sbrr"r1r#triton_scaled_mmsn ..     r z*sgl_kernel::sgl_per_token_group_quant_8bitcCdSr,r"routput_qoutput_srXrHrrrcr"r"r#rrz)sgl_kernel::sgl_per_token_group_quant_fp8cCr r,r"r r"r"r#rrz#sgl_kernel::sgl_per_token_quant_fp8cCr r,r")rrrr"r"r#rs)r_FFFFN)r_FFFFNN)F)rm)NNFrp)Nr-r-r-T)u functoolsrloggingrrtypingrrrrrrrtriton.languagelanguager6triton.tools.tensor_descriptorrsglang.srt.layersr sglang.srt.utilsr r r r rrrrsglang.srt.utils.custom_oprr_is_cuda_is_cpur sgl_kernelr&sglang.jit_kernel.per_tensor_quant_fp8rrrr ImportErrorrraiterrrrvllm._Cvllm getLogger__name__rr4r$float8_e4m3fnuzrr float8_e4m3fnr}r=rTensorr.jit constexprrRr^rr"r@rper_token_group_quant_fp8rrrrrrrrrrrrrrr'r)r7r8rSr]rhrlrsrrrrrtuplerautotune fp8_autotunerrrrrrrr library register_fakerr"r"r"r#sd   (               * 5 l N  ! / :  ,  - 8 Z 1 .    0  Z       L   ?   . 3 .   4 %  ,    &    V