File size: 10,239 Bytes
4a28d4d
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
import pytest
import torch


def do_test():
    try:
        import causal_conv1d  # noqa: F401
        import tilelang  # noqa: F401
        causal_conv1d_fn = causal_conv1d.causal_conv1d_fn  # noqa: F841
        causal_conv1d_update = causal_conv1d.causal_conv1d_update  # noqa: F841
        return True
    except Exception:
        return False


@pytest.mark.skipif(not do_test(), reason='tilelang or causal_conv1d is not available')
class TestCausalConv1dUpdate:

    @pytest.fixture
    def device(self):
        yield 'cuda'

    @pytest.fixture
    def batch(self):
        yield 128

    @pytest.fixture
    def hidden_size(self):
        yield 2048

    @pytest.fixture
    def width(self):
        yield 4

    @pytest.fixture(params=[1, 4])
    def seqlen(self, request):
        yield request.param

    @pytest.fixture
    def x(self, batch, hidden_size, seqlen, device):
        yield torch.randn(batch, hidden_size, seqlen, device=device)

    @pytest.fixture
    def weight(self, hidden_size, width, device):
        yield torch.randn(hidden_size, width, device=device)

    @pytest.fixture
    def bias(self, hidden_size, device):
        yield torch.randn(hidden_size, device=device)

    @pytest.fixture(params=[True, False])
    def cache_seqlens(self, request, batch, device):
        if request.param:
            yield torch.randint(0, 4096, (batch, ), dtype=torch.int32, device=device)
        else:
            yield None

    @pytest.fixture(params=[True, False])
    def conv_state_indices(self, request, batch, device):
        if request.param:
            conv_state_indices = batch * 2 - 1 - torch.arange(0, batch * 2, 2, device=device)
            yield conv_state_indices.to(torch.int32)
        else:
            yield None

    @pytest.fixture
    def conv_state(self, batch, hidden_size, width, device, conv_state_indices):
        if conv_state_indices is not None:
            conv_state = torch.randn(batch * 4, hidden_size, width, device=device)
            conv_state = conv_state[::2]
        else:
            conv_state = torch.randn(batch, hidden_size, width, device=device)
        yield conv_state

    @pytest.fixture(params=[None, 'silu'])
    def activation(self, request):
        yield request.param

    def test_causal_conv1d_update(self, x, conv_state, weight, bias, activation, cache_seqlens, conv_state_indices):
        from causal_conv1d import causal_conv1d_update as causal_conv1d_update_gt

        from lmdeploy.pytorch.kernels.cuda.causal_conv1d import causal_conv1d_update

        conv_state_clone = conv_state.clone()
        out = causal_conv1d_update(x=x,
                                   conv_state=conv_state_clone,
                                   weight=weight,
                                   bias=bias,
                                   activation=activation,
                                   cache_seqlens=cache_seqlens,
                                   conv_state_indices=conv_state_indices)
        out_gt = causal_conv1d_update_gt(x=x,
                                         conv_state=conv_state,
                                         weight=weight,
                                         bias=bias,
                                         activation=activation,
                                         cache_seqlens=cache_seqlens,
                                         conv_state_indices=conv_state_indices)
        torch.testing.assert_close(out, out_gt, rtol=1e-3, atol=1e-3)
        torch.testing.assert_close(conv_state_clone, conv_state, rtol=1e-3, atol=1e-3)


@pytest.mark.skipif(not do_test(), reason='tilelang or causal_conv1d is not available')
class TestCausalConv1dFn:

    @pytest.fixture
    def device(self):
        yield 'cuda'

    @pytest.fixture
    def hidden_size(self):
        yield 2048

    @pytest.fixture
    def seqlen(self):
        yield 4096

    @pytest.fixture
    def seq_idx(self, seqlen, device):
        seq_idx = torch.zeros(seqlen, dtype=torch.int32, device=device)
        seq_idx[seqlen // 4 * 3:] = 1
        seq_idx = seq_idx.view(1, -1)
        yield seq_idx

    @pytest.fixture
    def x(self, hidden_size, seqlen, device):
        yield torch.randn(1, hidden_size, seqlen, device=device).transpose(1, 2).contiguous().transpose(1, 2)

    @pytest.fixture
    def weight(self, hidden_size, device):
        yield torch.randn(hidden_size, 4, device=device)

    @pytest.fixture
    def bias(self, hidden_size, device):
        yield torch.randn(hidden_size, device=device)

    @pytest.fixture(params=[None, 'silu'])
    def activation(self, request):
        yield request.param

    def test_causal_conv1d_fn(self, x, weight, bias, activation, seq_idx):
        from causal_conv1d import causal_conv1d_fn as causal_conv1d_fn_gt

        from lmdeploy.pytorch.kernels.cuda.causal_conv1d import causal_conv1d_fn

        out = causal_conv1d_fn(x=x,
                               weight=weight,
                               bias=bias,
                               activation=activation,
                               return_final_states=False,
                               seq_idx=seq_idx)
        out_gt = causal_conv1d_fn_gt(x=x,
                                     weight=weight,
                                     bias=bias,
                                     activation=activation,
                                     return_final_states=False,
                                     seq_idx=seq_idx)
        torch.testing.assert_close(out, out_gt, rtol=1e-3, atol=1e-3)


@pytest.mark.skipif(not do_test(), reason='tilelang or causal_conv1d is not available')
class TestCausalConv1dFnInitStates:
    """Test causal_conv1d_fn with per-sequence initial_states."""

    @pytest.fixture
    def device(self):
        yield 'cuda'

    @pytest.fixture
    def hidden_size(self):
        yield 2048

    @pytest.fixture
    def width(self):
        yield 4

    @pytest.fixture(params=[None, 'silu'])
    def activation(self, request):
        yield request.param

    @pytest.fixture
    def weight(self, hidden_size, width, device):
        yield torch.randn(hidden_size, width, device=device)

    @pytest.fixture
    def bias(self, hidden_size, device):
        yield torch.randn(hidden_size, device=device)

    def _ref_conv1d(self, x_seq, weight, bias, init_state, activation):
        """Reference: prepend init_state and run F.conv1d per sequence."""
        dim = weight.shape[0]
        x_cat = torch.cat([init_state, x_seq], dim=-1).float()
        w = weight.float().unsqueeze(1)
        b = bias.float() if bias is not None else None
        out = torch.nn.functional.conv1d(x_cat, w, b, padding=0, groups=dim)
        out = out[..., :x_seq.shape[-1]]
        if activation in ('silu', 'swish'):
            out = torch.nn.functional.silu(out)
        return out.to(x_seq.dtype)

    def test_single_seq(self, hidden_size, width, weight, bias, activation, device):
        """Single sequence with non-zero initial states."""
        seqlen = 128
        x = torch.randn(1, hidden_size, seqlen, device=device).transpose(1, 2).contiguous().transpose(1, 2)
        seq_idx = torch.zeros(1, seqlen, dtype=torch.int32, device=device)
        init_states = torch.randn(1, hidden_size, width - 1, device=device, dtype=x.dtype)

        from lmdeploy.pytorch.kernels.cuda.causal_conv1d import causal_conv1d_fn

        out = causal_conv1d_fn(x=x, weight=weight, bias=bias, seq_idx=seq_idx, initial_states=init_states,
                               activation=activation)
        out_ref = self._ref_conv1d(x, weight, bias, init_states, activation)
        torch.testing.assert_close(out, out_ref, rtol=1e-3, atol=1e-3)

    def test_multi_seq(self, hidden_size, width, weight, bias, activation, device):
        """Multiple sequences packed, each with its own initial state."""
        seqlens = [100, 200, 50]
        n_seqs = len(seqlens)

        x_parts = [torch.randn(1, hidden_size, sl, device=device) for sl in seqlens]
        # Make channel-last stride like real lmdeploy usage.
        x_packed = torch.cat(x_parts, dim=-1).transpose(1, 2).contiguous().transpose(1, 2)

        seq_idx_parts = [torch.full((sl, ), i, dtype=torch.int32, device=device) for i, sl in enumerate(seqlens)]
        seq_idx = torch.cat(seq_idx_parts).unsqueeze(0)

        init_states = torch.randn(n_seqs, hidden_size, width - 1, device=device, dtype=x_packed.dtype)

        from lmdeploy.pytorch.kernels.cuda.causal_conv1d import causal_conv1d_fn

        out = causal_conv1d_fn(x=x_packed, weight=weight, bias=bias, seq_idx=seq_idx, initial_states=init_states,
                               activation=activation)

        # Build reference per-sequence.
        offset = 0
        for i, sl in enumerate(seqlens):
            x_seq = x_packed[:, :, offset:offset + sl]
            ref = self._ref_conv1d(x_seq, weight, bias, init_states[i:i + 1], activation)
            torch.testing.assert_close(out[:, :, offset:offset + sl], ref, rtol=1e-3, atol=1e-3)
            offset += sl

    def test_zero_init_matches_no_init(self, hidden_size, width, weight, bias, activation, device):
        """Zero initial_states should produce the same result as no
        initial_states."""
        seqlens = [80, 120]
        n_seqs = len(seqlens)

        x_packed = torch.randn(1, hidden_size, sum(seqlens), device=device).transpose(1, 2).contiguous().transpose(
            1, 2)
        seq_idx_parts = [torch.full((sl, ), i, dtype=torch.int32, device=device) for i, sl in enumerate(seqlens)]
        seq_idx = torch.cat(seq_idx_parts).unsqueeze(0)

        zero_states = torch.zeros(n_seqs, hidden_size, width - 1, device=device, dtype=x_packed.dtype)

        from lmdeploy.pytorch.kernels.cuda.causal_conv1d import causal_conv1d_fn

        out_with_init = causal_conv1d_fn(x=x_packed, weight=weight, bias=bias, seq_idx=seq_idx,
                                         initial_states=zero_states, activation=activation)
        out_no_init = causal_conv1d_fn(x=x_packed, weight=weight, bias=bias, seq_idx=seq_idx, initial_states=None,
                                       activation=activation)
        torch.testing.assert_close(out_with_init, out_no_init, rtol=1e-3, atol=1e-3)