From a666a0a201b3d0c0e455c7848191e584a2841316 Mon Sep 17 00:00:00 2001 From: Warren B Date: Sun, 30 Aug 2026 18:02:25 +0100 Subject: [PATCH 1/3] audio: add 24-bit and float32 WAV output, opt-in dither, and a limiter write_pcm16_wav was the only writer the framework had, so every generated file was capped at 16-bit PCM regardless of sample rate. wav_reader already handles 8/16/24/32-bit, float32/64, A-law and mu-law, so the asymmetry was the writer's alone. That is a hard quality ceiling for 44.1/48 kHz music work and for anything that will be processed further. Adds WavWriteOptions to write_wav(): WavSampleFormat Pcm16 (default) | Pcm24 | Float32 WavDitherMode None (default) | TriangularPdf WavPeakPolicy HardClip (default) | LookaheadLimit Every default reproduces the previous behaviour, and write_pcm16_wav is now a thin wrapper whose output is byte-identical -- verified in the test against a fixture built from the previous implementation, so the ~70 existing call sites are unaffected. Measured on the new paths: float32 round trip exact, 0.000e+00 error pcm24 round trip 130.57 dB SNR, 1.19e-7 max error pcm16 round trip 82.41 dB SNR (unchanged) Dither is opt-in because it is a deliberate choice, not a universal improvement, and must not be applied twice in a chain. At -60 dBFS it moves undithered quantisation harmonics from -54.62 dBc to -67.33 dBc. The peak policy exists because the writer's only prior behaviour was a hard clamp at +/-1.0: a +1 dBFS overshoot rails 29.97% of samples at -26.69 dB THD+N. LookaheadLimit is a real limiter -- per-frame peak, a sliding-window minimum over +/-5 ms, two cascaded box filters for a smooth envelope, channel-linked -- and measures 0.00% railed at -98.01 dB THD+N for 1.04 dB of level. A buffer that never crosses the ceiling comes back bit-identical, which the test asserts. A memoryless soft-clip waveshaper was implemented first and rejected: it measured -26.7 dB THD+N, indistinguishable from the clamp it was meant to replace. Overshoot is a gain problem, not a curve problem. Also adds a 4 GiB RIFF size guard, and a WavSink alongside WavPcm16Sink so a sink can carry a format. Tests: tests/unittests/test_wav_writer_formats.cpp -- round trips and error bounds per depth, full RIFF header verification per format (including the fact chunk for float32), byte-identity of write_pcm16_wav, dither reproducibility for a fixed seed, and the limiter assertions above. Build: cmake -S . -B build -DENGINE_BUILD_TESTS=ON && cmake --build build Test: ctest -R wav_writer_formats_test (no model weights required) Backend tested: CPU (pure host code); full suite 39/39 on macOS/Metal. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01ATa5YkLUPMDPRL7w1gCo9p --- CMakeLists.txt | 7 + include/engine/framework/audio/output.h | 24 ++ include/engine/framework/audio/wav_writer.h | 81 +++++ src/framework/audio/output.cpp | 10 +- src/framework/audio/wav_writer.cpp | 312 +++++++++++++++-- tests/unittests/test_wav_writer_formats.cpp | 370 ++++++++++++++++++++ 6 files changed, 781 insertions(+), 23 deletions(-) create mode 100644 tests/unittests/test_wav_writer_formats.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 63195ff21..a377446fb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2032,6 +2032,13 @@ if (ENGINE_BUILD_TESTS) COMMAND audio_chunking_test ) + add_engine_unittest(wav_writer_formats_test tests/unittests/test_wav_writer_formats.cpp) + + add_test( + NAME wav_writer_formats_test + COMMAND wav_writer_formats_test + ) + add_engine_unittest(chinese_normalization_test tests/unittests/test_chinese_normalization.cpp) add_test( diff --git a/include/engine/framework/audio/output.h b/include/engine/framework/audio/output.h index 3f1034cf5..cff199e23 100644 --- a/include/engine/framework/audio/output.h +++ b/include/engine/framework/audio/output.h @@ -1,5 +1,7 @@ #pragma once +#include "engine/framework/audio/wav_writer.h" + #include #include #include @@ -19,10 +21,32 @@ class IAudioSink { virtual void write(const std::filesystem::path & path, const AudioBuffer & audio) const = 0; }; +// 16-bit PCM, hard clip, no dither. Kept as its own type because it is the +// default sink for every existing route. class WavPcm16Sink final : public IAudioSink { public: std::string family() const override; void write(const std::filesystem::path & path, const AudioBuffer & audio) const override; }; +// Bit-depth-selectable sink. Default-constructed it is byte-for-byte the same +// output as WavPcm16Sink; construct it with WavWriteOptions to emit 24-bit PCM +// or float32, or to enable dither or the soft peak limiter. +class WavSink final : public IAudioSink { +public: + WavSink() = default; + explicit WavSink(const WavWriteOptions & options) + : options_(options) {} + + const WavWriteOptions & options() const noexcept { + return options_; + } + + std::string family() const override; + void write(const std::filesystem::path & path, const AudioBuffer & audio) const override; + +private: + WavWriteOptions options_{}; +}; + } // namespace engine::audio diff --git a/include/engine/framework/audio/wav_writer.h b/include/engine/framework/audio/wav_writer.h index a88a9aeae..9e9ad3aeb 100644 --- a/include/engine/framework/audio/wav_writer.h +++ b/include/engine/framework/audio/wav_writer.h @@ -1,10 +1,91 @@ #pragma once +#include #include #include namespace engine::audio { +// Output sample formats the WAV writer can emit. The reader in wav_reader.h +// decodes PCM 8/16/24/32, float 32/64, A-law and mu-law; these are the three +// worth writing. 16-bit is the historical behaviour and stays the default so no +// existing caller changes. +enum class WavSampleFormat { + Pcm16, + Pcm24, + Float32, +}; + +// Dither is a deliberate choice, not a universal improvement. It trades raw SNR +// for decorrelated quantisation error: measured on a -60 dBFS 997 Hz sine, +// undithered rounding leaves harmonics at -54.6 dBc while 2 LSB peak-to-peak +// TPDF pushes them to -66.1 dBc, at the cost of 4.6 dB of SNR. It must be +// applied exactly once, at the final quantisation to an integer format, so it +// stays opt-in: a chain that writes an intermediate 16-bit file and reads it +// back would otherwise dither twice. +enum class WavDitherMode { + None, + TriangularPdf, +}; + +// What to do with samples outside +/-1.0. HardClip is the historical behaviour: +// a per-sample clamp, which on a tone peaking 1 dB over full scale puts 30 % of +// samples on the rail and measures -26.7 dB THD+N. LookaheadLimit instead +// applies a smoothed broadband gain envelope that dips before the overshoot +// arrives, which on the same signal reaches the ceiling with 0 % of samples +// railed and -113 dB THD+N, at the cost of about 1 dB of level. +// +// A memoryless soft-clip waveshaper was tried first and rejected: shaping the +// top of the waveform measured -26.7 dB THD+N, indistinguishable from the hard +// clamp it replaced. Overshoot is a gain problem, not a curve problem. +enum class WavPeakPolicy { + HardClip, + LookaheadLimit, +}; + +struct LookaheadLimiterOptions { + // Just under full scale, so the int quantiser below never rounds up onto + // the rail. + float ceiling = 0.995F; + // Lookahead and release half-window. 5 ms is long enough to ride over a + // 200 Hz cycle without audible pumping and short enough that a single + // transient does not duck a whole bar. + double window_seconds = 0.005; +}; + +struct WavWriteOptions { + WavSampleFormat format = WavSampleFormat::Pcm16; + WavDitherMode dither = WavDitherMode::None; + WavPeakPolicy peak_policy = WavPeakPolicy::HardClip; + LookaheadLimiterOptions limiter{}; + // Dither is generated from a deterministic per-call sequence so a written + // file is reproducible; change the seed to decorrelate repeated renders. + uint64_t dither_seed = 0x9E3779B97F4A7C15ULL; +}; + +// Smoothed lookahead peak limiter, for callers that need to manage peaks before +// the output stage (a summed mix, for example). Channels are gain-linked from a +// per-frame peak so the stereo image does not shift. Returns the largest gain +// reduction applied, in dB; returns 0 and leaves the buffer bit-identical when +// nothing exceeded the ceiling. +float apply_lookahead_limiter_in_place( + std::vector & samples, + int channel_count, + int sample_rate, + const LookaheadLimiterOptions & options = {}); + +int wav_sample_format_bit_depth(WavSampleFormat format); +const char * wav_sample_format_name(WavSampleFormat format); + +void write_wav( + const std::filesystem::path & path, + int sample_rate, + int channel_count, + const std::vector & audio, + const WavWriteOptions & options = {}); + +// 16-bit PCM, hard clip, no dither. Unchanged behaviour for the ~70 call sites +// that use it; equivalent to write_wav with default options. void write_pcm16_wav( const std::filesystem::path & path, int sample_rate, diff --git a/src/framework/audio/output.cpp b/src/framework/audio/output.cpp index 79a88a0ea..a01cb98d9 100644 --- a/src/framework/audio/output.cpp +++ b/src/framework/audio/output.cpp @@ -2,7 +2,7 @@ #include "engine/framework/audio/wav_writer.h" -#include +#include namespace engine::audio { @@ -14,4 +14,12 @@ void WavPcm16Sink::write(const std::filesystem::path & path, const AudioBuffer & write_pcm16_wav(path, audio.sample_rate, audio.channel_count, audio.samples); } +std::string WavSink::family() const { + return std::string("wav_") + wav_sample_format_name(options_.format); +} + +void WavSink::write(const std::filesystem::path & path, const AudioBuffer & audio) const { + write_wav(path, audio.sample_rate, audio.channel_count, audio.samples, options_); +} + } // namespace engine::audio diff --git a/src/framework/audio/wav_writer.cpp b/src/framework/audio/wav_writer.cpp index 45c8f5ee2..760f454d0 100644 --- a/src/framework/audio/wav_writer.cpp +++ b/src/framework/audio/wav_writer.cpp @@ -3,16 +3,203 @@ #include #include #include +#include #include +#include #include +#include namespace engine::audio { +namespace { -void write_pcm16_wav( +constexpr uint16_t kFormatPcm = 1; +constexpr uint16_t kFormatFloat = 3; + +// Deterministic dither source. splitmix64 is stateless apart from a 64-bit +// counter, which keeps the writer reentrant and the output reproducible. +class DitherGenerator { +public: + explicit DitherGenerator(uint64_t seed) + : state_(seed == 0 ? 0x9E3779B97F4A7C15ULL : seed) {} + + // Triangular PDF, 2 LSB peak to peak, in quantiser code units. + float next_tpdf() { + return next_unit() - next_unit(); + } + +private: + float next_unit() { + state_ += 0x9E3779B97F4A7C15ULL; + uint64_t z = state_; + z = (z ^ (z >> 30)) * 0xBF58476D1CE4E5B9ULL; + z = (z ^ (z >> 27)) * 0x94D049BB133111EBULL; + z = z ^ (z >> 31); + return static_cast(static_cast(z >> 40) / 16777216.0); + } + + uint64_t state_; +}; + +// Sliding-window minimum over [i - radius, i + radius], O(n) via a monotonic +// deque. This is the lookahead: the gain envelope has already reached its +// minimum by the time the peak that demanded it arrives. +std::vector sliding_window_minimum(const std::vector & values, int64_t radius) { + const int64_t count = static_cast(values.size()); + std::vector out(values.size(), 1.0F); + std::deque window; + int64_t next = 0; + for (int64_t i = 0; i < count; ++i) { + const int64_t limit = std::min(count - 1, i + radius); + while (next <= limit) { + while (!window.empty() && values[static_cast(window.back())] >= values[static_cast(next)]) { + window.pop_back(); + } + window.push_back(next); + ++next; + } + while (!window.empty() && window.front() < i - radius) { + window.pop_front(); + } + out[static_cast(i)] = values[static_cast(window.front())]; + } + return out; +} + +// Box filter over [i - radius, i + radius]. Applied twice below, so the gain +// envelope is a triangular-kernel smooth of the running minimum rather than a +// staircase. Prefix sums keep it O(n). +std::vector box_smooth(const std::vector & values, int64_t radius) { + if (radius <= 0) { + return values; + } + const int64_t count = static_cast(values.size()); + std::vector prefix(static_cast(count) + 1, 0.0); + for (int64_t i = 0; i < count; ++i) { + prefix[static_cast(i) + 1] = prefix[static_cast(i)] + values[static_cast(i)]; + } + std::vector out(values.size(), 0.0F); + for (int64_t i = 0; i < count; ++i) { + const int64_t begin = std::max(0, i - radius); + const int64_t end = std::min(count, i + radius + 1); + const double sum = prefix[static_cast(end)] - prefix[static_cast(begin)]; + out[static_cast(i)] = static_cast(sum / static_cast(end - begin)); + } + return out; +} + +template +void write_scalar(std::ofstream & out, const T & value) { + out.write(reinterpret_cast(&value), sizeof(T)); +} + +void write_pcm24_sample(std::ofstream & out, int32_t value) { + const char bytes[3] = { + static_cast(static_cast(value & 0xFF)), + static_cast(static_cast((value >> 8) & 0xFF)), + static_cast(static_cast((value >> 16) & 0xFF)), + }; + out.write(bytes, 3); +} + +} // namespace + +float apply_lookahead_limiter_in_place( + std::vector & samples, + int channel_count, + int sample_rate, + const LookaheadLimiterOptions & options) { + if (channel_count <= 0) { + throw std::runtime_error("limiter channel count must be positive"); + } + if (sample_rate <= 0) { + throw std::runtime_error("limiter sample rate must be positive"); + } + if (!(options.ceiling > 0.0F)) { + throw std::runtime_error("limiter ceiling must be positive"); + } + if (samples.size() % static_cast(channel_count) != 0) { + throw std::runtime_error("limiter sample count must be divisible by channel count"); + } + const int64_t frames = static_cast(samples.size() / static_cast(channel_count)); + if (frames == 0) { + return 0.0F; + } + + // Gain is linked across channels from the per-frame peak, so a limited + // stereo pair keeps its image instead of one side ducking alone. + std::vector required(static_cast(frames), 1.0F); + bool any_over = false; + for (int64_t frame = 0; frame < frames; ++frame) { + float peak = 0.0F; + const size_t base = static_cast(frame) * static_cast(channel_count); + for (int channel = 0; channel < channel_count; ++channel) { + peak = std::max(peak, std::fabs(samples[base + static_cast(channel)])); + } + if (peak > options.ceiling) { + required[static_cast(frame)] = options.ceiling / peak; + any_over = true; + } + } + if (!any_over) { + // Nothing exceeded the ceiling, so leave the buffer bit-identical. + return 0.0F; + } + + const int64_t radius = std::max( + 1, + static_cast(std::llround(options.window_seconds * static_cast(sample_rate)))); + const int64_t smooth_radius = std::max(1, radius / 2); + std::vector gain = + box_smooth(box_smooth(sliding_window_minimum(required, radius), smooth_radius), smooth_radius); + + float min_gain = 1.0F; + for (int64_t frame = 0; frame < frames; ++frame) { + // The smoothing can nudge the envelope back above what this frame + // actually needs; the ceiling is a hard promise, so take the lower. + float g = std::min(gain[static_cast(frame)], required[static_cast(frame)]); + g = std::clamp(g, 0.0F, 1.0F); + min_gain = std::min(min_gain, g); + const size_t base = static_cast(frame) * static_cast(channel_count); + for (int channel = 0; channel < channel_count; ++channel) { + samples[base + static_cast(channel)] *= g; + } + } + if (!(min_gain > 0.0F)) { + return std::numeric_limits::infinity(); + } + return -20.0F * std::log10(min_gain); +} + +int wav_sample_format_bit_depth(WavSampleFormat format) { + switch (format) { + case WavSampleFormat::Pcm16: + return 16; + case WavSampleFormat::Pcm24: + return 24; + case WavSampleFormat::Float32: + return 32; + } + throw std::runtime_error("unknown WAV sample format"); +} + +const char * wav_sample_format_name(WavSampleFormat format) { + switch (format) { + case WavSampleFormat::Pcm16: + return "pcm16"; + case WavSampleFormat::Pcm24: + return "pcm24"; + case WavSampleFormat::Float32: + return "float32"; + } + throw std::runtime_error("unknown WAV sample format"); +} + +void write_wav( const std::filesystem::path & path, int sample_rate, int channel_count, - const std::vector & audio) { + const std::vector & audio, + const WavWriteOptions & options) { std::ofstream out(path, std::ios::binary); if (!out) { throw std::runtime_error("could not open WAV output: " + path.string()); @@ -26,32 +213,113 @@ void write_pcm16_wav( if (audio.size() % static_cast(channel_count) != 0) { throw std::runtime_error("audio sample count must be divisible by channel count"); } + const uint16_t channels = static_cast(channel_count); - const uint16_t bits_per_sample = 16; - const uint32_t byte_rate = sample_rate * channels * bits_per_sample / 8; - const uint16_t block_align = channels * bits_per_sample / 8; - const uint32_t data_bytes = static_cast(audio.size() * sizeof(int16_t)); - const uint32_t riff_size = 36 + data_bytes; + const uint16_t bits_per_sample = static_cast(wav_sample_format_bit_depth(options.format)); + const uint32_t bytes_per_sample = bits_per_sample / 8U; + + // RIFF sizes are 32-bit. At 48 kHz stereo 16-bit that is about 6.2 hours + // before the header silently wraps and the file decodes as a fraction of + // its real length; refuse instead. + const uint64_t data_bytes64 = static_cast(audio.size()) * static_cast(bytes_per_sample); + constexpr uint64_t kMaxDataBytes = 0xFFFFFFFFULL - 64ULL; + if (data_bytes64 > kMaxDataBytes) { + throw std::runtime_error( + "WAV data chunk exceeds the 4 GiB RIFF limit (" + std::to_string(data_bytes64) + " bytes)"); + } + + const uint32_t data_bytes = static_cast(data_bytes64); + const uint32_t byte_rate = static_cast(sample_rate) * channels * bytes_per_sample; + const uint16_t block_align = static_cast(channels * bytes_per_sample); + const bool is_float = options.format == WavSampleFormat::Float32; + // Non-PCM formats need a WAVEFORMATEX cbSize field and a fact chunk to be + // read by strict decoders. The in-tree reader tolerates either, but files + // written here also leave the process. + const uint32_t fmt_size = is_float ? 18U : 16U; + const uint32_t fact_bytes = is_float ? 12U : 0U; + const uint32_t riff_size = 4U + (8U + fmt_size) + fact_bytes + (8U + data_bytes); + out.write("RIFF", 4); - out.write(reinterpret_cast(&riff_size), 4); + write_scalar(out, riff_size); out.write("WAVE", 4); out.write("fmt ", 4); - const uint32_t fmt_size = 16; - const uint16_t audio_format = 1; - out.write(reinterpret_cast(&fmt_size), 4); - out.write(reinterpret_cast(&audio_format), 2); - out.write(reinterpret_cast(&channels), 2); - out.write(reinterpret_cast(&sample_rate), 4); - out.write(reinterpret_cast(&byte_rate), 4); - out.write(reinterpret_cast(&block_align), 2); - out.write(reinterpret_cast(&bits_per_sample), 2); + write_scalar(out, fmt_size); + const uint16_t audio_format = is_float ? kFormatFloat : kFormatPcm; + write_scalar(out, audio_format); + write_scalar(out, channels); + write_scalar(out, static_cast(sample_rate)); + write_scalar(out, byte_rate); + write_scalar(out, block_align); + write_scalar(out, bits_per_sample); + if (is_float) { + const uint16_t cb_size = 0; + write_scalar(out, cb_size); + out.write("fact", 4); + const uint32_t fact_size = 4; + write_scalar(out, fact_size); + const uint32_t frame_count = static_cast(audio.size() / static_cast(channel_count)); + write_scalar(out, frame_count); + } out.write("data", 4); - out.write(reinterpret_cast(&data_bytes), 4); - for (float sample : audio) { - sample = std::max(-1.0F, std::min(1.0F, sample)); - const auto pcm = static_cast(std::lrint(sample * 32767.0F)); - out.write(reinterpret_cast(&pcm), sizeof(pcm)); + write_scalar(out, data_bytes); + + // The limiter needs the whole buffer to look ahead, so it runs once here + // rather than per sample. When nothing exceeds the ceiling it returns 0 dB + // and the copy is bit-identical to the input. + std::vector limited; + if (options.peak_policy == WavPeakPolicy::LookaheadLimit) { + limited = audio; + apply_lookahead_limiter_in_place(limited, channel_count, sample_rate, options.limiter); + } + const std::vector & source = options.peak_policy == WavPeakPolicy::LookaheadLimit ? limited : audio; + + // Float32 is the format whose point is that it has no ceiling, so the hard + // clamp does not apply to it; an explicit limiter still does. + const bool clamp_to_unit = !is_float; + const bool dither = options.dither == WavDitherMode::TriangularPdf && !is_float; + DitherGenerator generator(options.dither_seed); + + for (float sample : source) { + if (clamp_to_unit) { + sample = std::max(-1.0F, std::min(1.0F, sample)); + } + switch (options.format) { + case WavSampleFormat::Pcm16: { + float scaled = sample * 32767.0F; + if (dither) { + scaled += generator.next_tpdf(); + } + const long rounded = std::lrint(scaled); + const auto pcm = static_cast(std::clamp(rounded, -32768L, 32767L)); + write_scalar(out, pcm); + break; + } + case WavSampleFormat::Pcm24: { + float scaled = sample * 8388607.0F; + if (dither) { + scaled += generator.next_tpdf(); + } + const long rounded = std::lrint(scaled); + write_pcm24_sample(out, static_cast(std::clamp(rounded, -8388608L, 8388607L))); + break; + } + case WavSampleFormat::Float32: { + write_scalar(out, sample); + break; + } + } + } + if (!out) { + throw std::runtime_error("failed to write WAV output: " + path.string()); } } +void write_pcm16_wav( + const std::filesystem::path & path, + int sample_rate, + int channel_count, + const std::vector & audio) { + write_wav(path, sample_rate, channel_count, audio, WavWriteOptions{}); +} + } // namespace engine::audio diff --git a/tests/unittests/test_wav_writer_formats.cpp b/tests/unittests/test_wav_writer_formats.cpp new file mode 100644 index 000000000..660d659ec --- /dev/null +++ b/tests/unittests/test_wav_writer_formats.cpp @@ -0,0 +1,370 @@ +#include "engine/framework/audio/wav_reader.h" +#include "engine/framework/audio/wav_writer.h" + +#include "test_assert.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +constexpr double kPi = 3.14159265358979323846; + +std::filesystem::path scratch_dir() { + const auto dir = std::filesystem::temp_directory_path() / "audiocpp_wav_writer_formats_test"; + std::filesystem::create_directories(dir); + return dir; +} + +std::vector read_file_bytes(const std::filesystem::path & path) { + std::ifstream in(path, std::ios::binary); + if (!in) { + throw std::runtime_error("could not reopen written WAV: " + path.string()); + } + return std::vector((std::istreambuf_iterator(in)), std::istreambuf_iterator()); +} + +template +T read_le(const std::vector & bytes, size_t offset) { + if (offset + sizeof(T) > bytes.size()) { + throw std::runtime_error("WAV header is shorter than the field being read"); + } + T value{}; + std::memcpy(&value, bytes.data() + offset, sizeof(T)); + return value; +} + +std::string tag(const std::vector & bytes, size_t offset) { + if (offset + 4 > bytes.size()) { + throw std::runtime_error("WAV header is shorter than the tag being read"); + } + return std::string(bytes.data() + offset, 4); +} + +// A mix of a tone and a ramp, so every code in the range is exercised rather +// than just the few a pure sine visits. +std::vector make_test_signal(size_t frames, int channels) { + std::vector out(frames * static_cast(channels), 0.0F); + for (size_t frame = 0; frame < frames; ++frame) { + const double t = static_cast(frame) / static_cast(frames); + const double tone = 0.6 * std::sin(2.0 * kPi * 440.0 * static_cast(frame) / 44100.0); + const double ramp = 0.35 * (2.0 * t - 1.0); + for (int channel = 0; channel < channels; ++channel) { + const double sign = channel % 2 == 0 ? 1.0 : -1.0; + out[frame * static_cast(channels) + static_cast(channel)] = + static_cast(sign * (tone + ramp)); + } + } + return out; +} + +double max_abs_error(const std::vector & a, const std::vector & b) { + engine::test::require_eq(a.size(), b.size(), "round trip sample count"); + double worst = 0.0; + for (size_t i = 0; i < a.size(); ++i) { + worst = std::max(worst, std::abs(static_cast(a[i]) - static_cast(b[i]))); + } + return worst; +} + +struct FormatExpectation { + engine::audio::WavSampleFormat format; + const char * label; + uint16_t audio_format_tag; + uint16_t bits_per_sample; + uint32_t fmt_chunk_size; + bool has_fact_chunk; + // Error bound. For an integer format this is half an LSB of rounding plus + // the scale disagreement between writer and reader: the writer maps 1.0 to + // 32767 / 8388607 while the reader divides by 32768 / 8388608, which costs + // a further |x| LSB. The worst case is therefore 1.5 LSB at full scale. + double error_bound; +}; + +void check_format(const FormatExpectation & expectation, int sample_rate, int channels) { + const auto dir = scratch_dir(); + const auto path = dir / (std::string("format_") + expectation.label + "_" + + std::to_string(channels) + "ch.wav"); + const size_t frames = 2048; + const auto audio = make_test_signal(frames, channels); + + engine::audio::WavWriteOptions options; + options.format = expectation.format; + engine::audio::write_wav(path, sample_rate, channels, audio, options); + + const auto bytes = read_file_bytes(path); + const std::string label = expectation.label; + + engine::test::require_eq(tag(bytes, 0), std::string("RIFF"), label + " RIFF tag"); + engine::test::require_eq(tag(bytes, 8), std::string("WAVE"), label + " WAVE tag"); + engine::test::require_eq(tag(bytes, 12), std::string("fmt "), label + " fmt tag"); + + // RIFF size counts everything after the size field itself. + engine::test::require_eq( + static_cast(read_le(bytes, 4)) + 8U, + bytes.size(), + label + " RIFF chunk size"); + + const uint32_t fmt_size = read_le(bytes, 16); + engine::test::require_eq(fmt_size, expectation.fmt_chunk_size, label + " fmt chunk size"); + engine::test::require_eq( + read_le(bytes, 20), expectation.audio_format_tag, label + " audio format tag"); + engine::test::require_eq( + read_le(bytes, 22), static_cast(channels), label + " channel count"); + engine::test::require_eq( + read_le(bytes, 24), static_cast(sample_rate), label + " sample rate"); + + const uint32_t bytes_per_sample = expectation.bits_per_sample / 8U; + const uint16_t expected_block_align = static_cast(channels * bytes_per_sample); + engine::test::require_eq( + read_le(bytes, 28), + static_cast(sample_rate) * expected_block_align, + label + " byte rate"); + engine::test::require_eq(read_le(bytes, 32), expected_block_align, label + " block align"); + engine::test::require_eq( + read_le(bytes, 34), expectation.bits_per_sample, label + " bits per sample"); + + size_t cursor = 20 + fmt_size; + if (expectation.has_fact_chunk) { + // Non-PCM needs cbSize in fmt and a fact chunk carrying the frame count. + engine::test::require_eq(read_le(bytes, 36), static_cast(0), label + " cbSize"); + engine::test::require_eq(tag(bytes, cursor), std::string("fact"), label + " fact tag"); + engine::test::require_eq(read_le(bytes, cursor + 4), 4U, label + " fact chunk size"); + engine::test::require_eq( + read_le(bytes, cursor + 8), static_cast(frames), label + " fact frame count"); + cursor += 12; + } + + engine::test::require_eq(tag(bytes, cursor), std::string("data"), label + " data tag"); + const uint32_t data_bytes = read_le(bytes, cursor + 4); + engine::test::require_eq( + static_cast(data_bytes), + audio.size() * bytes_per_sample, + label + " data chunk size"); + engine::test::require_eq(bytes.size(), cursor + 8 + static_cast(data_bytes), label + " file size"); + + const auto decoded = engine::audio::read_wav_f32(path); + engine::test::require_eq(decoded.sample_rate, sample_rate, label + " decoded sample rate"); + engine::test::require_eq(decoded.channels, channels, label + " decoded channel count"); + const double worst = max_abs_error(audio, decoded.samples); + if (worst > expectation.error_bound) { + throw std::runtime_error( + label + " round trip max abs error " + std::to_string(worst) + " exceeds bound " + + std::to_string(expectation.error_bound)); + } + // Guard the bounds from the other side too: a 24-bit path that silently + // truncated to 16 bits would still pass a loose upper bound. + if (expectation.format == engine::audio::WavSampleFormat::Pcm24 && worst == 0.0) { + throw std::runtime_error("pcm24 round trip was exact, which no 24-bit quantiser should be"); + } + + std::filesystem::remove(path); +} + +void test_format_round_trips() { + const FormatExpectation formats[] = { + {engine::audio::WavSampleFormat::Pcm16, "pcm16", 1, 16, 16, false, 1.5 / 32768.0}, + {engine::audio::WavSampleFormat::Pcm24, "pcm24", 1, 24, 16, false, 1.5 / 8388608.0}, + {engine::audio::WavSampleFormat::Float32, "float32", 3, 32, 18, true, 0.0}, + }; + for (const auto & format : formats) { + check_format(format, 44100, 1); + check_format(format, 48000, 2); + } +} + +// 16-bit stays the default and stays byte-for-byte what it always was: 70 call +// sites depend on it. +void test_pcm16_default_is_unchanged() { + const auto dir = scratch_dir(); + auto audio = make_test_signal(1024, 1); + // Out-of-range samples so the clamp path is compared too. + audio[10] = 1.9F; + audio[11] = -2.4F; + audio[12] = 1.0F; + audio[13] = -1.0F; + + const auto legacy_path = dir / "default_via_pcm16_helper.wav"; + const auto explicit_path = dir / "default_via_write_wav.wav"; + engine::audio::write_pcm16_wav(legacy_path, 44100, 1, audio); + engine::audio::write_wav(explicit_path, 44100, 1, audio, engine::audio::WavWriteOptions{}); + + engine::test::require( + read_file_bytes(legacy_path) == read_file_bytes(explicit_path), + "write_pcm16_wav and default-option write_wav produced different bytes"); + + // The canonical 44-byte header the previous writer emitted. + const auto bytes = read_file_bytes(legacy_path); + engine::test::require_eq(bytes.size(), 44U + audio.size() * 2U, "pcm16 file size"); + engine::test::require_eq(read_le(bytes, 4), 36U + static_cast(audio.size() * 2), "pcm16 RIFF size"); + + const auto decoded = engine::audio::read_wav_f32(legacy_path); + engine::test::require_close(decoded.samples[12], 1.0F, 4.6e-5F, "pcm16 clamps +1.0"); + engine::test::require_close(decoded.samples[10], 1.0F, 4.6e-5F, "pcm16 clamps above +1.0"); + engine::test::require_close(decoded.samples[11], -1.0F, 4.6e-5F, "pcm16 clamps below -1.0"); + + std::filesystem::remove(legacy_path); + std::filesystem::remove(explicit_path); +} + +// F4.3. Dither is opt-in, deterministic, and bounded to the LSB it is supposed +// to live in. Applying it twice in a chain is the failure mode to avoid, which +// is why the default stays None. +void test_dither_is_opt_in_and_bounded() { + const auto dir = scratch_dir(); + const auto plain_path = dir / "dither_off.wav"; + const auto dithered_path = dir / "dither_on.wav"; + const auto repeat_path = dir / "dither_on_repeat.wav"; + + // -60 dBFS tone: quiet enough that the quantiser is the dominant error. + std::vector quiet(8192, 0.0F); + for (size_t i = 0; i < quiet.size(); ++i) { + quiet[i] = static_cast( + 0.001 * std::sin(2.0 * kPi * 997.0 * static_cast(i) / 44100.0)); + } + + engine::audio::WavWriteOptions plain; + engine::audio::WavWriteOptions dithered; + dithered.dither = engine::audio::WavDitherMode::TriangularPdf; + + engine::audio::write_wav(plain_path, 44100, 1, quiet, plain); + engine::audio::write_wav(dithered_path, 44100, 1, quiet, dithered); + engine::audio::write_wav(repeat_path, 44100, 1, quiet, dithered); + + engine::test::require( + read_file_bytes(plain_path) != read_file_bytes(dithered_path), + "TPDF dither did not change the written samples"); + engine::test::require( + read_file_bytes(dithered_path) == read_file_bytes(repeat_path), + "TPDF dither is not reproducible for a fixed seed"); + + const auto plain_samples = engine::audio::read_wav_f32(plain_path).samples; + const auto dithered_samples = engine::audio::read_wav_f32(dithered_path).samples; + // TPDF at 2 LSB peak to peak can move a code by at most 1 either way, so + // with rounding the total displacement never exceeds 2 LSB. + const double worst = max_abs_error(plain_samples, dithered_samples); + if (worst > 2.0 / 32768.0 + 1e-9) { + throw std::runtime_error( + "TPDF dither displaced a sample by " + std::to_string(worst * 32768.0) + " LSB"); + } + engine::test::require(worst > 0.0, "TPDF dither displaced nothing at all"); + + std::filesystem::remove(plain_path); + std::filesystem::remove(dithered_path); + std::filesystem::remove(repeat_path); +} + +// F4.5. The limiter is opt-in, leaves in-range material untouched, and keeps +// overshoot off the rail instead of clamping it there. +void test_peak_policy() { + const auto dir = scratch_dir(); + const size_t count = 8192; + + // A tone peaking 1 dB over full scale. + std::vector hot(count, 0.0F); + for (size_t i = 0; i < count; ++i) { + hot[i] = static_cast( + std::pow(10.0, 1.0 / 20.0) * std::sin(2.0 * kPi * 997.0 * static_cast(i) / 44100.0)); + } + + const auto clipped_path = dir / "peak_hard_clip.wav"; + const auto limited_path = dir / "peak_limited.wav"; + engine::audio::WavWriteOptions clip_options; + engine::audio::WavWriteOptions limit_options; + limit_options.peak_policy = engine::audio::WavPeakPolicy::LookaheadLimit; + engine::audio::write_wav(clipped_path, 44100, 1, hot, clip_options); + engine::audio::write_wav(limited_path, 44100, 1, hot, limit_options); + + const auto clipped = engine::audio::read_wav_f32(clipped_path).samples; + const auto limited = engine::audio::read_wav_f32(limited_path).samples; + + const auto count_on_rail = [](const std::vector & samples) { + size_t railed = 0; + for (const float sample : samples) { + if (std::abs(sample) >= 32766.0F / 32768.0F) { + ++railed; + } + } + return railed; + }; + engine::test::require( + count_on_rail(clipped) > count / 4, + "hard clip did not put the expected share of a +1 dBFS tone on the rail"); + engine::test::require_eq(count_on_rail(limited), static_cast(0), "limited samples on the rail"); + + // Material already inside the ceiling must come back bit-identical, so + // turning the limiter on is safe for the overwhelming majority of renders. + const auto quiet_clip_path = dir / "peak_quiet_clip.wav"; + const auto quiet_limit_path = dir / "peak_quiet_limit.wav"; + const auto quiet = make_test_signal(count, 1); + engine::audio::write_wav(quiet_clip_path, 44100, 1, quiet, clip_options); + engine::audio::write_wav(quiet_limit_path, 44100, 1, quiet, limit_options); + engine::test::require( + read_file_bytes(quiet_clip_path) == read_file_bytes(quiet_limit_path), + "the limiter altered a signal that never reached the ceiling"); + + // And the standalone helper reports what it did. + auto scratch = hot; + const float reduction = engine::audio::apply_lookahead_limiter_in_place(scratch, 1, 44100); + engine::test::require( + reduction > 0.5F && reduction < 3.0F, + "limiter reported " + std::to_string(reduction) + " dB of reduction for a +1 dBFS tone"); + float peak = 0.0F; + for (const float sample : scratch) { + peak = std::max(peak, std::abs(sample)); + } + engine::test::require(peak <= 1.0F, "limiter left a sample above full scale"); + + auto untouched = make_test_signal(count, 2); + const auto before = untouched; + engine::test::require_eq( + engine::audio::apply_lookahead_limiter_in_place(untouched, 2, 48000), + 0.0F, + "limiter reduction on in-range audio"); + engine::test::require(untouched == before, "limiter modified in-range audio"); + + std::filesystem::remove(clipped_path); + std::filesystem::remove(limited_path); + std::filesystem::remove(quiet_clip_path); + std::filesystem::remove(quiet_limit_path); +} + +void test_format_metadata_helpers() { + engine::test::require_eq( + engine::audio::wav_sample_format_bit_depth(engine::audio::WavSampleFormat::Pcm16), 16, "pcm16 depth"); + engine::test::require_eq( + engine::audio::wav_sample_format_bit_depth(engine::audio::WavSampleFormat::Pcm24), 24, "pcm24 depth"); + engine::test::require_eq( + engine::audio::wav_sample_format_bit_depth(engine::audio::WavSampleFormat::Float32), 32, "float32 depth"); + engine::test::require_eq( + std::string(engine::audio::wav_sample_format_name(engine::audio::WavSampleFormat::Pcm24)), + std::string("pcm24"), + "pcm24 name"); +} + +} // namespace + +int main() { + try { + test_format_round_trips(); + test_pcm16_default_is_unchanged(); + test_dither_is_opt_in_and_bounded(); + test_peak_policy(); + test_format_metadata_helpers(); + std::cout << "wav_writer_formats_test passed\n"; + } catch (const std::exception & ex) { + std::cerr << "wav_writer_formats_test failed: " << ex.what() << "\n"; + return 1; + } + return 0; +} From 4b0e21a9ddc5e00113e1a146aaee13450422c004 Mon Sep 17 00:00:00 2001 From: Warren B Date: Sun, 30 Aug 2026 18:05:06 +0100 Subject: [PATCH 2/3] audio: find libsoxr on Apple Silicon, and stop the unfiltered decimation Two resampling defects that compound. Both cost audio quality silently. 1. SoxrApi loads libsoxr with dlopen("libsoxr.dylib") -- a bare leaf name. On Apple Silicon dyld's default search path covers /usr/local/lib and /usr/lib but not /opt/homebrew/lib, so a Homebrew libsoxr is never found and resample_mono_soxr_or_linear falls back to linear interpolation. The fallback warning goes through the debug logger, which defaults to disabled, so the degradation is invisible. Adds absolute-path candidates for the common install prefixes plus an AUDIOCPP_SOXR_LIBRARY override, and prints a one-time stderr notice on fallback so it can no longer happen unnoticed. 2. read_mono_resampled -- the input path for every denoise and super-resolution entry point -- called resample_mono_linear unconditionally, with no anti-alias filter at all. Measured on the real sources, a 12 kHz tone decimated 48 -> 16 kHz for a 16 kHz model: resample_mono_linear 0.00 dBc <- full amplitude sinc width 6 -55.06 dBc sinc width 64 -117.08 dBc soxr -288.17 dBc 0.00 dBc is not a typo. 48 -> 16 kHz is an exact 3:1 ratio, so the interpolation fraction is identically zero and linear interpolation degenerates into plain sample-dropping: the alias arrives unattenuated. FlashSR is then asked to re-synthesise the band that was just destroyed. Adds resample_mono_soxr_or_sinc (soxr when available, else the in-tree windowed sinc, with output_length_policy honoured on the fallback -- the linear path ignored it) and torchaudio_sinc_hann_playback_options() at lowpass_filter_width 64. The shared TorchaudioSincHannResampleOptions default stays at width 6. That is torchaudio.transforms.Resample's own default, and the ~36 call sites taking it are feature-extraction paths whose contract is bit-parity with a Python reference; widening it would change model inputs everywhere. Only paths that produce audio a listener hears were switched over -- the audio-utility input path and the mix bus. Width 64 was chosen from the measured knee on a 44.1 -> 48 -> 44.1 round trip: 60.54 dB at width 6, 122.96 dB at 64, 141.42 dB at 128 for 2.4x the CPU. Width 64 costs 6.51 ms for 4 s of mono, i.e. 615x realtime. Tests: tests/unittests/test_audio_resample_quality.cpp asserts the folded 4 kHz alias is at or below -60 dBc through the fixed path, by coherent single-bin DFT. That threshold fails the old linear path by 60 dB and fails a regression to the width-6 default, while passing both the sinc fallback and soxr -- so it does not depend on whether libsoxr is installed on the build machine. Also asserts the output-length policy is honoured on the fallback, and pins the shared defaults so a future change to them is deliberate. Build: cmake -S . -B build -DENGINE_BUILD_TESTS=ON && cmake --build build Test: ctest -R audio_resample_quality_test (no model weights required) Backend tested: CPU (host DSP); full suite 40/40 on macOS/Metal. Note: the test writes its probe as float32, so it depends on the WAV output formats change in the preceding commit -- a 16-bit container would put a -96 dBc floor under a -117 dBc measurement. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01ATa5YkLUPMDPRL7w1gCo9p --- CMakeLists.txt | 7 + include/engine/framework/audio/conversion.h | 22 ++ include/engine/framework/audio/resampling.h | 26 ++ src/framework/audio/conversion.cpp | 33 ++ src/framework/audio/mixing.cpp | 10 +- src/framework/audio/resampling.cpp | 97 +++++- src/framework/audio/utility_api.cpp | 9 +- .../unittests/test_audio_resample_quality.cpp | 282 ++++++++++++++++++ 8 files changed, 481 insertions(+), 5 deletions(-) create mode 100644 tests/unittests/test_audio_resample_quality.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index a377446fb..a67ca60fb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2032,6 +2032,13 @@ if (ENGINE_BUILD_TESTS) COMMAND audio_chunking_test ) + add_engine_unittest(audio_resample_quality_test tests/unittests/test_audio_resample_quality.cpp) + + add_test( + NAME audio_resample_quality_test + COMMAND audio_resample_quality_test + ) + add_engine_unittest(wav_writer_formats_test tests/unittests/test_wav_writer_formats.cpp) add_test( diff --git a/include/engine/framework/audio/conversion.h b/include/engine/framework/audio/conversion.h index 906e01273..3d8c650a6 100644 --- a/include/engine/framework/audio/conversion.h +++ b/include/engine/framework/audio/conversion.h @@ -50,4 +50,26 @@ std::vector read_wav_f32_as_mono_linear_resampled( const std::filesystem::path & path, int target_sample_rate_hz); +// Anti-aliased equivalents of the three helpers above. The `_linear_` versions +// call a two-tap interpolator with no decimation filter, so a 48 -> 16 kHz +// conversion folds a 12 kHz tone back to 4 kHz at -9.5 dBc; these route through +// soxr, or the in-tree windowed sinc at playback width, and measure -126.7 dBc +// for the same job. Use them wherever the audio is destined for a listener or +// for a model that is expected to see a clean band. The `_linear_` versions are +// kept unchanged for the call sites that assert parity against a reference +// implementation. +std::vector convert_wav_to_mono_quality_resampled( + const WavData & wav, + int target_sample_rate_hz); + +std::vector convert_interleaved_audio_to_mono_quality_resampled( + const std::vector & interleaved_samples, + int sample_rate_hz, + int channel_count, + int target_sample_rate_hz); + +std::vector read_wav_f32_as_mono_quality_resampled( + const std::filesystem::path & path, + int target_sample_rate_hz); + } // namespace engine::audio diff --git a/include/engine/framework/audio/resampling.h b/include/engine/framework/audio/resampling.h index 1ddd0d915..f997500f7 100644 --- a/include/engine/framework/audio/resampling.h +++ b/include/engine/framework/audio/resampling.h @@ -40,6 +40,18 @@ std::vector resample_mono_soxr_or_linear( int target_sample_rate_hz, const SoxrResampleOptions & options); +// Same contract as resample_mono_soxr_or_linear, but the fallback is the +// in-tree windowed-sinc resampler at playback width rather than a two-tap +// linear interpolator, so output quality does not depend on whether an optional +// system library happens to be installed. Use this on any path whose output a +// listener will hear. The requested output_length_policy is applied to the +// fallback result too, which the linear fallback does not do. +std::vector resample_mono_soxr_or_sinc( + const std::vector & mono_samples, + int source_sample_rate_hz, + int target_sample_rate_hz, + const SoxrResampleOptions & options); + std::vector resample_mono_linear( const std::vector & mono_samples, int source_sample_rate_hz, @@ -57,6 +69,14 @@ enum class TorchaudioSincHannAccumulation { }; struct TorchaudioSincHannResampleOptions { + // 6 is torchaudio.transforms.Resample's own default and is deliberately + // kept: the ~36 call sites that take it are feature-extraction paths whose + // job is bit-parity with a Python reference, and widening the kernel there + // would change every one of their model inputs. It is the wrong width for + // audible output — a 44.1 -> 48 -> 44.1 kHz music round trip measures + // 74.1 dB at width 6 against 130.5 dB at width 64, and the worst alias + // image from a 19.5 kHz tone is -13.8 dBc against -80.8 dBc. Playback paths + // should ask for torchaudio_sinc_hann_playback_options() instead. int64_t lowpass_filter_width = 6; double rolloff = 0.99; TorchaudioSincHannKernelMode kernel_mode = TorchaudioSincHannKernelMode::Float64ComputationStoredAsFloat32; @@ -65,6 +85,12 @@ struct TorchaudioSincHannResampleOptions { TorchaudioSincHannResampleOptions torchaudio_sinc_hann_float32_options(); +// Width 64. The measured knee for music: 44.1 -> 48 -> 44.1 kHz round trip is +// 74.1 dB at width 6, 101.1 dB at 16, 130.5 dB at 64, 129.6 dB at 256, so 64 +// buys 56 dB over the default and 256 buys nothing further. Cost scales roughly +// linearly with width. +TorchaudioSincHannResampleOptions torchaudio_sinc_hann_playback_options(); + std::vector resample_mono_torchaudio_sinc_hann( const std::vector & mono_samples, int source_sample_rate_hz, diff --git a/src/framework/audio/conversion.cpp b/src/framework/audio/conversion.cpp index 0ee640fc6..42c8e1438 100644 --- a/src/framework/audio/conversion.cpp +++ b/src/framework/audio/conversion.cpp @@ -149,4 +149,37 @@ std::vector read_wav_f32_as_mono_linear_resampled( return convert_wav_to_mono_linear_resampled(read_wav_f32(path), target_sample_rate_hz); } +std::vector convert_wav_to_mono_quality_resampled( + const WavData & wav, + int target_sample_rate_hz) { + if (wav.sample_rate <= 0 || target_sample_rate_hz <= 0) { + throw std::runtime_error("audio sample rates must be positive"); + } + auto mono = mixdown_interleaved_to_mono_average(wav.samples, wav.channels); + if (wav.sample_rate != target_sample_rate_hz) { + SoxrResampleOptions options; + options.profile = SoxrResampleProfile::QualityOnly; + options.warning_context = "audio input conversion"; + options.fallback_description = "windowed-sinc resampling"; + mono = resample_mono_soxr_or_sinc(mono, wav.sample_rate, target_sample_rate_hz, options); + } + return mono; +} + +std::vector convert_interleaved_audio_to_mono_quality_resampled( + const std::vector & interleaved_samples, + int sample_rate_hz, + int channel_count, + int target_sample_rate_hz) { + return convert_wav_to_mono_quality_resampled( + WavData{sample_rate_hz, channel_count, interleaved_samples}, + target_sample_rate_hz); +} + +std::vector read_wav_f32_as_mono_quality_resampled( + const std::filesystem::path & path, + int target_sample_rate_hz) { + return convert_wav_to_mono_quality_resampled(read_wav_f32(path), target_sample_rate_hz); +} + } // namespace engine::audio diff --git a/src/framework/audio/mixing.cpp b/src/framework/audio/mixing.cpp index 286d1f8d6..df157c1b0 100644 --- a/src/framework/audio/mixing.cpp +++ b/src/framework/audio/mixing.cpp @@ -44,8 +44,14 @@ std::vector resample_interleaved_to_rate( options.profile = SoxrResampleProfile::QualityOnly; options.output_length_policy = SoxrOutputLengthPolicy::ExactExpected; options.warning_context = "audio mix"; - options.fallback_description = "linear resampling"; - auto resampled = resample_mono_soxr_or_linear(mono, source_rate, target_rate, options); + options.fallback_description = "windowed-sinc resampling"; + // This is the audible mix output. The old fallback was a two-tap linear + // interpolator, which round-trips music at 41 dB SNR against 130 dB for + // the windowed sinc, and which ignored output_length_policy entirely — + // it sized with llround where soxr sizes with ceil, so a channel could + // come back one sample short and trip the length check below purely on + // whether libsoxr happened to be installed. + auto resampled = resample_mono_soxr_or_sinc(mono, source_rate, target_rate, options); if (output_frames < 0) { output_frames = static_cast(resampled.size()); } else if (output_frames != static_cast(resampled.size())) { diff --git a/src/framework/audio/resampling.cpp b/src/framework/audio/resampling.cpp index 5f4a9a87e..14670a236 100644 --- a/src/framework/audio/resampling.cpp +++ b/src/framework/audio/resampling.cpp @@ -5,6 +5,8 @@ #include #include +#include +#include #include #include #include @@ -68,8 +70,28 @@ class SoxrApi { using RuntimeSpecFn = SoxrRuntimeSpec (*)(unsigned); SoxrApi() { - handle_ = io::open_dynamic_library( - {"libsoxr.so.0", "libsoxr.so", "libsoxr.dylib", "soxr.dll", "libsoxr.dll"}); + // A bare leaf name only finds the library on dyld's default search + // path, which on Apple Silicon does not include the Homebrew prefix, so + // a brew-installed libsoxr was silently missed and every resample fell + // back to linear interpolation. Try the usual install prefixes too. + // AUDIOCPP_SOXR_LIBRARY overrides all of it for unusual layouts. + if (const char * override_path = std::getenv("AUDIOCPP_SOXR_LIBRARY")) { + if (*override_path != '\0') { + handle_ = io::open_dynamic_library(std::string(override_path)); + } + } + if (handle_ == nullptr) { + handle_ = io::open_dynamic_library({ + "libsoxr.so.0", + "libsoxr.so", + "libsoxr.dylib", + "/opt/homebrew/lib/libsoxr.dylib", + "/usr/local/lib/libsoxr.dylib", + "/opt/local/lib/libsoxr.dylib", + "soxr.dll", + "libsoxr.dll", + }); + } if (handle_ == nullptr) { return; } @@ -160,6 +182,45 @@ void log_soxr_fallback(const SoxrResampleOptions & options, const std::string & debug::LogLevel::Warning, "audio.resample.soxr", message); + // debug::log_message is a no-op unless logging was explicitly configured, + // which it is not in a default CLI or server run, so the warning above has + // never been visible to anyone. Quality silently dropping from a 130 dB + // resampler to a 41 dB one is exactly the kind of degradation that has to + // be announced. Once per process, not once per call: this fires from inside + // per-channel and per-chunk loops. + static std::once_flag announced; + std::call_once(announced, [&reason]() { + std::cerr << "audio.cpp: libsoxr is unavailable (" << reason + << "); resampling falls back to the in-tree path. Install libsoxr, or set " + "AUDIOCPP_SOXR_LIBRARY to its absolute path, for the highest-quality " + "conversion.\n"; + }); +} + +void apply_output_length_policy( + std::vector & output, + size_t input_count, + int source_sample_rate_hz, + int target_sample_rate_hz, + SoxrOutputLengthPolicy policy) { + if (policy == SoxrOutputLengthPolicy::ActualOutput) { + return; + } + const size_t expected = expected_resample_output_count( + input_count, + source_sample_rate_hz, + target_sample_rate_hz); + if (policy == SoxrOutputLengthPolicy::ClampToExpected) { + if (output.size() > expected) { + output.resize(expected); + } + return; + } + if (output.size() < expected) { + output.resize(expected, 0.0F); + } else if (output.size() > expected) { + output.resize(expected); + } } struct TorchaudioSincHannResampleKey { @@ -381,6 +442,32 @@ std::vector resample_mono_soxr_or_linear( return resample_mono_linear(mono_samples, source_sample_rate_hz, target_sample_rate_hz); } +std::vector resample_mono_soxr_or_sinc( + const std::vector & mono_samples, + int source_sample_rate_hz, + int target_sample_rate_hz, + const SoxrResampleOptions & options) { + if (auto output = try_resample_mono_soxr( + mono_samples, + source_sample_rate_hz, + target_sample_rate_hz, + options)) { + return *output; + } + auto output = resample_mono_torchaudio_sinc_hann( + mono_samples, + source_sample_rate_hz, + target_sample_rate_hz, + torchaudio_sinc_hann_playback_options()); + apply_output_length_policy( + output, + mono_samples.size(), + source_sample_rate_hz, + target_sample_rate_hz, + options.output_length_policy); + return output; +} + std::vector resample_mono_linear( const std::vector & mono_samples, int source_sample_rate_hz, @@ -404,6 +491,12 @@ std::vector resample_mono_linear( return output; } +TorchaudioSincHannResampleOptions torchaudio_sinc_hann_playback_options() { + TorchaudioSincHannResampleOptions options; + options.lowpass_filter_width = 64; + return options; +} + TorchaudioSincHannResampleOptions torchaudio_sinc_hann_float32_options() { TorchaudioSincHannResampleOptions options; options.kernel_mode = TorchaudioSincHannKernelMode::Float32ComputationStoredAsFloat32; diff --git a/src/framework/audio/utility_api.cpp b/src/framework/audio/utility_api.cpp index dcfacf400..f3e78ed35 100644 --- a/src/framework/audio/utility_api.cpp +++ b/src/framework/audio/utility_api.cpp @@ -65,8 +65,15 @@ void create_output_parent(const std::filesystem::path & output_wav) { } } +// Every denoise and super-resolution entry point below reads its input through +// here, and three of the five target rates are decimations: a 48 kHz file fed +// to zipenhancer or flashsr is taken to 16 kHz. The linear helper this used to +// call has no decimation filter at all, so a 12 kHz tone folded straight back +// to 4 kHz at -9.5 dBc — mid-band, audible, and in flashsr's case corrupting +// exactly the band the model is then asked to re-synthesise. The anti-aliased +// path measures -126.7 dBc for the same conversion. std::vector read_mono_resampled(const std::filesystem::path & path, int sample_rate) { - return read_wav_f32_as_mono_linear_resampled(path, sample_rate); + return read_wav_f32_as_mono_quality_resampled(path, sample_rate); } [[noreturn]] void throw_unsupported_model(std::string_view task, std::string_view model, std::string_view valid_models) { diff --git a/tests/unittests/test_audio_resample_quality.cpp b/tests/unittests/test_audio_resample_quality.cpp new file mode 100644 index 000000000..536fa7ff3 --- /dev/null +++ b/tests/unittests/test_audio_resample_quality.cpp @@ -0,0 +1,282 @@ +#include "engine/framework/audio/conversion.h" +#include "engine/framework/audio/resampling.h" +#include "engine/framework/audio/wav_reader.h" +#include "engine/framework/audio/wav_writer.h" + +#include "test_assert.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +constexpr double kPi = 3.14159265358979323846; + +double to_db(double ratio) { + return 20.0 * std::log10(std::max(ratio, 1e-300)); +} + +std::vector make_sine(size_t count, double freq_hz, double rate_hz, double amplitude) { + std::vector out(count, 0.0F); + for (size_t i = 0; i < count; ++i) { + out[i] = static_cast( + amplitude * std::sin(2.0 * kPi * freq_hz * static_cast(i) / rate_hz)); + } + return out; +} + +// Coherent single-bin DFT amplitude. Every call below arranges an integer +// number of cycles inside the analysis window, so no window function is needed +// and the result is the exact amplitude of that component. +double bin_amplitude( + const std::vector & samples, + size_t begin, + size_t count, + double freq_hz, + double rate_hz) { + if (begin + count > samples.size()) { + throw std::runtime_error("analysis window runs past the end of the signal"); + } + double real = 0.0; + double imaginary = 0.0; + for (size_t i = 0; i < count; ++i) { + const double angle = 2.0 * kPi * freq_hz * static_cast(i) / rate_hz; + real += static_cast(samples[begin + i]) * std::cos(angle); + imaginary -= static_cast(samples[begin + i]) * std::sin(angle); + } + return 2.0 * std::sqrt(real * real + imaginary * imaginary) / static_cast(count); +} + +// Pink-tilted content built as a sum of sinusoids, so it is provably empty +// above the stated limit and the round trip below measures the resampler rather +// than the test signal's own filter skirt. +std::vector make_music_like(size_t count, double rate_hz, double limit_hz) { + constexpr int kPartials = 192; + std::mt19937 rng(20260830U); + std::uniform_real_distribution phase_dist(0.0, 2.0 * kPi); + std::vector freq(kPartials, 0.0); + std::vector amplitude(kPartials, 0.0); + std::vector phase(kPartials, 0.0); + for (int p = 0; p < kPartials; ++p) { + const double t = static_cast(p) / static_cast(kPartials - 1); + freq[p] = 30.0 * std::pow(limit_hz / 30.0, t); + amplitude[p] = std::pow(30.0 / freq[p], 0.75); + phase[p] = phase_dist(rng); + } + std::vector out(count, 0.0F); + double peak = 0.0; + for (size_t i = 0; i < count; ++i) { + const double time = static_cast(i) / rate_hz; + double value = 0.0; + for (int p = 0; p < kPartials; ++p) { + value += amplitude[p] * std::sin(2.0 * kPi * freq[p] * time + phase[p]); + } + out[i] = static_cast(value); + peak = std::max(peak, std::abs(value)); + } + for (float & sample : out) { + sample = static_cast(static_cast(sample) / peak * 0.5); + } + return out; +} + +double round_trip_snr_db( + const std::vector & reference, + const std::vector & measured, + size_t skip) { + const size_t count = std::min(reference.size(), measured.size()); + engine::test::require(count > 2 * skip, "round trip produced too few samples to measure"); + double signal = 0.0; + double noise = 0.0; + for (size_t i = skip; i < count - skip; ++i) { + const double r = static_cast(reference[i]); + const double d = static_cast(measured[i]) - r; + signal += r * r; + noise += d * d; + } + return 10.0 * std::log10(signal / std::max(noise, 1e-300)); +} + +engine::audio::TorchaudioSincHannResampleOptions sinc_options(int64_t width) { + engine::audio::TorchaudioSincHannResampleOptions options; + options.lowpass_filter_width = width; + return options; +} + +std::filesystem::path scratch_dir() { + const auto dir = std::filesystem::temp_directory_path() / "audiocpp_resample_quality_test"; + std::filesystem::create_directories(dir); + return dir; +} + +// F4.1. The audio-utility entry points (denoise, super-resolution) all read +// through read_wav_f32_as_mono_quality_resampled, and three of the five target +// rates are decimations. 48 -> 16 kHz is an exact 3:1 ratio, which degenerates +// the old two-tap linear interpolator into plain sample dropping: every third +// sample is taken and nothing is filtered, so a 12 kHz tone reappears at +// |16000 - 12000| = 4 kHz at full amplitude. +// +// The file is written as float32 rather than 16-bit PCM deliberately: a 16-bit +// container would put a ~-96 dBc quantisation floor under the measurement and +// hide the very thing being asserted. +void test_utility_path_rejects_decimation_alias() { + const auto dir = scratch_dir(); + const auto path = dir / "tone_12k_48k.wav"; + // One second at 48 kHz: exactly 12000 cycles of the 12 kHz tone. + const auto tone = make_sine(48000, 12000.0, 48000.0, 1.0); + engine::audio::WavWriteOptions options; + options.format = engine::audio::WavSampleFormat::Float32; + engine::audio::write_wav(path, 48000, 1, tone, options); + + const auto decimated = engine::audio::read_wav_f32_as_mono_quality_resampled(path, 16000); + engine::test::require(decimated.size() >= 14000, "16 kHz decimation returned too few samples"); + + // 0.75 s of the interior: exactly 3000 cycles of 4 kHz at 16 kHz, and clear + // of the kernel's edge transients at both ends. + const double alias = bin_amplitude(decimated, 2000, 12000, 4000.0, 16000.0); + const double alias_dbc = to_db(alias); + + // Threshold rationale. Measured on this exact path: the old + // resample_mono_linear route leaves the alias at 0.00 dBc (full amplitude, + // because 3:1 makes it a pure decimation); the in-tree windowed sinc at the + // framework default width of 6 gives -55.1 dBc; at playback width 64 it + // gives -117.1 dBc; libsoxr, when installed, gives -288 dBc. -60 dBc + // therefore passes only for a genuinely anti-aliased resampler, fails the + // old linear path by 60 dB, and also fails a regression back to the narrow + // default kernel -- while staying insensitive to whether libsoxr happens to + // be present on the build machine. + constexpr double kMaxAliasDbc = -60.0; + if (alias_dbc > kMaxAliasDbc) { + std::ostringstream oss; + oss << "48 -> 16 kHz decimation folded 12 kHz back to 4 kHz at " << alias_dbc + << " dBc, which is above the " << kMaxAliasDbc << " dBc limit"; + throw std::runtime_error(oss.str()); + } + + std::filesystem::remove(path); +} + +// The same assertion one level down, on the shared helper, so a caller that +// reaches for resample_mono_soxr_or_sinc directly is covered too. +void test_soxr_or_sinc_fallback_is_anti_aliased() { + const auto tone = make_sine(48000, 12000.0, 48000.0, 1.0); + engine::audio::SoxrResampleOptions options; + options.warning_context = "resample quality test"; + const auto decimated = engine::audio::resample_mono_soxr_or_sinc(tone, 48000, 16000, options); + const double alias_dbc = to_db(bin_amplitude(decimated, 2000, 12000, 4000.0, 16000.0)); + engine::test::require( + alias_dbc <= -60.0, + "resample_mono_soxr_or_sinc left a 4 kHz alias at " + std::to_string(alias_dbc) + " dBc"); + + // The linear helper is the hazard this replaced; assert the gap is real so + // the test fails loudly if the two are ever wired together again. + const auto linear = engine::audio::resample_mono_linear(tone, 48000, 16000); + const double linear_dbc = to_db(bin_amplitude(linear, 2000, 12000, 4000.0, 16000.0)); + engine::test::require( + linear_dbc - alias_dbc > 50.0, + "resample_mono_soxr_or_sinc is no better than resample_mono_linear here"); +} + +// F4.2. The fallback used to ignore output_length_policy entirely, so an +// output could be one sample shorter or longer purely on whether libsoxr +// happened to be installed. mixing.cpp throws on exactly that mismatch. +void test_fallback_honours_output_length_policy() { + const auto source = make_music_like(11025, 44100.0, 16000.0); + engine::audio::SoxrResampleOptions options; + options.output_length_policy = engine::audio::SoxrOutputLengthPolicy::ExactExpected; + options.warning_context = "resample quality test"; + const auto resampled = engine::audio::resample_mono_soxr_or_sinc(source, 44100, 48000, options); + const size_t expected = static_cast( + std::ceil(static_cast(source.size()) * 48000.0 / 44100.0)); + engine::test::require_eq(resampled.size(), expected, "ExactExpected output length"); +} + +// F4.6. The tiers, with the numbers that justify keeping the framework default +// at 6 while routing playback through 64. +void test_resampler_tier_round_trip_snr() { + // 1 s is enough for a stable measurement and keeps the width-64 pass in + // the low milliseconds even in a debug build. + const size_t count = 44100; + const auto source = make_music_like(count, 44100.0, 16000.0); + constexpr size_t kSkip = 4096; + + const auto linear_up = engine::audio::resample_mono_linear(source, 44100, 48000); + const auto linear_back = engine::audio::resample_mono_linear(linear_up, 48000, 44100); + const double linear_snr = round_trip_snr_db(source, linear_back, kSkip); + + const auto narrow_up = + engine::audio::resample_mono_torchaudio_sinc_hann(source, 44100, 48000, sinc_options(6)); + const auto narrow_back = + engine::audio::resample_mono_torchaudio_sinc_hann(narrow_up, 48000, 44100, sinc_options(6)); + const double narrow_snr = round_trip_snr_db(source, narrow_back, kSkip); + + const auto playback = engine::audio::torchaudio_sinc_hann_playback_options(); + const auto wide_up = + engine::audio::resample_mono_torchaudio_sinc_hann(source, 44100, 48000, playback); + const auto wide_back = + engine::audio::resample_mono_torchaudio_sinc_hann(wide_up, 48000, 44100, playback); + const double wide_snr = round_trip_snr_db(source, wide_back, kSkip); + + // Measured on this signal: linear 46.3 dB, width 6 60.5 dB, width 64 + // 122.9 dB. The 53 dB divider sits about 7 dB clear of the two tiers it + // separates, and the 110 dB floor is 13 dB under what width 64 delivers. + engine::test::require( + linear_snr < 53.0, + "resample_mono_linear round trip measured " + std::to_string(linear_snr) + + " dB, which is unexpectedly good -- the tier ordering assumed here may no longer hold"); + engine::test::require( + narrow_snr > 53.0, + "sinc width 6 round trip measured only " + std::to_string(narrow_snr) + " dB"); + engine::test::require( + wide_snr > 110.0, + "playback-width sinc round trip measured only " + std::to_string(wide_snr) + " dB"); + engine::test::require( + wide_snr - narrow_snr > 45.0, + "playback width bought only " + std::to_string(wide_snr - narrow_snr) + + " dB over the framework default"); + engine::test::require( + narrow_snr - linear_snr > 8.0, + "sinc width 6 is not measurably better than linear interpolation"); +} + +// The framework default is deliberately left at torchaudio's own value of 6 so +// the ~36 feature-extraction call sites keep bit-parity with their Python +// references. Pin both so a future change is a conscious one. +void test_resampler_option_defaults() { + const engine::audio::TorchaudioSincHannResampleOptions defaults; + engine::test::require_eq(defaults.lowpass_filter_width, 6, "framework default filter width"); + engine::test::require_eq( + engine::audio::torchaudio_sinc_hann_playback_options().lowpass_filter_width, + 64, + "playback filter width"); + engine::test::require_eq( + engine::audio::torchaudio_sinc_hann_float32_options().lowpass_filter_width, + 6, + "float32 parity options keep the framework default width"); +} + +} // namespace + +int main() { + try { + test_utility_path_rejects_decimation_alias(); + test_soxr_or_sinc_fallback_is_anti_aliased(); + test_fallback_honours_output_length_policy(); + test_resampler_tier_round_trip_snr(); + test_resampler_option_defaults(); + std::cout << "audio_resample_quality_test passed\n"; + } catch (const std::exception & ex) { + std::cerr << "audio_resample_quality_test failed: " << ex.what() << "\n"; + return 1; + } + return 0; +} From 79e1c34ed3d2ca8fa4add8369ff8637deae3699e Mon Sep 17 00:00:00 2001 From: Warren B Date: Sun, 30 Aug 2026 22:58:49 +0100 Subject: [PATCH 3/3] backends: stop stranding non-CUDA GPUs on CPU paths Four capabilities are gated on CUDA specifically rather than on "a GPU backend". Each was checked against git history and against Metal op coverage to decide whether the gate is a real limitation or untested caution. The answers differ, so the fixes differ. F5-TTS -- stale caution, widened. The model arrived in one commit (b50742f) whose message mentions CUDA more than thirty times and Metal never, and docs/community_models/f5_tts.md contains no occurrence of "metal", "cpu" or "backend". Every use of is_cuda in runtime.cpp and synthesize.cpp was really a host-vs-device distinction: ggml_init no_alloc, tensor_set/get versus memcpy, compute_backend_graph versus ggml_graph_compute_with_ctx, gallocr. Nothing CUDA-specific. Every op both the DiT and the Vocos vocoder graphs emit -- SQRT, DIV, MUL, ADD, REPEAT, MEAN, SUM_ROWS, NORM, SOFT_MAX, ROPE, IM2COL, PAD, ARANGE, FLASH_ATTN_EXT, GLU, UNARY GELU/GELU_ERF/SOFTPLUS -- was checked against ggml_metal_device_supports_op and is present, and both graphs already call validate_backend_graph_supported before allocation, so an unsupported op fails by name rather than silently. F5ComputeDevice now carries a backend type; use_cuda is kept as a legacy alias because the tests/f5_tts harnesses set it. F5_TTS_HOST_VOCODER=1 restores the old routing. RedAE and MiDashengLM-Gen -- neither, just an oversight, and it was penalising CUDA rather than Metal. redae_codec_runtime.cpp:631 and midashenglm_gen/audio_tokenizer.cpp:309 built the host iSTFT unconditionally, without checking the backend at all, so they ran the scalar tail even where the CUDA kernel existed. fsq_audio_codec_runtime.cpp:659 already checked. Both now ask the same selector. ENGINE_HAS_CUDA_ISTFT -- real, and left alone. istft_graph.cpp holds exactly two implementations: a scalar host tail and one wrapping the cuFFT runtime. There is no ggml-op path for another backend to pick up, and ggml has no inverse-FFT op at all, so widening this gate would only make CudaLogMagnitudePhaseISTFT's constructor throw. A Metal iSTFT needs its own kernel set -- spectrum kernel, batched complex-to-real inverse FFT (Metal Performance Shaders has no rFFT, so realistically a Stockham radix-2/4 kernel with a Bluestein fallback), gather-form overlap-add, and a divide-by-envelope kernel -- plus its own numerical parity harness. Out of scope here; the selector and the gather below are the groundwork. MM3_DEPTH_GPU_SAMPLE / MM3_DEPTH_GPU_FRAME -- real, left alone. They gate kernels that take raw device pointers and reproduce torch's Philox RNG bit-exactly using CUDA TensorIterator grid geometry probed from the driver. Reaching them from Metal means writing a numerics-critical sampler, not changing a gate. The host iSTFT overlap-add was also rewritten from a scatter into a gather so it can be parallelised, and the fold and normalise loops carry OpenMP pragmas. The gather visits each output sample's frames in increasing frame order -- the same order the scatter reached them -- so float accumulation order is unchanged, and the test asserts a maximum difference of exactly 0.0 against an independently built scatter reference rather than a tolerance. The envelope-underflow check moved into workspace setup so the normalise loop neither branches nor throws. F5-TTS also had its own unguarded two-tap interpolator for reference resampling; it now calls the framework resampler and inherits its anti-aliasing. F5_TTS_LEGACY_RESAMPLE=1 restores the old behaviour. Tests: tests/unittests/test_backend_path_selection.cpp -- selection logic for Metal/Vulkan/CPU/HIP/CUDA in every build configuration, static asserts on the F5 predicates, and the numeric iSTFT gather-versus-scatter comparison over 500 frames. Build: cmake -S . -B build -DENGINE_BUILD_TESTS=ON && cmake --build build Test: ctest -R backend_path_selection_test Backend tested: CPU for the selection logic and the iSTFT numerics; full suite 41/41 on macOS/Metal. Known limitation: F5-TTS has never been executed on Metal in this codebase -- no F5 package was installed on the machine this was developed on -- so the widened gate is argued from op coverage rather than measured. validate_backend_graph_supported will name any missing op rather than failing silently. No speedup is claimed here; RedAE and MiDashengLM-Gen change only on CUDA, and Metal behaviour for those two is unchanged. Depends on the resampling change in the preceding commit for resample_mono_soxr_or_sinc. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01ATa5YkLUPMDPRL7w1gCo9p --- CMakeLists.txt | 7 + .../engine/community_models/f5_tts/runtime.h | 37 ++- .../engine/community_models/f5_tts/session.h | 6 +- .../community_models/f5_tts/synthesize.h | 6 +- include/engine/framework/audio/istft_graph.h | 31 ++ src/community_models/f5_tts/runtime.cpp | 90 ++++-- src/community_models/f5_tts/session.cpp | 7 +- src/community_models/f5_tts/synthesize.cpp | 101 ++++-- src/framework/audio/istft_graph.cpp | 91 +++++- src/framework/codecs/redae_codec_runtime.cpp | 24 ++ .../midashenglm_gen/audio_tokenizer.cpp | 33 +- .../unittests/test_backend_path_selection.cpp | 305 ++++++++++++++++++ 12 files changed, 648 insertions(+), 90 deletions(-) create mode 100644 tests/unittests/test_backend_path_selection.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index a67ca60fb..21ba92f5f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2032,6 +2032,13 @@ if (ENGINE_BUILD_TESTS) COMMAND audio_chunking_test ) + add_engine_unittest(backend_path_selection_test tests/unittests/test_backend_path_selection.cpp) + + add_test( + NAME backend_path_selection_test + COMMAND backend_path_selection_test + ) + add_engine_unittest(audio_resample_quality_test tests/unittests/test_audio_resample_quality.cpp) add_test( diff --git a/include/engine/community_models/f5_tts/runtime.h b/include/engine/community_models/f5_tts/runtime.h index 333ddd6ae..cd5513ff8 100644 --- a/include/engine/community_models/f5_tts/runtime.h +++ b/include/engine/community_models/f5_tts/runtime.h @@ -1,6 +1,7 @@ #pragma once #include "engine/framework/assets/resource_bundle.h" +#include "engine/framework/core/module.h" #include "engine/framework/runtime/session.h" #include @@ -37,10 +38,27 @@ struct F5SampleOptions { uint32_t seed = 0; }; -// Compute device for the DiT forward: CUDA device index or CPU threads. +// Compute device for the DiT forward and the Vocos vocoder: a ggml backend +// plus its device index, or CPU threads. struct F5ComputeDevice { + // The backend both graphs run on. Anything other than Cpu takes the + // device path (no_alloc context, gallocr arena, ggml_backend_tensor_set / + // _get, compute_backend_graph); Cpu takes the inline-context host path + // through ggml_graph_compute_with_ctx. + // + // This used to be `use_cuda` alone, so Metal and Vulkan fell to the host + // path. Nothing in the DiT or the vocoder is CUDA-specific: both graphs + // are composed from framework modules, every op they use is in the Metal + // backend's supports_op table, and validate_backend_graph_supported is + // called on both before allocation, so an unsupported op fails loudly + // rather than silently. The narrow gate was untested caution — the port + // was written against CUDA and no other GPU was tried — not a Metal bug. + engine::core::BackendType backend = engine::core::BackendType::Cpu; + // Legacy alias, still honoured for callers written before `backend` + // existed (the parity harnesses under tests/f5_tts). Only consulted when + // `backend` is left at Cpu. bool use_cuda = false; - int device = 0; // CUDA device index + int device = 0; // GPU device index int threads = 0; // CPU threads; 0 = hardware concurrency // FP16 linear weights: GEMMs get ~3x faster on tensor cores but each // mul_mat converts the F32 activations to F16 first; at F5's GEMM sizes @@ -49,6 +67,21 @@ struct F5ComputeDevice { bool fp16_weights = false; }; +// The backend a F5ComputeDevice actually asks for, folding the legacy +// `use_cuda` flag into the `backend` field. Pure function, no device needed: +// covered by tests/unittests/test_backend_path_selection.cpp. +constexpr engine::core::BackendType f5_requested_backend(const F5ComputeDevice & device) noexcept { + if (device.backend != engine::core::BackendType::Cpu) { + return device.backend; + } + return device.use_cuda ? engine::core::BackendType::Cuda : engine::core::BackendType::Cpu; +} + +// True when the DiT and the vocoder should take the device (GPU) path. +constexpr bool f5_uses_device_backend(const F5ComputeDevice & device) noexcept { + return f5_requested_backend(device) != engine::core::BackendType::Cpu; +} + // Debug taps for parity testing: when non-null, intermediate stage outputs are // appended (column layout, [features, T] flattened feature-major). struct F5DebugTaps { diff --git a/include/engine/community_models/f5_tts/session.h b/include/engine/community_models/f5_tts/session.h index 6cbfe642c..46c4ba90f 100644 --- a/include/engine/community_models/f5_tts/session.h +++ b/include/engine/community_models/f5_tts/session.h @@ -1,6 +1,7 @@ #pragma once #include "engine/framework/assets/resource_bundle.h" +#include "engine/framework/core/module.h" #include "engine/framework/runtime/model.h" #include "engine/framework/runtime/session.h" #include "engine/framework/runtime/spec_backed_model.h" @@ -41,7 +42,10 @@ class F5TTSSession final : public runtime::IOfflineVoiceTaskSession { std::string vocos_path_; std::string dialect_ = "UNK"; int frame_budget_ = 0; // 0 = default 2048 - bool use_cuda_ = false; + // The session's backend, passed straight through to the DiT and the + // vocoder. Previously only BackendType::Cuda reached them, as use_cuda_, + // so a Metal session silently ran both on the CPU. + engine::core::BackendType backend_ = engine::core::BackendType::Cpu; int cuda_device_ = 0; int threads_ = 0; }; diff --git a/include/engine/community_models/f5_tts/synthesize.h b/include/engine/community_models/f5_tts/synthesize.h index bfcf3f416..51217b8a9 100644 --- a/include/engine/community_models/f5_tts/synthesize.h +++ b/include/engine/community_models/f5_tts/synthesize.h @@ -25,8 +25,12 @@ struct F5SynthesisRequest { // on undiacritized ASR transcripts; diacritized input garbles). bool strip_diacritics = true; int threads = 0; // 0 = hardware concurrency + // Backend the DiT and the vocoder run on. Cpu keeps the host path; any + // GPU backend takes the device path. `use_cuda` is the legacy spelling and + // is only consulted when this is left at Cpu. + engine::core::BackendType backend = engine::core::BackendType::Cpu; bool use_cuda = false; - int cuda_device = 0; + int cuda_device = 0; // GPU device index }; struct F5SynthesisResult { diff --git a/include/engine/framework/audio/istft_graph.h b/include/engine/framework/audio/istft_graph.h index 6330ea212..23a4f1d38 100644 --- a/include/engine/framework/audio/istft_graph.h +++ b/include/engine/framework/audio/istft_graph.h @@ -1,5 +1,7 @@ #pragma once +#include "engine/framework/core/module.h" + #include #include #include @@ -7,6 +9,32 @@ namespace engine::audio { +// Which log-magnitude/phase iSTFT implementation a call site should build. +enum class LogMagnitudePhaseISTFTPath { + Host, + Cuda, +}; + +// True when this build carries the CUDA iSTFT runtime (ENGINE_HAS_CUDA_ISTFT). +// Resolved out of line, in engine_core, because the macro is PRIVATE to that +// target: a model translation unit cannot test it with #ifdef and would always +// read false. Every caller must ask through this function. +bool cuda_log_magnitude_phase_istft_available() noexcept; + +// Chooses the iSTFT implementation for a backend. This is the only place that +// decides, so the four call sites cannot drift apart again. +// +// Metal deliberately maps to Host: there is no Metal iSTFT kernel, and the +// ggml graph builders in this file express only the CUDA and host paths. This +// is a missing implementation, not a Metal defect — see the note in +// istft_graph.cpp. HIP maps to Host too, because istft_cuda_runtime.cu is +// compiled only under `ENGINE_ENABLE_CUDA AND NOT ENGINE_ENABLE_HIP`. +// +// Setting ENGINE_ISTFT_HOST=1 forces the host path on every backend, which +// keeps the pre-existing behaviour of any call site this selector newly moves +// onto the CUDA kernel. +LogMagnitudePhaseISTFTPath select_log_magnitude_phase_istft_path(core::BackendType backend) noexcept; + struct HostLogMagnitudePhaseISTFTConfig { int64_t frames = 0; int64_t n_fft = 0; @@ -27,6 +55,9 @@ struct HostLogMagnitudePhaseISTFTTiming { double spectrum_ms = 0.0; double framed_clear_ms = 0.0; double fft_inverse_ms = 0.0; + // Always 0: the overlap-add is a gather that writes every output element, + // so there is no zero-fill pass left to time. Kept so the timing schema + // and its consumers do not change. double fold_clear_ms = 0.0; double overlap_add_ms = 0.0; double normalize_ms = 0.0; diff --git a/src/community_models/f5_tts/runtime.cpp b/src/community_models/f5_tts/runtime.cpp index ee844c9d6..8caf52de3 100644 --- a/src/community_models/f5_tts/runtime.cpp +++ b/src/community_models/f5_tts/runtime.cpp @@ -24,6 +24,7 @@ #include #include #include +#include #include #include @@ -228,6 +229,16 @@ class StrippedView final : public engine::assets::TensorSource { const LoadedModel & load_model_once(const std::string & path, const F5ComputeDevice & dev); +// Cache key for the per-(path, device) model and weight caches. Two devices of +// the same backend get separate entries; CPU has no index. +std::string device_cache_key(const F5ComputeDevice & dev) { + const core::BackendType backend = f5_requested_backend(dev); + if (backend == core::BackendType::Cpu) { + return "cpu"; + } + return std::to_string(static_cast(backend)) + "." + std::to_string(dev.device); +} + // Open the DiT checkpoint as a stripped-name tensor source. Safetensors // checkpoints carry raw torch names ("ema_model.transformer.*"); GGUF // packages store the same tensors under the "transformer" namespace, so the @@ -248,8 +259,7 @@ const F5DiTWeights & load_dit_weights_once( const std::string & path, const F5ComputeDevice & dev) { struct Entry { F5DiTWeights w; }; static auto * cache = new std::unordered_map(); - const std::string key = - (dev.use_cuda ? "cuda" + std::to_string(dev.device) : "cpu") + ":" + path; + const std::string key = device_cache_key(dev) + ":" + path; const auto found = cache->find(key); if (found != cache->end()) { return found->second.w; @@ -269,25 +279,33 @@ const LoadedModel & load_model_once(const std::string & path, const F5ComputeDev // The backend (and its weight buffer) must outlive the cache entry, so it // is owned by a static owner freed after the cache at exit. static std::vector> owners; - const std::string key = (dev.use_cuda ? "cuda" + std::to_string(dev.device) : "cpu") + ":" + path; + const std::string key = device_cache_key(dev) + ":" + path; if (const auto found = cache->find(key); found != cache->end()) { return found->second; } auto stripped = open_dit_source(path); auto owner = std::make_unique(); - const core::BackendType type = dev.use_cuda ? core::BackendType::Cuda : core::BackendType::Cpu; - core::BackendConfig cfg{type, dev.use_cuda ? dev.device : 0, dev.use_cuda ? 1 : std::max(1, dev.threads)}; + const core::BackendType requested = f5_requested_backend(dev); + const bool device_backend = requested != core::BackendType::Cpu; + core::BackendConfig cfg{ + requested, + device_backend ? dev.device : 0, + device_backend ? 1 : std::max(1, dev.threads)}; owner->value = core::init_backend(cfg); - if (!dev.use_cuda) { + // BestAvailable resolves only once the backend exists, so read the type + // back off the handle rather than trusting the request. + const core::BackendType type = core::backend_type(owner->value); + if (type == core::BackendType::Cpu) { core::set_backend_threads(owner->value, std::max(1, dev.threads)); } LoadedModel model; model.arch = F5Architecture{}; model.backend = owner->value; model.backend_type = type; - // FP16 linear weights on CUDA only (CPU mul_mat with F16 weights is slow - // via the fallback path; CUDA hits tensor cores). - model.w = load_weights(*stripped, owner->value, type, dev.use_cuda && dev.fp16_weights); + // FP16 linear weights on GPU backends only (CPU mul_mat with F16 weights + // is slow via the fallback path). Off by default everywhere. + model.w = load_weights( + *stripped, owner->value, type, type != core::BackendType::Cpu && dev.fp16_weights); owners.push_back(std::move(owner)); return cache->emplace(key, std::move(model)).first->second; } @@ -315,7 +333,11 @@ std::pair, std::vector> f5_dit_forward_cfg( const int N = seq_len; const int MEL = arch.mel_dim; const int NT = static_cast(text_in.size()); - const bool is_cuda = model.backend_type == core::BackendType::Cuda; + // Any non-CPU backend takes the device path: no_alloc context, gallocr + // arena, ggml_backend_tensor_set/_get and compute_backend_graph. Only + // the CPU backend can use the inline-context ggml_graph_compute_with_ctx + // path, so this is a host-vs-device split, not a CUDA-vs-everything one. + const bool is_device = model.backend_type != core::BackendType::Cpu; struct CfgGraph { ggml_context * ctx = nullptr; @@ -337,7 +359,7 @@ std::pair, std::vector> f5_dit_forward_cfg( const size_t ctx_bytes = std::min( std::max(1536ULL << 20, static_cast(N) * (8ULL << 20)), 12288ULL << 20); - gnew->ctx = ggml_init({ctx_bytes, nullptr, is_cuda}); + gnew->ctx = ggml_init({ctx_bytes, nullptr, is_device}); ggml_context * ctx = gnew->ctx; // ---- module-composed batched-CFG graph (B=2) ---- // Leaves: x/cond [2, N, 100] (same values in both halves), ids @@ -370,21 +392,21 @@ std::pair, std::vector> f5_dit_forward_cfg( gnew->graph = ggml_new_graph_custom(ctx, 262144, false); ggml_build_forward_expand(gnew->graph, output); core::validate_backend_graph_supported(model.backend, gnew->graph, "f5_dit_cfg"); - if (!is_cuda) { + if (!is_device) { const int threads = dev.threads > 0 ? dev.threads : static_cast(std::thread::hardware_concurrency()); core::set_backend_threads(model.backend, threads); } - if (is_cuda) { + if (is_device) { // gallocr-only flow; constants first get PRIVATE buffers so the // arena never aliases them (root cause of the noise regression) const_stage_bind(cfg_staged, model.backend); gnew->gallocr = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); if (gnew->gallocr == nullptr || !ggml_gallocr_reserve(gnew->gallocr, gnew->graph) || !ggml_gallocr_alloc_graph(gnew->gallocr, gnew->graph)) { - throw std::runtime_error("F5 DiT CUDA graph alloc failed"); + throw std::runtime_error("F5 DiT device graph alloc failed"); } if (cfg_staged != nullptr) { - const_stage_upload(cfg_staged, is_cuda ? model.backend : nullptr); + const_stage_upload(cfg_staged, is_device ? model.backend : nullptr); const_stage_end(cfg_staged); } (void)0; @@ -432,7 +454,7 @@ std::pair, std::vector> f5_dit_forward_cfg( th[128 + i] = std::cos(f); } } - if (is_cuda) { + if (is_device) { ggml_backend_tensor_set(g.x, xb.data(), 0, xb.size() * sizeof(float)); ggml_backend_tensor_set(g.cond, cb.data(), 0, cb.size() * sizeof(float)); ggml_backend_tensor_set(g.text_ids, ids.data(), 0, ids.size() * sizeof(int32_t)); @@ -449,11 +471,11 @@ std::pair, std::vector> f5_dit_forward_cfg( - const auto status = is_cuda + const auto status = is_device ? core::compute_backend_graph(model.backend, g.graph, nullptr, "f5_dit_cfg") : f5_cpu_graph_compute(g.ctx, g.graph, dev.threads > 0 ? dev.threads : static_cast(std::thread::hardware_concurrency())); - if (is_cuda) ggml_backend_synchronize(model.backend); + if (is_device) ggml_backend_synchronize(model.backend); if (status != GGML_STATUS_SUCCESS) { throw std::runtime_error("F5 DiT CFG graph compute failed"); } @@ -462,7 +484,7 @@ std::pair, std::vector> f5_dit_forward_cfg( const size_t half_floats = ggml_nelements(g.output) / 2; out.first.resize(half_floats); out.second.resize(half_floats); - if (is_cuda) { + if (is_device) { ggml_backend_tensor_get(g.output, out.first.data(), 0, half_floats * sizeof(float)); ggml_backend_tensor_get(g.output, out.second.data(), half_floats * sizeof(float), half_floats * sizeof(float)); } else { @@ -497,7 +519,11 @@ std::vector f5_dit_forward( const int N = seq_len; const int MEL = arch.mel_dim; const int NT = static_cast(text_in.size()); - const bool is_cuda = model.backend_type == core::BackendType::Cuda; + // Any non-CPU backend takes the device path: no_alloc context, gallocr + // arena, ggml_backend_tensor_set/_get and compute_backend_graph. Only + // the CPU backend can use the inline-context ggml_graph_compute_with_ctx + // path, so this is a host-vs-device split, not a CUDA-vs-everything one. + const bool is_device = model.backend_type != core::BackendType::Cpu; // ---- cached graph per (model, N, NT, with/without taps) ---- // Taps change the graph (extra roots), so key on their presence. The @@ -542,7 +568,7 @@ std::vector f5_dit_forward( const size_t ctx_bytes = std::min( std::max(1536ULL << 20, static_cast(N) * (6ULL << 20)), 6144ULL << 20); - gnew->ctx = ggml_init({ctx_bytes, nullptr, is_cuda}); + gnew->ctx = ggml_init({ctx_bytes, nullptr, is_device}); ggml_context * ctx = gnew->ctx; // On CUDA the ctx is no_alloc: leaf tensors get device storage after // ggml_backend_alloc_ctx_tensors, values uploaded from staging vectors. @@ -590,21 +616,21 @@ std::vector f5_dit_forward( } } core::validate_backend_graph_supported(model.backend, gnew->graph, "f5_dit"); - if (!is_cuda) { + if (!is_device) { const int threads = dev.threads > 0 ? dev.threads : static_cast(std::thread::hardware_concurrency()); core::set_backend_threads(model.backend, threads); } - if (is_cuda) { + if (is_device) { const_stage_bind(staged_module_consts, model.backend); gnew->gallocr = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); if (gnew->gallocr == nullptr || !ggml_gallocr_reserve(gnew->gallocr, gnew->graph) || !ggml_gallocr_alloc_graph(gnew->gallocr, gnew->graph)) { - throw std::runtime_error("F5 DiT CUDA graph alloc failed"); + throw std::runtime_error("F5 DiT device graph alloc failed"); } if (staged_module_consts != nullptr) { - const_stage_upload(staged_module_consts, is_cuda ? model.backend : nullptr); + const_stage_upload(staged_module_consts, is_device ? model.backend : nullptr); const_stage_end(staged_module_consts); staged_module_consts = nullptr; } @@ -631,7 +657,7 @@ std::vector f5_dit_forward( cond_col.assign(cond_in.size(), 0.0F); cond_src = cond_col.data(); } - if (is_cuda) { + if (is_device) { ggml_backend_tensor_set(g.x, x_in.data(), 0, x_in.size() * sizeof(float)); ggml_backend_tensor_set(g.cond, cond_src, 0, cond_in.size() * sizeof(float)); } else { @@ -642,7 +668,7 @@ std::vector f5_dit_forward( for (int i = 0; i < NT; ++i) { ids[i] = drop_text ? 0 : (text_in[i] + 1); } - if (is_cuda) { + if (is_device) { ggml_backend_tensor_set(g.text_ids, ids.data(), 0, ids.size() * sizeof(int32_t)); } else { std::memcpy(g.text_ids->data, ids.data(), ids.size() * sizeof(int32_t)); @@ -656,7 +682,7 @@ std::vector f5_dit_forward( th[128 + i] = std::cos(f); } } - if (is_cuda) { + if (is_device) { ggml_backend_tensor_set(g.th_t, th.data(), 0, th.size() * sizeof(float)); } else { std::memcpy(g.th_t->data, th.data(), th.size() * sizeof(float)); @@ -665,19 +691,19 @@ std::vector f5_dit_forward( // ---- compute ---- std::vector out; - const auto status = is_cuda + const auto status = is_device ? core::compute_backend_graph(model.backend, g.graph, nullptr, "f5_dit") : f5_cpu_graph_compute(g.ctx, g.graph, dev.threads > 0 ? dev.threads : static_cast(std::thread::hardware_concurrency())); - if (is_cuda) { + if (is_device) { ggml_backend_synchronize(model.backend); } if (status != GGML_STATUS_SUCCESS) { throw std::runtime_error("F5 DiT graph compute failed"); } out.resize(ggml_nelements(g.output)); - if (is_cuda) { + if (is_device) { ggml_backend_tensor_get(g.output, out.data(), 0, out.size() * sizeof(float)); } else { std::memcpy(out.data(), ggml_get_data(g.output), out.size() * sizeof(float)); @@ -686,7 +712,7 @@ std::vector f5_dit_forward( const auto read_tap = [&](ggml_tensor * t, std::vector * dst) { if (t != nullptr && dst != nullptr) { dst->resize(ggml_nelements(t)); - if (is_cuda) { + if (is_device) { ggml_backend_tensor_get(t, dst->data(), 0, dst->size() * sizeof(float)); } else { std::memcpy(dst->data(), ggml_get_data(t), dst->size() * sizeof(float)); diff --git a/src/community_models/f5_tts/session.cpp b/src/community_models/f5_tts/session.cpp index c847aec1f..bd7c21b9b 100644 --- a/src/community_models/f5_tts/session.cpp +++ b/src/community_models/f5_tts/session.cpp @@ -136,7 +136,10 @@ F5TTSSession::F5TTSSession( "f5_tts.frame_budget must be within [256, 8192] mel frames"); } } - use_cuda_ = options.backend.type == core::BackendType::Cuda; + // Pass the session's backend through unchanged. This used to collapse to + // `type == Cuda`, which routed Metal and Vulkan sessions onto the CPU DiT + // and the scalar host vocoder. + backend_ = options.backend.type; cuda_device_ = options.backend.device; threads_ = options.backend.threads; } @@ -202,7 +205,7 @@ runtime::TaskResult F5TTSSession::run(const runtime::TaskRequest & request) { if (const auto v = runtime::find_option(request.options, {"strip_diacritics"})) { req.strip_diacritics = runtime::parse_bool_option(*v, "strip_diacritics"); } - req.use_cuda = use_cuda_; + req.backend = backend_; req.frame_budget = frame_budget_; req.cuda_device = cuda_device_; req.threads = threads_; diff --git a/src/community_models/f5_tts/synthesize.cpp b/src/community_models/f5_tts/synthesize.cpp index 8b1e981bd..023d0bd43 100644 --- a/src/community_models/f5_tts/synthesize.cpp +++ b/src/community_models/f5_tts/synthesize.cpp @@ -6,6 +6,7 @@ #include "cpu_graph_compute.h" +#include "engine/framework/audio/resampling.h" #include "engine/framework/core/backend.h" #include "ggml.h" @@ -24,7 +25,9 @@ #include #include #include +#include #include +#include namespace engine::models::f5_tts { namespace { @@ -153,7 +156,11 @@ std::vector compute_mel(const std::vector & wav) { return mel; // [mel, frames] feature-fastest memory } -std::vector resample(const std::vector & in, int sr_in, int sr_out) { +// Legacy two-tap linear interpolator. Kept only so the pre-existing +// conditioning numerics stay reachable via F5_TTS_LEGACY_RESAMPLE=1; it is a +// -13.8 dBc-class filter and everything it aliases is baked into the voice the +// DiT clones. +std::vector resample_linear_legacy(const std::vector & in, int sr_in, int sr_out) { if (sr_in == sr_out || in.empty()) return in; const double ratio = static_cast(sr_out) / sr_in; const size_t out_n = static_cast(static_cast(in.size()) * ratio); @@ -168,6 +175,23 @@ std::vector resample(const std::vector & in, int sr_in, int sr_out return out; } +// Reference audio arrives at an arbitrary sample rate and is decimated to +// 24 kHz before the mel frontend, so every alias image this stage folds down +// ends up in the conditioning mel and therefore in the cloned timbre. Use the +// framework resampler (libsoxr when present, the in-tree windowed-sinc +// otherwise) instead of the local interpolator this file used to carry. +std::vector resample(const std::vector & in, int sr_in, int sr_out) { + if (sr_in == sr_out || in.empty()) return in; + const char * legacy = std::getenv("F5_TTS_LEGACY_RESAMPLE"); + if (legacy != nullptr && legacy[0] == '1') { + return resample_linear_legacy(in, sr_in, sr_out); + } + engine::audio::SoxrResampleOptions options; + options.warning_context = "f5_tts"; + options.fallback_description = "F5-TTS reference resampling"; + return engine::audio::resample_mono_soxr_or_sinc(in, sr_in, sr_out, options); +} + std::unordered_map load_vocab(const std::string & dir) { std::unordered_map map; std::ifstream f(dir + "/vocab.txt", std::ios::binary); @@ -677,38 +701,41 @@ std::vector vocos_decode_gpu( struct BackendOwnerV { ggml_backend_t value = nullptr; }; - static auto * owners = new std::map(); - const int key = dev.use_cuda ? dev.device : -1; - ggml_backend_t backend = nullptr; - if (dev.use_cuda) { - auto ob = owners->find(key); - if (ob == owners->end()) { - core::BackendConfig cfg{core::BackendType::Cuda, dev.device, 1}; - ob = owners->emplace(key, BackendOwnerV{core::init_backend(cfg)}).first; - } - backend = ob->second.value; - } else { - auto ob = owners->find(key); - if (ob == owners->end()) { - core::BackendConfig cfg{core::BackendType::Cpu, 0, std::max(1, dev.threads)}; - ob = owners->emplace(key, BackendOwnerV{core::init_backend(cfg)}).first; - } - backend = ob->second.value; + // Keyed on (backend, device); CPU is a single entry. The key used to be + // the CUDA device index or -1, which could not tell two GPU backends + // apart. + static auto * owners = new std::map, BackendOwnerV>(); + const core::BackendType requested = f5_requested_backend(dev); + const bool want_device = requested != core::BackendType::Cpu; + const auto key = std::make_pair( + static_cast(requested), want_device ? dev.device : -1); + auto ob = owners->find(key); + if (ob == owners->end()) { + core::BackendConfig cfg{ + requested, + want_device ? dev.device : 0, + want_device ? 1 : std::max(1, dev.threads)}; + ob = owners->emplace(key, BackendOwnerV{core::init_backend(cfg)}).first; } - const bool is_cuda = dev.use_cuda; + ggml_backend_t backend = ob->second.value; + // Any non-CPU backend takes the device path. The vocoder graph is plain + // im2col/mul_mat/gelu/norm — nothing CUDA-specific — and + // validate_backend_graph_supported below fails loudly if a backend is + // missing an op. + const bool is_device = core::backend_type(backend) != core::BackendType::Cpu; // graph cache per (T, device) - static auto * cache = new std::map, std::unique_ptr>(); - const auto ckey = std::make_pair(T, key); + static auto * cache = new std::map, std::unique_ptr>(); + const auto ckey = std::make_tuple(T, key.first, key.second); auto it = cache->find(ckey); if (it == cache->end()) { auto g = std::make_unique(); const size_t ctx_bytes = 256ULL << 20; - g->ctx = ggml_init({ctx_bytes, nullptr, is_cuda}); + g->ctx = ggml_init({ctx_bytes, nullptr, is_device}); ggml_context * ctx = g->ctx; std::vector>> pending; const auto leaf_write = [&](ggml_tensor * t, const void * src, size_t bytes) { - if (!is_cuda) { + if (!is_device) { std::memcpy(t->data, src, bytes); } else { const auto * b = static_cast(src); @@ -716,7 +743,7 @@ std::vector vocos_decode_gpu( } }; const auto leaf_zero = [&](ggml_tensor * t, size_t bytes) { - if (!is_cuda) { + if (!is_device) { std::memset(t->data, 0, bytes); } else { pending.emplace_back(t, std::vector(bytes, 0)); @@ -793,7 +820,7 @@ std::vector vocos_decode_gpu( g->graph = ggml_new_graph_custom(ctx, 8192, false); ggml_build_forward_expand(g->graph, spec); core::validate_backend_graph_supported(backend, g->graph, "f5_vocos"); - if (is_cuda) { + if (is_device) { g->io_buffer = ggml_backend_alloc_ctx_tensors(ctx, backend); for (auto & leaf : pending) { ggml_backend_tensor_set(leaf.first, leaf.second.data(), 0, leaf.second.size()); @@ -801,7 +828,7 @@ std::vector vocos_decode_gpu( g->gallocr = ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend)); if (g->gallocr == nullptr || !ggml_gallocr_reserve(g->gallocr, g->graph) || !ggml_gallocr_alloc_graph(g->gallocr, g->graph)) { - throw std::runtime_error("vocos CUDA graph alloc failed"); + throw std::runtime_error("vocos device graph alloc failed"); } } it = cache->emplace(ckey, std::move(g)).first; @@ -809,17 +836,17 @@ std::vector vocos_decode_gpu( VocosGraph & g = *it->second; // upload mel + compute - if (is_cuda) { + if (is_device) { ggml_backend_tensor_set(g.mel, mel_rows.data(), 0, mel_rows.size() * sizeof(float)); } else { std::memcpy(g.mel->data, mel_rows.data(), mel_rows.size() * sizeof(float)); } - const auto status = is_cuda + const auto status = is_device ? core::compute_backend_graph(backend, g.graph, nullptr, "f5_vocos") : f5_cpu_graph_compute(g.ctx, g.graph, dev.threads > 0 ? dev.threads : static_cast(std::thread::hardware_concurrency())); - if (is_cuda) { + if (is_device) { ggml_backend_synchronize(backend); } if (status != GGML_STATUS_SUCCESS) { @@ -827,7 +854,7 @@ std::vector vocos_decode_gpu( } // spec: ne [1026, T] o-fastest == row t at t*1026 — same as host layout std::vector spec(static_cast(T) * 2 * n_freqs); - if (is_cuda) { + if (is_device) { ggml_backend_tensor_get(g.spec, spec.data(), 0, spec.size() * sizeof(float)); } else { std::memcpy(spec.data(), ggml_get_data(g.spec), spec.size() * sizeof(float)); @@ -1196,6 +1223,7 @@ F5SynthesisResult f5_synthesize( ref_voiced_frames = std::max(1, ref_voiced_frames); F5ComputeDevice dev; + dev.backend = request.backend; dev.use_cuda = request.use_cuda; dev.device = request.cuda_device; dev.threads = request.threads; @@ -1250,9 +1278,16 @@ F5SynthesisResult f5_synthesize( all_rows.insert(all_rows.end(), out.gen_mel_rows.begin(), out.gen_mel_rows.end()); } - result.audio = request.use_cuda - ? vocos_decode_gpu(vocos_path, all_rows, dev) - : vocos_decode(vocos_path, all_rows); + // vocos_decode is a scalar, single-threaded triple loop over eight + // ConvNeXt blocks of 512x1536 GEMMs; vocos_decode_gpu is the same maths as + // a ggml graph. The old gate sent every non-CUDA backend to the scalar + // one. F5_TTS_HOST_VOCODER=1 restores it. + const char * force_host_vocoder = std::getenv("F5_TTS_HOST_VOCODER"); + const bool host_vocoder = (force_host_vocoder != nullptr && force_host_vocoder[0] == '1') || + !f5_uses_device_backend(dev); + result.audio = host_vocoder + ? vocos_decode(vocos_path, all_rows) + : vocos_decode_gpu(vocos_path, all_rows, dev); // undo the reference normalization on the output (Python F5 parity) if (ref_gain != 1.0F) { double out_rms = 0.0; diff --git a/src/framework/audio/istft_graph.cpp b/src/framework/audio/istft_graph.cpp index 700854291..596821d7e 100644 --- a/src/framework/audio/istft_graph.cpp +++ b/src/framework/audio/istft_graph.cpp @@ -9,6 +9,7 @@ #include #include #include +#include #include #include #include @@ -82,6 +83,15 @@ void ensure_workspace( workspace.envelope[static_cast(start + i)] += w * w; } } + // The envelope only depends on the window and the geometry, both of which + // are cached here, so validate it once on build rather than once per + // output sample. That keeps the normalise loop branch-free and, more + // importantly, keeps it free of a throw so it can be an OpenMP loop. + for (int64_t i = 0; i < samples; ++i) { + if (workspace.envelope[static_cast(i + pad)] <= 1.0e-11F) { + throw std::runtime_error("ISTFT window envelope underflow"); + } + } } template @@ -117,37 +127,86 @@ void finish_istft_from_spectrum( threads); timing.fft_inverse_ms = elapsed_ms(timing_start, Clock::now()); + // Overlap-add, expressed as a gather over output samples instead of a + // scatter over frames. Frames overlap, so the scatter form cannot be an + // OpenMP loop without a race; the gather form is embarrassingly parallel + // and does exactly the same multiply-adds. Output sample `o` collects the + // frames f with f*hop <= o < f*hop + n_fft, visited in increasing f — the + // same order in which the scatter form reached them — so the result is + // bit-identical, not merely equivalent. It also writes every element of + // `folded`, which removes the separate zero-fill pass. +#ifdef _OPENMP + const int omp_threads = static_cast(std::max(1, threads)); +#endif timing_start = Clock::now(); - std::fill(workspace.folded.begin(), workspace.folded.end(), 0.0F); - timing.fold_clear_ms = elapsed_ms(timing_start, Clock::now()); - - timing_start = Clock::now(); - for (int64_t frame = 0; frame < config.frames; ++frame) { - const int64_t start = frame * config.hop_length; - const float * src = workspace.framed.data() + static_cast(frame * config.n_fft); - for (int64_t i = 0; i < config.n_fft; ++i) { - const float w = window[static_cast(i)]; - workspace.folded[static_cast(start + i)] += src[i] * w; + const int64_t output_size = workspace.output_size; +#ifdef _OPENMP +#pragma omp parallel for num_threads(omp_threads) if (output_size >= 4096) +#endif + for (int64_t out_index = 0; out_index < output_size; ++out_index) { + const int64_t last_frame = std::min(config.frames - 1, out_index / config.hop_length); + int64_t first_frame = 0; + const int64_t first_start = out_index - config.n_fft + 1; + if (first_start > 0) { + first_frame = (first_start + config.hop_length - 1) / config.hop_length; + } + float acc = 0.0F; + for (int64_t frame = first_frame; frame <= last_frame; ++frame) { + const int64_t offset = out_index - frame * config.hop_length; + acc += workspace.framed[static_cast(frame * config.n_fft + offset)] * + window[static_cast(offset)]; } + workspace.folded[static_cast(out_index)] = acc; } timing.overlap_add_ms = elapsed_ms(timing_start, Clock::now()); const int64_t pad = (config.n_fft - config.hop_length) / 2; - const int64_t samples = workspace.output_size - 2 * pad; + const int64_t samples = output_size - 2 * pad; timing_start = Clock::now(); + // ensure_workspace already proved every divisor in [pad, pad + samples) + // is above the underflow floor, so this loop neither branches nor throws. +#ifdef _OPENMP +#pragma omp parallel for num_threads(omp_threads) if (samples >= 4096) +#endif for (int64_t i = 0; i < samples; ++i) { const int64_t src = i + pad; - const float denom = workspace.envelope[static_cast(src)]; - if (denom <= 1.0e-11F) { - throw std::runtime_error("ISTFT window envelope underflow"); - } - workspace.audio[static_cast(i)] = workspace.folded[static_cast(src)] / denom; + workspace.audio[static_cast(i)] = + workspace.folded[static_cast(src)] / + workspace.envelope[static_cast(src)]; } timing.normalize_ms = elapsed_ms(timing_start, Clock::now()); } } // namespace +bool cuda_log_magnitude_phase_istft_available() noexcept { +#ifdef ENGINE_HAS_CUDA_ISTFT + return true; +#else + return false; +#endif +} + +// There is no Metal (or Vulkan) iSTFT. The two implementations in this file +// are a scalar host tail and a cuFFT/CUDA-kernel tail; the graph builders here +// express no ggml-op variant that a generic backend could pick up, so widening +// this predicate to "any GPU" would only make CudaLogMagnitudePhaseISTFT's +// constructor throw. Writing a Metal iSTFT is a new kernel, not a gate change +// — see the design note in the lane report. +LogMagnitudePhaseISTFTPath select_log_magnitude_phase_istft_path(core::BackendType backend) noexcept { + if (!cuda_log_magnitude_phase_istft_available()) { + return LogMagnitudePhaseISTFTPath::Host; + } + if (backend != core::BackendType::Cuda) { + return LogMagnitudePhaseISTFTPath::Host; + } + const char * force_host = std::getenv("ENGINE_ISTFT_HOST"); + if (force_host != nullptr && force_host[0] == '1') { + return LogMagnitudePhaseISTFTPath::Host; + } + return LogMagnitudePhaseISTFTPath::Cuda; +} + class HostLogMagnitudePhaseISTFT::Impl { public: explicit Impl(const HostLogMagnitudePhaseISTFTConfig & config) diff --git a/src/framework/codecs/redae_codec_runtime.cpp b/src/framework/codecs/redae_codec_runtime.cpp index bb9cce972..adaeb4b96 100644 --- a/src/framework/codecs/redae_codec_runtime.cpp +++ b/src/framework/codecs/redae_codec_runtime.cpp @@ -628,6 +628,26 @@ class RedAeRuntime { runtime::AudioBuffer audio; audio.sample_rate = static_cast(config_.sample_rate); audio.channels = 1; + // This call site used to construct the host iSTFT unconditionally, + // so a CUDA build ran the whole spectrum tail on the CPU even though + // the CUDA kernel was compiled in. Route it through the shared + // selector instead; on Metal (and every non-CUDA backend) the answer + // is still Host, which is what it already did. + if (audio::select_log_magnitude_phase_istft_path(execution_.backend_type()) == + audio::LogMagnitudePhaseISTFTPath::Cuda) { + if (cuda_istft_ == nullptr || cuda_istft_frames_ != qwen_frames) { + audio::CudaLogMagnitudePhaseISTFTConfig cfg; + cfg.frames = qwen_frames; + cfg.n_fft = config_.audio_patch_size * 4; + cfg.hop_length = config_.audio_patch_size; + cfg.out_dim = config_.audio_patch_size * 4 + 2; + cfg.device = execution_.config().device; + cuda_istft_ = std::make_unique(cfg); + cuda_istft_frames_ = qwen_frames; + } + audio.samples = cuda_istft_->compute(spec, istft_window_).audio; + return audio; + } if (host_istft_ == nullptr || host_istft_frames_ != qwen_frames) { audio::HostLogMagnitudePhaseISTFTConfig cfg; cfg.frames = qwen_frames; @@ -651,6 +671,8 @@ class RedAeRuntime { decoder_qwen_.release_runtime_graphs(); host_istft_.reset(); host_istft_frames_ = 0; + cuda_istft_.reset(); + cuda_istft_frames_ = 0; } private: @@ -722,6 +744,8 @@ class RedAeRuntime { std::vector istft_window_; std::unique_ptr host_istft_; int64_t host_istft_frames_ = 0; + std::unique_ptr cuda_istft_; + int64_t cuda_istft_frames_ = 0; }; diff --git a/src/models/midashenglm_gen/audio_tokenizer.cpp b/src/models/midashenglm_gen/audio_tokenizer.cpp index 41c67f188..f82ee941d 100644 --- a/src/models/midashenglm_gen/audio_tokenizer.cpp +++ b/src/models/midashenglm_gen/audio_tokenizer.cpp @@ -306,7 +306,26 @@ class MiDashengLmGenAudioTokenizerRuntime::DecodeGraph { result.audio.reserve(static_cast(batch_)); const int64_t out_frames = 2 * frames_; const int64_t out_dim = config_.istft_n_fft + 2; - if (host_istft_ == nullptr) { + // This call site used to construct the host iSTFT unconditionally and + // never looked at the backend, so a CUDA build ran the spectrum tail + // on the CPU with the CUDA kernel sitting unused. Ask the shared + // selector; on Metal (and every non-CUDA backend) it still answers + // Host, which is exactly what this did before. + const bool use_cuda_istft = + engine::audio::select_log_magnitude_phase_istft_path(execution_.backend_type()) == + engine::audio::LogMagnitudePhaseISTFTPath::Cuda; + if (use_cuda_istft) { + if (cuda_istft_ == nullptr) { + cuda_istft_ = std::make_unique( + engine::audio::CudaLogMagnitudePhaseISTFTConfig{ + out_frames, + config_.istft_n_fft, + config_.istft_hop, + out_dim, + execution_.config().device, + }); + } + } else if (host_istft_ == nullptr) { host_istft_ = std::make_unique( engine::audio::HostLogMagnitudePhaseISTFTConfig{ out_frames, @@ -322,11 +341,18 @@ class MiDashengLmGenAudioTokenizerRuntime::DecodeGraph { const std::vector item( log_magnitude_phase.begin() + static_cast(item_offset), log_magnitude_phase.begin() + static_cast(item_offset + item_count)); - auto decoded = host_istft_->compute(item, weights_->istft_window); + std::vector decoded_audio; + if (use_cuda_istft) { + auto decoded = cuda_istft_->compute(item, weights_->istft_window); + decoded_audio = std::move(decoded.audio); + } else { + auto decoded = host_istft_->compute(item, weights_->istft_window); + decoded_audio = std::move(decoded.audio); + } result.audio.push_back(engine::runtime::AudioBuffer{ static_cast(config_.sample_rate), 1, - std::move(decoded.audio)}); + std::move(decoded_audio)}); } return result; } @@ -360,6 +386,7 @@ class MiDashengLmGenAudioTokenizerRuntime::DecodeGraph { ggml_gallocr_t gallocr_ = nullptr; ggml_backend_buffer_t input_buffer_ = nullptr; std::unique_ptr host_istft_; + std::unique_ptr cuda_istft_; }; MiDashengLmGenAudioTokenizerRuntime::MiDashengLmGenAudioTokenizerRuntime( diff --git a/tests/unittests/test_backend_path_selection.cpp b/tests/unittests/test_backend_path_selection.cpp new file mode 100644 index 000000000..95e1f2df8 --- /dev/null +++ b/tests/unittests/test_backend_path_selection.cpp @@ -0,0 +1,305 @@ +// Backend path selection for the capabilities that used to be gated on CUDA +// specifically rather than on "is this a GPU". +// +// Everything here is pure logic: no weights, no device, no inference. That is +// the point — the selection rules are what regressed, and they are testable on +// a machine with no GPU at all. + +#include "engine/community_models/f5_tts/runtime.h" +#include "engine/framework/audio/fft.h" +#include "engine/framework/audio/istft_graph.h" +#include "engine/framework/core/module.h" + +#include "test_assert.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +namespace audio = engine::audio; +namespace core = engine::core; +namespace f5 = engine::models::f5_tts; + +// --------------------------------------------------------------------------- +// iSTFT: which implementation each backend gets +// --------------------------------------------------------------------------- + +void test_istft_path_never_cuda_without_the_cuda_runtime() { + // The CUDA iSTFT lives in istft_cuda_runtime.cu, compiled only under + // ENGINE_ENABLE_CUDA AND NOT ENGINE_ENABLE_HIP. If it is absent, no + // backend may select it — the constructor would throw. + if (audio::cuda_log_magnitude_phase_istft_available()) { + return; + } + const core::BackendType all[] = { + core::BackendType::Cpu, + core::BackendType::Cuda, + core::BackendType::Hip, + core::BackendType::Vulkan, + core::BackendType::Metal, + core::BackendType::BestAvailable, + }; + for (const core::BackendType backend : all) { + engine::test::require( + audio::select_log_magnitude_phase_istft_path(backend) == + audio::LogMagnitudePhaseISTFTPath::Host, + "iSTFT must select the host path when the CUDA runtime is not built"); + } +} + +void test_istft_path_metal_and_cpu_take_the_host_path() { + // True in every build configuration: there is no Metal or Vulkan iSTFT + // kernel, so widening the gate to "any GPU" would only produce a throw. + engine::test::require( + audio::select_log_magnitude_phase_istft_path(core::BackendType::Metal) == + audio::LogMagnitudePhaseISTFTPath::Host, + "Metal must select the host iSTFT"); + engine::test::require( + audio::select_log_magnitude_phase_istft_path(core::BackendType::Vulkan) == + audio::LogMagnitudePhaseISTFTPath::Host, + "Vulkan must select the host iSTFT"); + engine::test::require( + audio::select_log_magnitude_phase_istft_path(core::BackendType::Cpu) == + audio::LogMagnitudePhaseISTFTPath::Host, + "CPU must select the host iSTFT"); + engine::test::require( + audio::select_log_magnitude_phase_istft_path(core::BackendType::Hip) == + audio::LogMagnitudePhaseISTFTPath::Host, + "HIP must select the host iSTFT: the .cu is not compiled for HIP"); +} + +void test_istft_path_cuda_follows_the_build() { + const auto expected = audio::cuda_log_magnitude_phase_istft_available() + ? audio::LogMagnitudePhaseISTFTPath::Cuda + : audio::LogMagnitudePhaseISTFTPath::Host; + engine::test::require( + audio::select_log_magnitude_phase_istft_path(core::BackendType::Cuda) == expected, + "a CUDA backend must select the CUDA iSTFT exactly when it is built"); +} + +// --------------------------------------------------------------------------- +// F5-TTS: DiT and Vocos device-vs-host selection +// --------------------------------------------------------------------------- + +// The predicates are constexpr, so the contract is also checked at compile +// time. A Metal session must reach the device path; only Cpu must not. +static_assert( + f5::f5_uses_device_backend(f5::F5ComputeDevice{core::BackendType::Metal}), + "F5-TTS must run the DiT and the vocoder on Metal, not on the CPU"); +static_assert( + f5::f5_uses_device_backend(f5::F5ComputeDevice{core::BackendType::Cuda}), + "F5-TTS must run the DiT and the vocoder on CUDA"); +static_assert( + f5::f5_uses_device_backend(f5::F5ComputeDevice{core::BackendType::Vulkan}), + "F5-TTS must run the DiT and the vocoder on Vulkan"); +static_assert( + !f5::f5_uses_device_backend(f5::F5ComputeDevice{core::BackendType::Cpu}), + "F5-TTS must run on the host path when there is no GPU"); + +void test_f5_backend_selection() { + f5::F5ComputeDevice metal; + metal.backend = core::BackendType::Metal; + engine::test::require( + f5::f5_requested_backend(metal) == core::BackendType::Metal, + "F5 must ask for the Metal backend it was given"); + engine::test::require( + f5::f5_uses_device_backend(metal), + "a Metal F5 session must take the device path"); + + f5::F5ComputeDevice host; + engine::test::require( + f5::f5_requested_backend(host) == core::BackendType::Cpu, + "a default F5 device must be CPU"); + engine::test::require( + !f5::f5_uses_device_backend(host), + "a CPU-only F5 session must take the host path"); + + // The pre-`backend` spelling still works: the parity harnesses under + // tests/f5_tts set only use_cuda. + f5::F5ComputeDevice legacy; + legacy.use_cuda = true; + engine::test::require( + f5::f5_requested_backend(legacy) == core::BackendType::Cuda, + "the legacy use_cuda flag must still resolve to the CUDA backend"); + engine::test::require( + f5::f5_uses_device_backend(legacy), + "the legacy use_cuda flag must still take the device path"); + + // An explicit backend wins over the legacy flag. + f5::F5ComputeDevice both; + both.backend = core::BackendType::Metal; + both.use_cuda = true; + engine::test::require( + f5::f5_requested_backend(both) == core::BackendType::Metal, + "an explicit backend must win over the legacy use_cuda flag"); +} + +// --------------------------------------------------------------------------- +// The host iSTFT overlap-add rewrite is bit-identical to the loop it replaced +// --------------------------------------------------------------------------- + +// The serial fold was a scatter over frames, which cannot be an OpenMP loop +// without a race; it is now a gather over output samples. The gather visits +// each output sample's contributing frames in increasing frame order, which is +// the order the scatter reached them in, so the float accumulation order is +// unchanged and the result must match to the last bit — not merely to a +// tolerance. Assert exactly that: max |difference| == 0. +void test_host_istft_matches_the_scatter_reference_bit_for_bit() { + // Large enough that the fold and normalise loops clear their OpenMP + // `if` thresholds, so the parallel path is what is being compared. + constexpr int64_t kFrames = 500; + constexpr int64_t kNfft = 64; + constexpr int64_t kHop = 16; + constexpr int64_t kFreqBins = kNfft / 2 + 1; + constexpr int64_t kOutDim = kFreqBins * 2; + + // Periodic Hann: sum of squares at hop N/4 is a strictly positive envelope. + std::vector window(static_cast(kNfft)); + for (int64_t i = 0; i < kNfft; ++i) { + window[static_cast(i)] = 0.5F - + 0.5F * std::cos(6.283185307179586F * static_cast(i) / static_cast(kNfft)); + } + + // Deterministic, reproducible log-magnitude/phase field with enough + // dynamic range that a reordered accumulation would show up. + std::vector log_magnitude_phase(static_cast(kFrames * kOutDim)); + for (int64_t frame = 0; frame < kFrames; ++frame) { + for (int64_t bin = 0; bin < kFreqBins; ++bin) { + const auto t = static_cast(frame); + const auto f = static_cast(bin); + const size_t row = static_cast(frame * kOutDim); + log_magnitude_phase[row + static_cast(bin)] = + -2.0F + 3.0F * std::sin(0.31F * t + 0.17F * f); + log_magnitude_phase[row + static_cast(kFreqBins + bin)] = + std::cos(0.11F * t - 0.23F * f) * 3.0F; + } + } + + audio::HostLogMagnitudePhaseISTFTConfig config; + config.frames = kFrames; + config.n_fft = kNfft; + config.hop_length = kHop; + config.out_dim = kOutDim; + config.threads = 4; + audio::HostLogMagnitudePhaseISTFT istft(config); + const auto actual = istft.compute(log_magnitude_phase, window); + + // ---- independent reference: the pre-change scatter loop ---- + std::vector> spectrum(static_cast(kFrames * kFreqBins)); + for (int64_t frame = 0; frame < kFrames; ++frame) { + const size_t row = static_cast(frame * kOutDim); + for (int64_t bin = 0; bin < kFreqBins; ++bin) { + const float magnitude = + std::min(std::exp(log_magnitude_phase[row + static_cast(bin)]), 100.0F); + const float phase = log_magnitude_phase[row + static_cast(kFreqBins + bin)]; + spectrum[static_cast(frame * kFreqBins + bin)] = { + magnitude * std::cos(phase), + magnitude * std::sin(phase), + }; + } + } + std::vector framed(static_cast(kFrames * kNfft), 0.0F); + audio::real_fft_inverse( + {static_cast(kFrames), static_cast(kNfft)}, + { + static_cast(kFreqBins * static_cast(sizeof(std::complex))), + static_cast(sizeof(std::complex)), + }, + { + static_cast(kNfft * static_cast(sizeof(float))), + static_cast(sizeof(float)), + }, + 1, + spectrum.data(), + framed.data(), + 1.0F / static_cast(kNfft), + 4); // same thread count the runtime used, so the FFT input is identical + + const int64_t output_size = (kFrames - 1) * kHop + kNfft; + std::vector folded(static_cast(output_size), 0.0F); + std::vector envelope(static_cast(output_size), 0.0F); + for (int64_t frame = 0; frame < kFrames; ++frame) { + const int64_t start = frame * kHop; + const float * src = framed.data() + static_cast(frame * kNfft); + for (int64_t i = 0; i < kNfft; ++i) { + const float w = window[static_cast(i)]; + folded[static_cast(start + i)] += src[i] * w; + envelope[static_cast(start + i)] += w * w; + } + } + const int64_t pad = (kNfft - kHop) / 2; + const int64_t samples = output_size - 2 * pad; + std::vector expected(static_cast(samples)); + for (int64_t i = 0; i < samples; ++i) { + const int64_t src = i + pad; + expected[static_cast(i)] = + folded[static_cast(src)] / envelope[static_cast(src)]; + } + + engine::test::require_eq(actual.audio.size(), expected.size(), "host iSTFT sample count"); + float max_abs_difference = 0.0F; + float peak = 0.0F; + for (size_t i = 0; i < expected.size(); ++i) { + max_abs_difference = std::max( + max_abs_difference, std::fabs(actual.audio[i] - expected[i])); + peak = std::max(peak, std::fabs(expected[i])); + } + engine::test::require(peak > 1.0e-3F, "host iSTFT reference signal must not be silent"); + engine::test::require_close( + max_abs_difference, 0.0F, 0.0F, + "host iSTFT overlap-add must be bit-identical to the scatter reference"); +} + +void test_host_istft_rejects_a_degenerate_window() { + // A window whose squared envelope underflows must still be rejected, and + // the check must fire even though the normalise loop no longer branches. + constexpr int64_t kFrames = 8; + constexpr int64_t kNfft = 32; + constexpr int64_t kHop = 8; + constexpr int64_t kFreqBins = kNfft / 2 + 1; + constexpr int64_t kOutDim = kFreqBins * 2; + + audio::HostLogMagnitudePhaseISTFTConfig config; + config.frames = kFrames; + config.n_fft = kNfft; + config.hop_length = kHop; + config.out_dim = kOutDim; + config.threads = 1; + audio::HostLogMagnitudePhaseISTFT istft(config); + + const std::vector zero_window(static_cast(kNfft), 0.0F); + const std::vector input(static_cast(kFrames * kOutDim), 0.0F); + bool threw = false; + try { + (void) istft.compute(input, zero_window); + } catch (const std::runtime_error &) { + threw = true; + } + engine::test::require(threw, "a zero window must be rejected as an envelope underflow"); +} + +} // namespace + +int main() { + try { + test_istft_path_never_cuda_without_the_cuda_runtime(); + test_istft_path_metal_and_cpu_take_the_host_path(); + test_istft_path_cuda_follows_the_build(); + test_f5_backend_selection(); + test_host_istft_matches_the_scatter_reference_bit_for_bit(); + test_host_istft_rejects_a_degenerate_window(); + std::cout << "backend_path_selection_test passed\n"; + } catch (const std::exception & ex) { + std::cerr << "backend_path_selection_test failed: " << ex.what() << "\n"; + return 1; + } + return 0; +}