Repository navigation
flint: add conv_transpose1d, snake, pad1d, clamp, stft and istft operators; remove waifu::audio's compositions - #127
Conversation
…ove waifu::audio's compositions conv_transpose1d, snake, stft and istft were declared through the C API and Graph, but no backend ran them; waifu::audio composed each out of 6-25 graph nodes instead. They are now flint operators, along with two new ones, pad1d (zero, reflect, replicate) and clamp. The CPU and CUDA run conv_transpose1d (a GEMM, then a col2im pass), snake, pad1d and clamp as kernels (flint/cpu/signal.cc, flint/cuda/signal.cu). Every other device, and stft and istft everywhere, runs a composition in the Operators base class (flint/functional.cc), so Metal and Vulkan keep working. waifu::audio loses pad1d, Padding, reversal, conv_transpose1d, snake, stft, istft, stft_basis, istft_basis and window_envelope, and gains halves(). Its callers in layers.rs, BigVGAN, HiFT, both feature front-ends, w2v-bert, campplus and the CosyVoice3 flow use the graph ops. The at_least, at_most and clamp_unit helpers become Graph::clamp, except Whisper's, whose floor is a graph value. Functions that only carried dtype and device for those helpers drop them. BigVGAN's GPU test compares half with the processor's float32 over 200 frames, on the mean error, instead of the worst of 24 samples against a fixed bound: that statistic ranged 2.5e-3 to 7.6e-3 across mels for the same code, and the old compositions exceeded the old bound on 13 of 20. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WkfaWVmr9m5uxoJ4grPuLm
| for (int n = 0; n < nFft; ++n) envelope[frame * hop + n] += taps[n] * taps[n]; | ||
| } | ||
| for (float &value : envelope) value = std::max(value, 1e-11f); | ||
| Tensor overlap = hostConstant(op, device, dtype, {1, 1, fullLength}, envelope); |
There was a problem hiding this comment.
[P2] Keep the ISTFT envelope floor representable in float16
The envelope is floored at 1e-11f and then cast to the spectrum dtype by hostConstant, but 1e-11 rounds to zero in float16. For a Hann window with centered=false, the first envelope sample therefore remains zero and the subsequent division produces 0/0 (NaN); with nFft=1024, nearby small squared-window values also underflow. Please normalize in float32 or use a dtype-appropriate positive floor, and add a float16, non-centered ISTFT test. The new CUDA STFT/ISTFT test currently only exercises float32.
There was a problem hiding this comment.
Fixed in f80a4f5. composedIstft now divides by the overlap in float32 whatever the spectrum type, and casts only the kept samples back, so the floor and the small end values (w[1]^2 ≈ 9e-11 for a 1024-point Hann) never pass through half. Added a half, non-centred 1024-point stft→istft round trip on CUDA (flint/cuda/signal_test.cc) and on Vulkan. It checks every sample is finite and the covered region matches to 2e-2. With the old normalization it fails at sample 0.
| uint32_t numel) { | ||
| for (uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; idx < numel; | ||
| idx += blockDim.x * gridDim.x) { | ||
| out[idx] = store<T>(fminf(fmaxf(load(input, idx), lo), hi)); |
There was a problem hiding this comment.
[P2] Preserve NaN inputs in the CUDA clamp kernel
fmaxf(NaN, lo) returns lo, so this expression turns a NaN input into a finite bound (for example, clamp(NaN, -0.5, 0.25) returns -0.5). The CPU implementation using std::min/std::max preserves the input NaN, so the operator behaves differently across backends and CUDA silently hides upstream numerical failures. Please explicitly propagate NaN inputs and cover this case in the CPU/CUDA comparison.
There was a problem hiding this comment.
Fixed in f80a4f5, and not only on CUDA. The Vulkan shader (GLSL leaves min/max of a NaN undefined) and the base-class composition (relu makes 0 of a NaN) dropped it too. The CUDA, Vulkan and CPU kernels now test isnan and pass a NaN through. The composition adds exp(-|x|) * 0, which is 0 for every number and for ±inf, and NaN only for a NaN. Tests cover NaN and ±inf for kernel and composition on the CPU, CUDA and Vulkan.
Vulkan ran the four through the Operators base class's compositions. It now has shaders of its own, in f32 and f16: snake.comp, clamp.comp, pad1d.comp and overlap_add1d.comp, the col2im half of a transposed convolution behind the device's own matmul. Metal is the one backend left on the compositions. The snake reduces its angle into [-pi, pi] before the sine, since GLSL bounds sin's error only there and alpha * x is routinely outside it. flint/vulkan/operators_test.cc checks the kernels in both float types against the CPU, including a snake at twenty radians, and still runs the compositions on the same device for Metal's sake. BigVGAN gains it_runs_on_vulkan, sharing it_runs_on_a_card's comparison with the processor; its constants become CARD_*. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WkfaWVmr9m5uxoJ4grPuLm
Two review findings on #127. composedIstft divided the overlap-added wave by the window's squared overlap in the spectrum's type. In half, the overlap's floor of 1e-11 and the ends of a long Hann window (w[1]^2 of 1024 points is 9e-11) round to zero, so a non-centred half istft made NaNs at its first sample. The division is now done in float32 and the kept samples cast back. clamp turned a NaN into a bound on CUDA (fmaxf, fminf), on Vulkan (GLSL leaves min and max of a NaN undefined) and in the composition (a relu makes 0 of a NaN), where the CPU kept it, as torch.clamp does. The kernels test for a NaN and pass it through; the composition adds exp(-|x|) * 0, which is 0 for every number and both infinities and NaN for a NaN. Tests: clamp of NaN and +-inf, kernel and composition, on the CPU, CUDA and Vulkan; a half, non-centred 1024-point istft on CUDA and on Vulkan, which fails at sample 0 with the old normalization. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WkfaWVmr9m5uxoJ4grPuLm
|
Thanks for the review. Both findings were real, and both are fixed in f80a4f5. Details are in the inline threads.
The same push also adds Vulkan kernels for
|
ling0322
left a comment
There was a problem hiding this comment.
Re-reviewed f80a4f5, including the Vulkan kernels and fixes to both previous findings. The previous ISTFT denominator-underflow and CUDA clamp NaN issues are addressed in the implementation and regression tests. Built with Metal enabled on macOS and ran the signal suite: 6 cases, 2519 assertions passed. The process could not access a Metal device, so GPU execution was not validated locally. One additional issue in the shared clamp composition is reproduced below by running that same composition on the CPU in float16.
| // always the right operand. | ||
| if (hasHi) { | ||
| Tensor bound = scalar(op, device, dtype, hi, shape); | ||
| op->sub(x, bound, b); |
There was a problem hiding this comment.
[P2] Avoid cancellation when clamping half-precision values already inside the interval
The arithmetic composition stores each subtraction/addition in the input dtype. For a float16 input containing 0.0001, composedClamp(..., -0.5f, 0.25f, ...) returns 0 rather than the original representable value (~0.000100017): x - lo rounds to exactly 0.5, losing x before the bound is added back. The upper-only case composedClamp(..., -inf, 100.0f, ...) also returns 0. I reproduced both by executing this exact composition with CPU operators in float16; the CPU clamp kernel retains the input. Metal inherits this composition, so values inside the interval can be erased depending on the backend. Please avoid the subtract-and-add cancellation (for example with a comparison/select kernel) and test small in-range half values; the current comparisons predominantly exercise float32 or allow 1e-3 error, which hides this case.
There was a problem hiding this comment.
Fixed in c962b67. This was real, and it isn't limited to half precision: in float32, 1e-8 against a lower bound of -0.5 was lost the same way.
composedClamp no longer clamps by arithmetic. It puts each element beside its bound in an (n, 2) tensor and takes the max and then the min reduction of the last dimension, operators every backend (Metal included) already has. A reduction returns one of its two inputs unchanged, so:
- a value inside the interval comes out exactly itself;
- a value outside it comes out exactly the bound, including ±inf;
- the
exp(-|x|) * 0term still puts back a NaN if a device's reduction drops one.
The tests now require exact equality between the composition and the kernel:
- CPU, float32:
1e-8and-1e30. - CUDA and Vulkan, float16: your
1e-4, a half subnormal, values next to each bound, and ±1e4, for both[-0.5, 0.25]and(-inf, 100].
With the previous composition all three fail: 0 against 1e-08 on the CPU, and 0 against 0.000100017 on CUDA and Vulkan.
…ide the interval stay exact The composition clamped by arithmetic, `lo + relu(x - lo)` and its mirror, which rounds `x - lo` to the input's type: 1e-4 in half against a lower bound of -0.5 came back 0, and so did 1e-8 in float32. Metal runs this composition. Each element now sits beside its bound in an (n, 2) tensor, and the max and then the min reduction of the last dimension pick one of the two as it is: a value inside the interval comes out exactly itself, one beyond it exactly the bound. The NaN a reduction may drop is still put back by exp(-|x|) * 0. Tests compare the composition with the kernel exactly: 1e-8 and -1e30 in float32 on the CPU, and 1e-4, a half subnormal and +-1e4 in half on CUDA and Vulkan. Each fails with the arithmetic composition. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WkfaWVmr9m5uxoJ4grPuLm
Brings in #125 (anima: adapter.rs and dit.rs from diffusers) and #126 (cargo fmt --all over the workspace). The four conflicts -- waifu/src/flint/{functional,mod}.rs and waifu/tests/audio{,_cuda}.rs -- were formatting only on main's side: each is the merge base run through rustfmt. They take this branch's content, and the workspace is then run through cargo fmt --all as CLAUDE.md now asks. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WkfaWVmr9m5uxoJ4grPuLm
…ls (#129) Metal ran these four as compositions in flint/functional.cc, which had never been run on Metal. They now have kernels of their own, ported from the Vulkan shaders, in kernels/signal.metal and signal.cc; stft and istft pick them up through the operators they are composed of. Also fixes waifu-tools, which #127 broke by changing w2v_bert::graph's signature without updating parameters.rs. Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…uild (#130) - waifu-cli: draws_from_a_picture judges a published model with no manifest on the disk by its family (only sdxl can start from a picture), so Anima, Krea 2 and Qwen-Image are left out of the img2img list before they are fetched. - waifu-modelmanager: the config round-trip test uses a path that is absolute on Windows too. - waifu-tools: parameters calls w2v_bert::graph with its signature since #127. Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
This PR moves the audio operators out of
waifu::audioand into flint:conv_transpose1d,snake,stftandistftwere already declared but nothing ran them; they are now real flint operators.pad1d(zero/reflect/replicate) andclampare new operators.conv_transpose1d,snake,pad1dandclampas kernels.stft/istfton every device, run a composition in theOperatorsbase class (flint/functional.cc).The compositions that stood in for these in
waifu::audioare removed (audio.rs: 981 → 473 lines), and every caller uses the graph ops.Why these
Four of them were already plumbed through the C API,
OpandGraph, but no backend ran them, sowaifu::audiobuilt each out of 6–25 graph nodes. The snake alone was six passes over the tensor, and BigVGAN runs it in all 109 of its activations.pad1dhad 13 call sites, mirroring by a matmul against an anti-diagonal.clampreplacesat_least(three copies),at_mostandclamp_unit.What moved
Operators::pad1d/clamp, withFL_PAD_*,fl_pad1dandfl_clampin the C API.conv_transpose1dis one GEMM through the device's own matmul (F::transposedConvolutionColumns), then a col2im kernel that adds the bias.flint/cpu/signal.cc,flint/cuda/signal.cu, andflint/vulkan/signal.ccwithflint/vulkan/shaders/{snake,clamp,pad1d,overlap_add1d}.comp, in f32 and f16.Op::Pad1d,Op::Clamp,Graph::pad1d,Graph::clamp,PadModeandaudio::halves.waifu::audio:pad1d,Padding,reversal,conv_transpose1d,snake,stft,istft,stft_basis,istft_basisandwindow_envelope.dtype/devicefor those helpers drop them, as in audio: remove conv1d and depthwise_conv1d; callers use Graph::conv1d #122.at_least, because its bound is a graph value andclamptakes numbers.Numerics
sin, because GLSL boundssin's error only there.istftdivides by the window's overlap in float32 whatever the spectrum's type.clampkeeps a NaN on every device, astorch.clampdoes. Inside the interval it returns the input exactly on every path, including the composition.BigVGAN's GPU test
The old test compared the card's half-precision output with BigVGAN on the worst of 24 samples, under 4e-3. That statistic is noise:
It now compares the card in half with the processor in float32 over 200 frames: mean error ≤ 2.5e-3, max ≤ 1.5e-2. It runs as
it_runs_on_a_card(CUDA) andit_runs_on_vulkan. Both measure a mean of about 1.5e-3.Review fixes (f80a4f5, c962b67)
istftin half made NaNs at the ends when not centred. The 1e-11 floor and w[1]² ≈ 9e-11 rounded to zero. Fixed by dividing in float32, and tested with a half, non-centred 1024-point round trip on CUDA and Vulkan, which fails at sample 0 without the fix.clampturned a NaN into a bound on CUDA (fmaxf), on Vulkan (GLSLmin/max) and in the composition (relu). The kernels now pass a NaN through; the composition addsexp(-|x|) * 0.clampcomposition erased small values inside the interval.lo + relu(x - lo)roundedx - loto the input type, so1e-4in half (and1e-8in float32) came back as 0. It now uses themaxandminreductions of each element beside its bound, which return one of their inputs unchanged. The tests require exact equality with the kernel, and they fail with the old composition.Tests
./build/unittest: all 115193 assertions pass.cargo test: 393 passed, 0 failed. The CUDA- and Vulkan-gatedaudio_cudaandbigvgantests pass too.Not done
New activations (gelu-tanh, mish, leaky_relu, elu) and a graph rotary op are left for later. They are duplicated in waifu, but each needs Metal shader work.
🤖 Generated with Claude Code
https://claude.ai/code/session_01WkfaWVmr9m5uxoJ4grPuLm