File size: 4,815 Bytes
c335050 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144 145 146 147 148 149 |
import os
import pytest
import torch
import triton
from fla.ops.nsa.naive import naive_nsa
from fla.ops.nsa.parallel import parallel_nsa
from fla.ops.utils import prepare_token_indices
from fla.utils import assert_close, device
# FIXME
@pytest.mark.parametrize(
('B', 'T', 'H', 'HQ', 'D', 'S', 'block_size', 'scale', 'dtype'),
[
pytest.param(*test, id="B{}-T{}-H{}-HQ{}-D{}-S{}-block_size{}-scale{}-{}".format(*test))
for test in [
(1, 63, 1, 16, 64, 16, 32, 1.0, torch.float16),
(3, 111, 1, 32, 100, 16, 32, 1.0, torch.float16),
(3, 1024, 2, 32, 60, 16, 32, 0.1, torch.float16),
(3, 1024, 2, 32, 128, 16, 32, 0.1, torch.float16),
(4, 2048, 2, 32, 64, 16, 32, 0.1, torch.float16),
]
],
)
def test_parallel(
B: int,
T: int,
H: int,
HQ: int,
D: int,
S: int,
block_size: int,
scale: float,
dtype: torch.dtype,
):
torch.manual_seed(42)
os.environ['TRITON_F32_DEFAULT'] = 'ieee'
q = torch.randn((B, T, HQ, D), dtype=dtype, device=device).requires_grad_(True)
k = torch.randn((B, T, H, D), dtype=dtype, device=device).requires_grad_(True)
v = torch.randn((B, T, H, D), dtype=dtype, device=device).requires_grad_(True)
do = torch.randn((B, T, HQ, D), dtype=dtype, device=device)
block_indices = torch.full((B, T, H, S), T, dtype=torch.long, device=device)
for b in range(B):
for t in range(T):
for h in range(H):
i_i = torch.randperm(max(1, triton.cdiv(t, block_size)))[:S]
block_indices[b, t, h, :len(i_i)] = i_i
block_indices = block_indices.sort(-1)[0]
ref = naive_nsa(q=q, k=k, v=v, block_indices=block_indices, block_size=block_size, scale=scale)
ref.backward(do)
ref_dq, q.grad = q.grad.clone(), None
ref_dk, k.grad = k.grad.clone(), None
ref_dv, v.grad = v.grad.clone(), None
tri = parallel_nsa(q=q, k=k, v=v, block_indices=block_indices, block_size=block_size, scale=scale)
tri.backward(do)
tri_dq, q.grad = q.grad.clone(), None
tri_dk, k.grad = k.grad.clone(), None
tri_dv, v.grad = v.grad.clone(), None
assert_close(" o", ref, tri, 0.005)
assert_close("dq", ref_dq, tri_dq, 0.005)
assert_close("dk", ref_dk, tri_dk, 0.005)
assert_close("dv", ref_dv, tri_dv, 0.005)
@pytest.mark.parametrize(
('H', 'HQ', 'D', 'S', 'block_size', 'cu_seqlens', 'dtype'),
[
pytest.param(*test, id="H{}-HQ{}-D{}-S{}-block_size{}-cu_seqlens{}-{}".format(*test))
for test in [
(1, 16, 64, 16, 32, [0, 15], torch.float16),
(2, 32, 64, 16, 32, [0, 256, 500, 1000], torch.float16),
(2, 32, 100, 16, 32, [0, 15, 100, 300, 1200, 2000], torch.float16),
]
],
)
@pytest.mark.skipif(
os.getenv('SKIP_TEST_CHUNK_VARLEN') == '1',
reason='Skipping test because SKIP_TEST_CHUNK_VARLEN is set',
)
def test_parallel_varlen(
H: int,
HQ: int,
D: int,
S: int,
block_size: int,
cu_seqlens: list[int],
dtype: torch.dtype,
):
torch.manual_seed(42)
os.environ['TRITON_F32_DEFAULT'] = 'ieee'
T = cu_seqlens[-1]
cu_seqlens = torch.tensor(cu_seqlens, dtype=torch.int32, device=device)
# seq-first required for inputs with variable lengths
q = torch.randn((1, T, HQ, D), dtype=dtype, device=device).requires_grad_()
k = torch.randn((1, T, H, D), dtype=dtype, device=device).requires_grad_()
v = torch.randn((1, T, H, D), dtype=dtype, device=device).requires_grad_()
do = torch.randn((1, T, HQ, D), dtype=dtype, device=device)
block_indices = torch.full((1, T, H, S), T, dtype=torch.long, device=device)
seq_indices = prepare_token_indices(cu_seqlens).tolist()
for i in range(T):
_, t = seq_indices[i]
for h in range(H):
i_i = torch.randperm(max(1, triton.cdiv(t, block_size)))[:S]
block_indices[0, i, h, :len(i_i)] = i_i
block_indices = block_indices.sort(-1)[0]
ref = naive_nsa(
q=q,
k=k,
v=v,
block_indices=block_indices,
block_size=block_size,
cu_seqlens=cu_seqlens,
)
ref.backward(do)
ref_dq, q.grad = q.grad.clone(), None
ref_dk, k.grad = k.grad.clone(), None
ref_dv, v.grad = v.grad.clone(), None
tri = parallel_nsa(
q=q,
k=k,
v=v,
block_indices=block_indices,
block_size=block_size,
cu_seqlens=cu_seqlens,
)
tri.backward(do)
tri_dq, q.grad = q.grad.clone(), None
tri_dk, k.grad = k.grad.clone(), None
tri_dv, v.grad = v.grad.clone(), None
assert_close('o', ref, tri, 0.004)
assert_close('dq', ref_dq, tri_dq, 0.005)
assert_close('dk', ref_dk, tri_dk, 0.005)
assert_close('dv', ref_dv, tri_dv, 0.005)
|