
    VhY                        d Z ddlmZ ddlZddlZddlmZ ddlmZm	Z	 ddl	m
Z
mZ ddlmZmZmZ dd	lmZmZ dd
lmZmZmZ ddlmZmZmZmZmZmZmZmZm Z  ejB                  jD                  Z"ejB                  jF                  Z#ed        Z$ ede$dez   ez   ez   ez   ez   ez         Z%de&de&de&de&fdZ'de(e&e&e&f   fdZ)d Z*y)zXTriton Implementation of the flex_attention Kernel for short query length (FlexDecoding)    )AnyN)V   )configir)FixedLayoutFlexibleLayout)emptyempty_strided	lowerings)is_power_of_2next_power_of_2)autotune_select_algorithmSymbolicGridFnTritonTemplate   )	compute_forward_block_mncompute_forward_innercompute_next_offset_funccreate_indices_fake create_num_blocks_fake_generatorget_bounded_indices_funcload_checked_2dload_checked_blockmaybe_realizec                     | |z  |d   dfS )ad  How is this kernel parallelized?
    We create a grid of (batch_size * kv_heads, SPLIT_KV, 1)
    Each block is responsible for iterating over blocks of keys and values calculating
    the local output for their tile of keys and values over all full length of query.
    groups of SPLIT_KV blocks then combine their output to produce the final result.
    SPLIT_KVr    )
batch_sizekv_headsgqa_group_sizen_keysd_modelmetas         T/home/dcms/DCMS/lib/python3.12/site-packages/torch/_inductor/kernel/flex_decoding.pyflex_decoding_gridr&   !   s     !4
#3Q77    flex_decodinga)  
    {{def_kernel("Q", "K", "V", "M", "L", "KV_NUM_BLKS", "KV_IDX", "FULL_KV_NUM_BLKS", "FULL_KV_IDX")}}
    # Sub notation for this kernel:
    # Q: Query, K: Key, V: Value
    # reduction buffers: M rowmax across local KV split, L local sumexp across local KV split
    # M: Number of queries, N: Number of keys/values
    # QK_HEAD_DIM: The dimension of the query and key embeddings
    # V_HEAD_DIM: The dimension of the value embeddings
    # BLOCK_M, QK_HEAD_DIM: M, and D dimemsion are always assigned to the same block
    # z: Batch size, h: Number of heads, m: Number of queries per head, k: Number of keys per head t: Number of kv splits
    # (Modifiable) Config options:
    # SPLIT_KV: number of blocks K & V are split into
    # TILE_KV: length of each local KV split
    # BLOCK_M: block size that Q is padded along seqlen dim.
    # BLOCK_N: block size of K & V along N dimension.
    # GQA_SHARED_HEADS: number of query heads sharing one kv head in GQA setups.
    #
    # change of base out of the loop
    # ROWS_GUARANTEED_SAFE: Is it guaranteed that at least one value in each row
    # is not masked out? If so, we can skip an extra safety check
    # SAFE_M_BOUNDARY: Is Q seqlen a multiple of BLOCK_M? If so, we can skip an extra boundary check for loading query.
    # SAFE_N_BOUNDARY: Is KV seqlen a multiple of BLOCK_N? If so, we can skip an extra boundary check for loading key/value.

    # PRESCALE_QK: Whether to pre-scale QK by 1/sqrt(d) and change of base.
    #
    # SPARSE_KV_BLOCK_SIZE: sparse mask block size along KV seqlen dim.
    # KV_NUM_BLKS: The number of KV blocks (that may or may not require masking) for each query.
    # KV_IDX: The indices of KV blocks (that may or may not require masking) for each query.
    #
    #
    # Output: ACC output accumulated across local KV split.

    tl.static_assert(SPARSE_KV_BLOCK_SIZE >= BLOCK_N and SPARSE_KV_BLOCK_SIZE % BLOCK_N == 0)

    # Define Q Strides
    stride_qz, stride_qh, stride_qg, stride_qm, stride_qk = {{stride("Q")}}
    stride_kz, stride_kh, stride_kn, stride_kk = {{stride("K")}}
    stride_vz, stride_vh, stride_vn, stride_vk = {{stride("V")}}
    stride_mz, stride_mt, stride_mh, stride_mm = {{stride("M")}}
    stride_lz, stride_lt, stride_lh, stride_lm = {{stride("L")}}


    Z = {{size("Q", 0)}}
    ZKV = {{size("K", 0)}}
    HKV = {{size("Q", 1)}}
    G: tl.constexpr = GQA_SHARED_HEADS
    HQ = HKV * G
    Q_LEN = {{size("Q", 3)}}
    KV_LEN = {{size("K", 2)}}

    MATMUL_PRECISION = Q.dtype.element_ty

    # Make sure each split is a multiple of BLOCK_N
    TILE_KV_OG = tl.cdiv(KV_LEN, SPLIT_KV)
    TILE_KV = tl.cdiv(TILE_KV_OG, BLOCK_N) * BLOCK_N
    TILE_KV_MULTIPLE: tl.constexpr = (TILE_KV // BLOCK_N)

    off_z = tl.program_id(0) // HKV
    off_zkv = off_z % ZKV
    off_hkv = tl.program_id(0) % HKV
    off_t = tl.program_id(1)

    q_offset = off_z * stride_qz + off_hkv * stride_qh
    k_offset = off_zkv * stride_kz + off_hkv * stride_kh
    v_offset = off_zkv * stride_vz + off_hkv * stride_vh

    SPARSE_Z = {{size("KV_NUM_BLKS", 0)}}
    SPARSE_HQ = {{size("KV_NUM_BLKS", 1)}}

    sparse_idx_z = off_z % SPARSE_Z
    sparse_idx_h = off_hkv % SPARSE_HQ

    SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N)
    SPARSE_KV_BLOCK_CNT = tl.cdiv(KV_LEN, SPARSE_KV_BLOCK_SIZE)

    # initialize pointer to m and l
    m_i = tl.zeros([BLOCK_M], dtype=tl.float32) - float("inf")
    l_i = tl.zeros([BLOCK_M], dtype=tl.float32)
    acc = tl.zeros([BLOCK_M, V_HEAD_DIM_ROUNDED], dtype=tl.float32)

    # initialize offsets
    tl.device_assert(BLOCK_M % G == 0)
    BLOCK_M_PER_HQ: tl.constexpr = BLOCK_M // G
    off_g = tl.arange(0, G)                                                 # [G]
    offs_g = tl.ravel(tl.broadcast_to(off_g[:, None], [G, BLOCK_M_PER_HQ])) # [BLOCK_M]
    offs_hq = offs_g + off_hkv * G
    off_m = tl.arange(0, BLOCK_M_PER_HQ)                                    # [BLOCK_M_PER_HQ]
    offs_m = tl.ravel(tl.broadcast_to(off_m[None, :], [G, BLOCK_M_PER_HQ])) # [BLOCK_M]
    offs_d = tl.arange(0, QK_HEAD_DIM_ROUNDED)
    offs_vd = tl.arange(0, V_HEAD_DIM_ROUNDED)

    # Get HZ offsets for KV_NUM_BLKS and KV_IDX
    stride_block_z, stride_block_h, stride_block_row = {{stride("KV_NUM_BLKS")}}
    sparse_block_hz_offset = sparse_idx_z * stride_block_z + sparse_idx_h * stride_block_h
    stride_kv_z, stride_kv_h, stride_kv_row, stride_kv_col = {{stride("KV_IDX")}}
    sparse_idx_hz_offset = sparse_idx_z * stride_kv_z + sparse_idx_h * stride_kv_h

    # Calculate KV blocks that belong this CTA.
    block_n_start = off_t * TILE_KV_MULTIPLE                        # n_offset inside sparse block
    block_n_end = block_n_start + TILE_KV_MULTIPLE                  # end BLOCK_N

    q_range = stride_qg * off_g[:, None, None] + stride_qm * off_m[None, :, None] + stride_qk * offs_d[None, None, :]

    if not SAFE_M_BOUNDARY and not SAFE_HEAD_DIM:
        q = tl.load(Q + q_offset + q_range, mask=(offs_d[None, None, :] < QK_HEAD_DIM) & (off_m[None, :, None] < Q_LEN))
    elif SAFE_M_BOUNDARY and not SAFE_HEAD_DIM:
        q = tl.load(Q + q_offset + q_range, mask=offs_d[None, None, :] < QK_HEAD_DIM)
    elif not SAFE_M_BOUNDARY and SAFE_HEAD_DIM:
        q = tl.load(Q + q_offset + q_range, mask=off_m[None, :, None] < Q_LEN)
    else:
        q = tl.load(Q + q_offset + q_range)

    q = tl.reshape(q, [BLOCK_M, QK_HEAD_DIM_ROUNDED])


    # ~~~~~~~~~~~~~~ normal blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
    # Apply both score_mod and mask_mod

    # find first kv block we are loading and the number of blocks we are loading
    # Offset the kv_indices tensor by the correct batch and head
    kv_indices = KV_IDX + sparse_idx_hz_offset
    kv_num_blocks = tl.load(KV_NUM_BLKS + sparse_block_hz_offset)
    indices_idx = block_n_start // SPARSE_KV_MULTIPLE
    off_n_block_in_sparse = block_n_start % SPARSE_KV_MULTIPLE
    off_n = tl.load(kv_indices + indices_idx) * SPARSE_KV_BLOCK_SIZE + off_n_block_in_sparse * BLOCK_N
    # first kv block we're loading

    # last valid block according to sparse mask
    block_n_last_valid = tl.minimum(kv_num_blocks * SPARSE_KV_MULTIPLE, tl.maximum(tl.cdiv(KV_LEN, BLOCK_N), 1))

    K_block_ptr = tl.make_block_ptr(
        base=K + k_offset,
        shape=(QK_HEAD_DIM, KV_LEN),                # (d, N)
        strides=(stride_kk, stride_kn),
        offsets=(0, off_n),
        block_shape=(QK_HEAD_DIM_ROUNDED, BLOCK_N),
        order=(0, 1)
    )
    V_block_ptr = tl.make_block_ptr(
        base=V + v_offset,
        shape=(KV_LEN, V_HEAD_DIM),
        strides=(stride_vn, stride_vk),
        offsets=(off_n, 0),
        block_shape=(BLOCK_N, V_HEAD_DIM_ROUNDED),
        order=(1, 0)
    )
    offs_n = tl.arange(0, BLOCK_N) + off_n

    acc, l_i, m_i = forward_inner(
        {{gen_argdefs()}},
        q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN,
        # accumulatd values
        acc, l_i, m_i,
        #offsets
        off_z, offs_hq[:, None], offs_m[:, None], offs_n[None, :],
        #block sparse data
        kv_indices, kv_num_blocks,
        block_n_start, block_n_end if block_n_end <= block_n_last_valid else block_n_last_valid,
        MATMUL_PRECISION,
        IS_FULL_BLOCKS=False,
    )


    # ~~~~~~~~~~~~~~ "full" blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
    # We know these blocks are guaranteed to be "full", so we don't need to
    # apply mask_mod to them - only score_mod
    if HAS_FULL_BLOCKS:
        kv_indices = FULL_KV_IDX + sparse_idx_hz_offset
        kv_num_blocks = tl.load(FULL_KV_NUM_BLKS + sparse_block_hz_offset)
        # Assign full block in a reverse order for off_t. Prioritize the last CTA.
        block_n_start = (SPLIT_KV - off_t - 1) * TILE_KV_MULTIPLE
        block_n_end = block_n_start + TILE_KV_MULTIPLE
        indices_idx = block_n_start // SPARSE_KV_MULTIPLE
        off_n_block_in_sparse = block_n_start % SPARSE_KV_MULTIPLE
        off_n = tl.load(kv_indices + indices_idx) * SPARSE_KV_BLOCK_SIZE + off_n_block_in_sparse * BLOCK_N

        # last valid block according to sparse mask
        block_n_last_valid = tl.minimum(kv_num_blocks * SPARSE_KV_MULTIPLE, tl.maximum(tl.cdiv(KV_LEN, BLOCK_N), 1))

        K_block_ptr = tl.make_block_ptr(
            base=K + k_offset,
            shape=(QK_HEAD_DIM, KV_LEN),                # (d, N)
            strides=(stride_kk, stride_kn),
            offsets=(0, off_n),
            block_shape=(QK_HEAD_DIM_ROUNDED, BLOCK_N),
            order=(0, 1)
        )
        V_block_ptr = tl.make_block_ptr(
            base=V + v_offset,
            shape=(KV_LEN, V_HEAD_DIM),
            strides=(stride_vn, stride_vk),
            offsets=(off_n, 0),
            block_shape=(BLOCK_N, V_HEAD_DIM_ROUNDED),
            order=(1, 0)
        )
        offs_n = tl.arange(0, BLOCK_N) + off_n

        acc, l_i, m_i = forward_inner(
            {{gen_argdefs()}},
            q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN,
            # accumulatd values
            acc, l_i, m_i,
            #offsets
            off_z, offs_hq[:, None], offs_m[:, None], offs_n[None, :],
            #block sparse data
            kv_indices, kv_num_blocks,
            block_n_start, block_n_end if block_n_end <= block_n_last_valid else block_n_last_valid,
            MATMUL_PRECISION,
            IS_FULL_BLOCKS=True,
        )

    m_offset = off_t * stride_mt + off_z * stride_mz
    l_offset = off_t * stride_lt + off_z * stride_lz

    M_block_ptr = tl.make_block_ptr(
        base=M + m_offset,
        shape=(G, Q_LEN),                   # (G, M)
        strides=(stride_mh, stride_mm),
        offsets=(off_hkv*G, 0),
        block_shape=(G, BLOCK_M_PER_HQ),
        order=(1, 0)
    )
    L_block_ptr = tl.make_block_ptr(
        base=L + l_offset,
        shape=(G, Q_LEN),                   # (G, M)
        strides=(stride_lh, stride_lm),
        offsets=(off_hkv*G, 0),
        block_shape=(G, BLOCK_M_PER_HQ),
        order=(1, 0)
    )

    # Store output, logsumexp and rowmax for cross CTA reduction. (all in float32, even when input data are in fp16)
    m_i = m_i.reshape(G, BLOCK_M_PER_HQ)
    l_i = l_i.reshape(G, BLOCK_M_PER_HQ)
    if SAFE_M_BOUNDARY:
        tl.store(M_block_ptr, m_i)
        tl.store(L_block_ptr, l_i)
    else:
        tl.store(M_block_ptr, m_i, boundary_check=(1,))
        tl.store(L_block_ptr, l_i, boundary_check=(1,))

    # -- store output
    idx_z = off_z
    idx_t = off_t
    idx_hq = off_hkv*G + off_g[:, None, None]
    idx_m = off_m[None, :, None]
    idx_d = offs_vd[None, None, :]

    mask = (idx_m < Q_LEN) & (idx_d < V_HEAD_DIM)
    acc = acc.reshape(G, BLOCK_M_PER_HQ, V_HEAD_DIM)
    {{store_output(("idx_z", "idx_t", "idx_hq", "idx_m", "idx_d"), "acc", "mask")}}
 )namegridsourceBHMkreturnc                     t         j                  j                  d      j                  }t	        | |z  d      }t        |t        t        j                  f      sJ d       ||z  dz  }t	        |d      }|S )Ncudar   z!B and H must be concrete integersr   )	torchr1   get_device_propertiesmulti_processor_countmax
isinstanceintsympyInteger)r,   r-   r.   num_SMbhsplit_ks         r%   get_split_kr=   5  sj    ZZ--f5KKF	QUABb3./T1TT/lQG '1oGNr'   c                    | j                         }| j                         d   }t        j                  j	                         }d}|dk\  r6|dkD  r|t        j
                  k(  r|S t        j                  j                  yy|S )N)@   r   r   )	   r      )r@   r      )	get_dtypeget_sizer2   r1   get_device_capabilityfloat32versionhip)keydtypehead_dim
sm_versiondefault_configs        r%   _get_decoding_default_configrO   A  sp    MMOE||~b!H113JNVc>eu}}4!!==$r'   c                    @ ddl m} | \
  @}}}}}}}	}
}|\  }}}}}}}}}}}}}@j                         \  }}}}|j                         \  }}}}t        j                  j
                  j                  t        j                  ||      t        j                  |d      z        sJ d| d|        |}t        |      }|j                         D ci c]K  \  }}|t        |t        j                        r)t        j                  j
                  j                  |      n|M }}}|dz  dk7  s|dz  dk7  r|j                  dd       n|j                  dd	       ||z  }t        |      st!        d
      |j                  d|       |d u}|j                  d|       |s@fdt#        d      D        \  }}t%        @||||||g      \  @}}}}}}t%        |
      }
t%        |      }g }g } | j'                  t)        |             t*        j,                  r9| g dz  } t.        j0                  j2                  r| D !cg c]  }!|!d   |!d   df } }!|j                  d|       |j                  dt5        |||             |d   }"||"|||g}#|#d d }$t7        |$d t.        j8                  @j;                               }%t7        |$d t.        j8                  @j;                               }&t=        @j;                         t.        j8                  |#t?        j@                  |#            }' ||||t        j                  j
                         |j                  dtC        tE        t        j                  j
                  jG                  |t.        jH                  j*                  jJ                        |z        d             tL        jN                  jQ                  @      @@jS                         \  }(})}*}+|||||f},|(|)|z  |)|*|+f}-tU        tV        jX                     @|,|-      @t        j                  j
                  j[                  ||z  t        j\                  |d                |j                  d||z  |d   z  dk(         |j                  dd	       t        j                  j
                  j                  |      }|j_                         }.| D ]
  \  }/}0}1||/z  dk7  r|.j_                         }2ta        |2jc                               D ]O  }|je                  d      r|2jg                  |      }||2|dd  <   |je                  d      s?|2jg                  |       Q |2j                  d|/       |2j                  d|       |2j                  d|0       |2j                  d|1       ti        jj                  d*|@|||%|&||||g	|'||	g|%|&g@j                         d |2  @|||%|&||||g	ta        |
      z   ta        |      z   }3tm        |      tn        tm        |      tn        d!}4tq        d"||3|'|4#      }5tU        tV        jB                     |%dd	$      d   }6tU        tV        jr                     |6tu        d%             }7tU        tV        jv                     |%|6      }8tU        tV        jx                     |7d|8      }8tU        tV        jz                     |8      }9tU        tV        j|                     |&|9      }&tU        tV        j~                     |&d&      }:tU        tV        j                     |7d'      };tU        tV        jx                     |;d(|:      }:tU        tV        j                     |:      }<tU        tV        j                     |<tU        tV        j                     |6d'            }<tU        tV        j                     |9d      }=tU        tV        j|                     |5|=      }5tU        tV        j~                     |5d&      }>tU        tV        j                     |:d)      }?tU        tV        j                     |>|?      }>tU        t        j                     |>@j                               }>|>|<fS c c}}w c c}!w )+Nr   )set_head_dim_valuesz&Bq and Bkv must broadcastable. Got Bq=z	 and Bkv=rB   r   IS_DIVISIBLEFTzJNumber of shared query heads sharing the same KV head must be power of 2. GQA_SHARED_HEADSHAS_FULL_BLOCKSc              3   T   K   | ]  }t        d j                                ! yw)r   )deviceN)r
   
get_device).0_querys     r%   	<genexpr>z.create_flex_decoding_kernel.<locals>.<genexpr>  s)      /
45E!E,,.///
s   %(r   ))r@   r   r   )    r   rC   )rB   r   rC   SM_SCALEr   r?   )rK   rV   BLOCK_M)fallback   SAFE_M_BOUNDARYSAFE_N_BOUNDARYfwd_   bwd_BLOCK_NSPARSE_KV_BLOCK_SIZE	num_warps
num_stages)choicesinput_nodeslayout	subgraphsmutated_inputs
call_sizes)            r(   )input_gen_fns)dimkeepdiminf)axis)ru   g      ?rC   r   )Hflex_attentionrQ   rE   r   graphsizevarsevaluate_exprr8   Eqdictitemsr6   Symbolevaluate_static_shape
setdefaultr   
ValueErrorranger   appendrO   r   max_autotuner2   rH   rI   r=   r   rG   rW   r   r	   contiguous_stridesr5   r   	size_hint	_inductorunbacked_symint_fallbackr   ExternKernelrealize_input
get_strider   aten
as_strided	guard_leqr9   copylistkeys
startswithpopflex_decoding_templatemaybe_append_choicer   r   r   eqfloatsubwhereexp2mulsumsqueezelog2add	unsqueezedivprimsconvert_element_typerD   )AargskwargsrQ   rJ   value
block_maskscalekernel_optionsscore_mod_subgraphmask_mod_subgraphscore_mod_other_buffersmask_mod_other_buffersrY   kv_num_blocks
kv_indicesfull_kv_num_blocksfull_kv_indicesrg   BqHq	seq_len_qqk_head_dimBkvHkv
seq_len_kv
v_head_dimr,   kvgqa_shared_headshas_full_blocksrj   configscMAX_SPLIT_KVbuf_ACC_shapebuf_ML_shapebuf_Mbuf_L
layout_accstride_b	stride_hqstride_seq_len_qstride_qk_head_dimgqa_query_shapegqa_query_strideoriginal_kernel_optionsrf   rh   ri   cur_kernel_optionsinputs_for_flex_decodingrt   buf_ACCg_Mmasked_rowsadj_Malphag_Lmasked_rows_squeezed	logsumexpalpha_unseqoutputL_unseqrZ   sA                                                                   @r%   create_flex_decoding_kernelr   P  s   3 	  									 &+^^%5"BI{',~~'7$Cj*77))%((2s*;ehhsA>N*NO 
0IcUCO 	A.)N #((*	 Aq 	
a& 7711!4	N  3!zC/14!!.%8!!.$7 Sy)*X
 	
 02BC )4O/A/
9>q/
+O 		

	 ,,CD*+ABG*,GNN/45 
 	
 ==0781!adA8G8 j%0j+aj*IJ!*-L b)Z@M "%Lmm!	E mm!	E ))-8	J ZAQAQR
 GG$$..!!&!7!7!P!P /  '	' 	& OO))%0E@E@P@P@R=Hi)+= #/KHO$$ doo&uo?OPEGG$$emmN94M&N 
&
&.*C	CI
 /677++AABVW,113 +2 '
&J')Q.499;(--/0 	*A||F#&**1-,-"1QR5)||F#"&&q)	* 	%%i9%%&<>RS%%k9=%%lJ?22 	
"
 "! "5>~~''	
( !)	
%'
V 
	
 &
'	( %
&	' " ,J7+O<	M ( #G DHH
eD
9!
<C DGG$S5<-8Kdhhs+Edjj!+q%8Edii 'Edhhu-E
DHH
e!
,C$T\\2;AF
DJJ
 4c3
?C$))$S)I$((#Iy/FsPQ/RSIDNN+E15K!';7Gtxx q1F'Q/Gtxx 1Fu11265??;LMF 	 iB 9s   Aa7a)+__doc__typingr   r8   r2   torch._inductor.virtualizedr    r   r   r   r	   loweringr
   r   r   runtime.runtime_utilsr   r   select_algorithmr   r   r   ry   r   r   r   r   r   r   r   r   r   opsr   r   r&   r   r7   r=   tuplerO   r   r   r'   r%   <module>r      s   ^    )  , 6 6 B X X
 
 
 yy~~		 8 8 (		{x y|z {}| }~~ @ A@B CAE P	3 	3 	C 	C 	sC})= `r'   