o پi @sddlmZmZddlZddlZddlmZejdej dej dej dej fddZ dd ej d e d ee d eej ej ffddZdS))OptionalTupleN group_size round_scaleBLOCK_MBLOCK_Nc "Cstd} td} d} d} d| } | |}| |}|td|}|td|}||k}||k}|dddf|dddf@}||dddf||dddf}tj||ddtj}t|}tj|dd }t|d }|rt || }t |}t |}n|| }|dddf}||}t t|| | }||dddf||dddf}tj |||d | }|||||} |}!tj | ||!d dS) zs Triton kernel for activation quantization. Each block processes BLOCK_M rows and group_size columns. rg|g|@g?Ng)maskother)axisg-C6?)r )tl program_idarangeloadtofloat32absmaxmaximumlog2ceilexp2minimumstore)"X_ptrY_ptrS_ptrMNrrrrpid_mpid_nfp8_minfp8_max fp8_max_inv row_start col_startrowscolsrow_maskcol_maskr x_ptrsxx_absamaxlog_vallog_ceilscalescale_broadcastyy_ptrss_colss_ptrss_maskr7a/home/ubuntu/.local/lib/python3.10/site-packages/sglang/srt/layers/attention/nsa/triton_kernel.py_act_quant_kernel s>   (    (r9r+ block_size scale_fmtreturnc Cs|sJd|d|dksJd|d|d}|d|}|d}tj|tjd}|d|}|jg|dd||Rdtji}|d||} d } |} t || t ||f} |du} t | ||| |||| | | | rydnd d ||fS) a Quantizes the input tensor `x` using block-wise quantization with Triton. Args: x (torch.Tensor): The input tensor to be quantized. Must be contiguous and its last dimension size must be divisible by `block_size`. block_size (int, optional): The size of the blocks to be used for quantization. Default is 128. scale_fmt (Optional[str], optional): The format of the scale. Default is None. Returns: Tuple[torch.Tensor, torch.Tensor]: A tuple containing: - The quantized tensor with dtype `torch.float8_e4m3fn`. - A tensor of scaling factors with dtype `torch.float32`. zInput tensor must be contiguousrz@Last dimension size must be divisible by block_size (block_size=))dtypeNr@ )rrrr num_stages) is_contiguoussizeviewtorch empty_like float8_e4m3fn new_emptyrtritoncdivr9)r+r;r<rx_flatrr2y_flatss_flatrrgridrr7r7r8 act_quantVs8     ,  rR)r:N)typingrrrGrKtriton.languagelanguager jit constexprr9TensorintstrrRr7r7r7r8s0  M