@triton.jit
def _fwd_kernel_aligned(Q, K, V, B0, sm_scale, Out, stride_qh, stride_qm,
    stride_qk, stride_kh, stride_kn, stride_kk, stride_vh, stride_vk,
    stride_vn, stride_oh, stride_om, stride_on, stride_b0h, stride_b0m, Z,
    H, N_CTX, P_SEQ, OUT_DTYPE: tl.constexpr, BIAS_LAST_SIZE: tl.constexpr,
    B0_NUMEL: tl.constexpr, BLOCK_DMODEL: tl.constexpr, BLOCK_M: tl.
    constexpr, BLOCK_N: tl.constexpr):
    start_m = tl.program_id(0)
    off_hz = tl.program_id(1)
    q_offset = off_hz * stride_qh
    kv_offset = off_hz * stride_kh
    Q_block_ptr = tl.make_block_ptr(base=Q + q_offset, shape=(N_CTX,
        BLOCK_DMODEL), strides=(stride_qm, stride_qk), offsets=(start_m *
        BLOCK_M, 0), block_shape=(BLOCK_M, BLOCK_DMODEL), order=(1, 0))
    K_block_ptr = tl.make_block_ptr(base=K + kv_offset, shape=(BLOCK_DMODEL,
        N_CTX + P_SEQ), strides=(stride_kk, stride_kn), offsets=(0, 0),
        block_shape=(BLOCK_DMODEL, BLOCK_N), order=(0, 1))
    V_block_ptr = tl.make_block_ptr(base=V + kv_offset, shape=(N_CTX +
        P_SEQ, BLOCK_DMODEL), strides=(stride_vk, stride_vn), offsets=(0, 0
        ), block_shape=(BLOCK_N, BLOCK_DMODEL), order=(1, 0))
    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, BLOCK_DMODEL], dtype=tl.float32)
    qk_scale = sm_scale * 1.44269504
    q = tl.load(Q_block_ptr)
    q = (q * qk_scale).to(OUT_DTYPE)
    lo = 0
    hi = N_CTX + P_SEQ
    b_ptr_offsets_m = tl.arange(0, BLOCK_M)
    b_offset = off_hz * stride_b0h
    b_ptr_offsets_n_1 = tl.arange(0, BLOCK_N) % BIAS_LAST_SIZE + BIAS_LAST_SIZE
    b1 = tl.load(B0 + b_offset + ((start_m * BLOCK_M + b_ptr_offsets_m) *
        stride_b0m)[:, None] + b_ptr_offsets_n_1[None, :])
    for start_n in range(lo, hi, BLOCK_N):
        k = tl.load(K_block_ptr)
        v = tl.load(V_block_ptr)
        qk = tl.zeros([BLOCK_M, BLOCK_N], dtype=OUT_DTYPE)
        qk += tl.dot(q, k)
        b0 = tl.load(B0 + b_offset + ((start_m * BLOCK_M + b_ptr_offsets_m) *
            stride_b0m)[:, None] + start_n // BLOCK_N)
        qk += (b0 + b1) * 1.44269504
        m_i_new = tl.maximum(m_i, tl.max(qk, 1))
        alpha = tl.math.exp2(m_i - m_i_new)
        p = tl.math.exp2(qk - m_i_new[:, None])
        acc *= alpha[:, None]
        acc += tl.dot(p.to(OUT_DTYPE), v)
        l_i = l_i * alpha + tl.sum(p, 1)
        m_i = m_i_new
        K_block_ptr = tl.advance(K_block_ptr, (0, BLOCK_N))
        V_block_ptr = tl.advance(V_block_ptr, (BLOCK_N, 0))
    acc = acc / l_i[:, None]
    O_block_ptr = tl.make_block_ptr(base=Out + q_offset, shape=(N_CTX,
        BLOCK_DMODEL), strides=(stride_om, stride_on), offsets=(start_m *
        BLOCK_M, 0), block_shape=(BLOCK_M, BLOCK_DMODEL), order=(1, 0))
    tl.store(O_block_ptr, acc.to(OUT_DTYPE))
