Vh<dZddlZddlZddlZddlmZddlmZddlm Z m Z ddl m Z m Z mZddlZddlZddlmZddlmZdd lmZdd lmZdd lmZd d lmZd dlmZmZm Z m!Z!m"Z"m#Z#m$Z$m%Z%m&Z&m'Z'm(Z(m)Z)m*Z*d dl+m,Z,m-Z-m.Z.m/Z/m0Z0m1Z1m2Z2m3Z3m4Z4d dl5m6Z6m7Z7m8Z8m9Z9ejte;Zej~Z?dee@dee@dee@fdZAdee@dee@fdZBe8dZC dbdeDdejdejde eGe@de*f dZHdeGe e%fdZIdZJd eGe@fd!ZKeeGe ee efZLdeGe*d"ejjdeLfd#ZOdeGe*d$e)deLfd%ZPd&ZQd'ZRd(ZSd)ZTd*ZUd+ZVd,ZWe9d-eCeUeVzeQzeWzeSzeRz.ZXd/ZYejd0fd1ejd2fd3ejd4fd5ejd0fd6ejd2fd7ejd4fd8ejd0fd6ejd2fd9ejd4fd8i Z]ejd0fd1ejd2fd1ejd4fd:ejd0fd;ejd2fd7ejd4fd3ejd0fd;ejd2fd7ejd4fd3i Z^ejd0fdejd2fd>ejd4fd?ejd0fd>ejd2fd>ejd4fd@i Z_GdAdBe Z`dCe`deae@e@e@e@ffdDZbdCe`deae@e@e@e@ffdEZcdeae@e@e@e@ffdFZddeae@e@e@e@ffdGZedHZfdejfdIZhddJlimjZjd dKlkmlZldLZmdMZndNZodOZpdPZqdQereDe ffdRZse3ejzjjdSdTZudUZve9dVevdWeQzeRzeTz.ZwdXejjfdYZyedZ[Gd\d]Zzd^eLd_e@dezfd`Z{e3ejzjjdSdaZ|y)cz2Triton Implementation of the flex_attention KernelN)Sequence) dataclass)autoEnum)AnyOptionalUnion)V) OrderedSettree_map)int_oo) ValueRanges)config) BufferComputedBuffer ExternKernel FixedLayoutFlexibleLayoutget_fill_order InputBufferIRNodeMutationLayoutSHOULDREMOVEScatter StorageBoxSubgraph TensorBox) _fullcheck_and_broadcast_indicesempty empty_stridedexpandindex_output_size_and_inner_fn loweringsregister_loweringto_dtype)autotune_select_algorithmrealize_inputsSymbolicGridFnTritonTemplatesizes fill_orderreturnct|t|k(sJddgt|z}d}|D]}|||<|||z}|S)zTFrom a list of sizes and a fill order, construct the strides of the permuted tensor.z7Length of sizes must match the length of the fill orderr)len)r,r-stridescurrent_stridedims U/home/dcms/DCMS/lib/python3.12/site-packages/torch/_inductor/kernel/flex_attention.pyconstruct_stridesr6=sl u:Z (A (cCJGN%% %*$% Nsize orig_stridescvt|tjjj}t ||S)aThis is a mirror of the same function in aten/src/ATen/ExpandUtils.cpp Args: size: The size of the output tensor orig_strides: The strides of the input tensor Returns: List[int]: Dense non-overlapping strides that preserve the input tensor's layout permutation. The returned strides follow the same stride propagation rules as TensorIterator. This matches The behavior of empty_like() )rr graphsizevars shape_envr6)r8r9r-s r5infer_dense_stridesr>Ss- agg.>.>.H.HIJ T: ..r7c&|||d||zdfS)zHow is this kernel parallelized? We create a grid of (batch_size * num_heads, ceil_div(n_queries, query_block_size), 1) Each block is responsible for iterating over blocks of keys and values calculating the final attention output. BLOCK_Mr0) batch_sizeq_heads num_queriesd_modelmetacdivs r5flex_attention_gridrHbs! d9o . W0Da HHr7namedtypedevicec t|t|||r|ng|rtj|ng}t j |S)zBCreates a placeholder input buffers for producing subgraph_output.)rIlayout)rrrcontiguous_stridesrcreate)rIrJrKr8 input_buffers r5create_placeholderrQlsK   Db7;N - -d 3  L   L ))r7argsctd|S)zIAccepts a list of optional IRNodes and returns a list of realized IRNodescT|%t|tjs t|S|SN) isinstancesympySymbolr))xs r5zmaybe_realize..s-}Z5<<%@ 1 r7r rRs r5 maybe_realizer\s   r7ctjdk(s8tjjstjj ryy)Nhighestz'ieee'z'tf32')torchget_float32_matmul_precisionversionhipmtia is_availablerAr7r5get_float32_precisionres6 **, 9 ==   :: " " $r7shapec 8td|jtj|}t |t sJ|j |j}t||j}|Dcgc]}||jnd}}t||j\}}t||dj}tt|Dcgc]}|| } }t|||||| dd\} } t!|| }|j} | Jt#| |j|j| | d} t%|j&j&j(t+|| }|Scc}wcc}w)NrT)check atomic_add)rKrJinner_fnrangesoutput_indexer scatter_moderIrMdata)r get_devicer_float32rVrrealizeget_sizer' get_dtype make_loaderr listranger1r$r#rrrorIr)rfindicesvaluesgradx_sizeiindices_loaderstensor_indices tensor_size indexed_sizeexpected_vals_sizerjrKscatterbuffers r5zeros_and_scatter_loweringrs F%%' >D dI && &LLN ]]_F fdnn. /FKRSa!-q}}TASOS9'4??CTUG^w~a01::<=K',S\':;!F1I;L;#A  $ F. /F __ F   nn##%!! G YY^^ )$/ F MET.convert_output_node_to_buffers   m^ 4 -3 Z  6  3-,,j9 ]  <  9)!$))446#((224"''002 ##(( r7)subgraph_loweringrr r;r r_opsflex_libzeros_and_scatterdefaultrset_graph_handlerrunr1buffersregister_bufferrrr graph_outputs)rRrr pw_subgraphrrs r5build_subgraph_module_bufferrs>+GG$eii&8&8&J&J&R&R%ST II   0 0 8 8:T K  [ )   ;  !#!)) ,F GG # #F + ,8P2 1;3L3L MMEs D%%D.subgraphc.t||jSrU)rr)rRrs r5build_subgraph_bufferrs 'h.C.C DDr7a @triton.jit def get_offset_for_next_block( loop_iter, col_indices, total_blocks, SPARSE_BLOCK, SPARSE_BLOCK_MULTIPLE, BLOCK, BLOCKS_ARE_CONTIGUOUS: tl.constexpr ): if BLOCKS_ARE_CONTIGUOUS: return BLOCK cur_block_idx = loop_iter // SPARSE_BLOCK_MULTIPLE cur_block = tl.load(col_indices + cur_block_idx, eviction_policy="evict_last") next_block = tl.load(col_indices + cur_block_idx + 1, eviction_policy="evict_last", mask=cur_block_idx + 1 < total_blocks) needs_jump = (loop_iter + 1) % SPARSE_BLOCK_MULTIPLE == 0 jump_to_block = (next_block - cur_block ) * SPARSE_BLOCK - (SPARSE_BLOCK_MULTIPLE - 1) * BLOCK offset = jump_to_block * needs_jump + (1 - needs_jump) * BLOCK return offset z~ @triton.jit def get_bounded_indices(indices, max_len=None): return indices % max_len if max_len is not None else indices a @triton.jit def load_checked_block(block_ptr, IS_DIVISIBLE: tl.constexpr, SAFE_HEAD_DIM: tl.constexpr): if IS_DIVISIBLE and SAFE_HEAD_DIM: return tl.load(block_ptr) elif IS_DIVISIBLE and not SAFE_HEAD_DIM: return tl.load(block_ptr, boundary_check=(1,), padding_option="zero") elif not IS_DIVISIBLE and SAFE_HEAD_DIM: return tl.load(block_ptr, boundary_check=(0,), padding_option="zero") else: return tl.load(block_ptr, boundary_check=(0, 1), padding_option="zero") ak @triton.jit def load_checked_2d( ptr, offs_m, offs_n, stride_m, stride_n, IS_DIVISIBLE_M: tl.constexpr, IS_DIVISIBLE_N: tl.constexpr, M_LEN: tl.constexpr, N_DIM: tl.constexpr, ): # Calculate final pointer if strides are provided if stride_m is not None and stride_n is not None: ptr = ptr + offs_m[:, None] * stride_m + offs_n[None, :] * stride_n # Handle all masking cases if not IS_DIVISIBLE_M and not IS_DIVISIBLE_N: return tl.load(ptr, mask=(offs_m[:, None] < M_LEN) & (offs_n[None, :] < N_DIM), other=0.0) elif IS_DIVISIBLE_M and not IS_DIVISIBLE_N: return tl.load(ptr, mask=(offs_n[None, :] < N_DIM), other=0.0) elif not IS_DIVISIBLE_M and IS_DIVISIBLE_N: return tl.load(ptr, mask=(offs_m[:, None] < M_LEN), other=0.0) else: # Both divisible return tl.load(ptr) aj {{def_kernel("Q", "K", "V", "LSE", "KV_NUM_BLKS", "KV_IDX", "FULL_KV_NUM_BLKS", "FULL_KV_IDX")}} # Sub notation for this kernel: # # Q: Query, K: Key, V: Value # M: Number of queries, N: Number of keys/values, D: Model dimension # QK_HEAD_DIM: The dimension of the query and key embeddings # V_HEAD_DIM: The dimension of the value embeddings # z: Batch size, h: Number of heads, m: Number of queries per head, k: Number of keys per head # GQA_SHARED_HEADS: number of query heads sharing one kv head in GQA setups. # # The following FULL_* and PARTIAL_* is defined in the block sparse mask grid, rather than the thread block grid. # 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. # FULL_KV_NUM_BLKS: The number of fully unmasked KV blocks (so we don't need masking) for each query. # FULL_KV_IDX: The indices of fully unmasked KV blocks (so we don't need masking) for each query. # # OUTPUT_LOGSUMEXP: We only need to store the logsumexp if we require grad # # (Modifiable) Performance tuning options # BLOCK_M: The thread block size across the seqlen dim of Q. # BLOCK_N: Iterate over BLOCK_N across the seqlen dim of K/V in each thread block. # The below are kernel options that can be applied for certain score_mods, # or involve a numerics vs. perf tradeoff # PRESCALE_QK: Whether to pre-scale QK by 1/sqrt(d) and change of base. Has # about 20% more numerical error, but slightly faster. # 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 # BLOCKS_ARE_CONTIGUOUS: Is it guaranteed that all blocks in the mask are # contiguous? If so, we don't need to do an indirect jump for every block tl.static_assert(SPARSE_Q_BLOCK_SIZE >= BLOCK_M and SPARSE_Q_BLOCK_SIZE % BLOCK_M == 0) tl.static_assert(SPARSE_KV_BLOCK_SIZE >= BLOCK_N and SPARSE_KV_BLOCK_SIZE % BLOCK_N == 0) # Define strides of inputs stride_qz, stride_qh, 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")}} ZQ = {{size("Q", 0)}} HQ = {{size("Q", 1)}} Q_LEN = {{size("Q", 2)}} ZKV = {{size("K", 0)}} KV_LEN = {{size("K", 2)}} MATMUL_PRECISION = Q.dtype.element_ty q_start = tl.program_id(0) off_zq = tl.program_id(1) // HQ off_hq = tl.program_id(1) % HQ # We support two cases for batch dimension. a) (ZKV == ZQ) where off_zkv = off_zq. # b) (ZKV == 1 and ZQ > 1) where KV is broadcasted along the batch dimension and off_zkv=0. off_zkv = off_zq % ZKV off_hkv = off_hq // GQA_SHARED_HEADS off_g = off_hq % GQA_SHARED_HEADS q_offset = off_zq * stride_qz + off_hq * stride_qh k_offset = off_zkv * stride_kz + off_hkv * stride_kh v_offset = off_zkv * stride_vz + off_hkv * stride_vh Q = Q + q_offset K = K + k_offset V = V + v_offset SPARSE_Z = {{size("KV_NUM_BLKS", 0)}} SPARSE_HQ = {{size("KV_NUM_BLKS", 1)}} sparse_idx_z = off_zq % SPARSE_Z sparse_idx_hq = off_hq % SPARSE_HQ SPARSE_Q_MULTIPLE: tl.constexpr = (SPARSE_Q_BLOCK_SIZE // BLOCK_M) SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N) stride_kv_num_blks_h = {{stride("KV_NUM_BLKS", 1)}} stride_kv_idx_h = {{stride("KV_IDX", 1)}} stride_kv_idx_m = {{stride("KV_IDX", 2)}} # 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) offs_m = q_start * BLOCK_M + tl.arange(0, BLOCK_M) # KV_IDX and KV_NUM_BLKS are always contiguous. sparse_hz_offset = sparse_idx_z * SPARSE_HQ + sparse_idx_hq sparse_kv_num_blks_offset = sparse_hz_offset * stride_kv_num_blks_h + q_start // SPARSE_Q_MULTIPLE sparse_kv_idx_offset = sparse_hz_offset * stride_kv_idx_h + (q_start // SPARSE_Q_MULTIPLE) * stride_kv_idx_m # noqa: B950 Q_block_ptr = tl.make_block_ptr( base=Q, shape=(Q_LEN, QK_HEAD_DIM), strides=(stride_qm, stride_qk), offsets=(q_start * BLOCK_M, 0), block_shape=(BLOCK_M, QK_HEAD_DIM_ROUNDED), order=(1, 0) ) q = load_checked_block(Q_block_ptr, IS_DIVISIBLE, SAFE_HEAD_DIM) # ~~~~~~~~~~~~~~ normal blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # We don't know anything "special" about these blocks, so we need to apply # both score_mod and mask_mod to it kv_indices = KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading kv_num_blocks = tl.load(KV_NUM_BLKS + sparse_kv_num_blks_offset) block_n_end = 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, shape=(QK_HEAD_DIM, KV_LEN), strides=(stride_kk, stride_kn), offsets=(0, kv_start), block_shape=(QK_HEAD_DIM_ROUNDED, BLOCK_N), order=(0, 1) ) V_block_ptr = tl.make_block_ptr( base=V, shape=(KV_LEN, V_HEAD_DIM), strides=(stride_vn, stride_vk), offsets=(kv_start, 0), block_shape=(BLOCK_N, V_HEAD_DIM_ROUNDED), order=(1, 0) ) offs_n = kv_start + tl.arange(0, BLOCK_N) acc, l_i, m_i = forward_inner( {{gen_argdefs()}}, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, acc, l_i, m_i, off_zq, off_hq, offs_m[:, None], offs_n[None, :], kv_indices, kv_num_blocks, 0, block_n_end, 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: # FULL_KV_IDX and FULL_KV_NUM_BLKS are always contiguous. kv_indices = FULL_KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading kv_num_blocks = tl.load(FULL_KV_NUM_BLKS + sparse_kv_num_blks_offset) block_n_end = 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, shape=(QK_HEAD_DIM, KV_LEN), strides=(stride_kk, stride_kn), offsets=(0, kv_start), block_shape=(QK_HEAD_DIM_ROUNDED, BLOCK_N), order=(0, 1) ) V_block_ptr = tl.make_block_ptr( base=V, shape=(KV_LEN, V_HEAD_DIM), strides=(stride_vn, stride_vk), offsets=(kv_start, 0), block_shape=(BLOCK_N, V_HEAD_DIM_ROUNDED), order=(1, 0) ) offs_n = kv_start + tl.arange(0, BLOCK_N) acc, l_i, m_i = forward_inner( {{gen_argdefs()}}, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, acc, l_i, m_i, off_zq, off_hq, offs_m[:, None], offs_n[None, :], kv_indices, kv_num_blocks, 0, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS=True, ) # [Note] Handle fully masked out rows: # Li will be the sum(e^(-inf)) == 0.0 for masked out rows, mi will be -inf. # We set Li to 1.0 which will result in lse/out = 0.0 | after the log(li) + mi(0.0) step l_i = tl.where(l_i == 0.0, 1, l_i) acc = acc / l_i[:, None] idx_zq = tl.program_id(1) // HQ idx_hq = tl.program_id(1) % HQ idx_m = offs_m[:, None] idx_d = tl.arange(0, V_HEAD_DIM_ROUNDED)[None, :] mask = (idx_m < Q_LEN) & (idx_d < V_HEAD_DIM) {{store_output(("idx_zq", "idx_hq", "idx_m", "idx_d"), "acc", "mask")}} if OUTPUT_LOGSUMEXP: off_hz = tl.program_id(1) l_ptrs = LSE + off_hz * Q_LEN + offs_m lse = m_i + tl.math.log2(l_i) if IS_DIVISIBLE: tl.store(l_ptrs, lse) else: tl.store(l_ptrs, lse, mask=offs_m < Q_LEN) a @triton.jit def forward_inner( {{gen_argdefs()}}, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets used as inputs to score_mod & mask_mod # of size [BLOCK_M, BLOCK_N] or scalar. off_z, off_h, offs_m, offs_n, # blocksparse data kv_indices, kv_num_blocks, # start kv and end kv block block_n_start, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS, ): # Redefines all kernel parameters (BLOCK_M, etc.) so we don't need to plumb them all through {{gen_defines() | indent_except_first(1)}} SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N) RCP_LN2: tl.constexpr = 1.44269504 if PRESCALE_QK: q = (q * SM_SCALE * RCP_LN2).to(MATMUL_PRECISION) # loop over k, v and update accumulator until block_n_end for start_n in range(block_n_start, block_n_end): if IS_DIVISIBLE: acc, l_i, m_i = forward_block_mn( {{gen_argdefs()}}, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, ) else: # Benchmark shows even we applied mod & mask to each block for non divisible seqlen, # it's on par or slightly faster than only applying to the last block in fwd. # However, we choose different strategy for bwd, where we only apply mod & mask # to the last block because it's faster a lot. acc, l_i, m_i = forward_block_mn( {{gen_argdefs()}}, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=True, ) # update pointers offset = get_offset_for_next_block( start_n, kv_indices, kv_num_blocks, SPARSE_KV_BLOCK_SIZE, SPARSE_KV_MULTIPLE, BLOCK_N, BLOCKS_ARE_CONTIGUOUS ) V_block_ptr = tl.advance(V_block_ptr, (offset, 0)) K_block_ptr = tl.advance(K_block_ptr, (0, offset)) offs_n = offs_n + offset return acc, l_i, m_i aX @triton.jit def forward_block_mn( {{gen_argdefs()}}, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=False, ): # Redefines all kernel parameters (BLOCK_M, etc.) so we don't need to plumb them all through {{gen_defines() | indent_except_first(1)}} # -- load k -- # NB reversed order to since K is transposed k = load_checked_block(K_block_ptr, SAFE_HEAD_DIM, IS_DIVISIBLE) # -- compute qk --- qk = tl.dot(q, k, input_precision=FLOAT32_PRECISION) # TODO: use cuda matmul when q_len <= 2. if not PRESCALE_QK: qk *= SM_SCALE # ~~~~~~~~~~~~~~~~~~~ Apply score modification ~~~~~~~~~~~~~~~~~~~ # If this is the last block of a non divisible seqlen, we still need to load [BLOCK_M, BLOCK_N] elements, # which is larger than the actual number of elements. To avoid access memory out of bound, # we need to mask out the elements that are out of Q_LEN & KV_LEN. m = get_bounded_indices(offs_m, Q_LEN if CHECK_BLOCK_BOUNDARY else None) n = get_bounded_indices(offs_n, KV_LEN if CHECK_BLOCK_BOUNDARY else None) {{ modification( subgraph_number=0, output_name="post_mod_scores", score="qk", b="off_z", h="off_h", m="m", n="n", out="qk" ) | indent_except_first(1) }} if CHECK_BLOCK_BOUNDARY: # Mask out the elements that are out of the KV_LEN for non divisible seqlen. post_mod_scores = tl.where(offs_n < KV_LEN, post_mod_scores, float("-inf")) if not IS_FULL_BLOCKS: {{ modification( subgraph_number=1, output_name="mask_mod_output", score="qk", b="off_z", h="off_h", m="m", n="n", ) | indent_except_first(2) }} if CHECK_BLOCK_BOUNDARY: mask_mod_output = tl.where(offs_n < KV_LEN, mask_mod_output, False) # apply mask for partially unmasked blocks post_mod_scores = tl.where(mask_mod_output, post_mod_scores, float("-inf")) if not PRESCALE_QK: post_mod_scores *= RCP_LN2 # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # -- compute scaling constant --- m_ij = tl.maximum(m_i, tl.max(post_mod_scores, 1)) if not ROWS_GUARANTEED_SAFE: masked_out_rows = (m_ij == float("-inf")) m_ij_masked = tl.where(masked_out_rows, 0, m_ij) else: m_ij_masked = m_ij alpha = tl.math.exp2(m_i - m_ij_masked) p = tl.math.exp2(post_mod_scores - m_ij_masked[:, None]) # NB: l_i update is pulled up here since it's a bit faster # NB: For headdim=256, it's faster to move it back down to after m_i = # m_ij l_i = l_i * alpha + tl.sum(p, 1) # # -- scale and update acc -- acc = acc * alpha[:, None] v = load_checked_block(V_block_ptr, IS_DIVISIBLE, SAFE_HEAD_DIM) acc = tl.dot(p.to(MATMUL_PRECISION), v, acc, input_precision=FLOAT32_PRECISION) # -- update m_i m_i = m_ij return acc, l_i, m_i flex_attention)rIgridsourcec "|jdd}tjjj t j |jdd}tjjj t j|jdd}t|jdtt jf}t|jdtt jf}|rOtjjj t j|jdd} ntjjj t jt j|jddt j|jd|jd} | xr|xr|xr |xr|xr| S)a;Decide which kernel to use, return true if use flex decoding kernel. Note: Since the number of splits is calculated based of the the number of batch and head dims we need to ensure that the batch and head dims are statically known. Otherwise we just use the main flex_attention kernel. FORCE_USE_FLEX_ATTENTIONFrr0)getr r;r< evaluate_exprrWLtrsGtrVintIntegerEqOr) query kv_indiceskernel_options enable_gqa force_flexshort_query_lengthnon_zero_length static_batchstatic_num_headsvalid_block_mask_num_headss r5_use_flex_decodingrs ##$>FJ))77 !"%s+gg&&44UXXenn>Nr>RTU5VWOenn.q1C3GHL!%.."21"5U]]7KL &'WW%5%5%C%C HHZ((*1-q 1& "&'WW%5%5%C%C HH,,.q115,,.q15>>3CA3FG & "  '  '  '  '   ' ' r7@)r r)rrrr)rrrrrrrr)rrrrrrr)rrrr)rrrrrrr)rrrr0)rrrr0rrrr0)rrrr0)rrrr0c(eZdZeZeZy)ModeN)__name__ __module__ __qualname__rfwdbwdrAr7r5rr s &C &Cr7rmodec|j}tjjj |j d}d}|t jk(rR|dkr2|tjk(rd}nd}tj||f|}|S|tjk(rd}|Sd}|S|t jk(sJ|tjk(ry|dkr |dk(ry|d k(ry yy) Nr)rrrr0r)rrrr0)rrrr0rrrr0rr)rrrr0) rtr r;r<evaluate_static_shapersrrr_rq_rocm_default_configrr)rrrJhead_dim fwd_configs r5_get_rocm_configrs OO Eww55enn6Fr6JKHJ txx s? %+ , -115(2CZPJ   %+ , txx EMM !! _2~%S&%!r7c|j}tjjj |j d}d}d}t jj}|tjk(rv|dkrV|t jk(rd}nd}|dk\rtj||f|}|S|dk\rtj||f|}|S|t jk(rd}|Sd}|S|tjk(sJ|t jk(rd }|S|dkr|dk\r|d k(rd}|S|d k(rd }|Sd }|S|dk\r&|d k\rd}|S|d k(r|ddk(rdnd}d d d|f}|Sd }|Sd }|S)Nrrrrrrr) r)rr)rrrrrrrr)rrrr)rrrr)rrrrrrrr)rtr r;r<rrsr_cudaget_device_capabilityrrrq_h100_default_configr_a100_default_configr)rrrJrr bwd_config capability num_stagess r5_get_nv_configr2s OO Eww55enn6Fr6JKHJJ113J txx s? %+ , V#155uh6GT v%155uh6GT   %+ , txx EMM !'J&%_v!52~+ S, , 6 !2~, S",R.A"5Q1  "a4 , (Jr7ctjjt|tj St |tj SN)r)r_rarbrrrrrs r5_get_default_config_fwdrc4 }} e$((33DHH55r7ctjjt|tj St |tj Sr)r_rarbrrrrrs r5_get_default_config_bwdrjrr7c4dtjffd }|S)Nr.cltjjjjd}|j Dcgc]+}tjjj|-}}t j|||j|jScc}w)NrrJrK) r r;r< size_hintrfrsr_fullrtrp)rYnum_blocks_for_autotuningr|r8sparse_indicess r5create_num_blocks_fakez@create_num_blocks_fake_generator..create_num_blocks_fake|s$%GG$4$4$>$>~?S?STV?W$X!78zz|D!  **1-DDzz  %++-<<>   Es 0B1)r_Tensor)rrs` r5 create_num_blocks_fake_generatorrqs U\\  "!r7cF|jDcgc]+}tjjj |-}}t j d|d|j|j}|j|j}|Scc}w)Nrrr) rsr r;r<rr_arangertrpr# contiguous)rYr|r8rxs r5create_indices_fakersv34::< @aAGG   & &q ) @D @ll1d2hakkmALLNSGnnT"--/G N As0B)create_flex_decoding_kernel)CppFlexAttentionTemplatecddl}ddl}tjj xr|j ddk7}|xr1tj j xr|jdk( }|S)NrATEN_CPU_CAPABILITYrdarwin) ossysr_cpu_is_avx2_supportedgetenvxpurdplatform)rrrequires_avx2_on_cpu supporteds r5check_cpu_supportedrsq   $$&X2995J+Ky+X  ) &&( ( ) ( ( r7c |j}|rR|ddk7rJtttt |j }t j||S|S)zJEnsure that realized IR node has a contigous stride in the last dimension.rr0)maybe_get_stridervreversedrwr1rsrrequire_stride_order)rYr2contiguous_stride_orders r5contiguous_last_dimrsW  "G72;!#"&xc!**,6G0H'I"J004KLL Hr7c >?@|\ } } } } } } }}}}}}}|dr tdts tdg}tjjj j jj?tjjj j jj>tjjj }tdt|j?<tdt|j><tj@d@?>gfdtjgfdtjgfdtj?d gfd tjd >gffDcgc]"\}}}t|||j!|$}}}}t#|t%|z|}|;t'|t$r|D]}||j)n|j)d@?>gfdtjgfdtjgfdtj?d gfd tjd >gffDcgc]"\}}}t|||j!|$}}}}>?@fd }||}t+|t%|z|} tjjj j,}!|!D"cgc] }"|"?>fvs |" c}"tjjj _|t%|z|zt%|z}#|#D]8}$t'|$t.s|j1|$j2j2:t5t6|||g\}}}t9|||| | | | ||||g \ }}}} } } } }}}}t;t=|j?|j?|j?gd k7r td |jAtjtjBtjDfvrtd|jAdt9|}t9|}|jG\}%}&}'}(|jG\})}*}+},|%}-|-|&|'|,g}.tI|.|jK}/tM|j!|jA|-|&|'|,g|/D0cgc]}0tOjP|0c}0}1g}2|||| | g}3| sd}4nd}4|3| gz }3d}5i}6|s|rd}5d|fd|ffD]8\}7}8|6jStU|8D9:cic] \}9}:|7d|9|:c}:}9:|3|6jWDcgc]}t'|tNjXs| c}z }3|j[dd};tjjj]|}tjjj]|}tjjj_tOj`|'tOjb| jGd|sJdtjjj_tOj`|+tOjb| jGd|sJdtejf|2|3|1||;rdn||;rdn| ||5|4|t;|t;||6?>f||g}.convert_mask_graph_modules ]]:22 3  # #Dekk):$; < 8 ,,Y7L 8 KK Dww(""   &&&$$Q' #_5  " "9 - ++ %,/-,I  " "9 - ,, $$I|Y+O-J  '=  HH((U3 E 8 8     s$F1F 9F FF F)rz=Unsupported for now if query, key, value are the same buffer.z}`torch.float` , `torch.float16` and `torch.bfloat16` are supported in FlexAttention for CPU device. Found input tensors are `z`.strideTF score_others mask_others_SKIP_MASK_SCORErzkQ seqlen must be smaller than the block_mask size in the Q dimension, considering pass a larger block_mask.rzmKV seqlen must be smaller than the block_mask size in the KV dimension, considering pass a larger block_mask.)choices input_nodesrMscale score_modmask_mod kv_block_sizehas_other_bufferno_full_kv_block fake_bufferslen_score_otherlen_mask_otherkernel_input_name_to_buffer block_varsr)5NotImplementedErrorrr r;r<r=create_unbacked_symintr)exprrr var_to_ranger_rint64rQrprrvrV freeze_layoutrpending_fresh_unbacked_symbolsrappendromaprr\r1r get_namertbfloat16float16rsr> get_striderrWsympifyupdate enumerateryrXrrrLeMulr add_choicesr()Arkeyvaluer block_maskr;rscore_mod_other_buffersmask_mod_other_buffersr7 kv_num_blocksrfull_kv_num_blocksfull_kv_indices q_num_blocks q_indicesfull_q_num_blocksfull_q_indicesSPARSE_Q_BLOCK_SIZESPARSE_KV_BLOCK_SIZEr%rAr=rIrJr8placeholder_inpsr_bufmask_graph_placeholder_inpsr2converted_mask_graph_modulemask_graph_bufferpendingrY buffer_listitemBqHq seq_len_q qk_head_dimBkvHkv seq_len_kv v_head_dimBout_size out_stridessrM_choicesr:r@r?rDprefixrr|bufskip_mask_scoreinputs_for_autotuningresr/r0r1sA @@@r5 lower_cpurs2   ()! a    ! I  "$LWW%%//FFHMMRRNgg&&00GGINNSSO  **I .9F-CI>*.9!V.DI?+++KkNO#D E %++r " %++r " ekkNA#6 7 u{{Q$8 9 "    D% 4(8(8(:DA  ,4 788(O" ot ,' )#&&( )  ) ) +kNO#D E %++r " %++r " ekkNA#6 7 u{{Q$8 9 "  # # D% 4(8(8(:DA # #2&P#  ,,CD*+AB%*^^%5"BI{',~~'7$Cj* A2y*-H%h0@0@0BCK    B :&*56Q a 6 F H#umZ@K  *++ "$"84 5 2 3  OFG ( . .4=g4FG&!SF81QC#%G    4;;= eU\\2   %(():EBO77++AABVW''**@@ATU 77   ) ) EIIj&9&9&;B&?ATUV  v  77   ) ) UYYz':':'/0AB ()Nmz!flex_attention..s)/ 45E!E,,. / // %(r)rr)rrrr)rrrrrkpackzRQ and KV block size must be divisible by BLOCK_M and BLOCK_N. We got Q_BLOCK_SIZE=z and KV_BLOCK_SIZE=.fwd_rbwd_r num_warpsr@BLOCK_Nrerfr9r:rM subgraphsmutated_inputs call_sizes)rr input_gen_fnsrA)7rprrr r;r<rrWrrsrFrtr_int32rQrrvdictitemsrVrXrrreNerrr\rrrRr>rrSr"rqrwrrMrr max_autotunerarbrr1 ValueErrorkeys startswithpopflex_attention_templatemaybe_append_choicerrr()>rrYrZrr[r;rr\r] small_dqksmall_dvr7r^rr_r`rarbrcrdrerfr%rIrJrgrrirkkvrrorprqrrrsrtrurvrw q_stridesrxryrzrMlogsumexp_shape logsumexpgqa_shared_headshas_full_blocksr9configscoriginal_kernel_optionsr@rrrcur_kernel_optionserrorrrs>` r5rrs"  %'        # "    ..uxx8H8Lb/QRIww--ehhu~~7G7KR.PQHH!%%*^^%5b%9$:(5>>CSTVCWBX Z  &    eoo' ( %++  %++  %++  %++     D% 4(8(8(:;  ,4 788(O%++  %++  %++  %++   # D% 4(8(8(:;##.#d+A&BBJ.)N #((*  Aq a & 77   1 1! 4  N 13H3JK!!// !!$cllnQ&78J%^ZH*         # "  2                  ",,CD*+AB%*^^%5"BI{',~~'7$Cj* 77   ) )%((2s*;ehhsA>N*N O 0IcUC O 77   ) )%((9a*@ A- A 77   ) )%((:q*A B+ B A3!zC/14!!.%8!!.$7  "I2y*-H%h :K    B :&*56Q a 6 F"i(O mm! I j%0Sy02BC)4O/A / 9>q/ +O ZAQAQRG/1G NN*512    ==  6=>!adAaD!,>G>77++AABVW''**@@ATU }}"#w -1133:3/)Z ' )Q .2E2OST2T7|q  ((;'<RS';; " ! ~~') *!+ .  W!2Kg3l          & ' ( % & ' ,J 7  +O <  M "   ! '     c #j7R?s/%b8%bAb"b(9b-cxddl}|j||d||zz|j||dzd||zfS)a?How is this kernel parallelized? Currently this is only parallelizing over batch* kv_heads, but we can, and want to parallelize over ceil_div(q_heads//kv_heads * num_key_value, key_value_block_size). To do this will either require atomic updates to some grad values or to have a two pass kernel design. rNBLOCK_M2BLOCK_N1r0)tritonrG)rBrCrDrEkv_heads num_key_valuerFrs r5flex_attention_backward_gridr+sS  Kj!12g6IJ ++mT*%5 6 7 X  r7flex_attention_backwardam {{def_kernel("Q", "K", "V", "LSE", "DELTA", "DO", "DQ", "DV", "KV_NUM_BLKS", "KV_IDX", "Q_NUM_BLKS", "Q_IDX", "FULL_KV_NUM_BLKS", "FULL_KV_IDX", "FULL_Q_NUM_BLKS", "FULL_Q_IDX")}} # Sub notation for this kernel: # # Q: Query, K: Key, V: Value # LSE: logsumexp (logsumexp is always stored in fp32 regardless of the input dtype) # DELTA: Precomputed sum(OUT*DO, axis=-1) # DO: Derivative of Output, DQ: Derivative of Query, DV: Derivative of Value # DK: Derivative of Key, is the written to via the store_output call due to some limitations with # inductor codegen # 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 # z: Batch size, h: Number of heads, m: Number of queries or keys/values, d: Head dim # GQA_SHARED_HEADS: number of query heads sharing one kv head in GQA setups. # (Modifiable) Performance tuning options # BLOCK_M1: when calculating DK & DV, iterate over BLOCK_M1 across the seqlen dim of Q in each thread block. # BLOCK_N1: when calculating DK & DV, the thread block size across the seqlen dim of K/V. # BLOCK_M2: when calculating DQ, the thread block size across the seqlen dim of Q. # BLOCK_N2: when calculating DQ, iterate over BLOCK_N2 across the seqlen dim of K/V in each thread block. # # The following FULL_* and PARTIAL_* is defined in the block sparse mask grid, rather than the thread block grid. # 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. # Q_NUM_BLKS: The number of Q blocks (that may or may not require masking) for each query. # Q_IDX: The indices of Q blocks (that may or may not require masking) for each query. # FULL_KV_NUM_BLKS: The number of fully unmasked KV blocks (so we don't need masking) for each query. # FULL_KV_IDX: The indices of fully unmasked KV blocks (so we don't need masking) for each query. # FULL_Q_NUM_BLKS: The number of fully unmasked Q blocks (so we don't need masking) for each query. # FULL_Q_IDX: The indices of fully unmasked Q blocks (so we don't need masking) for each query. # The below are kernel options that can be applied for certain score_mods, # or involve a numerics vs. perf tradeoff # PRESCALE_QK: Whether to pre-scale QK by 1/sqrt(d) and change of base. Has # about 20% more numerical error, but slightly faster. # Define strides of inputs stride_qz, stride_qh, stride_qm, stride_qd = {{stride("Q")}} stride_kz, stride_kh, stride_kn, stride_kd = {{stride("K")}} stride_vz, stride_vh, stride_vn, stride_vd = {{stride("V")}} stride_doz, stride_doh, stride_dom, stride_dod = {{stride("DO")}} stride_dqz, stride_dqh, stride_dqm, stride_dqd = {{stride("DQ")}} stride_dvz, stride_dvh, stride_dvm, stride_dvd = {{stride("DV")}} ZQ = {{size("Q", 0)}} HQ = {{size("Q", 1)}} HKV = {{size("K", 1)}} Q_LEN = {{size("Q", 2)}} ZKV = {{size("K", 0)}} KV_LEN = {{size("K", 2)}} MATMUL_PRECISION = Q.dtype.element_ty pid = tl.program_id(0) NUM_KV_BLOCKS = tl.cdiv(KV_LEN, BLOCK_N1) NUM_Q_BLOCKS = tl.cdiv(Q_LEN, BLOCK_M2) off_hz = tl.program_id(2) off_zq = off_hz // HKV # q batch idx off_hkv = off_hz % HKV # kv head idx off_zkv = off_zq % ZKV # kv batch idx SPARSE_Z = {{size("KV_NUM_BLKS", 0)}} SPARSE_HQ = {{size("KV_NUM_BLKS", 1)}} sparse_idx_z = off_zq % SPARSE_Z k_adj = (stride_kh * off_hkv + stride_kz * off_zkv).to(tl.int64) v_adj = (stride_vh * off_hkv + stride_vz * off_zkv).to(tl.int64) # first compute broadcasted dv of shape [Bq, Hkv, KV_LEN, V_HEAD_DIM] # then reduce to dv of shape [Bkv, Hkv, KV_LEN, V_HEAD_DIM] dv_adj = (stride_dvh * off_hkv + stride_dvz * off_zq).to(tl.int64) # offset K, V, DV pointers for batch/kv-head K += k_adj V += v_adj DV += dv_adj RCP_LN2 = 1.44269504 offs_k = tl.arange(0, QK_HEAD_DIM_ROUNDED) offs_v = tl.arange(0, V_HEAD_DIM_ROUNDED) if pid >= NUM_KV_BLOCKS: off_pid = pid - NUM_KV_BLOCKS # THIS BLOCK DOES DQ SPARSE_Q_MULTIPLE = (SPARSE_Q_BLOCK_SIZE // BLOCK_M2) SPARSE_KV_MULTIPLE = (SPARSE_KV_BLOCK_SIZE // BLOCK_N2) off_hq2 = off_pid // NUM_Q_BLOCKS + off_hkv * GQA_SHARED_HEADS start_m2_block = off_pid % NUM_Q_BLOCKS off_pid_mask = start_m2_block // SPARSE_Q_MULTIPLE stride_kv_num_blks_h = {{stride("KV_NUM_BLKS", 1)}} stride_kv_idx_h = {{stride("KV_IDX", 1)}} stride_kv_idx_m = {{stride("KV_IDX", 2)}} sparse_idx_hq2 = off_hq2 % SPARSE_HQ sparse_hz_offset = sparse_idx_z * SPARSE_HQ + sparse_idx_hq2 sparse_kv_num_blks_offset = sparse_hz_offset * stride_kv_num_blks_h + off_pid_mask sparse_kv_idx_offset = sparse_hz_offset * stride_kv_idx_h + off_pid_mask * stride_kv_idx_m # noqa: B950 # Offset Q, DQ, DO, DELTA & LSE. These inputs are offseted by query heads. q_adj2 = (stride_qh * off_hq2 + stride_qz * off_zq).to(tl.int64) do_adj2 = (stride_doh * off_hq2 + stride_doz * off_zq).to(tl.int64) dq_adj2 = (stride_dqh * off_hq2 + stride_dqz * off_zq).to(tl.int64) off_chz2 = ((off_zq * HQ + off_hq2) * Q_LEN).to(tl.int64) Q2 = Q + q_adj2 DO2 = DO + do_adj2 # TODO: This does not work if DQ is not the same layout as Q (for example, # if Q is broadcasted) DQ2 = DQ + dq_adj2 LSE2 = LSE + off_chz2 DELTA2 = DELTA + off_chz2 # dq = tl.zeros([BLOCK_M2, QK_HEAD_DIM], dtype=tl.float32) dq = tl.zeros([BLOCK_M2, QK_HEAD_DIM_ROUNDED], dtype=tl.float32) start_m2 = start_m2_block * BLOCK_M2 offs_m2 = start_m2 + tl.arange(0, BLOCK_M2) # load Q and do: they stay in SRAM throughout the inner loop. q = load_checked_2d(Q2, offs_m2, offs_k, stride_qm, stride_qd, IS_DIVISIBLE, SAFE_HEAD_DIM, Q_LEN, QK_HEAD_DIM) do = load_checked_2d(DO2, offs_m2, offs_v, stride_dom, stride_dod, IS_DIVISIBLE, SAFE_HEAD_DIM, Q_LEN, V_HEAD_DIM) if PRESCALE_QK: q = (q * SM_SCALE * RCP_LN2).to(MATMUL_PRECISION) if IS_DIVISIBLE: Di = tl.load(DELTA2 + offs_m2) lse = tl.load(LSE2 + offs_m2) else: Di = tl.load(DELTA2 + offs_m2, mask=offs_m2 < Q_LEN) lse = tl.load(LSE2 + offs_m2, mask=offs_m2 < Q_LEN) lse = tl.where(lse == -float("inf"), 0.0, lse) lse = lse[:, None] # ~~~~~~~~~~~ fully unmasked blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # KV_IDX and KV_NUM_BLKS are always contiguous. kv_indices = KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading sparse_kv_num_blocks = tl.load(KV_NUM_BLKS + sparse_kv_num_blks_offset) offs_n2 = kv_start + tl.arange(0, BLOCK_N2) dq = bwd_dq_inner( {{gen_argdefs()}}, K, V, dq, q, do, Di, lse, off_zq, off_hq2, offs_m2, offs_n2, stride_kn, stride_kd, stride_vn, stride_vd, kv_indices, sparse_kv_num_blocks, MATMUL_PRECISION, IS_FULL_BLOCKS=False, ) if HAS_FULL_BLOCKS: # ~~~~~~~~~~~ partial unmasked blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # FULL_KV_IDX and FULL_KV_NUM_BLKS are always contiguous. kv_indices = FULL_KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading sparse_kv_num_blocks = tl.load(FULL_KV_NUM_BLKS + sparse_kv_num_blks_offset) offs_n2 = kv_start + tl.arange(0, BLOCK_N2) dq = bwd_dq_inner( {{gen_argdefs()}}, K, V, dq, q, do, Di, lse, off_zq, off_hq2, offs_m2, offs_n2, stride_kn, stride_kd, stride_vn, stride_vd, kv_indices, sparse_kv_num_blocks, MATMUL_PRECISION, IS_FULL_BLOCKS=True, ) # Write back dQ. dq_ptrs = DQ2 + offs_m2[:, None] * stride_dqm + offs_k[None, :] * stride_dqd dq *= SM_SCALE if IS_DIVISIBLE and SAFE_HEAD_DIM: tl.store(dq_ptrs, dq) else: tl.store(dq_ptrs, dq, mask=(offs_m2[:, None] < Q_LEN) & (offs_k[None, :] < QK_HEAD_DIM)) else: # THIS BLOCK DOES DK & DV SPARSE_Q_MULTIPLE = (SPARSE_Q_BLOCK_SIZE // BLOCK_M1) SPARSE_KV_MULTIPLE = (SPARSE_KV_BLOCK_SIZE // BLOCK_N1) pid_mask = pid // SPARSE_KV_MULTIPLE stride_q_num_blks_h = {{stride("Q_NUM_BLKS", 1)}} stride_q_idx_h = {{stride("Q_IDX", 1)}} stride_q_idx_n = {{stride("Q_IDX", 2)}} dv = tl.zeros([BLOCK_N1, V_HEAD_DIM_ROUNDED], dtype=tl.float32) dk = tl.zeros([BLOCK_N1, QK_HEAD_DIM_ROUNDED], dtype=tl.float32) start_n1 = pid * BLOCK_N1 offs_n1 = start_n1 + tl.arange(0, BLOCK_N1) # load K and V: they stay in SRAM throughout the inner loop. k = load_checked_2d(K, offs_n1, offs_k, stride_kn, stride_kd, IS_DIVISIBLE, SAFE_HEAD_DIM, KV_LEN, QK_HEAD_DIM) v = load_checked_2d(V, offs_n1, offs_v, stride_vn, stride_vd, IS_DIVISIBLE, SAFE_HEAD_DIM, KV_LEN, V_HEAD_DIM) if PRESCALE_QK: k = (k * SM_SCALE * RCP_LN2).to(MATMUL_PRECISION) for off_g in range(0, GQA_SHARED_HEADS): off_hq1 = off_hkv * GQA_SHARED_HEADS + off_g # Offset Q, DQ, DO, DELTA & LSE. These inputs are offseted by query heads. q_adj1 = (stride_qh * off_hq1 + stride_qz * off_zq).to(tl.int64) do_adj1 = (stride_doh * off_hq1 + stride_doz * off_zq).to(tl.int64) dq_adj1 = (stride_dqh * off_hq1 + stride_dqz * off_zq).to(tl.int64) off_chz1 = ((off_zq * HQ + off_hq1) * Q_LEN).to(tl.int64) Q1 = Q + q_adj1 DO1 = DO + do_adj1 # TODO: This does not work if DQ is not the same layout as Q (for example, # if Q is broadcasted) LSE1 = LSE + off_chz1 DELTA1 = DELTA + off_chz1 sparse_idx_hq1 = off_hq1 % SPARSE_HQ sparse_hz_offset = sparse_idx_z * SPARSE_HQ + sparse_idx_hq1 sparse_q_num_blks_offset = sparse_hz_offset * stride_q_num_blks_h + pid_mask sparse_q_idx_offset = sparse_hz_offset * stride_q_idx_h + pid_mask * stride_q_idx_n # noqa: B950 # ~~~~~~~~~~~~~~~ fully unmasked blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # Q_IDX and Q_NUM_BLKS are always contiguous. q_indices = Q_IDX + sparse_q_idx_offset q_start = tl.load(q_indices) * SPARSE_Q_BLOCK_SIZE # first q block we're loading sparse_q_num_blocks = tl.load(Q_NUM_BLKS + sparse_q_num_blks_offset) offs_m1 = q_start + tl.arange(0, BLOCK_M1) dk, dv = bwd_dkdv_inner( {{gen_argdefs()}}, Q1, DO1, DELTA1, LSE1, dk, dv, k, v, off_zq, off_hq1, offs_n1, offs_m1, stride_qm, stride_qd, stride_dom, stride_dod, q_indices, sparse_q_num_blocks, MATMUL_PRECISION, IS_FULL_BLOCKS=False, ) if HAS_FULL_BLOCKS: # ~~~~~~~~~~~~~~~ fully unmasked blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # FULL_Q_IDX and FULL_Q_NUM_BLKS are always contiguous. q_indices = FULL_Q_IDX + sparse_q_idx_offset q_start = tl.load(q_indices) * SPARSE_Q_BLOCK_SIZE # first q block we're loading sparse_q_num_blocks = tl.load(FULL_Q_NUM_BLKS + sparse_q_num_blks_offset) offs_m1 = q_start + tl.arange(0, BLOCK_M1) dk, dv = bwd_dkdv_inner( {{gen_argdefs()}}, Q1, DO1, DELTA1, LSE1, dk, dv, k, v, off_zq, off_hq1, offs_n1, offs_m1, stride_qm, stride_qd, stride_dom, stride_dod, q_indices, sparse_q_num_blocks, MATMUL_PRECISION, IS_FULL_BLOCKS=True, ) # Write back dV and dK. dv_ptrs = DV + offs_n1[:, None] * stride_dvm + offs_v[None, :] * stride_dvd index_n = offs_n1[:, None] index_k = offs_k[None, :] index_v = offs_v[None, :] if IS_DIVISIBLE and SAFE_HEAD_DIM: tl.store(dv_ptrs, dv) else: tl.store(dv_ptrs, dv, mask=(index_n < KV_LEN) & (index_v < V_HEAD_DIM)) dk *= SM_SCALE if SAFE_HEAD_DIM: mask = index_n < KV_LEN else: mask = (index_n < KV_LEN) & (index_k < QK_HEAD_DIM) # first compute broadcasted dk of shape [Bq, Hkv, KV_LEN, V_HEAD_DIM] # then reduce to dk of shape [Bkv, Hkv, KV_LEN, V_HEAD_DIM] {{store_output(("off_zq", "off_hkv", "index_n", "index_k"), "dk", "mask", indent_width=8)}} @triton.jit def bwd_dq_inner( {{gen_argdefs()}}, K, V, # pointers dq, q, do, Di, lse, off_z, off_hq, offs_m2, offs_n2, stride_kn, stride_kd, stride_vn, stride_vd, kv_indices, sparse_kv_num_blocks, MATMUL_PRECISION, IS_FULL_BLOCKS, ): {{gen_defines() | indent_except_first(1) }} SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N2) RCP_LN2: tl.constexpr = 1.44269504 Q_LEN = {{size("Q", 2)}} KV_LEN = {{size("K", 2)}} offs_k = tl.arange(0, QK_HEAD_DIM_ROUNDED) offs_v = tl.arange(0, V_HEAD_DIM_ROUNDED) kT_ptrs = K + offs_n2[None, :] * stride_kn + offs_k[:, None] * stride_kd vT_ptrs = V + offs_n2[None, :] * stride_vn + offs_v[:, None] * stride_vd # BLOCK_M2 must be a multiple of BLOCK_N2, otherwise the code wouldn't work. tl.static_assert(BLOCK_M2 % BLOCK_N2 == 0) hi = tl.minimum(sparse_kv_num_blocks * SPARSE_KV_MULTIPLE, tl.maximum(tl.cdiv(KV_LEN, BLOCK_N2), 1)) if not IS_DIVISIBLE: if hi >= 1: for start_n in range(0, hi - 1): dq = bwd_dq_block_mn( {{gen_argdefs()}}, dq, q, kT_ptrs, vT_ptrs, do, Di, lse, Q_LEN, KV_LEN, off_z, off_hq, offs_m2, offs_n2, offs_k, offs_v, stride_kn, stride_kd, stride_vn, stride_vd, kv_indices, sparse_kv_num_blocks, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, ) # Increment pointers. offset = get_offset_for_next_block( start_n, kv_indices, sparse_kv_num_blocks, SPARSE_KV_BLOCK_SIZE, SPARSE_KV_MULTIPLE, BLOCK_N2, BLOCKS_ARE_CONTIGUOUS ) kT_ptrs += offset * stride_kn vT_ptrs += offset * stride_vn offs_n2 += offset dq = bwd_dq_block_mn( {{gen_argdefs()}}, dq, q, kT_ptrs, vT_ptrs, do, Di, lse, Q_LEN, KV_LEN, off_z, off_hq, offs_m2, offs_n2, offs_k, offs_v, stride_kn, stride_kd, stride_vn, stride_vd, kv_indices, sparse_kv_num_blocks, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=True, ) else: for start_n in range(0, hi): dq = bwd_dq_block_mn( {{gen_argdefs()}}, dq, q, kT_ptrs, vT_ptrs, do, Di, lse, Q_LEN, KV_LEN, off_z, off_hq, offs_m2, offs_n2, offs_k, offs_v, stride_kn, stride_kd, stride_vn, stride_vd, kv_indices, sparse_kv_num_blocks, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, ) # Increment pointers. offset = get_offset_for_next_block( start_n, kv_indices, sparse_kv_num_blocks, SPARSE_KV_BLOCK_SIZE, SPARSE_KV_MULTIPLE, BLOCK_N2, BLOCKS_ARE_CONTIGUOUS ) kT_ptrs += offset * stride_kn vT_ptrs += offset * stride_vn offs_n2 += offset return dq @triton.jit def bwd_dq_block_mn( {{gen_argdefs()}}, dq, q, kT_ptrs, vT_ptrs, do, Di, lse, Q_LEN, KV_LEN, off_z, off_hq, offs_m2, offs_n2, offs_k, offs_v, stride_kn, stride_kd, stride_vn, stride_vd, kv_indices, sparse_kv_num_blocks, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=False, ): {{gen_defines() | indent_except_first(1)}} # NB reversed order to since K is transposed kT = load_checked_2d(kT_ptrs, offs_k, offs_n2, None, None, SAFE_HEAD_DIM, IS_DIVISIBLE, QK_HEAD_DIM, KV_LEN) qk = tl.dot(q, kT, input_precision=FLOAT32_PRECISION) if not PRESCALE_QK: qk *= SM_SCALE # ~~~~~~~~~~~~~~~~~~~ Apply score modification ~~~~~~~~~~~~~~~~~~~ pre_mod_scores = qk n = get_bounded_indices(offs_n2[None, :], KV_LEN if CHECK_BLOCK_BOUNDARY else None) # The boundary check is done for the outer loop, but here it's possible since we're iterating across N dim # that the M reads out of bounds prior to the last loop m = get_bounded_indices(offs_m2[:, None], Q_LEN if (not IS_DIVISIBLE or CHECK_BLOCK_BOUNDARY) else None) {{ modification( subgraph_number=0, output_name="post_mod_scores", score="qk", b="off_z", h="off_hq", m="m", n="n", out="qk" ) | indent_except_first(1) }} if CHECK_BLOCK_BOUNDARY: # Mask out the elements that are out of the KV_LEN for non divisible seqlen. post_mod_scores = tl.where(offs_n2[None, :] < KV_LEN, post_mod_scores, float("-inf")) if not IS_FULL_BLOCKS: {{ modification( subgraph_number=2, output_name="mask_mod_output", score="qk", b="off_z", h="off_hq", m="m", n="n", ) | indent_except_first(2) }} if CHECK_BLOCK_BOUNDARY: mask_mod_output = tl.where(offs_n2[None, :] < KV_LEN, mask_mod_output, False) # apply mask for partial masked block post_mod_scores = tl.where(mask_mod_output, post_mod_scores, float("-inf")) # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ if not PRESCALE_QK: post_mod_scores *= RCP_LN2 p = tl.math.exp2(post_mod_scores - lse) # Compute dP and dS. # NB reversed order to since V is transposed vT = load_checked_2d(vT_ptrs, offs_v, offs_n2, None, None, SAFE_HEAD_DIM, IS_DIVISIBLE, V_HEAD_DIM, KV_LEN) dp = tl.dot(do, vT, input_precision=FLOAT32_PRECISION) ds = p * (dp - Di[:, None]) # ~~~~~~~~~~~~~~~~~~~ Apply joint modification ~~~~~~~~~~~~~~~~~~~ {{ modification( subgraph_number=1, output_name = "grad_scores", score="pre_mod_scores", b="off_z", h="off_hq", m="m", n="n", grad_score_mod="ds" ) | indent_except_first(1) }} if CHECK_BLOCK_BOUNDARY: grad_scores = tl.where(offs_n2[None, :] < KV_LEN, grad_scores, 0.0) # ~~~~~~~~~~~~~~~~~~~ Apply other buffer grad writes ~~~~~~~~~~~~~ if WRITE_DQ: scatter_mask = offs_m2[:, None] < Q_LEN and offs_n2[None, :] < KV_LEN {{ modification( subgraph_number=3, output_name=None, mask="scatter_mask", score="pre_mod_scores", b="off_z", h="off_hq", m="m", n="n", grad_score_mod="ds" ) | indent_except_first(2) }} # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ ds = grad_scores if not IS_FULL_BLOCKS: if CHECK_BLOCK_BOUNDARY: mask_mod_output = tl.where(offs_n2[None, :] < KV_LEN, mask_mod_output, False) # (grads) apply mask for partially unmasked block ds = tl.where(mask_mod_output, ds, 0.0) # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ ds = ds.to(MATMUL_PRECISION) # Compute dQ. dq += tl.dot(ds, tl.trans(kT), input_precision=FLOAT32_PRECISION) return dq @triton.jit def bwd_dkdv_inner( {{gen_argdefs()}}, Q, DO, DELTA, LSE, # pointers dk, dv, k, v, off_z, off_hq, offs_n1, offs_m1, stride_qm, stride_qd, stride_dom, stride_dod, q_indices, sparse_q_num_blocks, MATMUL_PRECISION, IS_FULL_BLOCKS, ): {{gen_defines() | indent_except_first(1) }} SPARSE_Q_MULTIPLE: tl.constexpr = (SPARSE_Q_BLOCK_SIZE // BLOCK_M1) RCP_LN2: tl.constexpr = 1.44269504 Q_LEN = {{size("Q", 2)}} KV_LEN = {{size("K", 2)}} offs_k = tl.arange(0, QK_HEAD_DIM_ROUNDED) offs_v = tl.arange(0, V_HEAD_DIM_ROUNDED) qT_ptrs = Q + offs_m1[None, :] * stride_qm + offs_k[:, None] * stride_qd do_ptrs = DO + offs_m1[:, None] * stride_dom + offs_v[None, :] * stride_dod # BLOCK_N1 must be a multiple of BLOCK_M1, otherwise the code wouldn't work. tl.static_assert(BLOCK_N1 % BLOCK_M1 == 0) hi = tl.minimum(sparse_q_num_blocks * SPARSE_Q_MULTIPLE, tl.maximum(tl.cdiv(Q_LEN, BLOCK_M1), 1)) if not IS_DIVISIBLE: if hi >= 1: for start_m in range(0, hi - 1): dk, dv = bwd_dkdv_block_mn( {{gen_argdefs()}}, dk, dv, qT_ptrs, k, v, do_ptrs, DELTA, LSE, Q_LEN, KV_LEN, off_z, off_hq, offs_n1, offs_m1, offs_k, offs_v, stride_qm, stride_qd, stride_dom, stride_dod, q_indices, sparse_q_num_blocks, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, ) # Increment pointers. offset = get_offset_for_next_block( start_m, q_indices, sparse_q_num_blocks, SPARSE_Q_BLOCK_SIZE, SPARSE_Q_MULTIPLE, BLOCK_M1, BLOCKS_ARE_CONTIGUOUS ) qT_ptrs += offset * stride_qm do_ptrs += offset * stride_dom offs_m1 += offset dk, dv = bwd_dkdv_block_mn( {{gen_argdefs()}}, dk, dv, qT_ptrs, k, v, do_ptrs, DELTA, LSE, Q_LEN, KV_LEN, off_z, off_hq, offs_n1, offs_m1, offs_k, offs_v, stride_qm, stride_qd, stride_dom, stride_dod, q_indices, sparse_q_num_blocks, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=True, ) else: for start_m in range(0, hi): dk, dv = bwd_dkdv_block_mn( {{gen_argdefs()}}, dk, dv, qT_ptrs, k, v, do_ptrs, DELTA, LSE, Q_LEN, KV_LEN, off_z, off_hq, offs_n1, offs_m1, offs_k, offs_v, stride_qm, stride_qd, stride_dom, stride_dod, q_indices, sparse_q_num_blocks, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, ) # Increment pointers. offset = get_offset_for_next_block( start_m, q_indices, sparse_q_num_blocks, SPARSE_Q_BLOCK_SIZE, SPARSE_Q_MULTIPLE, BLOCK_M1, BLOCKS_ARE_CONTIGUOUS ) qT_ptrs += offset * stride_qm do_ptrs += offset * stride_dom offs_m1 += offset return dk, dv @triton.jit def bwd_dkdv_block_mn( {{gen_argdefs()}}, dk, dv, qT_ptrs, k, v, do_ptrs, DELTA, LSE, Q_LEN, KV_LEN, off_z, off_hq, offs_n1, offs_m1, offs_k, offs_v, stride_qm, stride_qd, stride_dom, stride_dod, q_indices, sparse_q_num_blocks, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=False, ): {{gen_defines() | indent_except_first(1) }} # NB reversed order since Q is transposed qT = load_checked_2d(qT_ptrs, offs_k, offs_m1, None, None, SAFE_HEAD_DIM, IS_DIVISIBLE, QK_HEAD_DIM, Q_LEN) # Load LSE before computing qk to reduce pipeline stall. if IS_DIVISIBLE: lse = tl.load(LSE + offs_m1) else: lse = tl.load(LSE + offs_m1, mask=offs_m1 < Q_LEN) lse = tl.where(lse == -float("inf"), 0.0, lse) qkT = tl.dot(k, qT, input_precision=FLOAT32_PRECISION) if not PRESCALE_QK: qkT *= SM_SCALE # ~~~~~~~~~~~~~~~~~~~ Apply score modification ~~~~~~~~~~~~~~~~~~~ m = get_bounded_indices(offs_m1[None, :], Q_LEN if CHECK_BLOCK_BOUNDARY else None) # The boundary check is done for the outer loop, but here it's possible since we're iterating across M dim # that the n reads out of bounds prior to the last loop n = get_bounded_indices(offs_n1[:, None], KV_LEN if (not IS_DIVISIBLE or CHECK_BLOCK_BOUNDARY) else None) pre_mod_scores = qkT {{ modification( subgraph_number=0, output_name="post_mod_scores", score="qkT", b="off_z", h="off_hq", m="m", n="n", out="qkT" ) | indent_except_first(1) }} if CHECK_BLOCK_BOUNDARY: # Mask out the elements that are out of the KV_LEN for non divisible seqlen. post_mod_scores = tl.where(offs_n1[:, None] < KV_LEN, post_mod_scores, float("-inf")) if not IS_FULL_BLOCKS: {{ modification( subgraph_number=2, output_name="mask_mod_output", score="qkT", b="off_z", h="off_hq", m="m", n="n", ) | indent_except_first(2) }} if CHECK_BLOCK_BOUNDARY: mask_mod_output = tl.where(offs_n1[:, None] < KV_LEN, mask_mod_output, False) # (grads) apply mask for fully masked block post_mod_scores = tl.where(mask_mod_output, post_mod_scores, float("-inf")) # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ if not PRESCALE_QK: post_mod_scores *= RCP_LN2 pT = tl.math.exp2(post_mod_scores - lse[None, :]) do = load_checked_2d(do_ptrs, offs_m1, offs_v, None, None, IS_DIVISIBLE, SAFE_HEAD_DIM, Q_LEN, V_HEAD_DIM) # Compute dV. ppT = pT dv += tl.dot(ppT.to(MATMUL_PRECISION), do, input_precision=FLOAT32_PRECISION) if IS_DIVISIBLE: Di = tl.load(DELTA + offs_m1) else: Di = tl.load(DELTA + offs_m1, mask=offs_m1 < Q_LEN) # Compute dP and dS. dpT = tl.dot(v, tl.trans(do), input_precision=FLOAT32_PRECISION) dsT = pT * (dpT - Di[None, :]) # ~~~~~~~~~~~~~~~~~~~ Apply joint modification ~~~~~~~~~~~~~~~~~~~ {{ modification( subgraph_number=1, output_name = "grad_scores", score="pre_mod_scores", b="off_z", h="off_hq", m="m", n="n", grad_score_mod="dsT" ) | indent_except_first(1) }} # ~~~~~~~~~~~~~~~~~~~ Apply other buffer grad writes ~~~~~~~~~~~~~ if not WRITE_DQ: idx_b = off_z idx_h = off_hq idx_m = m idx_n = n scatter_mask = offs_m1[None, :] < Q_LEN and offs_n1[:, None] < KV_LEN {{ modification( subgraph_number=3, output_name=None, mask="scatter_mask", score="pre_mod_scores", b="idx_b", h="idx_h", m="idx_m", n="idx_n", grad_score_mod="dsT" ) | indent_except_first(2) }} # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ if CHECK_BLOCK_BOUNDARY: grad_scores = tl.where(offs_n1[:, None] < KV_LEN, grad_scores, 0.0) dsT = grad_scores if not IS_FULL_BLOCKS: if CHECK_BLOCK_BOUNDARY: mask_mod_output = tl.where(offs_n1[:, None] < KV_LEN, mask_mod_output, False) # (grads) apply mask for partially unmasked block dsT = tl.where(mask_mod_output, dsT, 0.0) # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ dk += tl.dot(dsT.to(MATMUL_PRECISION), tl.trans(qT), input_precision=FLOAT32_PRECISION) return dk, dv joint_graphc|jD]y}|jdk(s|jtjj j jk(sO|jD]}|jdk7std{y)zKWe do some pre lowering graph checks in order to raise nicer error messagesrraSUsing multiple indexing operations on the same tensor that requires gradients in a score_mod function is not currently supported. This typically happens when indexing the same tensor multiple times, like: def score_mod(score, b, h, q_idx, kv_idx): return score + bias[q_idx] + bias[kv_idx] # bias used twice! A valid workaround is to clone() the tensors that will be indexed multiple times. For example: bias1 = bias.clone() def score_mod(score, b, h, q_idx, kv_idx): return score + bias[q_idx] + bias1[kv_idx] Note that this solution will use additional memory.N) rrtargetr_rrrrusersrF)rr)users r5validate_joint_graphrs|!! GG & uyy11CCKKK 77h&- N    & r7T)frozencReZdZUdZeed<eeed<eeeed<eeed<y)JointOutputResultz&Results from processing joint outputs. grad_inputcaptured_grads_computecaptured_grads mutated_gradsN) rrr__doc__r__annotations__rvrrrAr7r5rr s00 00),-- ?"r7rall_joint_outputsnum_placeholdersc t|tsJ|dJd|d}||dz d}|Dcgc]}|| }}d}|Dcgc] }|| }}|Dcgc]}|| } }t|||| Scc}wcc}wcc}w)a^Process joint outputs and extract various buffers needed for lowering Args: all_joint_outputs: List of all the outputs from build_subgraphs num_placeholders: The number of placeholder inputs, used to skip over unused backward compute buffers Returns: JointOutputResult containing processed buffers and gradients rNz.joint_subgraph_buffer is None - this is a bug!r0c|yt|tsJ|jJtjt j j|jSrU)rVrrIrrOr r; get_buffer)r}s r5get_outz&process_joint_outputs..get_out- sN ;#~...xx### 2 2388 <==r7)rrrr)rVrvr) rr joint_buffer other_gradsr} grads_computerrY grads_outrs r5process_joint_outputsr s ' .. . Q  +8 +%Q'L#$4q$8$:;K%0CS3?SCMC>&111I1$-ASSAMA , #  D2AsA7A7A<B"BcM|\M}}}}}}}} } } } } }| \ }}}}}}}}}}}}}tM|||||||||||g \ M}}}}}}}}}}}Mj}Mj}Mj\}}}} |j\}!}"}#}$tj j jtj||!tj|!dzs Jd|d|!t| } | jD%&cic]K\}%}&|%t|&tjr)tj j j|&n|&M} }%}&| jdt!|dzdk7s|#dzdk7r| jddn| jdd d |fd t"j$fd t"j$fd t"j$fdt"j$ffD'cgc]\}'}t'|'||}(}'}t)|(t+| z|})|(t'd|gz}*| j,j j/t1| j,j t)|*t+| z| }+t3|+t5|*},d t"j$fd t"j$fd t"j$fdt"j$ffD'cgc] \}'}t'|'|Mj"}-}'}t)|-t+|z|}.|.}.||"|#| g}/t7|/|j9}0t;|j|j|/|0D1cgc]}1tj<|1c}1}2t?t@jB|dtEjFdz }3t?t@jB||}4t?t@jH|4d}5t?t@jJ|5|3}5tMjN|5}5t|3|5g\}3}5|||| g}6t7|6Mj9}7tQ|6|7D1cgc]}1tj<|1c}1MjMj}8||"|#|$g}9t7|9|j9}:tQ|9|:D1cgc]}1tj<|1c}1|j|j};| jd| ||"z}<| jd|<|du}=| jd|=|=sMfdtSdD\}}}}tU| | |$tj j tj j j|}tj j j|}g}>g}?|?jWtYMtZj\rvt"j^j`gdndg}@|?jcdDABC1cgc]4}AdD]-}BAdk\sBdk\rddgndgD]}C@D]}1BAzdk(rABC|1f/6c}1}C}B}A| je}D|?D]\}A}B}E}F||Azdk7s|Azdk7s|Bzdk7s|Bzdk7r*Dje}Gt+|GjgD]O}%|%jidrGjk|%}&|&|G|%dd<|%jids?Gjk|%QGjd E|Gjd!F|Gjd"A|Gjd#B|Gjd$|B|Gjd%|A|Gjd&||Gjd'|tmjnd.|>M||||5||8|;||||||||g|2|)|,jp|.|,jrg|8|;g|,jtMj|jdd(zd)|GM||||5||8|;||||||||gt+| zt+|z|,jtz}Htw|txtw|txtw|txtw|txd*}It{d+|>|H|2|I,}Jtj j jtj||!rJ}K|;}Lntj j jtj||dtj|!dzs[Jdtj j j|dtj j j|!t?t@jHJdd -}Kt?t@jH|;dd -}L|8KLt|,jfScc}&}%wcc}}'wcc}}'wcc}1wcc}1wcc}1wcc}1}C}B}Aw)/Nr0rrrrrrFTr r r rrgrad_score_modr3rr)axis)r4rJrKrrrc3TK|]}tdj!ywrrrs r5rz*flex_attention_backward.. s+R 45E!E,,. / /R rr)r0rrr)rr)rrrrrrrrBLOCK_M1rrBLOCK_N2rerfrr)rr rr)rkeepdimsrA)Ar\rprtrsr r;r<rrWrrrrVrXrrrer_rrQrrvreliminate_dead_coderrr1r>rRrrSr%r mulrlogsumsubrrequire_contiguousr"rwrrMrrrrarbextendrrrr flex_attention_backward_templaterrrrrrr(rtupler)NrRrrYrZoutrgrad_outgrad_logsumexpfw_graphrr[r;rr\r]r7r^rr_r`rarbrcrdrerfr%rKrJrorprqrrrsrtrurvrrrIfwd_placeholder_inpsfw_subgraph_bufferjoint_placeholder_inpsr joint_outputsrirkkey_size key_stridesrzlayout_broadcasted_k grad_lse_exp2 mul_deltadelta query_sizegrad_query_strides grad_query value_size value_stridesbroadcasted_grad_valuerrr9rnum_stages_listBLOCK1BLOCK2wrrrrrrbroadcasted_grad_keygrad_key grad_valuersN @r5rr@ s &                            $   F OO E%*^^%5"BI{',~~'7$Cj* 77   ) )%((2s*;ehhsA>N*N O 0IcUC O.)N #((*  Aq a & 77   1 1! 4  N 13H3JK3!zC/14!!.%8!!.$7 e  %++  %++  %++  %++     D% 4/  /t$;<> *356M%++  %++  %++  %++   # D% 4(8(8(:;##.#d+A&BBJ*C[1H%h0@AK&  *56Q a 6 dhh'DHHQKHM$((#C2I dhh   3E dhh } 5E  + +E 2E(-)?@M5b)[1J,Z9I9I9KL*<=Q a =oo! Jc:z2J' E4D4D4FGM**78Q a 8oo! j%0Sy02BC)4O/A R 9>qR NO-> ZAQAQR''**@@ATU77++AABVWG/1G NN*512 */--*;*;*C,!'  + $*cMVs]1a&  (   F?a' A& & & &  -11318@ - : 6 )Q ."V+q0#f,1"V+q0 599;(--/0 *A||F#&**1-,-"1QR5)||F#"&&q)  * %%k9=%%lJ?%%j&9%%j&9%%j&9%%j&9%%&;=PQ%%&<>RS(<<" &"!!$("((!44 &,, ~~'#,,.1*==A" B!C" =@ F        "        ! $ & '% (& % &' '(  % %) &0 ,J 7  ,Y 7  ,_ =  ,^ <  M5!#  ww%%ehhr3&78'+ ww--ehhr1oa@P.PQ gg&&44R89:ww''55c:; = Q TXX&';!dStxx()?aRVW *eM4P4P.Q RRA B#.7">9> s+Ah2>h8:%h>ii .i #9irU)}rrloggingrcollections.abcr dataclassesrenumrrtypingrrr rWr_torch._inductor.virtualizedr torch.utils._ordered_setr torch.utils._pytreer torch.utils._sympy.numbersrtorch.utils._sympy.value_rangesrrirrrrrrrrrrrrrrloweringrr r!r"r#r$r%r&r'select_algorithmr(r)r*r+ getLoggerrrrr Exprrr6r>rHstrrJrKrvrQr\rerSubgraphResultsr#r$rrcompute_next_offset_funcget_bounded_indices_funcload_checked_blockload_checked_2dcompute_flex_attentioncompute_forward_innercompute_forward_block_mnrrrqrPrQrrrrrrrrrrrr$torch._inductor.kernel.flex_decodingr#codegen.cpp_flex_attention_templaterrrrrrrr higher_orderrrrGraphrrrrrAr7r5r?s8  $!'' )/(-7   g! yy~~ zz C= c], /hsm /8C= /II!% * * ;;* LL* 49  *  *& Xf-. )d3i)XXn568PPQ6N y/6N).)=)=6N6NrEYE8EE $ 8HVDNZz)   !         $P ]]B ]]C. ]]C. ^^R* ^^S? ^^S> ]]B) ]]C* ]]C.  ]]B ]]C/ ]]C. ^^R/ ^^S? ^^S> ]]B ]]C/ ]]C.  ]]B ]]C/ ]]C. ^^R/ ^^S? ^^S> ]]B ]]C/ ]]C. 4 "$"5c31C+D"B..sCc/A)B.b6eCc3,>&?66eCc3,>&?6"0ellMJ  \~+( cND599))88dSlTld $$2 " %l Z[m \]n ^_o s $ l ehhnn 0 $###'&':=''V II22JSJSr7