Coverage for src/flag_gems/fused/FLA/chunk_delta_h.py: 9%
163 statements
« prev ^ index » next coverage.py v7.6.9, created at 2026-03-20 02:31 +0800
1# This file contains code copied from the flash-linear-attention project.
2# The original source code was licensed under the MIT license and included
3# the following copyright notice:
4# Copyright (c) 2023-2025, Songlin Yang, Yu Zhang
5# ruff: noqa: E501
7import torch
8import triton
9import triton.language as tl
11from flag_gems.fused.FLA.index import prepare_chunk_indices, prepare_chunk_offsets
12from flag_gems.fused.FLA.triton_ops_helper import exp
13from flag_gems.fused.FLA.utils import use_cuda_graph
14from flag_gems.utils import libentry, libtuner
16NUM_WARPS = [2, 4, 8, 16]
19@libentry()
20@triton.heuristics(
21 {
22 "USE_G": lambda args: args["g"] is not None,
23 "USE_GK": lambda args: args["gk"] is not None,
24 "USE_INITIAL_STATE": lambda args: args["h0"] is not None,
25 "STORE_FINAL_STATE": lambda args: args["ht"] is not None,
26 "SAVE_NEW_VALUE": lambda args: args["v_new"] is not None,
27 "IS_VARLEN": lambda args: args["cu_seqlens"] is not None,
28 }
29)
30@libtuner(
31 configs=[
32 triton.Config({"BV": BV}, num_warps=num_warps, num_stages=num_stages)
33 for num_warps in [2, 4]
34 for num_stages in [2, 3, 4]
35 for BV in [32, 64]
36 ],
37 key=["H", "K", "V", "BT"],
38 use_cuda_graph=use_cuda_graph,
39)
40@triton.jit(do_not_specialize=["T"])
41def chunk_gated_delta_rule_fwd_kernel_h_blockdim64(
42 k,
43 v,
44 w,
45 v_new,
46 g,
47 gk,
48 h,
49 h0,
50 ht,
51 cu_seqlens,
52 chunk_offsets,
53 T,
54 H: tl.constexpr,
55 Hg: tl.constexpr,
56 K: tl.constexpr,
57 V: tl.constexpr,
58 BT: tl.constexpr,
59 BV: tl.constexpr,
60 USE_G: tl.constexpr,
61 USE_GK: tl.constexpr,
62 USE_INITIAL_STATE: tl.constexpr,
63 STORE_FINAL_STATE: tl.constexpr,
64 SAVE_NEW_VALUE: tl.constexpr,
65 IS_VARLEN: tl.constexpr,
66):
67 i_v, i_nh = tl.program_id(0), tl.program_id(1)
68 i_n, i_h = i_nh // H, i_nh % H
69 if IS_VARLEN:
70 bos, eos = (
71 tl.load(cu_seqlens + i_n).to(tl.int32),
72 tl.load(cu_seqlens + i_n + 1).to(tl.int32),
73 )
74 T = eos - bos
75 NT = tl.cdiv(T, BT)
76 boh = tl.load(chunk_offsets + i_n).to(tl.int32)
77 else:
78 bos, eos = i_n * T, i_n * T + T
79 NT = tl.cdiv(T, BT)
80 boh = i_n * NT
82 # [BK, BV]
83 b_h1 = tl.zeros([64, BV], dtype=tl.float32)
84 if K > 64:
85 b_h2 = tl.zeros([64, BV], dtype=tl.float32)
86 if K > 128:
87 b_h3 = tl.zeros([64, BV], dtype=tl.float32)
88 if K > 192:
89 b_h4 = tl.zeros([64, BV], dtype=tl.float32)
91 # calculate offset
92 h += ((boh * H + i_h) * K * V).to(tl.int64)
93 v += ((bos * H + i_h) * V).to(tl.int64)
94 k += ((bos * Hg + i_h // (H // Hg)) * K).to(tl.int64)
95 w += ((bos * H + i_h) * K).to(tl.int64)
96 if SAVE_NEW_VALUE:
97 v_new += ((bos * H + i_h) * V).to(tl.int64)
98 stride_v = H * V
99 stride_h = H * K * V
100 stride_k = Hg * K
101 stride_w = H * K
102 if USE_INITIAL_STATE:
103 h0 = h0 + i_nh * K * V
104 if STORE_FINAL_STATE:
105 ht = ht + i_nh * K * V
107 # load initial state
108 if USE_INITIAL_STATE:
109 p_h0_1 = tl.make_block_ptr(h0, (K, V), (V, 1), (0, i_v * BV), (64, BV), (1, 0))
110 b_h1 += tl.load(p_h0_1, boundary_check=(0, 1)).to(tl.float32)
111 if K > 64:
112 p_h0_2 = tl.make_block_ptr(
113 h0, (K, V), (V, 1), (64, i_v * BV), (64, BV), (1, 0)
114 )
115 b_h2 += tl.load(p_h0_2, boundary_check=(0, 1)).to(tl.float32)
116 if K > 128:
117 p_h0_3 = tl.make_block_ptr(
118 h0, (K, V), (V, 1), (128, i_v * BV), (64, BV), (1, 0)
119 )
120 b_h3 += tl.load(p_h0_3, boundary_check=(0, 1)).to(tl.float32)
121 if K > 192:
122 p_h0_4 = tl.make_block_ptr(
123 h0, (K, V), (V, 1), (192, i_v * BV), (64, BV), (1, 0)
124 )
125 b_h4 += tl.load(p_h0_4, boundary_check=(0, 1)).to(tl.float32)
127 # main recurrence
128 for i_t in range(NT):
129 p_h1 = tl.make_block_ptr(
130 h + i_t * stride_h, (K, V), (V, 1), (0, i_v * BV), (64, BV), (1, 0)
131 )
132 tl.store(p_h1, b_h1.to(p_h1.dtype.element_ty), boundary_check=(0, 1))
133 if K > 64:
134 p_h2 = tl.make_block_ptr(
135 h + i_t * stride_h, (K, V), (V, 1), (64, i_v * BV), (64, BV), (1, 0)
136 )
137 tl.store(p_h2, b_h2.to(p_h2.dtype.element_ty), boundary_check=(0, 1))
138 if K > 128:
139 p_h3 = tl.make_block_ptr(
140 h + i_t * stride_h, (K, V), (V, 1), (128, i_v * BV), (64, BV), (1, 0)
141 )
142 tl.store(p_h3, b_h3.to(p_h3.dtype.element_ty), boundary_check=(0, 1))
143 if K > 192:
144 p_h4 = tl.make_block_ptr(
145 h + i_t * stride_h, (K, V), (V, 1), (192, i_v * BV), (64, BV), (1, 0)
146 )
147 tl.store(p_h4, b_h4.to(p_h4.dtype.element_ty), boundary_check=(0, 1))
149 p_w = tl.make_block_ptr(
150 w, (T, K), (stride_w, 1), (i_t * BT, 0), (BT, 64), (1, 0)
151 )
152 b_w = tl.load(p_w, boundary_check=(0, 1))
153 b_v = tl.dot(b_w, b_h1.to(b_w.dtype))
154 if K > 64:
155 p_w = tl.make_block_ptr(
156 w, (T, K), (stride_w, 1), (i_t * BT, 64), (BT, 64), (1, 0)
157 )
158 b_w = tl.load(p_w, boundary_check=(0, 1))
159 b_v += tl.dot(b_w, b_h2.to(b_w.dtype))
160 if K > 128:
161 p_w = tl.make_block_ptr(
162 w, (T, K), (stride_w, 1), (i_t * BT, 128), (BT, 64), (1, 0)
163 )
164 b_w = tl.load(p_w, boundary_check=(0, 1))
165 b_v += tl.dot(b_w, b_h3.to(b_w.dtype))
166 if K > 192:
167 p_w = tl.make_block_ptr(
168 w, (T, K), (stride_w, 1), (i_t * BT, 192), (BT, 64), (1, 0)
169 )
170 b_w = tl.load(p_w, boundary_check=(0, 1))
171 b_v += tl.dot(b_w, b_h4.to(b_w.dtype))
172 p_v = tl.make_block_ptr(
173 v, (T, V), (stride_v, 1), (i_t * BT, i_v * BV), (BT, BV), (1, 0)
174 )
175 b_v = tl.load(p_v, boundary_check=(0, 1)) - b_v
177 if SAVE_NEW_VALUE:
178 p_v = tl.make_block_ptr(
179 v_new, (T, V), (stride_v, 1), (i_t * BT, i_v * BV), (BT, BV), (1, 0)
180 )
181 tl.store(p_v, b_v.to(p_v.dtype.element_ty), boundary_check=(0, 1))
183 last_idx = min((i_t + 1) * BT, T) - 1
184 if USE_G:
185 m_t = (i_t * BT + tl.arange(0, BT)) < T
186 b_g_last = tl.load(g + bos * H + last_idx * H + i_h)
187 p_g = tl.make_block_ptr(
188 g + bos * H + i_h, (T,), (H,), (i_t * BT,), (BT,), (0,)
189 )
190 b_g = tl.load(p_g, boundary_check=(0,))
191 b_v = b_v * tl.where(m_t, exp(b_g_last - b_g), 0)[:, None]
192 b_g_last = exp(b_g_last)
193 b_h1 *= b_g_last
194 if K > 64:
195 b_h2 *= b_g_last
196 if K > 128:
197 b_h3 *= b_g_last
198 if K > 192:
199 b_h4 *= b_g_last
201 if USE_GK:
202 o_k1 = tl.arange(0, 64)
203 b_gk_last1 = tl.load(
204 gk + (bos + last_idx) * H * K + i_h * K + o_k1,
205 mask=(o_k1 < K),
206 other=0.0,
207 )
208 b_h1 *= exp(b_gk_last1)[:, None]
209 if K > 64:
210 o_k2 = 64 + o_k1
211 b_gk_last2 = tl.load(
212 gk + (bos + last_idx) * H * K + i_h * K + o_k2,
213 mask=(o_k2 < K),
214 other=0.0,
215 )
216 b_h2 *= exp(b_gk_last2)[:, None]
217 if K > 128:
218 o_k3 = 128 + o_k1
219 b_gk_last3 = tl.load(
220 gk + (bos + last_idx) * H * K + i_h * K + o_k3,
221 mask=(o_k3 < K),
222 other=0.0,
223 )
224 b_h3 *= exp(b_gk_last3)[:, None]
225 if K > 192:
226 o_k4 = 192 + o_k1
227 b_gk_last4 = tl.load(
228 gk + (bos + last_idx) * H * K + i_h * K + o_k4,
229 mask=(o_k4 < K),
230 other=0.0,
231 )
232 b_h4 *= exp(b_gk_last4)[:, None]
233 b_v = b_v.to(k.dtype.element_ty)
235 p_k = tl.make_block_ptr(
236 k, (K, T), (1, stride_k), (0, i_t * BT), (64, BT), (0, 1)
237 )
238 b_k = tl.load(p_k, boundary_check=(0, 1))
239 b_h1 += tl.dot(b_k, b_v)
240 if K > 64:
241 p_k = tl.make_block_ptr(
242 k, (K, T), (1, stride_k), (64, i_t * BT), (64, BT), (0, 1)
243 )
244 b_k = tl.load(p_k, boundary_check=(0, 1))
245 b_h2 += tl.dot(b_k, b_v)
246 if K > 128:
247 p_k = tl.make_block_ptr(
248 k, (K, T), (1, stride_k), (128, i_t * BT), (64, BT), (0, 1)
249 )
250 b_k = tl.load(p_k, boundary_check=(0, 1))
251 b_h3 += tl.dot(b_k, b_v)
252 if K > 192:
253 p_k = tl.make_block_ptr(
254 k, (K, T), (1, stride_k), (192, i_t * BT), (64, BT), (0, 1)
255 )
256 b_k = tl.load(p_k, boundary_check=(0, 1))
257 b_h4 += tl.dot(b_k, b_v)
258 # epilogue
259 if STORE_FINAL_STATE:
260 p_ht = tl.make_block_ptr(ht, (K, V), (V, 1), (0, i_v * BV), (64, BV), (1, 0))
261 tl.store(p_ht, b_h1.to(p_ht.dtype.element_ty), boundary_check=(0, 1))
262 if K > 64:
263 p_ht = tl.make_block_ptr(
264 ht, (K, V), (V, 1), (64, i_v * BV), (64, BV), (1, 0)
265 )
266 tl.store(p_ht, b_h2.to(p_ht.dtype.element_ty), boundary_check=(0, 1))
267 if K > 128:
268 p_ht = tl.make_block_ptr(
269 ht, (K, V), (V, 1), (128, i_v * BV), (64, BV), (1, 0)
270 )
271 tl.store(p_ht, b_h3.to(p_ht.dtype.element_ty), boundary_check=(0, 1))
272 if K > 192:
273 p_ht = tl.make_block_ptr(
274 ht, (K, V), (V, 1), (192, i_v * BV), (64, BV), (1, 0)
275 )
276 tl.store(p_ht, b_h4.to(p_ht.dtype.element_ty), boundary_check=(0, 1))
279def chunk_gated_delta_rule_fwd_h(
280 k: torch.Tensor,
281 w: torch.Tensor,
282 u: torch.Tensor,
283 g: torch.Tensor | None = None,
284 gk: torch.Tensor | None = None,
285 initial_state: torch.Tensor | None = None,
286 output_final_state: bool = False,
287 chunk_size: int = 64, # SY: remove this argument and force chunk size 64?
288 save_new_value: bool = True,
289 cu_seqlens: torch.LongTensor | None = None,
290) -> tuple[torch.Tensor, torch.Tensor]:
291 # This kernel is slightly different from fla to support Q/K with different head numbers.
292 # In fla, Q/K always have the same head number, so Hg is always equal to H.
293 B, T, Hg, K, V = *k.shape, u.shape[-1]
294 H = u.shape[-2]
295 BT = chunk_size
297 chunk_indices = (
298 prepare_chunk_indices(cu_seqlens, chunk_size)
299 if cu_seqlens is not None
300 else None
301 )
302 # N: the actual number of sequences in the batch with either equal or variable lengths
303 if cu_seqlens is None:
304 N, NT, chunk_offsets = B, triton.cdiv(T, BT), None
305 else:
306 N, NT, chunk_offsets = (
307 len(cu_seqlens) - 1,
308 len(chunk_indices),
309 prepare_chunk_offsets(cu_seqlens, BT),
310 )
311 assert K <= 256, "current kernel does not support head dimension larger than 256."
313 h = k.new_empty(B, NT, H, K, V)
314 final_state = (
315 k.new_empty(N, H, K, V, dtype=torch.float32) if output_final_state else None
316 )
318 v_new = torch.empty_like(u) if save_new_value else None
320 def grid(meta):
321 return (triton.cdiv(V, meta["BV"]), N * H)
323 chunk_gated_delta_rule_fwd_kernel_h_blockdim64[grid](
324 k=k,
325 v=u,
326 w=w,
327 v_new=v_new,
328 g=g,
329 gk=gk,
330 h=h,
331 h0=initial_state,
332 ht=final_state,
333 cu_seqlens=cu_seqlens,
334 chunk_offsets=chunk_offsets,
335 T=T,
336 H=H,
337 Hg=Hg,
338 K=K,
339 V=V,
340 BT=BT,
341 )
342 return h, v_new, final_state