diff --git a/CMakeLists.txt b/CMakeLists.txt index 63195ff21..21ba92f5f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2032,6 +2032,27 @@ 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( + 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( + 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/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/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/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/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/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/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/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/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/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/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/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/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/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/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_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; +} 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; +} 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; +}