seq parallel for backwards nsa

Author: lucidrains

native_sparse_attention_pytorch/triton_native_sparse_attention.py

@@ -1526,93 +1526,91 @@ def backward_kernel(
         off_qh * seqlen_q_rounded
     )
 
-    num_block_n = tl.cdiv(seqlen_k, BLOCK)
+    start_n = tl.program_id(2)
 
     if IS_CAUSAL:
-        for start_n in range(0, num_block_n):
-            backward_kernel_one_col_block_causal(
-                start_n,
-                Q,
-                K,
-                V,
-                kv_block_indices,
-                kv_block_mask,
-                DO,
-                DQ,
-                DK,
-                DV,
-                LSE,
-                D,
-                softmax_scale,
-                stride_qm,
-                stride_kn,
-                stride_vn,
-                stride_dom,
-                stride_dqm,
-                stride_dkn,
-                stride_dvn,
-                stride_kvbl_m,
-                stride_qh,
-                stride_doh,
-                stride_dqh,
-                seqlen_q,
-                seqlen_k,
-                seqlen_q_rounded,
-                headdim,
-                BLOCK_HEADDIM = BLOCK_HEADDIM,
-                EVEN_M = EVEN_M,
-                EVEN_N = EVEN_N,
-                EVEN_HEADDIM = EVEN_HEADDIM,
-                BLOCK = BLOCK,
-                SEL_BLOCK = SEL_BLOCK,
-                QUERY_HEAD_GROUPS = QUERY_HEAD_GROUPS,
-                QUERY_EXPAND_DIM = QUERY_EXPAND_DIM,
-                SLIDING = SLIDING
-            )
+        backward_kernel_one_col_block_causal(
+            start_n,
+            Q,
+            K,
+            V,
+            kv_block_indices,
+            kv_block_mask,
+            DO,
+            DQ,
+            DK,
+            DV,
+            LSE,
+            D,
+            softmax_scale,
+            stride_qm,
+            stride_kn,
+            stride_vn,
+            stride_dom,
+            stride_dqm,
+            stride_dkn,
+            stride_dvn,
+            stride_kvbl_m,
+            stride_qh,
+            stride_doh,
+            stride_dqh,
+            seqlen_q,
+            seqlen_k,
+            seqlen_q_rounded,
+            headdim,
+            BLOCK_HEADDIM = BLOCK_HEADDIM,
+            EVEN_M = EVEN_M,
+            EVEN_N = EVEN_N,
+            EVEN_HEADDIM = EVEN_HEADDIM,
+            BLOCK = BLOCK,
+            SEL_BLOCK = SEL_BLOCK,
+            QUERY_HEAD_GROUPS = QUERY_HEAD_GROUPS,
+            QUERY_EXPAND_DIM = QUERY_EXPAND_DIM,
+            SLIDING = SLIDING
+        )
     else:
-        for start_n in range(0, num_block_n):
-            backward_kernel_one_col_block_sparse(
-                start_n,
-                Q,
-                K,
-                V,
-                kv_block_indices,
-                kv_block_mask,
-                kv_block_grads,
-                DO,
-                DQ,
-                DK,
-                DV,
-                LSE,
-                D,
-                softmax_scale,
-                stride_qm,
-                stride_kn,
-                stride_vn,
-                stride_dom,
-                stride_dqm,
-                stride_dkn,
-                stride_dvn,
-                stride_kvbl_m,
-                stride_qh,
-                stride_doh,
-                stride_dqh,
-                seqlen_q,
-                seqlen_k,
-                seqlen_q_rounded,
-                headdim,
-                BLOCK_HEADDIM = BLOCK_HEADDIM,
-                EVEN_M = EVEN_M,
-                EVEN_N = EVEN_N,
-                EVEN_HEADDIM = EVEN_HEADDIM,
-                BLOCK = BLOCK,
-                QUERY_HEAD_GROUPS = QUERY_HEAD_GROUPS,
-                QUERY_EXPAND_DIM = QUERY_EXPAND_DIM,
-                RETURN_SEL_GRADS = RETURN_SEL_GRADS,
-                OFF_SEL_KV_BLOCKS = OFF_SEL_KV_BLOCKS,
-                BLOCK_DV_USE_DOT = BLOCK_DV_USE_DOT,
-                BLOCK_DK_USE_DOT = BLOCK_DK_USE_DOT,
-            )
+        backward_kernel_one_col_block_sparse(
+            start_n,
+            Q,
+            K,
+            V,
+            kv_block_indices,
+            kv_block_mask,
+            kv_block_grads,
+            DO,
+            DQ,
+            DK,
+            DV,
+            LSE,
+            D,
+            softmax_scale,
+            stride_qm,
+            stride_kn,
+            stride_vn,
+            stride_dom,
+            stride_dqm,
+            stride_dkn,
+            stride_dvn,
+            stride_kvbl_m,
+            stride_qh,
+            stride_doh,
+            stride_dqh,
+            seqlen_q,
+            seqlen_k,
+            seqlen_q_rounded,
+            headdim,
+            BLOCK_HEADDIM = BLOCK_HEADDIM,
+            EVEN_M = EVEN_M,
+            EVEN_N = EVEN_N,
+            EVEN_HEADDIM = EVEN_HEADDIM,
+            BLOCK = BLOCK,
+            QUERY_HEAD_GROUPS = QUERY_HEAD_GROUPS,
+            QUERY_EXPAND_DIM = QUERY_EXPAND_DIM,
+            RETURN_SEL_GRADS = RETURN_SEL_GRADS,
+            OFF_SEL_KV_BLOCKS = OFF_SEL_KV_BLOCKS,
+            BLOCK_DV_USE_DOT = BLOCK_DV_USE_DOT,
+            BLOCK_DK_USE_DOT = BLOCK_DK_USE_DOT,
+        )
 
 def native_sparse_attn_backward(
     do,
@@ -1692,7 +1690,8 @@ def native_sparse_attn_backward(
 
     grid = lambda META: (
         num_sel_fine_blocks + int(include_block_causal),
-        batch * kv_heads
+        batch * kv_heads,
+        triton.cdiv(seqlen_k, META['BLOCK'])
     )
 
     backward_kernel[grid](

pyproject.toml

@@ -1,6 +1,6 @@
 [project]
 name = "native-sparse-attention-pytorch"
-version = "0.1.20"
+version = "0.1.21"
 description = "Native Sparse Attention"
 authors = [
     { name = "Phil Wang", email = "lucidrains@gmail.com" }