o پiD3@sddlZddlZddlZddlZddlmZmZmZmZm Z ddl Z ddl Z ddl m ZddlmZmZeZerUz ddlmZdZWneyTddlmZdZYnweeZe jdejd ejfd d Ze jdfd d Ze jd ejfddZ de j!fde j"de#de$de j%de e j"e j"ff ddZ&de j!dfde j"de#de$de j%dee'f ddZ(e jdejdejdejdejfdd Z)ej*d!e#d"e#d#e#d$e#deee#eff d%d&Z+e j,fd'e j"d(e j"d)e j"d*e j"d+ee#d,e j%de j"fd-d.Z-dS)/N)AnyDictListOptionalTuple)get_device_nameis_cuda)sgl_per_token_group_quant_8bitT)sgl_per_token_group_quant_int8FCAL_SUMBLOCKc Cstd} td|} | |k} tj|| || | ddtj} ttt| d} | d}| d| }tj j j |tj }|rYtj| dd}t|| ||jjtj|| || || dt|| ||jjdS)Nrmaskother绽|=axisr)tl program_idarangeloadtofloat32maximummaxabsextracuda libdeviceroundint8sumstoredtype element_ty)x_ptrxq_ptr scale_ptr x_sum_ptrstride_x stride_xqNr r row_idcolsrxabsmaxscale_xx_qx_sumr6^/home/ubuntu/.local/lib/python3.10/site-packages/sglang/srt/layers/quantization/int8_kernel.py_per_token_quant_int8s $ r8c Cs||jd}|jd}tj||jtjd}tj|jddd|j|d}|r:tj|jdd|j|jd}nd}t |}t t |ddd} | sQJt |f|||||d|d|||| dd |rp|||fS||fS) Ndevicer&)r<)r,r-r.r r num_warps num_stages)numelshapetorch empty_liker;r#emptyr&tritonnext_power_of_2minr is_contiguousr8stride) r1 scale_dtypecal_sumMr.r4scalesr5r r@r6r6r7per_token_quant_int8;s4      rPc Cstd} || |7}|| |7}|| 7}td|} | |k} tj|| | ddtj} ttt| |} | |}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 int8 values. rr rrN) rrrrrrrrrclampr&r'r%)y_ptry_q_ptry_s_ptry_strider.epsint8_minint8_maxr g_idr0ry_absmaxy_sy_qr6r6r7_per_token_group_quant_int8\s    r^rr1 group_sizerVr&returnc Cs|jd|dks Jd|sJdt|}|j}|j}tj||j|d}||}|} tj |jdd|jd|f|jtj d} t | } tt| ddd } d} t |f||| || |||| | | d || fS) aWFunction to perform per-token-group quantization on an input tensor `x`. It converts the tensor values into signed int8 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. Note that only `torch.int8` is supported for now. Returns: Tuple[torch.Tensor, torch.Tensor]: The quantized tensor and the scaling factor for quantization. r9r=the last dimension of `x` cannot be divisible by `group_size``x` is not contiguousr:Nr=r<r>)rWrXr r@rA)rCrJrDiinforrIrEr;rBrFrrGrHr^)r1r_rVr&rcrXrWr4rNr.x_sr r@rAr6r6r7per_token_group_quant_int8s@   re enable_v2c Cs|jd|dks Jd|sJdt|}|j}|j}tj||j|d}tj|jdd|jd|f|jtj d} t rRt ||| |||||d|| fS|rVJt ||| |||||| fS)Nr9rrarbr:)rf) rCrJrDrcrrIrEr;rFr%enable_sgl_per_token_group_quant_8bitr r ) r1r_rVr&rfrcrXrWr4rdr6r6r7!sglang_per_token_group_quant_int8s* rh BLOCK_SIZE_M BLOCK_SIZE_N BLOCK_SIZE_K GROUP_SIZE_Mc6Cstjdd}t||}t||}||}||}||}t|||}|||}|||} ||td||}!| |td||}"td|}#||!dddf| |#dddf| }$||#dddf| |"dddf| }%||!|}&|"|}'||'|}(tj||ftjd})tdt||D]l}*tj|$|#dddf||*|kdd}+tj|%|#dddf||*|kdd},|*|}-|-| }.t|&|.|}/t|(|.|}0|)t |+|, tj|/dddf|0dddf7})|$|| 7}$|%|| 7}%q|j j tj kr|) tj }1n|j j tjkr)|) tj}1n|) tj}1||td|}2| |td|}3|||2dddf||3dddf}4|2dddf|k|3dddf|k@}5tj|4|1|5ddS)zTriton-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`. rrNr&r rr)rrcdivrIrzerosrrangerdotrr&r'bfloat16float16r%)6ABCAsBsrNr.Kgroup_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_nrirjrkrlpid 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 accumulatorkabk_startoffs_ksa_sb_scoffs_cmoffs_cnc_ptrsc_maskr6r6r7_w8a8_block_int8_matmulsL %     ,,  ((8  ,(rr.ryblock_nblock_kc Cstdd}d|d|d|d|d|d }tjtjtjtd |}tj|rUt |}t d |d d t |DWd S1sPwYt d|d S)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.  _zN=z,K=z ,device_name=z,dtype=int8_w8a8,block_shape=[z, z].jsonconfigsz7Using configuration from %s for W8A8 Block INT8 kernel.cSsi|] \}}t||qSr6)int).0keyvalr6r6r7 Usz/get_w8a8_block_int8_configs..NzjUsing default W8A8 Block INT8 kernel config. Performance might be sub-optimal! Config file not found at %s)rreplaceospathjoindirnamerealpath__file__existsopenloggerinfojsonritemswarning)r.ryrr device_namejson_file_nameconfig_file_pathfr6r6r7get_w8a8_block_int8_configs9s$$    rrtrurwrx block_size output_dtypecst|dksJ|d|d}}|jd|jdksJ|jdd|jddkr/|s1Jt|jd||jdksAJ||jd|jdkrX|rX|jdksZJ|j\}t||jdkslJt|||jdksyJ|jddf} |j| |d} t||d|d} | r| t | fddd } n d |d|dd d d d} fdd} t | ||| |||||| d| d| d| d| d| d| d| d| d| dfi| | 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. rr<r9Nrmcs t|S)N)r)r1)rNr6r7s z(w8a8_block_int8_matmul..)r@ )rirjrkrlr@rAcs"t|dt|dfS)Nrirj)rGrn)METArNr.r6r7gridsz$w8a8_block_int8_matmul..gridr?) lenrCrJrGrnrBndim new_emptyrrIkeysrrK)rtrurwrxrrrrryC_shapervrconfigrr6rr7w8a8_block_int8_matmulcs`(    r). functoolsrloggingrtypingrrrrrrDrGtriton.languagelanguagersglang.srt.utilsrr_is_cuda sgl_kernelr rg ImportErrorr getLogger__name__rjit constexprr8rrPr^r#Tensorrfloatr&reboolrhr lru_cacherrsrr6r6r6r7s      ! , > # T /