from typing import Optional, Tuple import tilelang import tilelang.language as T import torch from sglang.srt.layers.quantization.fp8_kernel import is_fp8_fnuz from sglang.srt.utils import is_gfx95_supported, is_hip tilelang.set_log_level("WARNING") pass_configs = { tilelang.PassConfigKey.TL_DISABLE_WARP_SPECIALIZED: True, tilelang.PassConfigKey.TL_DISABLE_TMA_LOWER: True, } # TL_DISABLE_FAST_MATH has deprecated in v0.1.7.post1 tilelang if hasattr(tilelang.PassConfigKey, "TL_DISABLE_FAST_MATH"): pass_configs[tilelang.PassConfigKey.TL_DISABLE_FAST_MATH] = True elif hasattr(tilelang.PassConfigKey, "TL_ENABLE_FAST_MATH"): pass_configs[tilelang.PassConfigKey.TL_ENABLE_FAST_MATH] = False _is_hip = is_hip() _is_gfx95_supported = is_gfx95_supported() _is_fp8_fnuz = is_fp8_fnuz() BF16 = "bfloat16" FP8 = "float8_e4m3fnuz" if _is_fp8_fnuz else "float8_e4m3" FP32 = "float32" def fast_log2_ceil(x): bits_x = T.reinterpret("uint32", x) exp_x = (bits_x >> 23) & 0xFF man_bits = bits_x & ((1 << 23) - 1) return T.Cast("int32", exp_x - 127 + T.if_then_else(man_bits != 0, 1, 0)) def fast_pow2(x): bits_x = (x + 127) << 23 return T.reinterpret("float32", bits_x) def fast_round_scale(amax, fp8_max_inv): return fast_pow2(fast_log2_ceil(amax * fp8_max_inv)) @tilelang.jit(pass_configs=pass_configs) def act_quant_kernel( N, in_dtype=BF16, out_dtype=FP8, scale_dtype=FP32, round_scale=False ): M = T.symbolic("M") fp8_min = -224.0 if _is_fp8_fnuz else -448.0 fp8_max = 224.0 if _is_fp8_fnuz else 448.0 fp8_max_inv = 1 / fp8_max num_stages = 0 if round_scale else 2 blk_m = 32 group_size = 128 @T.prim_func def act_quant_kernel_( X: T.Tensor[(M, N), in_dtype], Y: T.Tensor[(M, N), out_dtype], S: T.Tensor[(M, T.ceildiv(N, group_size)), scale_dtype], ): with T.Kernel(T.ceildiv(M, blk_m), T.ceildiv(N, group_size), threads=128) as ( pid_m, pid_n, ): x_shared = T.alloc_shared((blk_m, group_size), in_dtype) x_local = T.alloc_fragment((blk_m, group_size), in_dtype) amax_local = T.alloc_fragment((blk_m,), scale_dtype) s_local = T.alloc_fragment((blk_m,), scale_dtype) y_local = T.alloc_fragment((blk_m, group_size), out_dtype) y_shared = T.alloc_shared((blk_m, group_size), out_dtype) for _ in T.Pipelined(1, num_stages=num_stages): T.copy(X[pid_m * blk_m, pid_n * group_size], x_shared) T.copy(x_shared, x_local) T.reduce_absmax(x_local, amax_local, dim=1) for i in T.Parallel(blk_m): amax_local[i] = T.max(amax_local[i], 1e-4) if round_scale: s_local[i] = fast_round_scale(amax_local[i], fp8_max_inv) else: s_local[i] = amax_local[i] * fp8_max_inv for i, j in T.Parallel(blk_m, group_size): y_local[i, j] = T.clamp( x_local[i, j] / s_local[i], fp8_min, fp8_max ) for i in T.Parallel(blk_m): S[pid_m * blk_m + i, pid_n] = s_local[i] T.copy(y_local, y_shared) T.copy(y_shared, Y[pid_m * blk_m, pid_n * group_size]) return act_quant_kernel_ def act_quant( x: torch.Tensor, block_size: int = 128, scale_fmt: Optional[str] = None ) -> Tuple[torch.Tensor, torch.Tensor]: """ 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`. """ assert x.is_contiguous(), "Input tensor must be contiguous" assert ( x.size(-1) % block_size == 0 ), f"Last dimension size must be divisible by block_size (block_size={block_size})" N = x.size(-1) if _is_fp8_fnuz: y = torch.empty_like(x, dtype=torch.float8_e4m3fnuz) else: y = torch.empty_like(x, dtype=torch.float8_e4m3fn) s = x.new_empty(*x.size()[:-1], N // block_size, dtype=torch.float32) kernel = act_quant_kernel(N, round_scale=scale_fmt is not None) kernel(x.view(-1, N), y.view(-1, N), s.view(-1, N // block_size)) return y, s @tilelang.jit(out_idx=[4], pass_configs=pass_configs) def fp8_index_kernel(h: int, d: int, clear_accum=True): b = T.symbolic("b") m = T.symbolic("m") n = T.symbolic("n") blk_n1 = 512 blk_n2 = 128 @T.prim_func def fp8_index_kernel_( q: T.Tensor[(b, m, h, d), FP8], q_s: T.Tensor[(b, m, h), FP32], k: T.Tensor[(b, n, d), FP8], k_s: T.Tensor[(b, n), FP32], o: T.Tensor[(b, m, n), FP32], ) -> None: with T.Kernel(b, m, T.ceildiv(n, blk_n1)) as (i_b, i_m, i1_n): q_smem = T.alloc_shared((h, d), FP8) T.copy(q[i_b, i_m, 0, 0], q_smem) q_s_frag = T.alloc_fragment(h, FP32) T.copy(q_s[i_b, i_m, 0], q_s_frag) for i2_n in T.Pipelined(blk_n1 // blk_n2, num_stages=2): k_smem = T.alloc_shared((blk_n2, d), FP8) T.copy(k[i_b, i1_n * blk_n1 + i2_n * blk_n2, 0], k_smem) k_s_frag = T.alloc_fragment(blk_n2, FP32) T.copy(k_s[i_b, i1_n * blk_n1 + i2_n * blk_n2], k_s_frag) logits = T.alloc_fragment((blk_n2, h), FP32) if not clear_accum: T.fill(logits, 0) T.gemm( k_smem, q_smem, logits, transpose_A=False, transpose_B=True, clear_accum=clear_accum, ) for i_h, i3_n in T.Parallel(h, blk_n2): logits[i3_n, i_h] = T.max(logits[i3_n, i_h], 0) * q_s_frag[i_h] logits_sum = T.alloc_fragment(blk_n2, FP32) T.reduce_sum(logits, logits_sum, dim=1) for i3_n in T.Parallel(blk_n2): logits_sum[i3_n] *= k_s_frag[i3_n] T.copy(logits_sum, o[i_b, i_m, i1_n * blk_n1 + i2_n * blk_n2]) return fp8_index_kernel_ def fp8_index( q: torch.Tensor, q_s: torch.Tensor, k: torch.Tensor, k_s: torch.Tensor, ) -> torch.Tensor: """ 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 """ if _is_hip: return fp8_index_kernel(q.shape[2], q.shape[3], False)(q, q_s, k, k_s) else: return fp8_index_kernel(q.shape[2], q.shape[3])(q, q_s, k, k_s) @tilelang.jit( out_idx=[-1], pass_configs={ tilelang.PassConfigKey.TL_DISABLE_TMA_LOWER: True, tilelang.PassConfigKey.TL_DISABLE_WARP_SPECIALIZED: True, }, ) def sparse_attention_fwd_kernel_v1( num_heads, dim, tail_dim, topk, *, kv_group=1, sm_scale=None, is_causal=True, block_I=64, num_stages=2, threads=256, ): assert dim == tilelang.math.next_power_of_2( dim ), f"haven't check padding correctness yet, dim={dim}" assert tail_dim == tilelang.math.next_power_of_2( tail_dim ), f"haven't check padding correctness yet, dim={tail_dim}" assert is_causal == True, "non-casual is not supported" assert ( topk % block_I == 0 ), "otherwise will load some index=0 thus causing wrong kv to be loaded" if sm_scale is None: sm_scale = (1.0 / (dim + tail_dim)) ** 0.5 * 1.44269504 # log2(e) else: sm_scale = sm_scale * 1.44269504 # log2(e) batch = T.symbolic("batch") seq_len = T.symbolic("seq_len") seq_len_kv = T.symbolic("seq_len_kv") head_kv = num_heads // kv_group q_shape = [batch, seq_len, num_heads, dim + tail_dim] kv_shape = [batch, seq_len_kv, kv_group, dim + tail_dim] o_shape = [batch, seq_len, num_heads, dim] indices_shape = [batch, seq_len, kv_group, topk] indices_dtype = "int32" dtype = "bfloat16" accum_dtype = "float" H = head_kv padded_H = max(tilelang.math.next_power_of_2(head_kv), 16) if padded_H != H: assert kv_group == 1 BI = block_I NI = tilelang.cdiv(topk, block_I) D = dim D_tail = tail_dim if head_kv > 64: assert head_kv % 64 == 0, "head_kv should be a multiple of 64" REPLICATE_H = head_kv // 64 else: REPLICATE_H = 1 H_per_block = padded_H if REPLICATE_H == 1 else 64 @T.prim_func def main( Q: T.Tensor(q_shape, dtype), # type: ignore KV: T.Tensor(kv_shape, dtype), # type: ignore Indices: T.Tensor(indices_shape, indices_dtype), # type: ignore Output: T.Tensor(o_shape, dtype), # type: ignore ): with T.Kernel(seq_len * REPLICATE_H, batch, kv_group, threads=threads) as ( bx, by, bz, ): Q_shared = T.alloc_shared([H_per_block, D], dtype) Q_tail_shared = T.alloc_shared([H_per_block, D_tail], dtype) KV_shared = T.alloc_shared([BI, D], dtype) K_tail_shared = T.alloc_shared([BI, D_tail], dtype) O_shared = T.alloc_shared([H_per_block, D], dtype) mask = T.alloc_fragment([BI], "bool") acc_o = T.alloc_fragment([H_per_block, D], accum_dtype) acc_s = T.alloc_fragment([H_per_block, BI], accum_dtype) S_shared = T.alloc_shared([H_per_block, BI], dtype) sumexp = T.alloc_fragment([H_per_block], accum_dtype) sumexp_i = T.alloc_fragment([H_per_block], accum_dtype) alpha = T.alloc_fragment([H_per_block], accum_dtype) m_i = T.alloc_fragment([H_per_block], accum_dtype) m_i_prev = T.alloc_fragment([H_per_block], accum_dtype) T.fill(acc_o, 0) T.fill(sumexp, 0) T.fill(m_i, -(2**30)) # avoid -inf - inf to cause nan b_i, g_i = by, bz s_i = bx if REPLICATE_H == 1 else (bx // REPLICATE_H) q_i = s_i max_kv_i = q_i H0 = g_i * padded_H + (0 if REPLICATE_H == 1 else (bx % REPLICATE_H) * 64) H1 = H0 + H_per_block T.copy(Q[b_i, s_i, H0:H1, :D], Q_shared) T.copy(Q[b_i, s_i, H0:H1, D:], Q_tail_shared) for i_i in T.Pipelined(NI, num_stages=num_stages): for bi_i in T.Parallel(BI): mask[bi_i] = Indices[b_i, s_i, g_i, i_i * BI + bi_i] >= 0 for bi_i, d_i in T.Parallel(BI, D): KV_shared[bi_i, d_i] = KV[ b_i, Indices[b_i, s_i, g_i, i_i * BI + bi_i], g_i, d_i ] for bi_i, d_i in T.Parallel(BI, D_tail): K_tail_shared[bi_i, d_i] = KV[ b_i, Indices[b_i, s_i, g_i, i_i * BI + bi_i], g_i, D + d_i ] for h_i, bi_i in T.Parallel(H_per_block, BI): acc_s[h_i, bi_i] = T.if_then_else( mask[bi_i], 0, -T.infinity(acc_s.dtype) ) T.gemm( Q_shared, KV_shared, acc_s, transpose_B=True, policy=T.GemmWarpPolicy.FullCol, ) T.gemm( Q_tail_shared, K_tail_shared, acc_s, transpose_B=True, policy=T.GemmWarpPolicy.FullCol, ) T.copy(m_i, m_i_prev) T.reduce_max(acc_s, m_i, dim=1, clear=False) for h_i in T.Parallel(H_per_block): alpha[h_i] = T.exp2((m_i_prev[h_i] - m_i[h_i]) * sm_scale) for h_i, bi_i in T.Parallel(H_per_block, BI): acc_s[h_i, bi_i] = T.exp2( acc_s[h_i, bi_i] * sm_scale - m_i[h_i] * sm_scale ) T.reduce_sum(acc_s, sumexp_i, dim=1) # is this a accumulate operator? for h_i in T.Parallel(H_per_block): sumexp[h_i] = sumexp[h_i] * alpha[h_i] + sumexp_i[h_i] for h_i, d_i in T.Parallel(H_per_block, D): acc_o[h_i, d_i] = acc_o[h_i, d_i] * alpha[h_i] T.copy(acc_s, S_shared) T.gemm(S_shared, KV_shared, acc_o, policy=T.GemmWarpPolicy.FullCol) # Rescale for h_i, d_i in T.Parallel(H_per_block, D): acc_o[h_i, d_i] /= sumexp[h_i] for h_i in T.Parallel(H_per_block): sumexp[h_i] = T.log2(sumexp[h_i]) + m_i[h_i] * sm_scale T.copy(acc_o, O_shared) T.copy(acc_o, Output[b_i, s_i, H0:H1, :]) return main @tilelang.jit( out_idx=[-1], compile_flags=[ "-O3", "-Wno-deprecated-declarations", "-U__CUDA_NO_HALF_OPERATORS__", "-U__CUDA_NO_HALF_CONVERSIONS__", "-U__CUDA_NO_HALF2_OPERATORS__", "-U__CUDA_NO_BFLOAT16_CONVERSIONS__", "--expt-relaxed-constexpr", "--expt-extended-lambda", "--ptxas-options=-v,--register-usage-level=10", "-DNDEBUG", ], ) # type: ignore def sparse_attention_fwd_kernel_v2( num_heads: int, dim: int, tail_dim: int, topk: int, *, kv_group: int = 1, sm_scale: Optional[float] = None, block_I: int = 64, ): assert dim == tilelang.math.next_power_of_2( dim ), f"haven't check padding correctness yet, dim={dim}" assert tail_dim == tilelang.math.next_power_of_2( tail_dim ), f"haven't check padding correctness yet, dim={tail_dim}" assert ( topk % block_I == 0 ), "otherwise will load some index=0 thus causing wrong kv to be loaded" if sm_scale is None: sm_scale = (1.0 / (dim + tail_dim)) ** 0.5 * 1.44269504 # log2(e) else: sm_scale = sm_scale * 1.44269504 # log2(e) threads = 384 batch = T.symbolic("batch") qo_len = T.symbolic("seq_len") num_pages = T.symbolic("num_pages") q_shape = [batch, qo_len, num_heads, dim + tail_dim] kv_shape = [batch, num_pages, kv_group, dim + tail_dim] o_shape = [batch, qo_len, num_heads, dim] indices_shape = [batch, qo_len, kv_group, topk] indices_dtype = "int32" dtype = "bfloat16" accum_dtype = "float" H = num_heads padded_H = max(tilelang.math.next_power_of_2(num_heads), 16) if padded_H != H: assert kv_group == 1 BI = block_I NI = tilelang.cdiv(topk, block_I) assert NI % 2 == 0, "NI should be a multiple of 2" D = dim D_tail = tail_dim if num_heads > 64: assert num_heads % 64 == 0, "head_kv should be a multiple of 64" REPLICATE_H = num_heads // 64 else: REPLICATE_H = 1 H_per_block = padded_H if REPLICATE_H == 1 else 64 @T.prim_func def main( Q: T.Tensor(q_shape, dtype), # type: ignore KV: T.Tensor(kv_shape, dtype), # type: ignore Indices: T.Tensor(indices_shape, indices_dtype), # type: ignore Output: T.Tensor(o_shape, dtype), # type: ignore ): """ 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) """ with T.Kernel(qo_len * REPLICATE_H, batch, 1, threads=threads) as (bx, by, bz): # type: ignore Q_shared_l = T.alloc_shared([H_per_block, D // 2], dtype) Q_shared_r = T.alloc_shared([H_per_block, D // 2], dtype) Q_tail_shared = T.alloc_shared([H_per_block, D_tail], dtype) KV_shared_0_l = T.alloc_shared([BI, D // 2], dtype) KV_shared_0_r = T.alloc_shared([BI, D // 2], dtype) KV_shared_1_l = T.alloc_shared([BI, D // 2], dtype) KV_shared_1_r = T.alloc_shared([BI, D // 2], dtype) K_tail_shared_0 = T.alloc_shared([BI, D_tail], dtype) K_tail_shared_1 = T.alloc_shared([BI, D_tail], dtype) O_shared_l = Q_shared_l O_shared_r = Q_shared_r is_kv_valid_0 = T.alloc_shared([BI], "bool", scope="shared") is_kv_valid_1 = T.alloc_shared([BI], "bool", scope="shared") acc_o_l = T.alloc_fragment([H_per_block, D // 2], accum_dtype) acc_o_r = T.alloc_fragment([H_per_block, D // 2], accum_dtype) acc_s = T.alloc_fragment([H_per_block, BI], accum_dtype) S_shared = T.alloc_shared([H_per_block, BI], dtype) sumexp = T.alloc_fragment([H_per_block], accum_dtype) sum_exp_shared = T.alloc_shared([H_per_block], accum_dtype) sumexp_i = T.alloc_fragment([H_per_block], accum_dtype) alpha_shared = T.alloc_shared([H_per_block], accum_dtype, scope="shared") alpha_local = T.alloc_fragment([H_per_block], accum_dtype) m_i = T.alloc_fragment([H_per_block], accum_dtype) m_i_prev = T.alloc_fragment([H_per_block], accum_dtype) indices_local = T.alloc_local([1], indices_dtype) indices_tmp = T.alloc_local([1], indices_dtype) bar_q = T.alloc_barrier(arrive_count=384) bar_k_0_ready = T.alloc_barrier(arrive_count=128) bar_k_1_ready = T.alloc_barrier(arrive_count=128) bar_k_0_free = T.alloc_barrier(arrive_count=256) bar_k_1_free = T.alloc_barrier(arrive_count=256) bar_sScale_and_sS_ready = T.alloc_barrier(arrive_count=256) bar_sScale_and_sS_free = T.alloc_barrier(arrive_count=256) bar_0_128 = T.alloc_barrier(arrive_count=128) bar_1_128 = T.alloc_barrier(arrive_count=128) bar_2_128 = T.alloc_barrier(arrive_count=128) bar_final = T.alloc_barrier(arrive_count=128) b_i, g_i = by, bz s_i = bx if REPLICATE_H == 1 else bx // REPLICATE_H H0 = g_i * padded_H + (0 if REPLICATE_H == 1 else (bx % REPLICATE_H) * 64) H1 = H0 + H_per_block tx = T.get_thread_binding() T.copy(Q[b_i, s_i, H0:H1, 0 : D // 2], Q_shared_l) T.copy(Q[b_i, s_i, H0:H1, D // 2 : D], Q_shared_r) T.copy(Q[b_i, s_i, H0:H1, D:], Q_tail_shared) T.barrier_arrive(bar_q) if tx < 128: T.set_max_nreg(240, 1) T.fill(sumexp, 0) T.fill(m_i, -(2**30)) # avoid -inf - inf to cause nan T.fill(acc_o_l, 0) T.barrier_wait(bar_q, 0) for i_i in T.serial(T.ceildiv(NI, 2)): # Buffer 0 # with sync_at(bar_0_128, 0): T.barrier_wait(bar_k_0_ready[0], (i_i & 1)) T.barrier_arrive(bar_0_128) T.barrier_wait(bar_0_128, 0) for h_i, bi_i in T.Parallel(H_per_block, BI): acc_s[h_i, bi_i] = T.if_then_else( is_kv_valid_0[bi_i], 0, -T.infinity(acc_s.dtype) ) T.gemm( Q_shared_l, KV_shared_0_l, acc_s, transpose_B=True, wg_wait=-1 ) T.gemm( Q_shared_r, KV_shared_0_r, acc_s, transpose_B=True, wg_wait=-1 ) T.gemm( Q_tail_shared, K_tail_shared_0, acc_s, transpose_B=True, wg_wait=-1, ) T.wait_wgmma(0) if i_i != 0: T.barrier_arrive(bar_sScale_and_sS_free) T.barrier_wait(bar_sScale_and_sS_free, ((i_i * 2) & 1) ^ 1) T.copy(m_i, m_i_prev) T.reduce_max(acc_s, m_i, dim=1, clear=False) for h_i in T.Parallel(H_per_block): alpha_local[h_i] = T.exp2((m_i_prev[h_i] - m_i[h_i]) * sm_scale) for h_i, bi_i in T.Parallel(H_per_block, BI): acc_s[h_i, bi_i] = T.exp2( acc_s[h_i, bi_i] * sm_scale - m_i[h_i] * sm_scale ) T.reduce_sum( acc_s, sumexp_i, dim=1 ) # is this a accumulate operator? for h_i in T.Parallel(H_per_block): sumexp[h_i] = sumexp[h_i] * alpha_local[h_i] + sumexp_i[h_i] for h_i, d_i in T.Parallel(H_per_block, D // 2): acc_o_l[h_i, d_i] *= alpha_local[h_i] T.copy(alpha_local, alpha_shared) T.copy(acc_s, S_shared) T.gemm(S_shared, KV_shared_0_l, acc_o_l) T.barrier_arrive(bar_sScale_and_sS_ready) T.barrier_arrive(bar_k_0_free[0]) # Buffer 1 T.barrier_wait(bar_k_1_ready[0], (i_i & 1)) T.barrier_arrive(bar_0_128) T.barrier_wait(bar_0_128, 1) for h_i, bi_i in T.Parallel(H_per_block, BI): acc_s[h_i, bi_i] = T.if_then_else( is_kv_valid_1[bi_i], 0, -T.infinity(acc_s.dtype) ) T.gemm( Q_shared_l, KV_shared_1_l, acc_s, transpose_B=True, wg_wait=-1 ) T.gemm( Q_shared_r, KV_shared_1_r, acc_s, transpose_B=True, wg_wait=-1 ) T.gemm( Q_tail_shared, K_tail_shared_1, acc_s, transpose_B=True, wg_wait=-1, ) T.wait_wgmma(0) T.barrier_arrive(bar_sScale_and_sS_free) T.barrier_wait(bar_sScale_and_sS_free, ((i_i * 2 + 1) & 1) ^ 1) T.copy(m_i, m_i_prev) T.reduce_max(acc_s, m_i, dim=1, clear=False) for h_i in T.Parallel(H_per_block): alpha_local[h_i] = T.exp2((m_i_prev[h_i] - m_i[h_i]) * sm_scale) for h_i, bi_i in T.Parallel(H_per_block, BI): acc_s[h_i, bi_i] = T.exp2( acc_s[h_i, bi_i] * sm_scale - m_i[h_i] * sm_scale ) T.reduce_sum( acc_s, sumexp_i, dim=1 ) # is this a accumulate operator? for h_i in T.Parallel(H_per_block): sumexp[h_i] = sumexp[h_i] * alpha_local[h_i] + sumexp_i[h_i] for h_i, d_i in T.Parallel(H_per_block, D // 2): acc_o_l[h_i, d_i] *= alpha_local[h_i] T.copy(alpha_local, alpha_shared) T.copy(acc_s, S_shared) T.gemm(S_shared, KV_shared_1_l, acc_o_l) T.barrier_arrive(bar_sScale_and_sS_ready) T.barrier_arrive(bar_k_1_free[0]) # Rescale for h_i in T.Parallel(H_per_block): sum_exp_shared[h_i] = sumexp[h_i] T.barrier_arrive(bar_final) for h_i, d_i in T.Parallel(H_per_block, D // 2): acc_o_l[h_i, d_i] /= sumexp[h_i] for h_i in T.Parallel(H_per_block): sumexp[h_i] = T.log2(sumexp[h_i]) + m_i[h_i] * sm_scale T.copy(acc_o_l, O_shared_l) T.copy(O_shared_l, Output[b_i, s_i, H0:H1, 0 : D // 2]) elif tx >= 128 and tx < 256: # T.set_max_nreg(168, 1) T.fill(acc_o_r, 0) for i_i in T.serial(T.ceildiv(NI, 2)): # Buffer 0 T.barrier_arrive(bar_sScale_and_sS_ready) T.barrier_wait(bar_sScale_and_sS_ready, ((i_i * 2) & 1)) T.barrier_arrive(bar_1_128) T.barrier_wait(bar_1_128, 0) for h_i, d_i in T.Parallel(H_per_block, D // 2): acc_o_r[h_i, d_i] *= alpha_shared[h_i] T.gemm(S_shared, KV_shared_0_r, acc_o_r) T.barrier_arrive(bar_k_0_free[0]) T.barrier_arrive(bar_sScale_and_sS_free) # Buffer 1 T.barrier_arrive(bar_sScale_and_sS_ready) T.barrier_wait(bar_sScale_and_sS_ready, ((i_i * 2 + 1) & 1)) T.barrier_arrive(bar_1_128) T.barrier_wait(bar_1_128, 1) for h_i, d_i in T.Parallel(H_per_block, D // 2): acc_o_r[h_i, d_i] *= alpha_shared[h_i] T.gemm(S_shared, KV_shared_1_r, acc_o_r) T.barrier_arrive(bar_k_1_free[0]) if i_i != T.ceildiv(NI, 2) - 1: T.barrier_arrive(bar_sScale_and_sS_free) # Rescale T.barrier_wait(bar_final, 0) for h_i, d_i in T.Parallel(H_per_block, D // 2): acc_o_r[h_i, d_i] /= sum_exp_shared[h_i] T.copy(acc_o_r, O_shared_r) T.copy(O_shared_r, Output[b_i, s_i, H0:H1, D // 2 : D]) elif tx >= 256: # producer T.set_max_nreg(80, 0) indices_local[0] = 0 for i_i in T.serial(T.ceildiv(NI, 2)): # Buffer 0 T.barrier_wait(bar_k_0_free[0], ((i_i & 1) ^ 1)) T.barrier_arrive(bar_2_128) T.barrier_wait(bar_2_128, 0) for r in T.serial(4): indices_tmp[0] = Indices[ b_i, s_i, g_i, (i_i * 2) * BI + r * 16 + (tx - 256) // 8 ] is_kv_valid_0[r * 16 + (tx - 256) // 8] = indices_tmp[0] >= 0 if is_kv_valid_0[r * 16 + (tx - 256) // 8]: indices_local[0] = indices_tmp[0] with T.attr("default", "async_scope", 1): # type: ignore for u in T.serial(4): for v in T.vectorized(8): KV_shared_0_l[ r * 16 + (tx - 256) // 8, 64 * u + (tx - 256) % 8 * 8 + v, ] = KV[ b_i, indices_local[0], g_i, 64 * u + (tx - 256) % 8 * 8 + v, ] KV_shared_0_r[ r * 16 + (tx - 256) // 8, 64 * u + (tx - 256) % 8 * 8 + v, ] = KV[ b_i, indices_local[0], g_i, D // 2 + 64 * u + (tx - 256) % 8 * 8 + v, ] with T.attr("default", "async_scope", 1): # type: ignore for v in T.vectorized(8): K_tail_shared_0[ r * 16 + (tx - 256) // 8, (tx - 256) % 8 * 8 + v ] = KV[ b_i, indices_local[0], g_i, D + (tx - 256) % 8 * 8 + v, ] T.cp_async_barrier_noinc(bar_k_0_ready[0]) # Buffer 1 T.barrier_wait(bar_k_1_free[0], ((i_i & 1) ^ 1)) T.barrier_arrive(bar_2_128) T.barrier_wait(bar_2_128, 1) for r in T.serial(4): indices_tmp[0] = Indices[ b_i, s_i, g_i, (i_i * 2 + 1) * BI + r * 16 + (tx - 256) // 8 ] is_kv_valid_1[r * 16 + (tx - 256) // 8] = indices_tmp[0] >= 0 if is_kv_valid_1[r * 16 + (tx - 256) // 8]: indices_local[0] = indices_tmp[0] with T.attr("default", "async_scope", 1): # type: ignore for u in T.serial(4): for v in T.vectorized(8): KV_shared_1_l[ r * 16 + (tx - 256) // 8, 64 * u + (tx - 256) % 8 * 8 + v, ] = KV[ b_i, indices_local[0], g_i, 64 * u + (tx - 256) % 8 * 8 + v, ] KV_shared_1_r[ r * 16 + (tx - 256) // 8, 64 * u + (tx - 256) % 8 * 8 + v, ] = KV[ b_i, indices_local[0], g_i, D // 2 + 64 * u + (tx - 256) % 8 * 8 + v, ] with T.attr("default", "async_scope", 1): # type: ignore for v in T.vectorized(8): K_tail_shared_1[ r * 16 + (tx - 256) // 8, (tx - 256) % 8 * 8 + v ] = KV[ b_i, indices_local[0], g_i, D + (tx - 256) % 8 * 8 + v, ] T.cp_async_barrier_noinc(bar_k_1_ready[0]) return main @tilelang.jit( out_idx=[-2, -1], pass_configs={ tilelang.PassConfigKey.TL_DISABLE_TMA_LOWER: True, tilelang.PassConfigKey.TL_DISABLE_WARP_SPECIALIZED: True, }, ) def sparse_mla_fwd_decode_partial( heads, dim, tail_dim, topk, *, kv_group=1, sm_scale=None, is_causal=True, block_I=64, threads=256, ): """ grid: (seq_len * REPLICATE_H, top_k_blocks). Each block does one topk block, writes partial_o, partial_lse. """ assert is_causal == True, "non-causal is not supported" assert kv_group == 1 assert topk % block_I == 0 # log2(e) = 1.44269504 if sm_scale is None: sm_scale = (1.0 / (dim + tail_dim)) ** 0.5 * 1.44269504 else: sm_scale = sm_scale * 1.44269504 batch = 1 seq_len = T.dynamic("seq_len") seq_len_kv = T.dynamic("seq_len_kv") head_kv = heads // kv_group padded_H = max(tilelang.math.next_power_of_2(head_kv), 16) REPLICATE_H = (head_kv // 64) if head_kv > 64 else 1 H_per_block = padded_H if REPLICATE_H == 1 else 64 BI = block_I NI = topk // block_I D = dim D_tail = tail_dim q_shape = [batch, seq_len, heads, dim + tail_dim] kv_shape = [batch, seq_len_kv, kv_group, dim + tail_dim] indices_shape = [batch, seq_len, kv_group, topk] partial_o_shape = [batch, seq_len, NI, heads, dim] partial_lse_shape = [batch, seq_len, NI, heads] indices_dtype = T.int32 dtype = T.bfloat16 accum_dtype = T.float32 @T.prim_func def main( Q: T.Tensor(q_shape, dtype), KV: T.Tensor(kv_shape, dtype), Indices: T.Tensor(indices_shape, indices_dtype), Partial_O: T.Tensor(partial_o_shape, dtype), Partial_Lse: T.Tensor(partial_lse_shape, accum_dtype), ): with T.Kernel(seq_len * REPLICATE_H, NI, threads=threads) as (bx, by): Q_shared = T.alloc_shared([H_per_block, D], dtype) Q_tail_shared = T.alloc_shared([H_per_block, D_tail], dtype) KV_shared = T.alloc_shared([BI, D], dtype) K_tail_shared = T.alloc_shared([BI, D_tail], dtype) mask = T.alloc_fragment([BI], T.bool) acc_o = T.alloc_fragment([H_per_block, D], accum_dtype) acc_s = T.alloc_fragment([H_per_block, BI], accum_dtype) S_shared = T.alloc_shared([H_per_block, BI], dtype) sumexp_i = T.alloc_fragment([H_per_block], accum_dtype) m_i = T.alloc_fragment([H_per_block], accum_dtype) T.fill(acc_o, 0) b_i, g_i = 0, 0 s_i = bx if REPLICATE_H == 1 else (bx // REPLICATE_H) topk_block_i = by q_i = s_i H0 = 0 if REPLICATE_H == 1 else (bx % REPLICATE_H) * 64 H1 = H0 + H_per_block T.copy(Q[b_i, s_i, H0:H1, :D], Q_shared) T.copy(Q[b_i, s_i, H0:H1, D:], Q_tail_shared) for bi_i in T.Parallel(BI): mask[bi_i] = Indices[b_i, s_i, g_i, topk_block_i * BI + bi_i] >= 0 for bi_i, d_i in T.Parallel(BI, D): KV_shared[bi_i, d_i] = KV[ b_i, Indices[b_i, s_i, g_i, topk_block_i * BI + bi_i], g_i, d_i ] for bi_i, d_i in T.Parallel(BI, D_tail): K_tail_shared[bi_i, d_i] = KV[ b_i, Indices[b_i, s_i, g_i, topk_block_i * BI + bi_i], g_i, D + d_i ] for h_i, bi_i in T.Parallel(H_per_block, BI): acc_s[h_i, bi_i] = T.if_then_else( mask[bi_i], 0, -T.infinity(acc_s.dtype) ) T.gemm( Q_shared, KV_shared, acc_s, transpose_B=True, policy=T.GemmWarpPolicy.FullCol, ) T.gemm( Q_tail_shared, K_tail_shared, acc_s, transpose_B=True, policy=T.GemmWarpPolicy.FullCol, ) T.reduce_max(acc_s, m_i, dim=1, clear=True) for h_i in T.Parallel(H_per_block): m_i[h_i] = T.max(m_i[h_i], -(2**30)) for h_i, bi_i in T.Parallel(H_per_block, BI): acc_s[h_i, bi_i] = T.exp2( acc_s[h_i, bi_i] * sm_scale - m_i[h_i] * sm_scale ) T.reduce_sum(acc_s, sumexp_i, dim=1) T.copy(acc_s, S_shared) T.gemm(S_shared, KV_shared, acc_o, policy=T.GemmWarpPolicy.FullCol) # sumexp_i==0 (all masked), divide by 1 to get 0 and avoid nan for h_i, d_i in T.Parallel(H_per_block, D): acc_o[h_i, d_i] = acc_o[h_i, d_i] / T.if_then_else( sumexp_i[h_i] == 0.0, 1.0, sumexp_i[h_i] ) # sumexp_i==0 (all masked), use large negative so combine ignores this split for h_i in T.Parallel(H_per_block): sumexp_i[h_i] = T.if_then_else( sumexp_i[h_i] == 0.0, -(2**30), T.log2(sumexp_i[h_i]) + m_i[h_i] * sm_scale, ) T.copy(acc_o, Partial_O[b_i, s_i, topk_block_i, H0:H1, :]) T.copy(sumexp_i, Partial_Lse[b_i, s_i, topk_block_i, H0:H1]) return main @tilelang.jit( out_idx=[-1], pass_configs={ tilelang.PassConfigKey.TL_DISABLE_TMA_LOWER: True, tilelang.PassConfigKey.TL_DISABLE_WARP_SPECIALIZED: True, }, ) def sparse_mla_fwd_decode_combine( heads, dim, topk, head_per_block, *, block_I=64, threads=256, ): """ grid: (seq_len * REPLICATE_H). batch=1, kv_group=1. Each block does one tile of heads (e.g. 4 or 8 for decode). """ assert heads % head_per_block == 0, f"head_per_block must divide heads" batch = 1 seq_len = T.dynamic("seq_len") NI = topk // block_I H_per_block = head_per_block REPLICATE_H = heads // H_per_block partial_o_shape = [batch, seq_len, NI, heads, dim] partial_lse_shape = [batch, seq_len, NI, heads] o_shape = [batch, seq_len, heads, dim] dtype = T.bfloat16 accum_dtype = T.float32 @T.prim_func def main( Partial_O: T.Tensor(partial_o_shape, dtype), Partial_Lse: T.Tensor(partial_lse_shape, accum_dtype), Output: T.Tensor(o_shape, dtype), ): with T.Kernel(seq_len * REPLICATE_H, threads=threads) as (bx,): shared_lse = T.alloc_shared([NI, H_per_block], accum_dtype) lse_max = T.alloc_fragment([H_per_block], accum_dtype) lse_sum = T.alloc_fragment([H_per_block], accum_dtype) scale = T.alloc_fragment([H_per_block, NI], accum_dtype) acc_o = T.alloc_fragment([H_per_block, dim], accum_dtype) b_i = 0 s_i = bx if REPLICATE_H == 1 else (bx // REPLICATE_H) H0 = 0 if REPLICATE_H == 1 else (bx % REPLICATE_H) * H_per_block H1 = H0 + H_per_block for k in T.serial(NI): T.copy(Partial_Lse[b_i, s_i, k, H0:H1], shared_lse[k, :]) T.fill(lse_max, -(2**30)) for k in T.serial(NI): for h_i in T.Parallel(H_per_block): lse_max[h_i] = T.max(lse_max[h_i], shared_lse[k, h_i]) T.fill(lse_sum, 0) for k in T.serial(NI): for h_i in T.Parallel(H_per_block): lse_sum[h_i] = lse_sum[h_i] + T.exp2( shared_lse[k, h_i] - lse_max[h_i] ) for k in T.serial(NI): for h_i in T.Parallel(H_per_block): scale[h_i, k] = T.exp2( shared_lse[k, h_i] - lse_max[h_i] - T.log2(lse_sum[h_i]) ) T.fill(acc_o, 0) for k in T.serial(NI): for h_i, d_i in T.Parallel(H_per_block, dim): acc_o[h_i, d_i] = acc_o[h_i, d_i] + scale[h_i, k] * Partial_O[ b_i, s_i, k, H0 + h_i, d_i ].astype(accum_dtype) T.copy(acc_o, Output[b_i, s_i, H0:H1, :]) return main def tilelang_sparse_fwd( q: torch.Tensor, kv: torch.Tensor, indices: torch.Tensor, sm_scale: float, d_v: int = 512, ) -> torch.Tensor: assert q.dim() == 3 and kv.dim() == 3 and indices.dim() == 3 num_heads = q.shape[1] dim = q.shape[2] tail_dim = dim - d_v topk = indices.shape[-1] assert topk == 2048 if _is_hip: if _is_gfx95_supported: # decode kernel if q.shape[0] <= 64: kernel_partial = sparse_mla_fwd_decode_partial( num_heads, d_v, tail_dim, topk, sm_scale=sm_scale, block_I=64, threads=256, ) kernel_combine = sparse_mla_fwd_decode_combine( num_heads, d_v, topk, head_per_block=4, block_I=64, threads=256 ) partial_o, partial_lse = kernel_partial( q.unsqueeze(0), kv.unsqueeze(0), indices.unsqueeze(0) ) out = kernel_combine(partial_o, partial_lse) return out # prefill kernel kernel = sparse_attention_fwd_kernel_v1( num_heads, d_v, tail_dim, topk, sm_scale=sm_scale, num_stages=1 ) else: # reduce LDS usage on gfx942 target kernel = sparse_attention_fwd_kernel_v1( num_heads, d_v, tail_dim, topk, sm_scale=sm_scale, block_I=32, num_stages=1, threads=128, ) else: kernel = sparse_attention_fwd_kernel_v2( num_heads, d_v, tail_dim, topk, sm_scale=sm_scale ) return kernel(q.unsqueeze(0), kv.unsqueeze(0), indices.unsqueeze(0)) # type: ignore