Support unequal sequence length of q and kv

[wenhaoli-xmu]

@triton.jit
def _attn_bwd(Q, K, V, mask, sm_scale,
              DO,
              DQ, DK, DV,
              M, D,
              stride_z, stride_h, stride_tok, stride_d,
              mask_stride_z, mask_stride_h, mask_stride_tok, mask_stride_tokk,
              H, N_CTX,
              BLOCK_M1: tl.constexpr,
              BLOCK_N1: tl.constexpr,
              BLOCK_M2: tl.constexpr,
              BLOCK_N2: tl.constexpr,
              BLK_SLICE_FACTOR: tl.constexpr,
              HEAD_DIM: tl.constexpr,
              USE_MASK: tl.constexpr):

The length of queries and keys/values should be equal now. Are you plan to support unequal input length?

[wenhaoli-xmu]

If it already support unequal input length, why does the qkv share data stride? I think the data stride should be different with various sequence length.