o پiH%@s ddlmZmZddlZddlmZddlmZdejdejdejdej fd d Z d}d d Z d dZ d}ddZ dejdejdejdejdejdejddfddZ d}dejdejdejdejdejdeejdejfddZ d}dejdejdeejdejfd d!Z " "  d~d#ejd$ejd%ejd&ed'ed(ed)ed*ed+ed,eejd-eeddfd.d/ZeZeZd#ejd$ejd%ejd0eddf d1d2Zd#ejd$ejd%ejddfd3d4Zd5ejd6ejd7ejd8ejd9ejd:ejdejfd;d<Zd#ejd=ejdeejejffd>d?Z d}d@ejdAejdBejdCejdDejdEejdFeejdejfdGdHZ d}d@ejdAejdIejdJejdBejdCejdFeejdejfdKdLZejfdMejdNejd:ejdejfdOdPZdQdRZ dSejd=ejdTejfdUdVZ!dSejd=ejdTejfdWdXZ" d}dSejd=ejdYejdZejd[ed\eejde#ejejffd]d^Z$ _ " " "dd5ejd`eejdaejdbejdceejddeejdeeejdfeejdgejdhediedjedkedledmednedoedejf$dpdqZ%d5ejdaejdrejdsejdtejduedvedejfdwdxZ&dyejdzejdveddfd{d|Z'dS))OptionalTupleN) ScalarType)_get_cache_bufqweightscalesqzerosreturncCstjjj|||SN)torchops sgl_kernelawq_dequantizedefault)rrrrC/home/ubuntu/.local/lib/python3.10/site-packages/sgl_kernel/gemm.pyrsrcCtjjj||||||Sr )r r r int8_scaled_mmrmat_amat_bscales_ascales_b out_dtypebiasrrrr rcCstjjj|||||Sr )r r r fp8_blockwise_scaled_mmr)rrrrrrrrrs rcCrr )r r r fp8_scaled_mmrrrrrr#rrworkspace_bufferABDA_scaleB_scalec Cs*tj}tjjj|||||||dSr )r cudacurrent_blas_handler r bmm_fp8r)rrr r!r"r# cublas_handlerrr_bmm_fp8_internal.s  r(dtypeoutcCsV|durtj|jd|jd|jdf|j|d}tdd|j}t|||||||S)Nrdevicer)bmm_fp8_workspacei)r emptyshaper.rr()rr r"r#r)r*rrrrr&Bsr&rroutputcCsD|durtj|jd|jdf|j|jd}tjjj||||S)Nrr+r-) r r0r1r.r)r r dsv3_fused_a_gemmr)rrr2rrrr3Usr3Finputoutput_qoutput_s group_sizeepsfp8_minfp8_max scale_ue8m0fuse_silu_and_mulmasked_m enable_v2c Cs| durddlm} | d} | r!tjjj|||||||||| S|r'Jd| dus/Jdtjjj||||||||dS)Nr)get_bool_env_var$SGLANG_PER_TOKEN_GROUP_QUANT_8BIT_V2z!only v2 support fuse_silu_and_mulzonly v2 support masked_m)sglang.srt.utilsr?r r r !sgl_per_token_group_quant_8bit_v2rsgl_per_token_group_quant_8bit) r4r5r6r7r8r9r:r;r<r=r>r?rrrrCds*    rC is_staticcCstjjj||||dSr )r r r sgl_per_tensor_quant_fp8r)r4r5r6rDrrrrEs rEcCstjjj|||dSr )r r r sgl_per_token_quant_fp8r)r4r5r6rrrrFsrFab block_scale_a block_scale_balpharc Csb|jdkr |jdks J|jd|jd}}tj||f||jd}tjjj|||||||S)Nr,rr)r.) ndimr1r r0r.r r cutlass_scaled_fp4_mmr) rGrHrIrJrKrmnr*rrrrNs  rNinput_global_scalec CsF|jdksJd|jd|jdkrdnd}|||jd}|j\}}d}|j}||dks8Jd|d|jtjtjfvsJJd|jdtj||d f|tj d }|d dd d }||} | d dd d } | | krtj || d f|tj d } n tj|| d f|tj d } tj j 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 a sizzled layout. r+z%input.ndim needs to be >= 1, but got .rz+last dim has to be multiple of 16, but got z-input.dtype needs to be fp16 or bf16 but got r,r-)rMreshaper1r.r)r float16bfloat16r0uint8zerosint32r r scaled_fp4_quantrview float8_e4m3fn) r4rQ other_dimsrOrP block_sizer.r2 rounded_mscale_n rounded_n output_scalerrrr]s:    r]in_featskernelwscalesascalesw_szsa_ssums out_featsc CL|durtj|jd|jdf|jtjd}tjjj||||||||SNrr-) r r0r1r.rXr r qserve_w4a8_per_chn_gemmr)rfrgrhrirjrkrlrrrro  ror[ scales_i8c Crmrn) r r0r1r.rXr r qserve_w4a8_per_group_gemmr)rfrgr[rqrhrirlrrrrrrprr hidden_statesrouter_weightscCs6tj|jd|jd|j|d}tjj||||Srn)r r0r1r.r r dsv3_router_gemm)rsrtrr2rrrrusrucCs,tj||j|jd}tjjj||||S)Nr-)r r0r.r)r r shuffle_rowsr) input_tensor dst2src_mapoutput_tensor_shape output_tensorrrrrv.srvrwmaskcCs$|j}|j\}}}d}||dksJd|d||}|ddd} | d} |ddd} tj|||d |tjd } tj|| | |tjd } tjjjj | |||d | || | | |||||d d | d d d} | tj  || d| dddd} | ddd dd d} | | fS)a Quantize input tensor to FP4 and return quantized tensor and scale, for grouped gemm inputs (e.g., grouped_gemm_nt_masked for flashinfer). Args: input: The input tensor to be quantized to FP4, with shape (l, m, k) l is number of groups, m is number of tokens per group, k is number of features. input_global_scale: A scalar scaling factor for the entire tensor, with shape (l,). Outputs: output: The quantized tensor in FP4, with shape (m, k // 2, l) but the physical layout is (l, m, k // 2). `// 2` is because two fp4 values are packed into an uint8. output_scales: The blockscale tensor in FP8-E4M3, with shape (32, 4, rm, 4, rk, l) but the physical layout is (l, rm, rk, 32, 4, 4). Note: For the shape of output_scales, `32 * 4 * rm` is a padded m to nearest multiple of 128. `4 * rk` is a padded `k // 16` to nearest multiple of 4. These layout constants are required by the NVIDIA Blackwell MMA operations. rTr"k must be multiple of 16, but got rRrVrUr,r-Fuse_silu_and_mulr+  r.r1r r0rZr\r r %silu_and_mul_scaled_fp4_experts_quantrr^permuter_)rwrQr{r.lrOk sf_vec_sizescale_kpadded_kpadded_k_int32padded_mr2 output_scalesrrrscaled_fp4_grouped_quant8s4    rcCs,|j}|j\}}}|d}d}||dksJd|d||} | ddd} | d} |dd d } tj|||d|tjd } tj|| | |tjd }tjjjj | |||d| || | | |||||d d | d dd} | tj  || d | dddd}| ddd ddd}| |fS)aQ Quantize input tensor to FP4 and return quantized tensor and scale, for grouped gemm inputs (e.g., grouped_gemm_nt_masked for flashinfer). Args: input: The input tensor to be quantized to FP4, with shape (l, m, k * 2) l is number of groups, m is number of tokens per group, k is number of features. input_global_scale: A scalar scaling factor for the entire tensor, with shape (l,). mask: The mask tensor, with shape (l,) Outputs: output: The quantized tensor in FP4, with shape (m, k // 2, l) but the physical layout is (l, m, k // 2). `// 2` is because two fp4 values are packed into an uint8. output_scales: The blockscale tensor in FP8-E4M3, with shape (32, 4, rm, 4, rk, l) but the physical layout is (l, rm, rk, 32, 4, 4). Note: For the shape of output_scales, `32 * 4 * rm` is a padded m to nearest multiple of 128. `4 * rk` is a padded `k // 16` to nearest multiple of 4. These layout constants are required by the NVIDIA Blackwell MMA operations. r,rTrr|rRr}rVr~rUr-Trr+rrr)rwrQr{r.rrOk_by_2rrrrrrr2rrrr%silu_and_mul_scaled_fp4_grouped_quantss6    rexpert_offsetsblockscale_offsetstopk expert_mapcCs&|jdksJd|jd|dur#|j\}}|||f}t|||}|j\} }ddl} t| jdd} | | |ksFJd| d | d |d } | d d } tj| |d|j tj d}| | krotj | || tj |j d}n tj| || tj |j d}tj jj|||||||tj}||fS)a Quantize input tensor to FP4 and return quantized tensor and scale, for packed MoE Inputs. Args: input: The input tensor to be quantized to FP4 expert_map: The expert map tensor 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 r,z%input.ndim needs to be == 2, but got rRNrMODELOPT_MAX_TOKENS_PER_EXPERTiz2m_numtopk must be less than MAX_TOKENS_PER_EXPERT(z,) for cutlass_moe_fp4, observed m_numtopk = z7. Use MODELOPT_MAX_TOKENS_PER_EXPERT to set this value.rTr}rVr-rL)rMr1rvosintenvirongetr r0r.rZr[r\r r scaled_fp4_experts_quantrr^r_)rwrQrrrrrOrry m_numtopkrMAX_TOKENS_PER_EXPERTscales_krr2rrrrrsZ         rTc b_q_weightb_scales global_scaleb_zerosg_idxperm workspaceb_q_typesize_msize_nsize_k is_k_fulluse_atomic_adduse_fp32_reduce is_zp_floatcCs0tjj|||||||||| j| | | | |||Sr )r r r gptq_marlin_gemmid)rGrrrrrrrrrrrrrrrrrrrrs&r b_gptq_qzeros b_gptq_scalesb_g_idx use_shufflebitc Cstjj|||||||Sr )r r r gptq_gemm)rGrrrrrrrrrr%s rq_weightq_permcCstjjj|||dSr )r r r gptq_shuffle)rrrrrrr3srr )FFNN)TFFF)(typingrrr sgl_kernel.scalar_typersgl_kernel.utilsrTensor ByteTensorrrrrr(r)r&r3rfloatboolrCsgl_per_token_group_quant_fp8sgl_per_token_group_quant_int8rErFrNr]rorrrYrurvrrtuplerrrrrrrrs            (     B     ; C [      ( "