o پi{@sddlmZmZddlZddlmZddlZddlm Z ddl m Z m Z e dejjdejjdiZeejdr@deejj<n eejdrLd eejj<e Ze Ze Zd Zer[d nd Zd ZddZddZddZejedeeed fddZ dKdej de!dee"deej ej ffddZ#ejdgeddLd e!d!e!fd"d#Z$d$ej d%ej d&ej d'ej dej f d(d)Z%ejd*gejjdejjdidd+ddd,d-d.d/d0d1Z&ejd*ggd2d3d+dd,d4d5e!d6e!d7e!d8e!d9e!d:ee'd;e!fdd*gejjdejjdidd+ddd,d.d?d@dAZ)ejd*gejjdejjdidd,d.dBdCdDZ* EdMd$ej dFej dGej d:e'dHe!dej f dIdJZ+dS)N)OptionalTupleN) is_fp8_fnuz)is_gfx95_supportedis_hipWARNINGTTL_DISABLE_FAST_MATHTL_ENABLE_FAST_MATHFbfloat16float8_e4m3fnuz float8_e4m3float32c CsBtd|}|d?d@}|d@}td|dt|dkddS) Nuint32iint32r)T reinterpretCast if_then_else)xbits_xexp_xman_bitsrc/home/ubuntu/.local/lib/python3.10/site-packages/sglang/srt/layers/attention/nsa/tilelang_kernel.pyfast_log2_ceils  "rcCs|dd>}td|S)Nrrr )rr)rrrrr fast_pow2&s  rcCstt||S)N)rr)amax fp8_max_invrrrfast_round_scale+sr") pass_configscstdtr dndtrdndd rdndd d tjd tjffd tjf fd tjtf ff f dd }|S)NMglg|gl@g|@rr XYSc stjttdd\}}tf}tf}tf }tf }tf } tf } tjddD]} t|||f|t||tj||ddtD] } t || d|| < rt || || <qk|| || <qktD]\} } t || | f|| | | | f<qtD]} || ||| |f<qt| | t| |||fqHWddS1swYdS)Nr'threadsr num_stagesdimg-C6?) rKernelceildiv alloc_sharedalloc_fragment Pipelinedcopy reduce_absmaxParallelmaxr"clamp)r(r)r*pid_mpid_nx_sharedx_local amax_locals_localy_localy_shared_ij r$Nblk_mfp8_maxr!fp8_min group_sizein_dtyper. out_dtype round_scale scale_dtyperract_quant_kernel_;s8"  "z+act_quant_kernel..act_quant_kernel_)rsymbolic _is_fp8_fnuz prim_funcTensorr2)rGrLrMrOrNrPrrFract_quant_kernel/s    &#rUr'r block_size scale_fmtreturncCs|sJd|d|dksJd|d|d}tr)tj|tjd}ntj|tjd}|jg|dd||Rdtji}t ||dud }|| d|| d|| d||||fS) av Quantizes the input tensor `x` using block-wise quantization. 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[)rN) is_contiguoussizerRtorch empty_liker float8_e4m3fn new_emptyr rUview)rrVrWrGyskernelrrr act_quantbs  ,(rf)out_idxr#hdcstdtdtdddtjdtjftfdtjftfdtjftfd tjftfd tjftfd df fd d }|S)Nbmnr'qq_skk_sorXc stt\}}}tft}t|||ddf|tt} t|||df| tjddD]} tft} t|||| df| tt} t|||| f| tft} s~t | dtj | || dddt D]\}}t | ||fd| || ||f<qtt}tj | |ddt D] }||| |9<qt|||||| fq=WddS1swYdS) Nrr%r-FT) transpose_A transpose_B clear_accumrr/)rr1r2r3FP8r6r4FP32r5fillgemmr8r9 reduce_sum)rorprqrrrsi_bi_mi1_nq_smemq_s_fragi2_nk_smemk_s_fraglogitsi_hi3_n logits_sumrkblk_n1blk_n2rvrjrirlrmrrfp8_index_kernel_s< "    & $"z+fp8_index_kernel..fp8_index_kernel_)rrQrSrTrwrx)rirjrvrrrrfp8_index_kernels(   ,rrorprqrrcCsFtrt|jd|jdd||||St|jd|jd||||S)a Perform index score using FP8 precision. Args: q (torch.Tensor): The Q tensor, must be contiguous. q_s (torch.Tensor): The scaling factor for Q (float), must be contiguous. k (torch.Tensor): The K tensor, must be contiguous. k_s (torch.Tensor): The scaling factor for K (e8m0 here), must be contiguous. fp8 q @ fp8 k -> fp32 logits relu(fp32 logits) * q_s (weights) -> fp32 logits fp32 logits -> fp32 logits_sum fp32 logits_sum * k_s (e8m0) -> fp32 index_score r%F)_is_hiprshape)rorprqrrrrr fp8_indexs" rrYr@r%)kv_groupsm_scale is_causalblock_Ir.r,cs|tj|ksJd||tj|ksJd||dks&Jd||dks0Jd dur?d||dd n d td td  td } | }  |||g} | ||g}  ||g} |g}d }d d| }ttj| d |kr dksJ|t||||| dkr| ddksJd| dnddkr ndtjdt| dt| dt||dt|f fdd }|S)N+haven't check padding correctness yet, dim=Tznon-casual is not supportedrCotherwise will load some index=0 thus causing wrong kv to be loaded??/ldG?batchseq_len seq_len_kvrr floatrr"head_kv should be a multiple of 64QKVIndicesOutputc stj  dE\}}}tg}tg}tg} tg} tg} tgd} tg} tg}tg}tg}tg}tg}tg}tg}t| dt|dt|d||}}dkr|n|}|}|}| dkrdn|d}|}t|||||df|t|||||df|tj dD]'}tD]}||||||fdk| |<qtD]\}}||||||||f||f| ||f<qtD]\}}||||||||f||f| ||f<qtD]\}}t| |dt |j  |||f<q@tj || |dtj j d tj || |dtj j d t||tj||dd d tD]}t|||| ||<qtD]\}}t|||f || |||f<qtj||dd tD]}||||||||<qtD]\}}| ||f||| ||f<qt||tj || | tj j d qtD]\}}| ||f||<q tD]}t|||| ||<q t| | t| |||||ddfWddS1sUwYdS)Nr+boolrrrr-TrupolicyFr0clearr/r)rr1r3r4ryr6r5r8rinfinityr[rzGemmWarpPolicyFullCol reduce_maxexp2r{log2) rrrrbxbybzQ_shared Q_tail_shared KV_shared K_tail_sharedO_sharedmaskacc_oacc_sS_sharedsumexpsumexp_ialpham_im_i_prevb_ig_is_iq_imax_kv_iH0H1i_ibi_id_ih_iBIDD_tail H_per_blockNI REPLICATE_H accum_dtyperr[rr.padded_Hrrr,rrmains       "" "   " "$z,sparse_attention_fwd_kernel_v1..main tilelangmathnext_power_of_2rrQr9cdivrSrT) num_headsr0tail_dimtopkrrrrr.r,rhead_kvq_shapekv_shapeo_shape indices_shape indices_dtypeHrrrrsparse_attention_fwd_kernel_v1sh            ,dr) z-O3z-Wno-deprecated-declarationsz-U__CUDA_NO_HALF_OPERATORS__z-U__CUDA_NO_HALF_CONVERSIONS__z-U__CUDA_NO_HALF2_OPERATORS__z"-U__CUDA_NO_BFLOAT16_CONVERSIONS__z--expt-relaxed-constexprz--expt-extended-lambdaz,--ptxas-options=-v,--register-usage-level=10z-DNDEBUG)rh compile_flags)rrrrr0rrrrrcs|tj|ksJd||tj|ksJd|||dks(Jd dur7d||dd n d d tdtd  td } |||g}||||g}  ||g}  ||g} d d d |} ttj|d | kr|dksJ|t||ddksJd|||dkr|ddksJd|dnddkr ndtjdt|dt| dt| dt| f fdd } | S)Nrrrrrrrr num_pagesrr rrrr%zNI should be a multiple of 2rrrrrrc9 stj d d\}}}tdg}tdg}tg} tdg} tdg} tdg} tdg} tg}tg}|}|}tjgddd}tjgddd}tdg}tdg}tg}tg}tg}tg}tg}tjgdd}tg}tg}tg}tdg }tdg } tjdd}!tjd d}"tjd d}#tjd d}$tjd d}%tjd d}&tjd d}'tjd d}(tjd d})tjd d}*tjd d}+||},}-dkr-|n|}.|- dkr;d n|d }/|/}0t}1t||,|.|/|0d df|t||,|.|/|0df|t||,|.|/|0d f| t|!|1d kr)t ddt |d t |dt |d t |!d t t dD]}2t |"d |2d@t|(t |(d tD]\}3}4t||4d t|j ||3|4f<qtj|| |dddtj|| |dddtj| ||dddtd |2d krt|'t |'|2dd@dAt||tj||dddtD]}3t||3||3 ||3<q2tD]\}3}4t||3|4f ||3 ||3|4f<qKtj||ddtD]}3||3||3||3||3<qstdD]\}3}5||3|5f||39<qt||t||t|| |t|&t|$d t |#d |2d@t|(t |(dtD]\}3}4t||4d t|j ||3|4f<qtj|| |dddtj|| |dddtj| ||dddtd t|'t |'|2ddd@dAt||tj||dddtD]}3t||3||3 ||3<q;tD]\}3}4t||3|4f ||3 ||3|4f<qTtj||ddtD]}3||3||3||3||3<q|tdD]\}3}5||3|5f||39<qt||t||t|| |t|&t|%d qtD] }3||3||3<qt|+tdD]\}3}5||3|5f||3<qtD]}3t||3||3 ||3<qt||t|||,|.|/|0d dfn}|1d kr|1d krt |d t t dD]}2t|&t |&|2dd@t|)t |)d tdD]\}3}5||3|5f||39<qft|| |t|$d t|'t|&t |&|2ddd@t|)t |)dtdD]\}3}5||3|5f||39<qt|| |t|%d |2t ddkrt|'qBt |+d tdD]\}3}5||3|5f||3<qt||t|||,|.|/|0dfn|1d krt dd d |d <t t dD]}2t |$d |2d@dAt|*t |*d t dD]}6||,|.|-|2d|6d|1d df| d <| d d k||6d|1d d<||6d|1d dr| d |d <tdddzt dD]l}7tdD]c}8||,|d |-d |7|1d dd|8f| |6d|1d dd |7|1d dd|8f<||,|d |-dd |7|1d dd|8f| |6d|1d dd |7|1d dd|8f<qqWd n 1swYtddd9tdD]+}8||,|d |-|1d dd|8f||6d|1d d|1d dd|8f<q'Wd n 1s^wYqRt|"d t |%d |2d@dAt|*t |*dt dD]}6||,|.|-|2dd|6d|1d df| d <| d d k||6d|1d d<||6d|1d dr| d |d <tdddzt dD]l}7tdD]c}8||,|d |-d |7|1d dd|8f| |6d|1d dd |7|1d dd|8f<||,|d |-dd |7|1d dd|8f| |6d|1d dd |7|1d dd|8f<qܐqWd n 1sMwYtddd9tdD]+}8||,|d |-|1d dd|8f||6d|1d d|1d dd|8f<q_Wd n 1swYqt|#d q3Wd d SWd d SWd d SWd d S1swYd S)z Q: [b, qo_len, H, D + D_tail] (bfloat16) KV: [b, num_pages, kv_group, D + D_tail] (bfloat16) Indices: [b, qo_len, kv_group, topk] (int32) rr+r%rshared)scoper) arrive_countr'rrrNrTrY)ruwg_waitFrr/Prgrdefault async_scope)rr1r3r4 alloc_local alloc_barrierget_thread_bindingr6barrier_arrive set_max_nregry barrier_waitserialr2r8rrr[rz wait_wgmmarrr{rattr vectorizedcp_async_barrier_noinc)9rrrrrrr Q_shared_l Q_shared_rr KV_shared_0_l KV_shared_0_r KV_shared_1_l KV_shared_1_rK_tail_shared_0K_tail_shared_1 O_shared_l O_shared_r is_kv_valid_0 is_kv_valid_1acc_o_lacc_o_rrrrsum_exp_sharedr alpha_shared alpha_localrr indices_local indices_tmpbar_q bar_k_0_ready bar_k_1_ready bar_k_0_free bar_k_1_freebar_sScale_and_sS_readybar_sScale_and_sS_free bar_0_128 bar_1_128 bar_2_128 bar_finalrrrrrtxrrrrruvrrrrrrrrr[rrqo_lenrr,rrrsl             "$$                "           "     " (           (    &   "&   *   "& TSX$z,sparse_attention_fwd_kernel_v2..mainr)rr0rrrrrrrrrrrrrr"rsparse_attention_fwd_kernel_v2|sl            *Dr$)rrrrr,cs||dksJd|dksJ||dksJ dur%d||dd n d d} td td } ||} ttj| d } | d krJ| d nddkrR| nd |||||| |||g} | | |||g}| ||g}| ||g}| |g}tj}tjtjtj d t | dt |dt ||dt |dt |f  f dd }|S)zy grid: (seq_len * REPLICATE_H, top_k_blocks). Each block does one topk block, writes partial_o, partial_lse. Tznon-causal is not supportedrrNrrrrrrrrrr Partial_O Partial_Lsec stj d\}}tg}tg}tg} tg} tgtj} tg} tg} tg}tg}tg}t| dd\}}dkrl|n|}|}|}dkrzdn|d}|}t|||||df|t|||||df|tD]}||||||fdk| |<qtD]\}}||||||||f||f| ||f<qtD]\}}||||||||f||f| ||f<qtD]\}}t| |dt | j  | ||f<qtj || | dtj j dtj || | dtj j dtj| |dddtD] }t||d ||<qDtD]\}}t| ||f || | ||f<qXtj| |dd t| |tj || | tj j d tD]\}}| ||ft||d kd ||| ||f<qtD]}t||d kd t|||| ||<qt| ||||||ddft|||||||fWddS1swYdS)Nr+r)rrrrTrrrr/rgr)rr1r3r4rryr6r8rrr[rzrrrr9rr{r)rrrr&r'rrrrrrrrrrrrrrr topk_block_irrrrrr rrrrrrrr[rrr,rrr@s   ""   "$z+sparse_mla_fwd_decode_partial..main) rdynamicr9rrrrr r rSrT)headsr0rrrrrrr,rrrrrrrpartial_o_shapepartial_lse_shaperrrr)rsparse_mla_fwd_decode_partialsL          $Zr.)rr,c s||dks Jdd}td||||||g}||g}||g} tj} tjtjdt|| dt|dt| | ffdd } | S) z} grid: (seq_len * REPLICATE_H). batch=1, kv_group=1. Each block does one tile of heads (e.g. 4 or 8 for decode). rz head_per_block must divide headsrrr&r'rc srtjd$\}tg}tg}tg}tg}tg}d} dkr:|n|} dkrDdn|} | } tD]} t|| | | | | f|| ddfqSt|dtD]} tD]}t|||| |f||<q|qut|dtD]} tD]}||t || |f||||<qqtD]"} tD]}t || |f||t ||||| f<qqt|dtD]+} tD]"\}}|||f||| f|| | | | ||f |||f<qqt||| | | | ddfWddS1s2wYdS)Nr+rrr) rr1r3r4rr6ryr8r9rrastype)r&r'rr shared_lselse_maxlse_sumscalerrrrrrqrrrrrrr0rr,rrrsR*      "$z+sparse_mla_fwd_decode_combine..main)rr*r r rSrT) r+r0rhead_per_blockrr,rr,r-rr[rrr4rsparse_mla_fwd_decode_combines(      .r6rnkvindicesd_vc Cs*|dkr|dkr|dksJ|jd}|jd}||}|jd}|dks-Jtr}trp|jddkret|||||ddd } t|||d ddd } | |d|d|d\} } | | | } | St|||||dd }nt|||||d ddd}n t|||||d}||d|d|dS)Nrrr%rYirrr)rrr,rg)r5rr,)rr.r&r')rrr.r,)r) r0rr_is_gfx95_supportedr.r6 unsqueezerr$)ror7r8rr9rr0rrkernel_partialkernel_combine partial_o partial_lseoutrerrrtilelang_sparse_fwdsV(        rA)r'N)T)rn),typingrrrtilelang.languagelanguagerr^)sglang.srt.layers.quantization.fp8_kernelrsglang.srt.utilsrr set_log_level PassConfigKeyTL_DISABLE_WARP_SPECIALIZEDTL_DISABLE_TMA_LOWERr#hasattrrr rr:rRBF16rwrxrrr"jitrUrTintstrrfrrrrr$r.r6rArrrrs         3 8  #   T