Skip to content

Commit 8221fd6

Browse files
committed
fix: apply remaining PR openvpi#312 fixes
- weight split at midpoint (support expansion_factor != 1) - variance nested glu_type from predictor configs - remove useless warmup from __init__ (model on CPU, no compilation) - use_fused_kernels default to false
1 parent c396635 commit 8221fd6

4 files changed

Lines changed: 53 additions & 62 deletions

File tree

configs/base.yaml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -85,7 +85,7 @@ nccl_p2p: true
8585
###########
8686
# fusing
8787
###########
88-
use_fused_kernels: true
88+
use_fused_kernels: false
8989

9090
###########
9191
# finetune

modules/kernels/fused_linear_softsign_glu.py

Lines changed: 43 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -39,7 +39,7 @@
3939
def _fused_linear_softsign_glu_fwd_kernel(
4040
x_ptr, w_left_ptr, w_right_ptr, b_left_ptr, b_right_ptr,
4141
y_ptr, left_ptr, gate_ptr,
42-
M, K,
42+
M, N, K,
4343
stride_x_b, stride_x_k,
4444
stride_wl_n, stride_wl_k,
4545
stride_wr_n, stride_wr_k,
@@ -50,13 +50,15 @@ def _fused_linear_softsign_glu_fwd_kernel(
5050
):
5151
"""
5252
y = (x @ W_left^T + b_left) * softsign(x @ W_right^T + b_right)
53-
where softsign(x) = x / (1 + |x|).
5453
55-
2D grid over (M // BLOCK_M, K // BLOCK_N).
54+
N = output dim per GLU half (= inner_dim = dim × expansion_factor)
55+
K = input feature dim (= dim for first Linear, inner_dim for second)
56+
57+
2D grid over (M // BLOCK_M, N // BLOCK_N).
5658
"""
5759
pid = tl.program_id(0)
5860
num_pid_m = tl.cdiv(M, BLOCK_M)
59-
num_pid_n = tl.cdiv(K, BLOCK_N)
61+
num_pid_n = tl.cdiv(N, BLOCK_N)
6062
pid_m = pid // num_pid_n
6163
pid_n = pid % num_pid_n
6264

@@ -65,9 +67,9 @@ def _fused_linear_softsign_glu_fwd_kernel(
6567
offs_k = tl.arange(0, BLOCK_K)
6668

6769
m_mask_2d = offs_m[:, None] < M
68-
n_mask_nk = offs_n[:, None] < K # [BLOCK_N, 1] for N×K weight access
69-
n_mask_mn = offs_n[None, :] < K # [1, BLOCK_N] for M×N output access
70-
n_mask_1d = offs_n < K
70+
n_mask_nk = offs_n[:, None] < N # [BLOCK_N, 1] for N×K weight access
71+
n_mask_mn = offs_n[None, :] < N # [1, BLOCK_N] for M×N output access
72+
n_mask_1d = offs_n < N
7173

7274
acc_left = tl.zeros([BLOCK_M, BLOCK_N], dtype=tl.float32)
7375
acc_gate = tl.zeros([BLOCK_M, BLOCK_N], dtype=tl.float32)
@@ -137,7 +139,7 @@ def _fused_linear_softsign_glu_bwd_kernel(
137139
left_ptr, gate_ptr, grad_y_ptr,
138140
grad_x_ptr,
139141
w_left_ptr, w_right_ptr,
140-
M, K,
142+
M, K, N,
141143
stride_l_b, stride_l_n,
142144
stride_g_b, stride_g_n,
143145
stride_gy_b, stride_gy_n,
@@ -167,10 +169,10 @@ def _fused_linear_softsign_glu_bwd_kernel(
167169
k_mask_nk = offs_k[None, :] < K
168170
acc = tl.zeros([BLOCK_M, BLOCK_K], dtype=tl.float32)
169171

170-
for n_start in range(0, K, BLOCK_N):
172+
for n_start in range(0, N, BLOCK_N):
171173
n_offs = n_start + offs_n
172-
n_mask_mn = n_offs[None, :] < K
173-
n_mask_nk = n_offs[:, None] < K
174+
n_mask_mn = n_offs[None, :] < N
175+
n_mask_nk = n_offs[:, None] < N
174176

175177
left = tl.load(
176178
left_ptr + offs_m[:, None] * stride_l_b + n_offs[None, :] * stride_l_n,
@@ -279,24 +281,25 @@ class FusedLinearSoftSignGLUFn(torch.autograd.Function):
279281
@staticmethod
280282
def forward(ctx, x, weight, bias):
281283
orig_shape = x.shape
282-
K = weight.shape[1]
284+
K = weight.shape[1] # input feature dim (contraction dim)
285+
N = weight.shape[0] // 2 # output dim per GLU half
283286
x_2d = x.reshape(-1, K)
284287
M = x_2d.shape[0]
285288

286-
w_left, w_right = weight.split(K, dim=0)
287-
b_left, b_right = bias.split(K, dim=0)
289+
w_left, w_right = weight.split(N, dim=0)
290+
b_left, b_right = bias.split(N, dim=0)
288291

289-
out = torch.empty(M, K, device=x.device, dtype=x.dtype)
290-
left = torch.empty(M, K, device=x.device, dtype=x.dtype)
291-
gate = torch.empty(M, K, device=x.device, dtype=x.dtype)
292+
out = torch.empty(M, N, device=x.device, dtype=x.dtype)
293+
left = torch.empty(M, N, device=x.device, dtype=x.dtype)
294+
gate = torch.empty(M, N, device=x.device, dtype=x.dtype)
292295

293296
def grid(meta):
294-
return (triton.cdiv(M, meta['BLOCK_M']) * triton.cdiv(K, meta['BLOCK_N']),)
297+
return (triton.cdiv(M, meta['BLOCK_M']) * triton.cdiv(N, meta['BLOCK_N']),)
295298

296299
_fused_linear_softsign_glu_fwd_kernel[grid](
297300
x_2d, w_left, w_right, b_left, b_right,
298301
out, left, gate,
299-
M, K,
302+
M, N, K,
300303
x_2d.stride(0), x_2d.stride(1),
301304
w_left.stride(0), w_left.stride(1),
302305
w_right.stride(0), w_right.stride(1),
@@ -306,34 +309,36 @@ def grid(meta):
306309
)
307310

308311
if x.dim() > 2:
309-
out = out.view(*orig_shape[:-1], K)
310-
left = left.view(M, K)
311-
gate = gate.view(M, K)
312+
out = out.view(*orig_shape[:-1], N)
313+
left = left.view(M, N)
314+
gate = gate.view(M, N)
312315

313316
ctx.save_for_backward(x_2d, weight, left, gate)
314317
ctx.orig_x_shape = orig_shape
318+
ctx.N = N
315319
return out
316320

317321
@staticmethod
318322
def backward(ctx, grad_y):
319323
x, weight, left, gate = ctx.saved_tensors
320324
M, K = x.shape
321-
w_left, w_right = weight.split(K, dim=0)
325+
N = ctx.N
326+
w_left, w_right = weight.split(N, dim=0)
322327

323328
if grad_y.dim() > 2:
324-
grad_y = grad_y.reshape(-1, K)
329+
grad_y = grad_y.reshape(-1, N)
325330

326331
# Step 1: Fused element-wise SoftSignGLU backward
327-
grad_left_pre = torch.empty(M, K, device=x.device, dtype=x.dtype)
328-
grad_gate = torch.empty(M, K, device=x.device, dtype=x.dtype)
332+
grad_left_pre = torch.empty(M, N, device=x.device, dtype=x.dtype)
333+
grad_gate = torch.empty(M, N, device=x.device, dtype=x.dtype)
329334

330335
def elem_grid(meta):
331-
return (triton.cdiv(M, meta['BLOCK_M']) * triton.cdiv(K, meta['BLOCK_K']),)
336+
return (triton.cdiv(M, meta['BLOCK_M']) * triton.cdiv(N, meta['BLOCK_K']),)
332337

333338
_softsign_glu_bwd_elem_kernel[elem_grid](
334339
left, gate, grad_y,
335340
grad_left_pre, grad_gate,
336-
M, K,
341+
M, N,
337342
left.stride(0), left.stride(1),
338343
gate.stride(0), gate.stride(1),
339344
grad_y.stride(0), grad_y.stride(1),
@@ -360,7 +365,7 @@ def bwd_grid(meta):
360365
left, gate, grad_y,
361366
grad_x,
362367
w_left, w_right,
363-
M, K,
368+
M, K, N,
364369
left.stride(0), left.stride(1),
365370
gate.stride(0), gate.stride(1),
366371
grad_y.stride(0), grad_y.stride(1),
@@ -381,20 +386,22 @@ def bwd_grid(meta):
381386
# ---------------------------------------------------------------------------
382387

383388
def fused_linear_softsign_glu(x, weight, bias):
384-
"""Fused Linear(2K, K) + SoftSignGLU.
389+
"""Fused Linear(C, 2*C) + SoftSignGLU, where C = dim (expansion_factor×dim).
385390
386391
y = left * gate / (1 + |gate|)
387392
393+
Supports expansion_factor != 1 by splitting weight at midpoint.
394+
388395
Args:
389-
x: Input [..., K]
390-
weight: [2*K, K]
391-
bias: [2*K]
396+
x: Input [..., K] where K = weight.shape[1] (input dim)
397+
weight: [2*N, K] where N = output dim per GLU half (= K × expansion_factor)
398+
bias: [2*N]
392399
393400
Returns:
394-
[..., K]
401+
[..., N]
395402
"""
396-
assert weight.shape[0] == 2 * weight.shape[1], \
397-
f"Expected [2*K, K], got {weight.shape}"
403+
N = weight.shape[0] // 2
404+
K = weight.shape[1]
398405
# Match weight/bias dtype to input (handles 16-mixed precision where
399406
# weights are fp32 but activations are autocast to fp16)
400407
if weight.dtype != x.dtype:

training/acoustic_task.py

Lines changed: 3 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -144,19 +144,13 @@ def __init__(self):
144144

145145
# ── Fuse LYNXNet2 backbone kernels (in-place) ──
146146
if hparams.get('use_fused_kernels', False):
147-
from modules.kernels.integration import patch_diffusion_module, warmup_fused_backbone
147+
from modules.kernels.integration import patch_diffusion_module
148148
from lightning.pytorch.utilities.rank_zero import rank_zero_info
149149
n = patch_diffusion_module(
150150
self.model.diffusion,
151-
glu_type=hparams['backbone_args'].get('glu_type', 'atanglu'),
151+
glu_type=hparams['backbone_args'].get('glu_type', 'softsign_glu'),
152152
)
153-
rank_zero_info('Fused kernels: patched %d LYNXNet2 blocks, warming up...', n)
154-
if n > 0:
155-
backbone = getattr(self.model.diffusion, 'denoise_fn',
156-
getattr(self.model.diffusion, 'velocity_fn', None))
157-
if backbone is not None:
158-
warmup_fused_backbone(backbone)
159-
rank_zero_info('Fused kernels: autotune complete')
153+
rank_zero_info('Fused kernels: patched %d LYNXNet2 blocks', n)
160154

161155
def _build_model(self):
162156
return DiffSingerAcoustic(

training/variance_task.py

Lines changed: 6 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -132,23 +132,13 @@ def __init__(self):
132132

133133
# ── Fuse LYNXNet2 backbone kernels (in-place) ──
134134
if hparams.get('use_fused_kernels', False):
135-
from modules.kernels.integration import patch_variance_model, warmup_fused_backbone
135+
from modules.kernels.integration import patch_variance_model
136136
from lightning.pytorch.utilities.rank_zero import rank_zero_info
137-
n = patch_variance_model(
138-
self.model,
139-
glu_type=hparams.get('backbone_args', {}).get('glu_type', 'atanglu'),
140-
)
141-
rank_zero_info('Fused kernels: patched %d LYNXNet2 blocks in variance model, warming up...', n)
142-
if n > 0:
143-
glu_type_patched = hparams.get('backbone_args', {}).get('glu_type', 'atanglu')
144-
for predictor_attr in ['pitch_predictor', 'variance_predictor']:
145-
predictor = getattr(self.model, predictor_attr, None)
146-
if predictor is None:
147-
continue
148-
backbone = getattr(predictor, 'denoise_fn', None) or getattr(predictor, 'velocity_fn', None)
149-
if backbone is not None:
150-
warmup_fused_backbone(backbone)
151-
rank_zero_info('Fused kernels: autotune complete')
137+
# Read glu_type from nested predictor configs
138+
pitch_glu = hparams.get('pitch_prediction_args', {}).get('backbone_args', {}).get('glu_type', 'softsign_glu')
139+
var_glu = hparams.get('variances_prediction_args', {}).get('backbone_args', {}).get('glu_type', 'softsign_glu')
140+
n = patch_variance_model(self.model, glu_type=pitch_glu)
141+
rank_zero_info('Fused kernels: patched %d LYNXNet2 blocks in variance model', n)
152142

153143
def _build_model(self):
154144
return DiffSingerVariance(

0 commit comments

Comments
 (0)