diff --git a/.coderabbit.yaml b/.coderabbit.yaml
index 5654f12..200ce48 100644
--- a/.coderabbit.yaml
+++ b/.coderabbit.yaml
@@ -12,3 +12,23 @@ reviews:
drafts: true
base_branches:
- ".*"
+ path_instructions:
+ - path: "**/*.{cu,cuh,cpp,c,h,hpp}"
+ instructions: |
+ Prioritize correctness, performance and portability:
+ - Correctness: races, uninitialized or stale state, edge cases (empty, masked,
+ boundary sizes), numerical stability, and fallback paths that differ from the fast
+ path.
+ - Performance: regressions on hot paths, redundant work, unnecessary allocations or
+ synchronization, and caches or reuse that do not actually take effect.
+ - Portability: assumptions tied to one GPU, architecture, backend or platform; hardware
+ limits must be queried or guarded, with a working fallback.
+ - path: "ggml-patches/**"
+ instructions: |
+ Patches to the pinned ggml submodule. Apply the same correctness, performance and
+ portability checks; also check that op preconditions match what the kernels support
+ and that the patch series stays consistent.
+ - path: "{BENCHMARK.md,README.md,docs/**,app/bench*}"
+ instructions: |
+ Check that measurements are sound and that documented numbers, defaults and flags match
+ the code.
diff --git a/BENCHMARK.md b/BENCHMARK.md
new file mode 100644
index 0000000..1d0f3b4
--- /dev/null
+++ b/BENCHMARK.md
@@ -0,0 +1,162 @@
+# Benchmarks
+
+## Speech recognition (Nemotron Speech Streaming)
+
+Nemotron Speech Streaming EN 0.6B, cache-aware streaming, LibriSpeech
+test-clean (2,620 utterances, 5.4 h).
+
+### GeForce RTX 4090
+
+| Engine | Chunk | Compute per chunk (ms)
```
+### Developer Certificate of Origin
+
+Signing off certifies that at least one of (a), (b), or (c) below applies to
+your contribution, and that you agree to (d).
The full, unmodified [Developer Certificate of Origin
1.1](https://developercertificate.org/) follows:
diff --git a/README.md b/README.md
index 453ea3b..498c37a 100644
--- a/README.md
+++ b/README.md
@@ -24,8 +24,36 @@
| Full-duplex voicechat | [Nemotron Labs VoiceChat](https://huggingface.co/nvidia/NVIDIA-NemotronLabs-VoiceChat-11B), including realtime audio, transcripts, and tool calling |
| Speech processing | [Silero VAD](https://github.com/snakers4/silero-vad), punctuation and capitalization, endpointing, text normalization, and subtitles |
+## Performance
+
+NeMo-Speech.cpp is blazing fast and built for real-time streaming. Speech recognition
+transcribes audio in 160 ms chunks up to 67× faster than real time, and speech synthesis
+generates speech up to 60× faster than real time with its first audio in under 10 ms. Both stay
+faster than real time even on a CPU.
+
+Streaming speech recognition with Nemotron Speech Streaming 0.6B (Q8_0), 160 ms chunks:
+
+| Device | Latency per chunk | Throughput | Speedup over NeMo (FP32) |
+|---|:---:|:---:|:---:|
+| GeForce RTX 4090 | **2.3 ms** | **67× real time** | **8.8×** |
+| CPU | **27 ms** | **6× real time** | **3.7×** |
+
+Streaming speech synthesis with MagpieTTS Multilingual (Q8_0), 186 ms audio chunks:
+
+| Device | Time to first audio | Inter-chunk latency | Throughput | Speedup over NeMo (FP32) |
+|---|:---:|:---:|:---:|:---:|
+| GeForce RTX 4090 | **9 ms** | **3 ms** | **60× real time** | **40×** |
+| CPU | **203 ms** | **64 ms** | **2.7× real time** | **9.7×** |
+
+See [BENCHMARK.md](BENCHMARK.md) for the methodology and more results.
+
## Installation
+> [!IMPORTANT]
+> **For the best performance and the latest features, build natively from source.** A native
+> build is compiled for your machine, and release tags can be out of sync with the
+> `main` branch. See [Build from source](#build-from-source).
+
Install the `nemo-speech` CLI for the detected platform and backend:
On Linux or macOS, run:
@@ -45,10 +73,12 @@ irm https://github.com/NVIDIA/NeMo-Speech.cpp/raw/main/scripts/install.ps1 | iex
Open a new PowerShell window after installation so the updated user `PATH`
takes effect.
-The installer prefers a verified native release and falls back to a source
-build when an artifact is unavailable. A source build requires Git, CMake 3.26
-or newer, Ninja, a C++17 compiler, SentencePiece development files, and the
-toolchain required by the selected backend, if any. See
+The installer downloads the prebuilt archive for the latest release, checks it
+against the SHA-256 checksum published with the release, and builds from source
+when no archive is available for your platform. **Pass `--source` (`-Source` on
+Windows) to always build from the `main` branch.** A source build requires Git,
+CMake 3.26 or newer, Ninja, a C++17 compiler, SentencePiece development files,
+and the toolchain required by the selected backend, if any. See
[Installation](docs/install.md) for platform-specific prerequisites and
options.
diff --git a/app/CMakeLists.txt b/app/CMakeLists.txt
index 3e9abd0..9b54fda 100644
--- a/app/CMakeLists.txt
+++ b/app/CMakeLists.txt
@@ -1,8 +1,10 @@
# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
+find_package(Threads REQUIRED)
add_executable(nemo_speech_cli
main.cpp
+ bench.cpp
cli_util.cpp
doctor.cpp
model.cpp
@@ -15,6 +17,7 @@ target_link_libraries(nemo_speech_cli PRIVATE
nemo_speech_engine_registry
ggml
ggml-base
+ Threads::Threads
)
target_include_directories(nemo_speech_cli PRIVATE ${CMAKE_SOURCE_DIR}/include)
if(WIN32)
@@ -22,14 +25,11 @@ if(WIN32)
endif()
if(NEMO_SPEECH_BUILD_ASR)
- find_package(Threads REQUIRED)
target_sources(nemo_speech_cli PRIVATE
- bench.cpp
+ bench_asr.cpp
transcribe.cpp)
target_compile_definitions(nemo_speech_cli PRIVATE NEMO_SPEECH_CLI_ASR=1)
- target_link_libraries(nemo_speech_cli PRIVATE
- nemo_speech_asr
- Threads::Threads)
+ target_link_libraries(nemo_speech_cli PRIVATE nemo_speech_asr)
if(NEMO_SPEECH_BUILD_MIC_CAPTURE)
if(NOT EXISTS "${CMAKE_SOURCE_DIR}/llama.cpp/vendor/miniaudio/miniaudio.h")
message(FATAL_ERROR
@@ -85,12 +85,12 @@ if(NEMO_SPEECH_WITH_NORM)
endif()
if(NEMO_SPEECH_BUILD_DIAR)
target_compile_definitions(nemo_speech_cli PRIVATE NEMO_SPEECH_CLI_DIAR=1)
- target_sources(nemo_speech_cli PRIVATE diarize.cpp)
+ target_sources(nemo_speech_cli PRIVATE bench_diarize.cpp diarize.cpp)
target_link_libraries(nemo_speech_cli PRIVATE nemo_speech_asr)
endif()
if(NEMO_SPEECH_BUILD_TTS)
target_compile_definitions(nemo_speech_cli PRIVATE NEMO_SPEECH_CLI_TTS=1)
- target_sources(nemo_speech_cli PRIVATE synthesize.cpp)
+ target_sources(nemo_speech_cli PRIVATE bench_tts.cpp synthesize.cpp)
target_link_libraries(nemo_speech_cli PRIVATE nemo_speech_tts_magpietts)
if(NEMO_SPEECH_TTS_WITH_JA)
target_compile_definitions(nemo_speech_cli PRIVATE NEMO_SPEECH_CLI_TTS_JA=1)
@@ -101,7 +101,7 @@ if(NEMO_SPEECH_BUILD_TTS)
endif()
if(NEMO_SPEECH_BUILD_NMT)
target_compile_definitions(nemo_speech_cli PRIVATE NEMO_SPEECH_CLI_NMT=1)
- target_sources(nemo_speech_cli PRIVATE translate.cpp)
+ target_sources(nemo_speech_cli PRIVATE bench_translate.cpp translate.cpp)
target_link_libraries(nemo_speech_cli PRIVATE nemo_speech_nmt)
endif()
diff --git a/app/bench.cpp b/app/bench.cpp
index 2aaf46f..40b443f 100644
--- a/app/bench.cpp
+++ b/app/bench.cpp
@@ -1,56 +1,95 @@
// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
// SPDX-License-Identifier: Apache-2.0
+#include "bench.h"
+
#include
#include
#include
#include
+#include
#include
+#include
+#include
#include
#include
#include
#include
-#include
#include
-#include "audio_file.h"
#include "cli_util.h"
#include "commands.h"
-#include "engine_registry.h"
-#include "json.h"
-#include "model_utils.h"
-#include "parameter_parser.h"
-namespace {
+namespace nemo_speech::bench {
namespace fs = std::filesystem;
-namespace asr = nemo_speech::asr;
-using Clock = std::chrono::steady_clock;
-using nemo_speech::json::Value;
-struct AudioInput {
- fs::path path;
- nemo_speech::audio::AudioFile audio;
-};
+double
+stat(
+ const Value& object, const std::string& group, const std::string& name,
+ const std::string& stat_name) {
+ const Value* node = object.find(group);
+ if (node && !name.empty())
+ node = node->find(name);
+ if (node)
+ node = node->find(stat_name);
+ return node && node->is_number() ? node->number() : 0.0;
+}
-struct Options {
- fs::path input;
- std::string model;
- std::string language;
- std::string config_file;
- asr::RecognizerConfig config;
- std::vector concurrency{1, 2, 4};
- int repetitions = 3;
- int warmup = 1;
- bool recursive = false;
- bool stream = false;
- bool json = cli_json();
-};
+std::vector
+collect_files(
+ const fs::path& input, bool recursive, const std::function& accept,
+ const std::string& kind) {
+ std::error_code error;
+ if (fs::is_regular_file(input, error))
+ return {fs::absolute(input)};
+ if (!fs::is_directory(input, error))
+ throw std::invalid_argument(input.string() + " is not a file or directory");
+ std::vector result;
+ auto add = [&](const auto& entry) {
+ if (entry.is_regular_file(error) && accept(entry.path()))
+ result.push_back(fs::absolute(entry.path()));
+ };
+ if (recursive)
+ for (const auto& entry : fs::recursive_directory_iterator(input)) add(entry);
+ else
+ for (const auto& entry : fs::directory_iterator(input)) add(entry);
+ std::sort(result.begin(), result.end());
+ if (result.empty())
+ throw std::invalid_argument(input.string() + " contains no " + kind + " files");
+ return result;
+}
+
+namespace {
+using Clock = std::chrono::steady_clock;
-std::string
-required_value(int& index, int argc, char** argv, const std::string& option) {
- if (++index >= argc)
- throw std::invalid_argument(option + " requires a value");
- return argv[index];
+const char* const kTasks[][2] = {
+ {"asr", "ASR"}, {"tts", "TTS"}, {"diarize", "DIAR"}, {"translate", "NMT"}};
+
+std::unique_ptr
+make_workload(const std::string& task) {
+#if defined(NEMO_SPEECH_CLI_ASR)
+ if (task == "asr")
+ return make_asr_workload();
+#endif
+#if defined(NEMO_SPEECH_CLI_TTS)
+ if (task == "tts")
+ return make_tts_workload();
+#endif
+#if defined(NEMO_SPEECH_CLI_DIAR)
+ if (task == "diarize")
+ return make_diarize_workload();
+#endif
+#if defined(NEMO_SPEECH_CLI_NMT)
+ if (task == "translate")
+ return make_translate_workload();
+#endif
+ for (const auto& known : kTasks)
+ if (task == known[0])
+ throw UnsupportedFeatureError(
+ "this build does not include the '" + task +
+ "' benchmark workload; rebuild with -DNEMO_SPEECH_BUILD_" + known[1] + "=ON");
+ throw std::invalid_argument(
+ "unknown bench workload '" + task + "' (expected asr, tts, diarize, or translate)");
}
std::vector
@@ -71,196 +110,218 @@ parse_concurrency(const std::string& value) {
return result;
}
-Options
-parse_options(int argc, char** argv) {
- Options options;
- options.config.backend.gpu = default_gpu_index();
- nemo_speech::common::ParameterParser parser;
- parser.Register("asr", options.config);
- for (int i = 0; i < argc; ++i)
- if (std::string(argv[i]) == "--config")
- options.config_file = required_value(i, argc, argv, "--config");
- if (!options.config_file.empty())
- parser.ApplyYaml(options.config_file);
+double
+percentile(std::vector values, double fraction) {
+ if (values.empty())
+ return 0.0;
+ std::sort(values.begin(), values.end());
+ const double position = fraction * static_cast(values.size() - 1);
+ const size_t lower = static_cast(position);
+ const size_t upper = std::min(lower + 1, values.size() - 1);
+ return values[lower] + (values[upper] - values[lower]) * (position - lower);
+}
+
+Value
+distribution(const std::vector& values) {
+ Value result(Value::Object{});
+ double sum = 0.0;
+ for (const double value : values) sum += value;
+ result["mean"] = values.empty() ? 0.0 : sum / values.size();
+ result["p50"] = percentile(values, 0.50);
+ result["p90"] = percentile(values, 0.90);
+ result["p95"] = percentile(values, 0.95);
+ result["p99"] = percentile(values, 0.99);
+ result["min"] = values.empty() ? 0.0 : *std::min_element(values.begin(), values.end());
+ result["max"] = values.empty() ? 0.0 : *std::max_element(values.begin(), values.end());
+ return result;
+}
+
+// Per-metric distributions in first-seen metric order; JSON objects sort keys.
+Value
+metric_distributions(const std::vector& items) {
+ std::vector order;
+ std::map> values;
+ for (const auto* item : items) {
+ for (const auto& [name, value] : item->metrics) {
+ auto [it, inserted] = values.try_emplace(name);
+ if (inserted)
+ order.push_back(name);
+ it->second.push_back(value);
+ }
+ for (const auto& [name, samples] : item->samples) {
+ auto [it, inserted] = values.try_emplace(name);
+ if (inserted)
+ order.push_back(name);
+ it->second.insert(it->second.end(), samples.begin(), samples.end());
+ }
+ }
+ Value result(Value::Object{});
+ for (const auto& name : order) result[name] = distribution(values[name]);
+ return result;
+}
+
+void
+print_table(
+ const std::vector& columns, const Value::Array& rows,
+ const std::vector* names = nullptr) {
+ size_t name_width = 6;
+ if (names)
+ for (const auto& name : *names) name_width = std::max(name_width, name.size() + 2);
+ std::vector widths;
+ for (const auto& column : columns)
+ widths.push_back(static_cast(std::max(10, column.header.size() + 2)));
+ if (names)
+ std::printf("%-*s", static_cast(name_width), "INPUT");
+ for (size_t c = 0; c < columns.size(); ++c)
+ std::printf("%-*s", widths[c], columns[c].header.c_str());
+ std::printf("\n");
+ for (size_t r = 0; r < rows.size(); ++r) {
+ if (names)
+ std::printf("%-*s", static_cast(name_width), (*names)[r].c_str());
+ for (size_t c = 0; c < columns.size(); ++c)
+ std::printf("%-*.*f", widths[c], columns[c].precision, columns[c].value(rows[r]));
+ std::printf("\n");
+ }
+}
+
+Column
+key_column(const std::string& header, const std::string& key, int precision) {
+ return {header, [key](const Value& row) { return row.number_or(key); }, precision};
+}
+
+Column
+stat_column(
+ const std::string& header, const std::string& group, const std::string& name,
+ const std::string& stat_name, int precision) {
+ return {header, [=](const Value& row) { return stat(row, group, name, stat_name); }, precision};
+}
+
+int
+run_bench(int argc, char** argv) {
+ if (argc == 0 || is_help_argument(argv[0])) {
+ print_bench_help("nemo-speech");
+ return 0;
+ }
+ auto workload = make_workload(argv[0]);
+ CommonOptions options;
+ options.gpu = default_gpu_index();
+ options.concurrency = workload->default_concurrency();
+ bool json = cli_json();
+
+ common::ParameterParser parser;
+ workload->register_parameters(parser);
+ std::string config_file;
+ for (int i = 1; i < argc; ++i)
+ if (std::string(argv[i]) == "--config") {
+ if (++i >= argc)
+ throw std::invalid_argument("--config requires a value");
+ config_file = argv[i];
+ }
+ if (!config_file.empty())
+ parser.ApplyYaml(config_file);
parser.ApplyEnv("NEMO_SPEECH");
- for (int i = 0; i < argc; ++i) {
+ for (int i = 1; i < argc; ++i) {
const std::string arg = argv[i];
- if (arg == "asr") {
- continue;
+ auto value = [&]() {
+ if (++i >= argc)
+ throw std::invalid_argument(arg + " requires a value");
+ return std::string(argv[i]);
+ };
+ if (is_help_argument(arg)) {
+ print_bench_help("nemo-speech");
+ return 0;
} else if (arg == "--config") {
++i;
- } else if (arg == "--model" || arg == "-m") {
- options.model = required_value(i, argc, argv, arg);
- } else if (arg == "--language" || arg == "-l") {
- options.language = required_value(i, argc, argv, arg);
- } else if (arg == "--device" || arg == "--backend") {
- options.config.backend.gpu = parse_device(required_value(i, argc, argv, arg), arg);
- } else if (arg == "--gpu") {
- options.config.backend.gpu =
- parse_int(required_value(i, argc, argv, arg), arg, -1, 1024);
} else if (arg == "--concurrency" || arg == "-c") {
- options.concurrency = parse_concurrency(required_value(i, argc, argv, arg));
+ options.concurrency = parse_concurrency(value());
} else if (arg == "--repetitions" || arg == "-n") {
- options.repetitions = parse_int(required_value(i, argc, argv, arg), arg, 1, 10000);
+ options.repetitions = parse_int(value(), arg, 1, 10000);
+ } else if (arg == "--per-stream") {
+ options.per_stream = parse_int(value(), arg, 1, 100000);
} else if (arg == "--warmup") {
- options.warmup = parse_int(required_value(i, argc, argv, arg), arg, 0, 1000);
- } else if (arg == "--mode") {
- const auto value = required_value(i, argc, argv, arg);
- if (value == "offline")
- options.stream = false;
- else if (value == "stream")
- options.stream = true;
- else
- throw std::invalid_argument("--mode must be offline or stream");
+ options.warmup = parse_int(value(), arg, 0, 1000);
+ } else if (arg == "--device" || arg == "--backend") {
+ options.device = value();
+ options.gpu = parse_device(options.device, arg);
+ options.device_set = true;
+ } else if (arg == "--gpu") {
+ options.gpu = parse_int(value(), arg, -1, 1024);
+ options.device = options.gpu < 0 ? "cpu" : "gpu:" + std::to_string(options.gpu);
+ options.device_set = true;
} else if (arg == "--recursive" || arg == "-r") {
options.recursive = true;
+ } else if (arg == "--save") {
+ options.save_dir = value();
} else if (arg == "--json") {
- options.json = true;
+ json = true;
+ } else if (workload->parse_option(arg, value)) {
+ continue;
} else if (!arg.empty() && arg.front() == '-') {
bool consumed = false;
- const char* next = i + 1 < argc ? argv[i + 1] : nullptr;
- if (!parser.ParseCliArg(arg, next, &consumed))
+ if (!parser.ParseCliArg(arg, i + 1 < argc ? argv[i + 1] : nullptr, &consumed))
throw std::invalid_argument("unknown option: " + arg);
if (consumed)
++i;
- } else if (options.input.empty()) {
- options.input = arg;
} else {
- throw std::invalid_argument("unexpected argument: " + arg);
+ options.positional.push_back(arg);
}
}
- if (options.input.empty())
- throw std::invalid_argument("bench asr requires a WAV file or directory");
- return options;
-}
-
-std::vector
-collect_files(const Options& options) {
- std::error_code error;
- if (fs::is_regular_file(options.input, error))
- return {fs::absolute(options.input)};
- if (!fs::is_directory(options.input, error))
- throw std::invalid_argument(options.input.string() + " is not a file or directory");
- std::vector result;
- auto add = [&](const auto& entry) {
- if (entry.is_regular_file(error) && nemo_speech::audio::is_wav_path(entry.path().string()))
- result.push_back(fs::absolute(entry.path()));
- };
- if (options.recursive)
- for (const auto& entry : fs::recursive_directory_iterator(options.input)) add(entry);
- else
- for (const auto& entry : fs::directory_iterator(options.input)) add(entry);
- std::sort(result.begin(), result.end());
- if (result.empty())
- throw std::invalid_argument(options.input.string() + " contains no WAV files");
- return result;
-}
-
-std::string
-append_result(const asr::Result& result) {
- if (result.alternatives.empty())
- return {};
- return result.alternatives.front().transcript;
-}
-
-std::string
-recognize(asr::Recognizer& recognizer, const AudioInput& input, const Options& options) {
- asr::AsrRequestOptions request;
- request.language_code = options.language;
- if (!options.stream) {
- return append_result(recognizer.recognize(
- input.audio.samples.data(), input.audio.samples.size(), request, options.language,
- input.audio.sample_rate));
- }
- auto stream = recognizer.streaming_recognize(request, options.language);
- const size_t chunk = std::max(1, input.audio.sample_rate * 160 / 1000);
- std::string transcript;
- for (size_t offset = 0; offset < input.audio.samples.size(); offset += chunk) {
- const size_t count = std::min(chunk, input.audio.samples.size() - offset);
- stream->push(input.audio.samples.data() + offset, count, input.audio.sample_rate);
- while (auto result = stream->next()) {
- if (result->is_final) {
- const auto text = append_result(*result);
- if (!text.empty()) {
- if (!transcript.empty())
- transcript += ' ';
- transcript += text;
- }
- } else {
- break;
- }
+ workload->prepare(options);
+ const size_t inputs = workload->input_count();
+ if (inputs == 0)
+ throw std::invalid_argument("bench " + workload->task() + " has no inputs");
+ if (!options.save_dir.empty()) {
+ // --save writes DIR/ saved;
+ for (size_t i = 0; i < inputs; ++i) {
+ const std::string name = workload->input_name(i);
+ const auto [it, inserted] = saved.emplace(fs::path(name).stem().string(), name);
+ if (!inserted)
+ throw std::invalid_argument(
+ "--save: inputs " + it->second + " and " + name + " would both be written to " +
+ it->first + ".txt; give them distinct file names");
}
}
- const auto final = append_result(stream->finish());
- if (!final.empty()) {
- if (!transcript.empty())
- transcript += ' ';
- transcript += final;
- }
- return transcript;
-}
-
-int
-run_bench(int argc, char** argv) {
- if (argc == 0 || is_help_argument(argv[0])) {
- print_bench_help("nemo-speech");
- return 0;
- }
- if (std::string(argv[0]) != "asr")
- throw std::invalid_argument(
- "supported workload is 'asr' (use 'nemo-speech bench asr ...')");
- Options options = parse_options(argc, argv);
- std::vector inputs;
- double corpus_seconds = 0.0;
- for (const auto& path : collect_files(options)) {
- auto audio = nemo_speech::audio::load_wav_file(path.string());
- corpus_seconds += static_cast(audio.samples.size()) / audio.sample_rate;
- inputs.push_back({path, std::move(audio)});
- }
const int max_concurrency =
*std::max_element(options.concurrency.begin(), options.concurrency.end());
- options.config.model.path =
- resolve_model_file(
- options.model.empty() ? options.config.model.path : options.model, "asr", "ASR model")
- .string();
- options.config.batching.enabled = max_concurrency > 1;
- options.config.batching.max_batch_size =
- std::max(options.config.batching.max_batch_size, max_concurrency);
- options.config.batching.max_queue_depth =
- std::max(options.config.batching.max_queue_depth, max_concurrency * 2);
- options.config.batching.state_arena_slots =
- std::max(options.config.batching.state_arena_slots, max_concurrency);
- options.config.log_status = !cli_quiet() && !cli_json();
- nemo_speech::EngineRegistry engines;
const auto load_start = Clock::now();
- auto recognizer = engines.load_asr(std::move(options.config));
+ workload->load(options, max_concurrency);
const double load_ms =
std::chrono::duration(Clock::now() - load_start).count();
const auto warmup_start = Clock::now();
- engines.warmup();
- for (int i = 0; i < options.warmup; ++i)
- (void)recognize(*recognizer, inputs[static_cast(i) % inputs.size()], options);
+ workload->warmup_engine();
+ for (int i = 0; i < options.warmup; ++i) (void)workload->run(static_cast(i) % inputs);
const double warmup_ms =
std::chrono::duration(Clock::now() - warmup_start).count();
Value output(Value::Object{});
- output["command"] = "bench asr";
- output["model"] = recognizer->model_name();
- output["mode"] = options.stream ? "stream" : "offline";
+ output["command"] = "bench " + workload->task();
+ output["task"] = workload->task();
+ workload->describe(output);
output["load_ms"] = load_ms;
output["warmup_ms"] = warmup_ms;
- output["files"] = static_cast(inputs.size());
- output["corpus_audio_seconds"] = corpus_seconds;
+ output["warmup"] = options.warmup;
output["repetitions"] = options.repetitions;
+ if (options.per_stream > 0)
+ output["per_stream"] = options.per_stream;
+ const auto input_columns = workload->input_columns();
+ std::vector names;
+ for (size_t i = 0; i < inputs; ++i) names.push_back(workload->input_name(i));
+
Value::Array runs;
- std::vector reference(inputs.size());
- std::vector reference_set(inputs.size(), false);
+ std::vector reference(inputs);
+ std::vector reference_set(inputs, false);
for (const int concurrency : options.concurrency) {
- const size_t work_count = inputs.size() * static_cast(options.repetitions);
+ const size_t work_count =
+ options.per_stream > 0
+ ? static_cast(options.per_stream) * static_cast(concurrency)
+ : inputs * static_cast(options.repetitions);
+ std::vector results(work_count);
+ std::vector latency(work_count);
std::atomic next{0};
- std::atomic mismatches{0};
- std::mutex error_mutex;
+ std::mutex failure_mutex;
std::exception_ptr failure;
const auto started = Clock::now();
std::vector workers;
@@ -272,22 +333,19 @@ run_bench(int argc, char** argv) {
const size_t work = next.fetch_add(1);
if (work >= work_count)
return;
- const size_t input_index = work % inputs.size();
- const auto transcript =
- recognize(*recognizer, inputs[input_index], options);
- std::lock_guard lock(error_mutex);
- if (!reference_set[input_index]) {
- reference[input_index] = transcript;
- reference_set[input_index] = true;
- } else if (reference[input_index] != transcript) {
- ++mismatches;
- }
+ const auto item_start = Clock::now();
+ results[work] = workload->run(work % inputs);
+ results[work].input = work % inputs;
+ latency[work] =
+ std::chrono::duration(Clock::now() - item_start)
+ .count();
}
}
catch (...) {
- std::lock_guard lock(error_mutex);
+ std::lock_guard lock(failure_mutex);
if (!failure)
failure = std::current_exception();
+ next = work_count;
}
});
}
@@ -295,77 +353,192 @@ run_bench(int argc, char** argv) {
if (failure)
std::rethrow_exception(failure);
const double wall_seconds = std::chrono::duration(Clock::now() - started).count();
- const double audio_seconds = corpus_seconds * options.repetitions;
+
+ std::vector mismatches(inputs, 0);
+ int total_mismatches = 0;
+ for (size_t work = 0; work < work_count; ++work) {
+ const size_t index = work % inputs;
+ if (!reference_set[index]) {
+ reference[index] = results[work].signature;
+ reference_set[index] = true;
+ } else if (reference[index] != results[work].signature) {
+ ++mismatches[index];
+ ++total_mismatches;
+ }
+ }
+ std::vector all_items;
+ for (const auto& item : results) all_items.push_back(&item);
+
Value run(Value::Object{});
run["concurrency"] = concurrency;
- run["utterances"] = static_cast(work_count);
- run["audio_seconds"] = audio_seconds;
+ run["items"] = static_cast(work_count);
run["wall_seconds"] = wall_seconds;
- run["rtfx"] = audio_seconds / wall_seconds;
- run["utterances_per_second"] = work_count / wall_seconds;
- run["transcript_mismatches"] = mismatches.load();
+ run["items_per_second"] = work_count / wall_seconds;
+ run["latency_ms"] = distribution(latency);
+ Value metrics = metric_distributions(all_items);
+ if (!metrics.object().empty())
+ run["metrics"] = std::move(metrics);
+ run[workload->mismatch_key()] = total_mismatches;
+ workload->summarize_run(run, results, wall_seconds);
+ if (!input_columns.empty()) {
+ Value::Array per_input;
+ for (size_t index = 0; index < inputs; ++index) {
+ std::vector items;
+ std::vector item_latency;
+ for (size_t work = index; work < work_count; work += inputs) {
+ items.push_back(&results[work]);
+ item_latency.push_back(latency[work]);
+ }
+ Value entry(Value::Object{});
+ entry["name"] = names[index];
+ workload->describe_input(index, entry);
+ entry["repetitions"] = static_cast(items.size());
+ entry["latency_ms"] = distribution(item_latency);
+ entry["metrics"] = metric_distributions(items);
+ entry[workload->mismatch_key()] = mismatches[index];
+ per_input.emplace_back(std::move(entry));
+ }
+ run["inputs"] = std::move(per_input);
+ }
runs.emplace_back(std::move(run));
}
output["runs"] = std::move(runs);
- if (options.json) {
+ if (!options.save_dir.empty()) {
+ fs::create_directories(options.save_dir);
+ for (size_t index = 0; index < inputs; ++index) {
+ if (!reference_set[index])
+ continue;
+ std::ofstream file(
+ fs::path(options.save_dir) / (fs::path(names[index]).stem().string() + ".txt"));
+ file << reference[index] << '\n';
+ if (!file)
+ throw std::runtime_error("cannot write outputs to " + options.save_dir);
+ }
+ }
+
+ if (json) {
std::printf("%s\n", output.dump(2).c_str());
- } else {
- std::printf(
- "Model: %s\nMode: %s\nLoad: %.1f ms Warmup: %.1f ms\n",
- recognizer->model_name().c_str(), options.stream ? "stream" : "offline", load_ms,
- warmup_ms);
- std::printf(
- "%-12s %-12s %-12s %-12s %-10s\n", "CONCURRENCY", "WALL (s)", "RTFx", "UTT/s",
- "MISMATCH");
- for (const auto& run : output.at("runs").array())
- std::printf(
- "%-12d %-12.3f %-12.2f %-12.2f %-10d\n",
- static_cast(run.at("concurrency").number()), run.at("wall_seconds").number(),
- run.at("rtfx").number(), run.at("utterances_per_second").number(),
- static_cast(run.at("transcript_mismatches").number()));
+ return 0;
}
+ for (const auto& line : workload->header_lines()) std::printf("%s\n", line.c_str());
+ std::printf("Load: %.1f ms Warmup: %.1f ms\n", load_ms, warmup_ms);
+ std::vector columns{
+ key_column("CONCURRENCY", "concurrency", 0),
+ key_column("ITEMS", "items", 0),
+ key_column("WALL (s)", "wall_seconds", 3),
+ key_column("ITEMS/s", "items_per_second", 2),
+ stat_column("P50 (ms)", "latency_ms", "", "p50", 1),
+ stat_column("P95 (ms)", "latency_ms", "", "p95", 1),
+ };
+ for (auto& column : workload->run_columns()) columns.push_back(std::move(column));
+ columns.push_back(key_column("MISMATCH", workload->mismatch_key(), 0));
+ print_table(columns, output.at("runs").array());
+ if (!input_columns.empty())
+ for (const auto& run : output.at("runs").array()) {
+ std::printf(
+ "\nPer input, concurrency %d (%d requests):\n",
+ static_cast(run.at("concurrency").number()),
+ static_cast(run.at("items").number()));
+ auto per_input = input_columns;
+ per_input.push_back(key_column("MISMATCH", workload->mismatch_key(), 0));
+ print_table(per_input, run.at("inputs").array(), &names);
+ }
return 0;
}
} // namespace
+} // namespace nemo_speech::bench
void
print_bench_help(const char* program) {
std::printf(
- "Usage: %s bench asr INPUT --model MODEL [options]\n\n"
- "Benchmark end-to-end ASR with one shared recognizer and concurrent utterances.\n"
- "Directory inputs are sorted and loaded once; GPU work uses the runtime's dynamic "
- "batching.\n\n"
- "Options:\n"
- " -m, --model MODEL Local ASR GGUF path\n"
- " -c, --concurrency LIST Comma-separated levels (default: 1,2,4)\n"
+ "Usage: %s bench [INPUT] [options]\n\n"
+ "Benchmark an end-to-end workload with one shared engine and concurrent requests.\n"
+ "Every task reports load/warmup time and, per concurrency level, wall time,\n"
+ "items/s, per-item latency (mean/p50/p95), task metrics, and output mismatches\n"
+ "against the first result seen for each input.\n\n"
+ "Tasks:\n"
+#if defined(NEMO_SPEECH_CLI_ASR)
+ " asr INPUT --model MODEL WAV file or directory\n"
+#endif
+#if defined(NEMO_SPEECH_CLI_TTS)
+ " tts [INPUT] [--text TEXT]... .txt file or directory (one utterance per file)\n"
+#endif
+#if defined(NEMO_SPEECH_CLI_DIAR)
+ " diarize INPUT [--model MODEL] WAV file or directory\n"
+#endif
+#if defined(NEMO_SPEECH_CLI_NMT)
+ " translate INPUT --model MODEL --from SRC --to DST\n"
+ " Text file, one input per line\n"
+#endif
+ "\nCommon options:\n"
+ " -c, --concurrency LIST Comma-separated levels (default: asr/diarize 1,2,4;\n"
+ " tts/translate 1)\n"
" -n, --repetitions N Corpus repetitions per level (default: 3)\n"
+ " --per-stream N Instead: each concurrent stream sends N requests,\n"
+ " cycling through the inputs\n"
" --warmup N Timed-input warmup iterations (default: 1)\n"
- " --mode offline|stream Recognition mode (default: offline)\n"
" --device, --backend DEVICE\n"
" auto, cpu, cuda[:N], metal, or vulkan[:N]\n"
- " -l, --language CODE Prompt language code\n"
" -r, --recursive Recurse into input directories\n"
- " --json Emit machine-readable results\n"
+ " --save DIR Write each input's output to DIR/
+#include
+#include
+#include
+#include
+#include
+#include
+
+#include "json.h"
+#include "parameter_parser.h"
+
+namespace nemo_speech::bench {
+
+using json::Value;
+
+struct CommonOptions {
+ std::vector concurrency;
+ int repetitions = 3;
+ int per_stream = 0; // > 0: requests per concurrent stream (cycling inputs)
+ int warmup = 1;
+ int gpu = 0;
+ std::string device = "auto";
+ bool device_set = false;
+ bool recursive = false;
+ std::string save_dir; // write each input's first output here as .txt
+ std::vector positional;
+};
+
+// Per-item output of Workload::run. `signature` is compared against the first
+// result for the same input to count mismatches across repetitions.
+struct ItemResult {
+ std::string signature;
+ std::vector> metrics;
+ // Per-item value lists (e.g. every chunk gap), pooled across items.
+ std::vector>> samples = {};
+ size_t input = 0; // index of the processed input (set by the harness)
+};
+
+// A table column; `value` reads a run object (summary table) or a per-input
+// object (input table).
+struct Column {
+ std::string header;
+ std::function value;
+ int precision = 2;
+};
+
+class Workload {
+ public:
+ virtual ~Workload() = default;
+
+ virtual std::string task() const = 0;
+ virtual std::vector default_concurrency() const { return {1, 2, 4}; }
+ virtual void register_parameters(common::ParameterParser& parser) = 0;
+ // Consume a task-specific option; `value` reads its argument.
+ virtual bool parse_option(const std::string& arg, const std::function& value) {
+ (void)arg;
+ (void)value;
+ return false;
+ }
+ // Validate options and load inputs (not timed).
+ virtual void prepare(const CommonOptions& options) = 0;
+ virtual size_t input_count() const = 0;
+ virtual std::string input_name(size_t index) const = 0;
+ // Build the engine sized for `max_concurrency` (timed as load_ms).
+ virtual void load(const CommonOptions& options, int max_concurrency) = 0;
+ // Engine-internal warmup before the timed-input warmup iterations.
+ virtual void warmup_engine() {}
+ // Process one input; called concurrently from worker threads.
+ virtual ItemResult run(size_t index) = 0;
+
+ // Top-level report keys (model(s), mode, corpus facts).
+ virtual void describe(Value& output) const = 0;
+ virtual std::vector header_lines() const = 0;
+ virtual std::string mismatch_key() const = 0;
+ // Task keys for one concurrency level (e.g. rtfx over the corpus).
+ virtual void summarize_run(
+ Value& run, const std::vector& items, double wall_seconds) const {
+ (void)run;
+ (void)items;
+ (void)wall_seconds;
+ }
+ virtual std::vector run_columns() const { return {}; }
+ // Non-empty input columns enable the per-input breakdown.
+ virtual std::vector input_columns() const { return {}; }
+ virtual void describe_input(size_t index, Value& entry) const {
+ (void)index;
+ (void)entry;
+ }
+};
+
+// Reads `object[group][name][stat]`, or 0 when absent.
+double stat(
+ const Value& object, const std::string& group, const std::string& name,
+ const std::string& stat_name);
+std::vector collect_files(
+ const std::filesystem::path& input, bool recursive,
+ const std::function& accept, const std::string& kind);
+
+#if defined(NEMO_SPEECH_CLI_ASR)
+std::unique_ptr make_asr_workload();
+#endif
+#if defined(NEMO_SPEECH_CLI_TTS)
+std::unique_ptr make_tts_workload();
+#endif
+#if defined(NEMO_SPEECH_CLI_DIAR)
+std::unique_ptr make_diarize_workload();
+#endif
+#if defined(NEMO_SPEECH_CLI_NMT)
+std::unique_ptr make_translate_workload();
+#endif
+
+} // namespace nemo_speech::bench
diff --git a/app/bench_asr.cpp b/app/bench_asr.cpp
new file mode 100644
index 0000000..49abf3f
--- /dev/null
+++ b/app/bench_asr.cpp
@@ -0,0 +1,252 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+
+#include "audio_file.h"
+#include "bench.h"
+#include "cli_util.h"
+#include "engine_registry.h"
+#include "model_utils.h"
+
+namespace nemo_speech::bench {
+namespace {
+namespace fs = std::filesystem;
+
+struct AudioInput {
+ fs::path path;
+ audio::AudioFile audio;
+};
+
+std::string
+first_transcript(const asr::Result& result) {
+ if (result.alternatives.empty())
+ return {};
+ return result.alternatives.front().transcript;
+}
+
+void
+append_text(std::string& transcript, const std::string& text) {
+ if (text.empty())
+ return;
+ if (!transcript.empty())
+ transcript += ' ';
+ transcript += text;
+}
+
+class AsrWorkload : public Workload {
+ public:
+ AsrWorkload() { config_.backend.gpu = default_gpu_index(); }
+
+ std::string task() const override { return "asr"; }
+ void register_parameters(common::ParameterParser& parser) override {
+ parser.Register("asr", config_);
+ }
+ bool parse_option(const std::string& arg, const std::function& value) override {
+ if (arg == "--model" || arg == "-m") {
+ model_ = value();
+ } else if (arg == "--language" || arg == "-l") {
+ language_ = value();
+ } else if (arg == "--mode") {
+ const auto mode = value();
+ if (mode != "offline" && mode != "stream")
+ throw std::invalid_argument("--mode must be offline or stream");
+ stream_ = mode == "stream";
+ } else if (arg == "--chunk-ms") {
+ chunk_ms_ = parse_int(value(), arg, 10, 60000);
+ } else if (arg == "--trace") {
+ trace_path_ = value();
+ } else {
+ return false;
+ }
+ return true;
+ }
+ void prepare(const CommonOptions& options) override {
+ if (options.positional.size() != 1)
+ throw std::invalid_argument(
+ options.positional.empty() ? "bench asr requires a WAV file or directory"
+ : "unexpected argument: " + options.positional[1]);
+ for (const auto& path : collect_files(
+ options.positional.front(), options.recursive,
+ [](const fs::path& file) { return audio::is_wav_path(file.string()); }, "WAV")) {
+ auto audio = audio::load_wav_file(path.string());
+ corpus_seconds_ += static_cast(audio.samples.size()) / audio.sample_rate;
+ inputs_.push_back({path, std::move(audio)});
+ }
+ if (options.device_set)
+ config_.backend.gpu = options.gpu;
+ if (!trace_path_.empty()) {
+ if (!stream_)
+ throw std::invalid_argument("--trace requires --mode stream");
+ trace_.open(trace_path_, std::ios::trunc);
+ if (!trace_)
+ throw std::runtime_error("cannot write " + trace_path_);
+ }
+ }
+ size_t input_count() const override { return inputs_.size(); }
+ std::string input_name(size_t index) const override {
+ return inputs_[index].path.filename().string();
+ }
+ void load(const CommonOptions&, int max_concurrency) override {
+ config_.model.path =
+ resolve_model_file(model_.empty() ? config_.model.path : model_, "asr", "ASR model")
+ .string();
+ config_.batching.enabled = max_concurrency > 1;
+ config_.batching.max_batch_size =
+ std::max(config_.batching.max_batch_size, max_concurrency);
+ config_.batching.max_queue_depth =
+ std::max(config_.batching.max_queue_depth, max_concurrency * 2);
+ config_.batching.state_arena_slots =
+ std::max(config_.batching.state_arena_slots, max_concurrency);
+ config_.log_status = !cli_quiet() && !cli_json();
+ recognizer_ = engines_.load_asr(config_);
+ }
+ void warmup_engine() override { engines_.warmup(); }
+ ItemResult run(size_t index) override {
+ std::vector chunk_latency_ms;
+ auto transcript = recognize(inputs_[index], chunk_latency_ms);
+ if (!stream_)
+ return {std::move(transcript), {}};
+ return {std::move(transcript), {}, {{"chunk_latency_ms", std::move(chunk_latency_ms)}}};
+ }
+
+ void describe(Value& output) const override {
+ output["model"] = recognizer_->model_name();
+ output["mode"] = stream_ ? "stream" : "offline";
+ if (stream_)
+ output["chunk_ms"] = chunk_ms();
+ output["files"] = static_cast(inputs_.size());
+ output["corpus_audio_seconds"] = corpus_seconds_;
+ }
+ std::vector header_lines() const override {
+ return {
+ "Model: " + recognizer_->model_name(),
+ std::string("Mode: ") +
+ (stream_ ? "stream, " + std::to_string(chunk_ms()) + " ms chunks" : "offline")};
+ }
+ std::string mismatch_key() const override { return "transcript_mismatches"; }
+ void summarize_run(
+ Value& run, const std::vector& items, double wall_seconds) const override {
+ double audio_seconds = 0.0;
+ for (const auto& item : items)
+ audio_seconds += static_cast(inputs_[item.input].audio.samples.size()) /
+ inputs_[item.input].audio.sample_rate;
+ run["utterances"] = static_cast(items.size());
+ run["audio_seconds"] = audio_seconds;
+ run["rtfx"] = audio_seconds / wall_seconds;
+ run["utterances_per_second"] = items.size() / wall_seconds;
+ }
+ std::vector run_columns() const override {
+ std::vector columns = {
+ {"RTFx", [](const Value& run) { return run.number_or("rtfx"); }, 2}};
+ if (stream_) {
+ // compute time per streamed chunk (push + drain), client-side
+ columns.push_back(
+ {"CHUNK avg (ms)",
+ [](const Value& run) { return stat(run, "metrics", "chunk_latency_ms", "mean"); },
+ 2});
+ columns.push_back(
+ {"CHUNK p99 (ms)",
+ [](const Value& run) { return stat(run, "metrics", "chunk_latency_ms", "p99"); },
+ 2});
+ }
+ return columns;
+ }
+
+ private:
+ // Streamed chunk length: the model's cache-aware chunk ((right context + 1) x 80 ms) unless
+ // --chunk-ms overrides it.
+ int chunk_ms() const {
+ if (chunk_ms_ > 0)
+ return chunk_ms_;
+ const int rc = config_.streaming.rnnt_right_context;
+ return rc >= 0 ? (rc + 1) * 80 : 160;
+ }
+ std::string recognize(const AudioInput& input, std::vector& chunk_latency_ms) {
+ asr::AsrRequestOptions request;
+ request.language_code = language_;
+ if (!stream_) {
+ return first_transcript(recognizer_->recognize(
+ input.audio.samples.data(), input.audio.samples.size(), request, language_,
+ input.audio.sample_rate));
+ }
+ auto stream = recognizer_->streaming_recognize(request, language_);
+ const size_t chunk =
+ std::max(1, static_cast(input.audio.sample_rate) * chunk_ms() / 1000);
+ const auto stream_started = std::chrono::steady_clock::now();
+ std::string transcript;
+ for (size_t offset = 0; offset < input.audio.samples.size(); offset += chunk) {
+ const size_t count = std::min(chunk, input.audio.samples.size() - offset);
+ const auto started = std::chrono::steady_clock::now();
+ stream->push(input.audio.samples.data() + offset, count, input.audio.sample_rate);
+ std::string interim;
+ while (auto result = stream->next()) {
+ if (!result->is_final) {
+ interim = first_transcript(*result);
+ break;
+ }
+ append_text(transcript, first_transcript(*result));
+ }
+ const auto now = std::chrono::steady_clock::now();
+ chunk_latency_ms.push_back(
+ std::chrono::duration(now - started).count());
+ if (trace_.is_open()) {
+ std::string text = transcript;
+ append_text(text, interim);
+ write_trace(
+ input, static_cast(offset + count) / input.audio.sample_rate,
+ std::chrono::duration(now - stream_started).count(), text);
+ }
+ }
+ append_text(transcript, first_transcript(stream->finish()));
+ if (trace_.is_open())
+ write_trace(
+ input, static_cast(input.audio.samples.size()) / input.audio.sample_rate,
+ std::chrono::duration(
+ std::chrono::steady_clock::now() - stream_started)
+ .count(),
+ transcript);
+ return transcript;
+ }
+
+ // One JSON line per streamed chunk: audio position, elapsed time since the stream started, and
+ // the transcript so far (finals plus the current interim). Used to render speed demos.
+ void write_trace(
+ const AudioInput& input, double audio_s, double elapsed_ms, const std::string& text) {
+ Value line;
+ line["input"] = input.path.filename().string();
+ line["audio_s"] = audio_s;
+ line["elapsed_ms"] = elapsed_ms;
+ line["text"] = text;
+ std::lock_guard lock(trace_mutex_);
+ trace_ << line.dump() << '\n';
+ }
+
+ asr::RecognizerConfig config_;
+ std::string model_;
+ std::string language_;
+ bool stream_ = false;
+ int chunk_ms_ = 0;
+ std::string trace_path_;
+ std::ofstream trace_;
+ std::mutex trace_mutex_;
+ std::vector inputs_;
+ double corpus_seconds_ = 0.0;
+ EngineRegistry engines_;
+ std::shared_ptr recognizer_;
+};
+
+} // namespace
+
+std::unique_ptr
+make_asr_workload() {
+ return std::make_unique();
+}
+
+} // namespace nemo_speech::bench
diff --git a/app/bench_diarize.cpp b/app/bench_diarize.cpp
new file mode 100644
index 0000000..a89f8bc
--- /dev/null
+++ b/app/bench_diarize.cpp
@@ -0,0 +1,132 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+
+#include
+#include
+#include
+#include
+#include
+
+#include "audio_file.h"
+#include "batching.h"
+#include "bench.h"
+#include "cli_util.h"
+#include "engine_registry.h"
+#include "model_utils.h"
+
+namespace nemo_speech::bench {
+namespace {
+namespace fs = std::filesystem;
+
+struct AudioInput {
+ fs::path path;
+ audio::AudioFile audio;
+};
+
+class DiarizeWorkload : public Workload {
+ public:
+ std::string task() const override { return "diarize"; }
+ void register_parameters(common::ParameterParser& parser) override {
+ parser.Register("diar", config_);
+ parser.Register("batching", batching_);
+ }
+ bool parse_option(const std::string& arg, const std::function& value) override {
+ if (arg == "--model" || arg == "-m")
+ config_.model_path = value();
+ else if (arg == "--offline")
+ offline_ = true;
+ else if (arg == "--preset")
+ config_.preset = value();
+ else if (arg == "--no-batching")
+ enable_batching_ = false;
+ else
+ return false;
+ return true;
+ }
+ void prepare(const CommonOptions& options) override {
+ if (options.positional.size() != 1)
+ throw std::invalid_argument(
+ options.positional.empty() ? "bench diarize requires a WAV file or directory"
+ : "unexpected argument: " + options.positional[1]);
+ for (const auto& path : collect_files(
+ options.positional.front(), options.recursive,
+ [](const fs::path& file) { return audio::is_wav_path(file.string()); }, "WAV")) {
+ auto audio = audio::load_wav_file(path.string());
+ corpus_seconds_ += static_cast(audio.samples.size()) / audio.sample_rate;
+ inputs_.push_back({path, std::move(audio)});
+ }
+ }
+ size_t input_count() const override { return inputs_.size(); }
+ std::string input_name(size_t index) const override {
+ return inputs_[index].path.filename().string();
+ }
+ void load(const CommonOptions& options, int max_concurrency) override {
+ batching_.enabled = enable_batching_ && max_concurrency > 1;
+ batching_.max_batch_size = std::min(batching_.max_batch_size, max_concurrency);
+ batching_.max_queue_depth = std::max(batching_.max_queue_depth, max_concurrency * 4);
+ batching_.state_arena_slots = std::max(batching_.state_arena_slots, max_concurrency);
+ config_.model_path =
+ resolve_model_file(config_.model_path, "diarization", "diarization model").string();
+ engine_ = engines_.load_diarization(
+ options.gpu, config_.model_path, config_.resolved_geometry(), batching_);
+ }
+ ItemResult run(size_t index) override {
+ const auto& audio = inputs_[index].audio;
+ const auto result = engine_->diarize(
+ audio.samples.data(), audio.samples.size(), audio.sample_rate,
+ offline_ ? asr::DiarizationMode::Offline : asr::DiarizationMode::Streaming);
+ std::string signature;
+ for (const auto& segment : result.segments) {
+ char line[96];
+ std::snprintf(
+ line, sizeof(line), "%.3f %.3f %d;", segment.t0, segment.t1, segment.speaker);
+ signature += line;
+ }
+ return {std::move(signature), {{"segments", static_cast(result.segments.size())}}};
+ }
+
+ void describe(Value& output) const override {
+ output["model"] = fs::path(config_.model_path).stem().string();
+ output["mode"] = mode();
+ output["batching"] = batching_.enabled;
+ output["files"] = static_cast(inputs_.size());
+ output["corpus_audio_seconds"] = corpus_seconds_;
+ }
+ std::vector header_lines() const override {
+ return {"Model: " + fs::path(config_.model_path).filename().string(), "Mode: " + mode()};
+ }
+ std::string mismatch_key() const override { return "segment_mismatches"; }
+ void summarize_run(
+ Value& run, const std::vector& items, double wall_seconds) const override {
+ double audio_seconds = 0.0;
+ for (const auto& item : items)
+ audio_seconds += static_cast(inputs_[item.input].audio.samples.size()) /
+ inputs_[item.input].audio.sample_rate;
+ run["audio_seconds"] = audio_seconds;
+ run["rtfx"] = audio_seconds / wall_seconds;
+ }
+ std::vector run_columns() const override {
+ return {{"RTFx", [](const Value& run) { return run.number_or("rtfx"); }, 2}};
+ }
+
+ private:
+ std::string mode() const { return offline_ ? "offline" : "streaming"; }
+
+ asr::DiarConfig config_;
+ asr::BatchingConfig batching_;
+ bool offline_ = false;
+ bool enable_batching_ = true;
+ std::vector inputs_;
+ double corpus_seconds_ = 0.0;
+ EngineRegistry engines_;
+ std::shared_ptr engine_;
+};
+
+} // namespace
+
+std::unique_ptr
+make_diarize_workload() {
+ return std::make_unique();
+}
+
+} // namespace nemo_speech::bench
diff --git a/app/bench_translate.cpp b/app/bench_translate.cpp
new file mode 100644
index 0000000..5776831
--- /dev/null
+++ b/app/bench_translate.cpp
@@ -0,0 +1,117 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+
+#include
+#include
+#include
+#include
+#include
+
+#include "bench.h"
+#include "cli_util.h"
+#include "engine_registry.h"
+#include "model_utils.h"
+#include "translator.h"
+
+namespace nemo_speech::bench {
+namespace {
+
+class TranslateWorkload : public Workload {
+ public:
+ TranslateWorkload() { config_.backend.gpu = default_gpu_index(); }
+
+ std::string task() const override { return "translate"; }
+ std::vector default_concurrency() const override { return {1}; }
+ void register_parameters(common::ParameterParser& parser) override {
+ parser.Register("nmt", config_);
+ }
+ bool parse_option(const std::string& arg, const std::function& value) override {
+ if (arg == "--model" || arg == "-m")
+ model_ = value();
+ else if (arg == "--from")
+ source_ = value();
+ else if (arg == "--to")
+ target_ = value();
+ else if (arg == "--text")
+ texts_.push_back(value());
+ else
+ return false;
+ return true;
+ }
+ void prepare(const CommonOptions& options) override {
+ if (source_.empty() || target_.empty())
+ throw std::invalid_argument("--from and --to are required");
+ for (const auto& positional : options.positional) {
+ std::istringstream lines(read_text_file(positional));
+ for (std::string line; std::getline(lines, line);)
+ if (!line.empty() && line.find_first_not_of(" \t\r") != std::string::npos)
+ texts_.push_back(std::move(line));
+ }
+ if (texts_.empty())
+ throw std::invalid_argument("bench translate requires --text TEXT or a text file");
+ for (const auto& text : texts_) corpus_bytes_ += text.size();
+ if (options.device_set)
+ config_.backend.gpu = options.gpu;
+ }
+ size_t input_count() const override { return texts_.size(); }
+ std::string input_name(size_t index) const override {
+ return "line" + std::to_string(index + 1);
+ }
+ void load(const CommonOptions&, int max_concurrency) override {
+ config_.pool.contexts = std::max(config_.pool.contexts, max_concurrency);
+ config_.verbose = cli_verbose();
+ config_.model.path =
+ require_model_file(model_.empty() ? config_.model.path : model_, "translation model")
+ .string();
+ translator_ = engines_.load_nmt(config_);
+ }
+ ItemResult run(size_t index) override {
+ const auto result = translator_->translate({texts_[index]}, source_, target_);
+ std::string text = result.empty() ? std::string() : result.front().text;
+ const double bytes = static_cast(text.size());
+ return {std::move(text), {{"output_bytes", bytes}}};
+ }
+
+ void describe(Value& output) const override {
+ output["model"] = translator_->model_name();
+ output["source_language"] = source_;
+ output["target_language"] = target_;
+ output["contexts"] = config_.pool.contexts;
+ output["lines"] = static_cast(texts_.size());
+ output["corpus_input_bytes"] = static_cast(corpus_bytes_);
+ }
+ std::vector header_lines() const override {
+ return {"Model: " + translator_->model_name(), "Pair: " + source_ + "-" + target_};
+ }
+ std::string mismatch_key() const override { return "translation_mismatches"; }
+ void summarize_run(
+ Value& run, const std::vector& items, double wall_seconds) const override {
+ double bytes = 0.0;
+ for (const auto& item : items) bytes += static_cast(texts_[item.input].size());
+ run["input_bytes_per_second"] = bytes / wall_seconds;
+ }
+ std::vector run_columns() const override {
+ return {
+ {"IN B/s", [](const Value& run) { return run.number_or("input_bytes_per_second"); },
+ 1}};
+ }
+
+ private:
+ nmt::TranslatorConfig config_;
+ std::string model_;
+ std::string source_;
+ std::string target_;
+ std::vector texts_;
+ size_t corpus_bytes_ = 0;
+ EngineRegistry engines_;
+ std::shared_ptr translator_;
+};
+
+} // namespace
+
+std::unique_ptr
+make_translate_workload() {
+ return std::make_unique();
+}
+
+} // namespace nemo_speech::bench
diff --git a/app/bench_tts.cpp b/app/bench_tts.cpp
new file mode 100644
index 0000000..8f85e12
--- /dev/null
+++ b/app/bench_tts.cpp
@@ -0,0 +1,300 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+
+#include
+#include
+#include
+#include
+#include
+
+#include "bench.h"
+#include "cli_util.h"
+#include "commands.h"
+#include "config.h"
+#include "engine_registry.h"
+
+namespace nemo_speech::bench {
+namespace {
+namespace fs = std::filesystem;
+
+constexpr int kDefaultSeed = 1;
+
+struct TextInput {
+ std::string name;
+ std::string text;
+};
+
+std::string
+trim(const std::string& value) {
+ const auto begin = value.find_first_not_of(" \t\r\n");
+ if (begin == std::string::npos)
+ return {};
+ return value.substr(begin, value.find_last_not_of(" \t\r\n") - begin + 1);
+}
+
+size_t
+utf8_length(const std::string& text) {
+ size_t count = 0;
+ for (const unsigned char c : text) count += (c & 0xC0) != 0x80;
+ return count;
+}
+
+Column
+metric_column(
+ const std::string& header, const std::string& metric, const char* stat_name, int precision) {
+ return {
+ header, [=](const Value& row) { return stat(row, "metrics", metric, stat_name); },
+ precision};
+}
+
+class TtsWorkload : public Workload {
+ public:
+ std::string task() const override { return "tts"; }
+ std::vector default_concurrency() const override { return {1}; }
+ void register_parameters(common::ParameterParser& parser) override {
+ parser.Register("tts", parsed_);
+ }
+ bool parse_option(const std::string& arg, const std::function& value) override {
+ if (arg == "--text") {
+ texts_.push_back(value());
+ } else if (arg == "--text-file") {
+ text_files_.push_back(value());
+ } else if (arg == "--magpie-model" || arg == "--model" || arg == "-m") {
+ parsed_.runtime.magpie_model = value();
+ } else if (arg == "--codec-model") {
+ parsed_.runtime.codec_model = value();
+ } else if (arg == "--tokenizer-dir") {
+ parsed_.tokenizer_model_dir = value();
+ } else if (arg == "--tn-model-dir") {
+ parsed_.tn_model_dir = value();
+ } else if (arg == "--language") {
+ language_ = value();
+ } else if (arg == "--voice") {
+ voice_ = value();
+ } else if (arg == "--speaker") {
+ options_.speaker = parse_int(value(), arg, 0, 100000);
+ } else if (arg == "--sample-rate") {
+ sample_rate_ = parse_int(value(), arg, 8000, 192000);
+ } else if (arg == "--seed") {
+ options_.seed = parse_int(value(), arg, -1, 2147483647);
+ seed_set_ = true;
+ } else if (arg == "--steps") {
+ options_.steps = parse_int(value(), arg, 1, 1000000);
+ } else if (arg == "--top-k") {
+ options_.top_k = parse_int(value(), arg, 1, 1000000);
+ } else if (arg == "--temperature") {
+ options_.temperature = static_cast(parse_double(value(), arg));
+ options_.override_temperature = true;
+ } else if (arg == "--cfg-scale") {
+ options_.cfg_scale = static_cast(parse_double(value(), arg));
+ options_.override_cfg_scale = true;
+ } else {
+ return false;
+ }
+ return true;
+ }
+ void prepare(const CommonOptions& options) override {
+ for (size_t i = 0; i < texts_.size(); ++i) {
+ if (trim(texts_[i]).empty())
+ throw std::invalid_argument("--text must not be empty");
+ inputs_.push_back({"text" + std::to_string(i + 1), trim(texts_[i])});
+ }
+ // one utterance per line; `audio|text` filelist lines use the text
+ for (const auto& file : text_files_) {
+ std::istringstream lines(read_text_file(file));
+ std::string line;
+ int number = 0;
+ while (std::getline(lines, line)) {
+ ++number;
+ std::string name = fs::path(file).stem().string() + ":" + std::to_string(number);
+ const auto bar = line.rfind('|');
+ if (bar != std::string::npos) {
+ const auto audio = fs::path(trim(line.substr(0, line.find('|'))));
+ if (!audio.empty())
+ name = audio.stem().string();
+ line = line.substr(bar + 1);
+ }
+ line = trim(line);
+ if (!line.empty())
+ inputs_.push_back({name, line});
+ }
+ }
+ for (const auto& positional : options.positional)
+ for (const auto& path : collect_files(
+ positional, options.recursive,
+ [](const fs::path& file) { return file.extension() == ".txt"; }, ".txt")) {
+ auto text = trim(read_text_file(path));
+ if (text.empty())
+ throw std::invalid_argument(path.string() + " is empty");
+ inputs_.push_back({path.stem().string(), std::move(text)});
+ }
+ if (inputs_.empty())
+ throw std::invalid_argument(
+ "bench tts requires --text TEXT, --text-file FILE, or a .txt file or directory");
+ if (!seed_set_ && parsed_.runtime.seed < 0)
+ options_.seed = kDefaultSeed;
+ seed_ = options_.seed >= 0 ? options_.seed : parsed_.runtime.seed;
+ }
+ size_t input_count() const override { return inputs_.size(); }
+ std::string input_name(size_t index) const override { return inputs_[index].name; }
+ void load(const CommonOptions& options, int) override {
+ auto config = make_synthesizer_config(parsed_, options.device, options.device_set);
+ magpie_model_ = config.runtime.magpie_model;
+ codec_model_ = config.runtime.codec_model;
+ tokenizer_dir_ = config.tokenizer_model_dir;
+ synthesizer_ = engines_.load_tts(std::move(config));
+ }
+ void warmup_engine() override { engines_.warmup(); }
+ ItemResult run(size_t index) override {
+ tts::SynthesisRequest request;
+ request.text = inputs_[index].text;
+ request.language_code = language_;
+ request.voice_name = voice_;
+ request.output_sample_rate = sample_rate_;
+ request.options = options_;
+ // client-side: submission to first audio chunk, and gaps between chunks
+ using Clock = std::chrono::steady_clock;
+ const auto submitted = Clock::now();
+ Clock::time_point last{};
+ double first_audio_ms = 0.0;
+ std::vector chunk_gaps_ms;
+ const auto result =
+ synthesizer_->synthesize(request, [&](const auto&, const std::string& chunk) {
+ if (chunk.empty())
+ return true;
+ const auto now = Clock::now();
+ if (last == Clock::time_point{})
+ first_audio_ms =
+ std::chrono::duration(now - submitted).count();
+ else
+ chunk_gaps_ms.push_back(
+ std::chrono::duration(now - last).count());
+ last = now;
+ return true;
+ });
+ if (result.output_samples == 0)
+ throw std::runtime_error("synthesizer returned no audio for " + inputs_[index].name);
+ const auto& stats = result.stats;
+ return {
+ std::to_string(result.output_samples),
+ {{"first_audio_ms", first_audio_ms},
+ {"audio_seconds", stats.audio_s},
+ {"e2e_ttfa_ms", stats.e2e_ttfa_ms},
+ {"e2e_rtfx", stats.e2e_rtfx},
+ {"decoder_rtfx", stats.decoder_rtfx},
+ {"decoder_itl_avg_ms", stats.decoder_itl_avg_ms},
+ {"decoder_itl_p95_ms", stats.decoder_itl_p95_ms},
+ {"codec_rtfx", stats.codec_rtfx},
+ {"encoder_ms", stats.encoder_ms},
+ {"decoder_frames", static_cast(stats.generated_frames)}},
+ {{"chunk_gap_ms", std::move(chunk_gaps_ms)}}};
+ }
+
+ void describe(Value& output) const override {
+ output["model"] = synthesizer_->model_name();
+ Value models(Value::Object{});
+ models["magpie"] = magpie_model_;
+ models["codec"] = codec_model_;
+ models["tokenizer"] = tokenizer_dir_;
+ output["models"] = std::move(models);
+ output["language"] = language();
+ output["voice"] = voice();
+ output["sample_rate"] = sample_rate_ > 0 ? sample_rate_ : synthesizer_->sample_rate();
+ output["seed"] = seed_;
+ if (options_.steps > 0)
+ output["steps"] = options_.steps;
+ if (options_.top_k > 0)
+ output["top_k"] = options_.top_k;
+ if (options_.override_temperature)
+ output["temperature"] = static_cast(options_.temperature);
+ if (options_.override_cfg_scale)
+ output["cfg_scale"] = static_cast(options_.cfg_scale);
+ output["utterances"] = static_cast(inputs_.size());
+ size_t chars = 0;
+ for (const auto& input : inputs_) chars += utf8_length(input.text);
+ output["corpus_chars"] = static_cast(chars);
+ }
+ std::vector header_lines() const override {
+ return {
+ "Model: " + synthesizer_->model_name() + " (" +
+ fs::path(magpie_model_).filename().string() + ", codec " +
+ fs::path(codec_model_).filename().string() + ")",
+ "Language: " + language() + " Voice: " + voice() + " Seed: " + std::to_string(seed_)};
+ }
+ std::string mismatch_key() const override { return "audio_length_mismatches"; }
+ void summarize_run(
+ Value& run, const std::vector& items, double wall_seconds) const override {
+ double audio_seconds = 0.0;
+ for (const auto& item : items)
+ for (const auto& [name, value] : item.metrics)
+ if (name == "audio_seconds")
+ audio_seconds += value;
+ run["audio_seconds"] = audio_seconds;
+ run["rtfx"] = audio_seconds / wall_seconds;
+ }
+ std::vector run_columns() const override {
+ // Riva/NIM TTS streaming metrics
+ return {
+ {"RTFx", [](const Value& run) { return run.number_or("rtfx"); }, 2},
+ metric_column("TTFA avg", "first_audio_ms", "mean", 1),
+ metric_column("TTFA p99", "first_audio_ms", "p99", 1),
+ metric_column("ICL avg", "chunk_gap_ms", "mean", 2),
+ metric_column("ICL p99", "chunk_gap_ms", "p99", 2)};
+ }
+ std::vector input_columns() const override {
+ return {
+ {"CHARS", [](const Value& row) { return row.number_or("chars"); }, 0},
+ metric_column("AUDIO (s)", "audio_seconds", "mean", 2),
+ metric_column("TTFA avg", "first_audio_ms", "mean", 1),
+ metric_column("TTFA p99", "first_audio_ms", "p99", 1),
+ metric_column("RTFx", "e2e_rtfx", "mean", 2),
+ metric_column("DEC RTFx", "decoder_rtfx", "mean", 2),
+ metric_column("ITL avg", "decoder_itl_avg_ms", "mean", 2),
+ metric_column("ITL p95", "decoder_itl_p95_ms", "mean", 2),
+ metric_column("CODEC RTFx", "codec_rtfx", "mean", 1),
+ metric_column("ENC (ms)", "encoder_ms", "mean", 1),
+ {"WALL (ms)", [](const Value& row) { return stat(row, "latency_ms", "", "mean"); }, 1}};
+ }
+ void describe_input(size_t index, Value& entry) const override {
+ entry["chars"] = static_cast(utf8_length(inputs_[index].text));
+ entry["text"] = inputs_[index].text;
+ }
+
+ private:
+ std::string language() const {
+ return language_.empty() ? synthesizer_->default_language_code() : language_;
+ }
+ std::string voice() const {
+ if (!voice_.empty())
+ return voice_;
+ const int speaker =
+ options_.speaker >= 0 ? options_.speaker : synthesizer_->default_speaker();
+ const auto& names = synthesizer_->speaker_names();
+ return speaker < static_cast(names.size()) ? names[speaker] : std::to_string(speaker);
+ }
+
+ tts::MagpieTtsServerConfig parsed_;
+ tts::MagpieSynthesisOptions options_;
+ std::vector texts_;
+ std::vector text_files_;
+ std::string language_;
+ std::string voice_;
+ int sample_rate_ = 0;
+ bool seed_set_ = false;
+ int seed_ = kDefaultSeed;
+ std::vector inputs_;
+ std::string magpie_model_;
+ std::string codec_model_;
+ std::string tokenizer_dir_;
+ EngineRegistry engines_;
+ std::shared_ptr synthesizer_;
+};
+
+} // namespace
+
+std::unique_ptr
+make_tts_workload() {
+ return std::make_unique();
+}
+
+} // namespace nemo_speech::bench
diff --git a/app/commands.h b/app/commands.h
index 4f3f3d3..9fd7ce8 100644
--- a/app/commands.h
+++ b/app/commands.h
@@ -2,6 +2,8 @@
// SPDX-License-Identifier: Apache-2.0
#pragma once
+#include
+
int command_transcribe(int argc, char** argv);
int command_diarize(int argc, char** argv);
int command_translate(int argc, char** argv);
@@ -22,3 +24,15 @@ void print_model_help(const char* program);
void print_doctor_help(const char* program);
void print_health_help(const char* program);
void print_serve_help(const char* program);
+
+#if defined(NEMO_SPEECH_CLI_TTS)
+namespace nemo_speech::tts {
+struct MagpieTtsServerConfig;
+struct SynthesizerConfig;
+} // namespace nemo_speech::tts
+
+// Resolve model references and apply a --device choice to a parsed TTS config.
+nemo_speech::tts::SynthesizerConfig make_synthesizer_config(
+ nemo_speech::tts::MagpieTtsServerConfig parsed, const std::string& device_name,
+ bool device_set);
+#endif
diff --git a/app/main.cpp b/app/main.cpp
index 51695b9..bec4db1 100644
--- a/app/main.cpp
+++ b/app/main.cpp
@@ -28,7 +28,7 @@ std::string
unavailable_command(const std::string& command) {
(void)command;
#if !defined(NEMO_SPEECH_CLI_ASR)
- if (command == "transcribe" || command == "bench")
+ if (command == "transcribe")
return "this build does not include ASR; rebuild with -DNEMO_SPEECH_BUILD_ASR=ON";
#endif
#if !defined(NEMO_SPEECH_CLI_DIAR)
@@ -93,9 +93,7 @@ print_help(const char* program) {
#if defined(NEMO_SPEECH_CLI_TTS)
" synthesize Synthesize speech to a WAV file\n"
#endif
-#if defined(NEMO_SPEECH_CLI_ASR)
- " bench Benchmark an end-to-end ASR workload\n"
-#endif
+ " bench Benchmark an end-to-end workload (asr, tts, diarize, ...)\n"
" pull Download a pinned model from Hugging Face\n"
" model List, pull, or inspect models\n"
" doctor Inspect runtime and device availability\n"
@@ -183,12 +181,10 @@ main(int argc, char** argv) {
return 0;
}
#endif
-#if defined(NEMO_SPEECH_CLI_ASR)
if (std::strcmp(argv[2], "bench") == 0) {
print_bench_help(argv[0]);
return 0;
}
-#endif
if (std::strcmp(argv[2], "pull") == 0) {
std::printf("Usage: %s pull REPO\n", argv[0]);
} else if (std::strcmp(argv[2], "model") == 0)
@@ -232,10 +228,8 @@ main(int argc, char** argv) {
return run_session(
"synthesize", argc, argv, [&] { return command_synthesize(argc - 2, argv + 2); });
#endif
-#if defined(NEMO_SPEECH_CLI_ASR)
if (std::strcmp(argv[1], "bench") == 0)
return run_session("bench", argc, argv, [&] { return command_bench(argc - 2, argv + 2); });
-#endif
if (std::strcmp(argv[1], "model") == 0)
return command_model(argc - 2, argv + 2);
if (std::strcmp(argv[1], "pull") == 0)
diff --git a/app/synthesize.cpp b/app/synthesize.cpp
index c104d74..b904cdd 100644
--- a/app/synthesize.cpp
+++ b/app/synthesize.cpp
@@ -6,6 +6,7 @@
#include
#include
#include
+#include
#include "commands.h"
#if defined(_WIN32)
@@ -57,6 +58,53 @@ write_audio(const std::filesystem::path& path, const std::string& audio, bool fo
} // namespace
+nemo_speech::tts::SynthesizerConfig
+make_synthesizer_config(
+ nemo_speech::tts::MagpieTtsServerConfig parsed, const std::string& device_name,
+ bool device_set) {
+ parsed.runtime.magpie_model =
+ resolve_model_file(parsed.runtime.magpie_model, "tts", "MagpieTTS model").string();
+ parsed.runtime.codec_model =
+ resolve_model_file(parsed.runtime.codec_model, "codec", "NanoCodec model").string();
+ parsed.tokenizer_model_dir =
+ resolve_model_directory(parsed.tokenizer_model_dir, "tokenizer", "tokenizer model")
+ .string();
+ if (!parsed.tn_model_dir.empty())
+ parsed.tn_model_dir =
+ require_model_directory(parsed.tn_model_dir, "text normalization model").string();
+ if (device_set) {
+ const int gpu = parse_device(device_name, "--device");
+ bool cuda_device = device_name == "cuda" || device_name.rfind("cuda:", 0) == 0;
+#if defined(NEMO_SPEECH_CLI_CUDA)
+ cuda_device = cuda_device || device_name == "auto" || device_name == "gpu" ||
+ device_name.rfind("gpu:", 0) == 0;
+#endif
+ if (gpu < 0) {
+ parsed.runtime.lt_backend = nemo_speech::tts::MagpieBackendPreference::Cpu;
+ parsed.runtime.sampling_backend = nemo_speech::tts::MagpieBackendPreference::Cpu;
+ parsed.runtime.magpie_cpu = true;
+ parsed.runtime.codec_cpu = true;
+ } else if (cuda_device) {
+ parsed.runtime.lt_backend = nemo_speech::tts::MagpieBackendPreference::Cuda;
+ parsed.runtime.codec_cpu = false;
+ } else {
+ parsed.runtime.lt_backend = nemo_speech::tts::MagpieBackendPreference::Cpu;
+ parsed.runtime.sampling_backend = nemo_speech::tts::MagpieBackendPreference::Cpu;
+ parsed.runtime.codec_cpu = false;
+ }
+ }
+ parsed.runtime.verbose = cli_verbose();
+
+ nemo_speech::tts::SynthesizerConfig config;
+ config.runtime = parsed.runtime;
+ config.tokenizer_model_dir = parsed.tokenizer_model_dir;
+ config.text_normalizer_model_dir = parsed.tn_model_dir;
+ config.tokenizer = parsed.tokenizer_config;
+ config.default_language_code = parsed.default_language_code;
+ config.default_voice_name = parsed.default_voice_name;
+ return config;
+}
+
void
print_synthesize_help(const char* program) {
std::printf(
@@ -115,7 +163,6 @@ command_synthesize(int argc, char** argv) {
int output_rate = 0;
bool device_set = false;
std::string device_name = "auto";
- int gpu = default_gpu_index();
nemo_speech::tts::MagpieSynthesisOptions request_options;
auto value = [&](int& i, const std::string& option) {
if (++i >= argc)
@@ -150,7 +197,7 @@ command_synthesize(int argc, char** argv) {
output_rate = parse_int(value(i, arg), arg, 8000, 192000);
else if (arg == "--device" || arg == "--backend") {
device_name = value(i, arg);
- gpu = parse_device(device_name, arg);
+ (void)parse_device(device_name, arg);
device_set = true;
} else if (arg == "--seed")
request_options.seed = parse_int(value(i, arg), arg, -1, 2147483647);
@@ -191,49 +238,11 @@ command_synthesize(int argc, char** argv) {
if (cli_json() && output_path == "-")
throw std::invalid_argument("--json cannot be combined with --output -");
- parsed.runtime.magpie_model =
- resolve_model_file(parsed.runtime.magpie_model, "tts", "MagpieTTS model").string();
- parsed.runtime.codec_model =
- resolve_model_file(parsed.runtime.codec_model, "codec", "NanoCodec model").string();
- parsed.tokenizer_model_dir =
- resolve_model_directory(parsed.tokenizer_model_dir, "tokenizer", "tokenizer model")
- .string();
- if (!parsed.tn_model_dir.empty())
- parsed.tn_model_dir =
- require_model_directory(parsed.tn_model_dir, "text normalization model").string();
- if (device_set) {
- bool cuda_device = device_name == "cuda" || device_name.rfind("cuda:", 0) == 0;
-#if defined(NEMO_SPEECH_CLI_CUDA)
- cuda_device = cuda_device || device_name == "auto" || device_name == "gpu" ||
- device_name.rfind("gpu:", 0) == 0;
-#endif
- if (gpu < 0) {
- parsed.runtime.lt_backend = nemo_speech::tts::MagpieBackendPreference::Cpu;
- parsed.runtime.sampling_backend = nemo_speech::tts::MagpieBackendPreference::Cpu;
- parsed.runtime.magpie_cpu = true;
- parsed.runtime.codec_cpu = true;
- } else if (cuda_device) {
- parsed.runtime.lt_backend = nemo_speech::tts::MagpieBackendPreference::Cuda;
- parsed.runtime.codec_cpu = false;
- } else {
- parsed.runtime.lt_backend = nemo_speech::tts::MagpieBackendPreference::Cpu;
- parsed.runtime.sampling_backend = nemo_speech::tts::MagpieBackendPreference::Cpu;
- parsed.runtime.codec_cpu = false;
- }
- }
- parsed.runtime.verbose = cli_verbose();
-
- nemo_speech::tts::SynthesizerConfig config;
- config.runtime = parsed.runtime;
- config.tokenizer_model_dir = parsed.tokenizer_model_dir;
- config.text_normalizer_model_dir = parsed.tn_model_dir;
- config.tokenizer = parsed.tokenizer_config;
- config.default_language_code = parsed.default_language_code;
- config.default_voice_name = parsed.default_voice_name;
+ auto config = make_synthesizer_config(std::move(parsed), device_name, device_set);
nemo_speech::EngineRegistry engines;
auto synthesizer = engines.load_tts(std::move(config));
if (warmup)
- synthesizer->warmup("Hello", 1);
+ synthesizer->warmup("Hello", 8);
nemo_speech::tts::SynthesisRequest request;
request.text = text;
diff --git a/conversion/registry.py b/conversion/registry.py
index 059dd69..84b6597 100644
--- a/conversion/registry.py
+++ b/conversion/registry.py
@@ -114,7 +114,7 @@ def _normalized_outtype(architecture: str, outtype: str) -> str:
"diarization": "f32",
"pnc": "q8_0",
"vad": "f32",
- "tts": "f16",
+ "tts": "q8_0",
"codec": "f16",
"nmt": "f16",
"s2s": "q4_k_m",
@@ -141,7 +141,7 @@ def _normalized_outtype(architecture: str, outtype: str) -> str:
},
"pnc": {"f16", "bf16", "q8_0"},
"vad": {"f32"},
- "tts": {"f16", "f32"},
+ "tts": {"f16", "f32", "q8_0"},
"codec": {"f16", "f32"},
"nmt": {"f32", "f16", "bf16", "q8_0", "auto"},
"s2s": {"bf16", "q4_k_m", "nvfp4"},
diff --git a/conversion/tts.py b/conversion/tts.py
index b58f859..cba60d0 100644
--- a/conversion/tts.py
+++ b/conversion/tts.py
@@ -15,6 +15,7 @@
from __future__ import annotations
import json
+import re
import tempfile
from pathlib import Path
from typing import Any
@@ -22,6 +23,7 @@
import gguf
import numpy as np
import torch
+from gguf.quants import quantize
from .source import extract_archive, find_checkpoint_files, load_state_dict, read_checkpoint_config
from .tts_tokenizer_profiles import tokenizer_profile
@@ -29,6 +31,21 @@
SPECIAL_AUDIO_TOKENS = 8
SPEAKER_NAMES = ["John", "Sofia", "Aria", "Jason", "Leo"]
+# 2-D projection weights stored as Q8_0 by ``--outtype q8_0``: the attention and feed-forward
+# projections (per-tap Conv1d slices) of the text encoder, the decoder and the local
+# transformer, the local-transformer output projections and the final projection. Norms and
+# biases stay f32; all other tensors keep the f16 layout. The CUDA runtime's fused decoder and
+# local-transformer kernels need these Q8_0 projections; other backends run them through
+# ggml's Q8_0 matrix multiplication.
+Q8_PROJECTION_WEIGHT = re.compile(
+ r"(self_attention\.(qkv_net|o_net)\.weight"
+ r"|cross_attention\.(q_net|kv_net|o_net)\.weight"
+ r"|pos_ff\.(proj|o_net)\.conv\.weight\.k\d+"
+ r"|local_transformer_out_projections\.\d+\.weight"
+ r"|final_proj\.weight)$"
+)
+Q8_BLOCK = 32
+
def extract_nemo(path: Path) -> tuple[Path, tempfile.TemporaryDirectory[str] | None]:
if path.is_dir():
@@ -276,10 +293,19 @@ def should_store_f32(name: str, tensor: torch.Tensor) -> bool:
return name.endswith(".bias")
+def is_q8_projection(name: str, tensor: torch.Tensor) -> bool:
+ return (
+ tensor.is_floating_point()
+ and tensor.ndim == 2
+ and int(tensor.shape[-1]) % Q8_BLOCK == 0
+ and Q8_PROJECTION_WEIGHT.search(name) is not None
+ )
+
+
def tensor_to_numpy(name: str, tensor: torch.Tensor, outtype: str) -> np.ndarray:
tensor = tensor.detach().cpu().contiguous()
if tensor.is_floating_point():
- if outtype == "f16" and not should_store_f32(name, tensor):
+ if outtype in ("f16", "q8_0") and not should_store_f32(name, tensor):
tensor = tensor.to(torch.float16)
else:
tensor = tensor.to(torch.float32)
@@ -289,6 +315,11 @@ def tensor_to_numpy(name: str, tensor: torch.Tensor, outtype: str) -> np.ndarray
def add_tensor(writer: gguf.GGUFWriter, name: str, tensor: torch.Tensor, outtype: str) -> None:
+ if outtype == "q8_0" and is_q8_projection(name, tensor):
+ f32 = tensor.detach().cpu().contiguous().to(torch.float32).numpy()
+ q8 = quantize(f32, gguf.GGMLQuantizationType.Q8_0)
+ writer.add_tensor(name, q8, raw_dtype=gguf.GGMLQuantizationType.Q8_0)
+ return
writer.add_tensor(name, tensor_to_numpy(name, tensor, outtype))
@@ -378,6 +409,7 @@ def convert(
summary["tensors_written"] = n_written
summary["tensors_skipped"] = skipped
summary["output"] = str(output)
+ summary["outtype"] = outtype
summary["local_transformer_outtype"] = local_transformer_outtype or outtype
if metadata_json:
metadata_json.parent.mkdir(parents=True, exist_ok=True)
diff --git a/docs/cli.md b/docs/cli.md
index 5bbe6bb..e25a878 100644
--- a/docs/cli.md
+++ b/docs/cli.md
@@ -238,11 +238,41 @@ configuration file](server.md#engine-and-listener-configuration).
## Benchmark
-Benchmark end-to-end ASR concurrency with one shared recognizer:
+`nemo-speech bench ` loads one engine, warms it up, and runs every input
+`--repetitions` times at each `--concurrency` level. The tasks compiled into
+the build are listed by `nemo-speech help bench`; a task missing from the build
+fails with `unsupported_feature`.
```bash
-nemo-speech bench asr recordings/ \
- --model asr.q8_0.gguf \
- --concurrency 1,2,4 \
- --json
+# ASR over a WAV file or directory
+nemo-speech bench asr recordings/ --model asr.q8_0.gguf --concurrency 1,2,4 --json
+
+# TTS over --text, a .txt file or directory (one utterance per file), or
+# --text-file (one utterance per line)
+nemo-speech bench tts --text-file test_files/tts/ljs_audio_text_test_filelist_small.txt \
+ --magpie-model magpie.q8_0.gguf --codec-model codec.gguf --tokenizer-dir tokenizer/ \
+ --per-stream 20 --json
+
+# Diarization over a WAV file or directory
+nemo-speech bench diarize meetings/ --model sortformer.gguf
+
+# Translation over a text file, one input per line (NMT builds)
+nemo-speech bench translate lines.txt --model nmt.gguf --from en --to de
```
+
+All tasks share `-n/--repetitions` (or `--per-stream N` requests per
+concurrent stream), `--warmup`, `-c/--concurrency`,
+`--device`, `--config`, the task's `--asr.*`/`--tts.*`/`--diar.*`/`--nmt.*`
+overrides, and `--json`. The report contains `command`, `task`, the model,
+`load_ms`, `warmup_ms`, and one `runs` entry per concurrency level with
+`concurrency`, `items`, `wall_seconds`, `items_per_second`, client-side
+`latency_ms` (`mean`/`p50`/`p90`/`p95`/`p99`/`min`/`max`), per-item `metrics`, task
+throughput such as `rtfx`, and an output-mismatch count against the first
+result observed for each input (`transcript_mismatches`,
+`audio_length_mismatches`, `segment_mismatches`, `translation_mismatches`).
+
+`bench tts` reports time to first audio (TTFA) and inter-chunk latency (ICL),
+measured at the client (`first_audio_ms`, `chunk_gap_ms`), and throughput RTFx,
+plus a per-input breakdown. It fixes `--seed`
+to 1 so repeated runs generate identical audio. The MagpieTTS runtime serializes
+requests, so TTS concurrency above 1 measures queueing.
diff --git a/docs/install.md b/docs/install.md
index 79596b6..c29d4ee 100644
--- a/docs/install.md
+++ b/docs/install.md
@@ -1,12 +1,18 @@
# Install NeMo-Speech.cpp
-The installer selects a backend-matched native release containing the ASR,
-diarization, translation, and TTS CLI, HTTP API, realtime WebSocket endpoint,
-browser playground, SDK, and notices. It builds from source when a matching
-archive is unavailable. Models are distributed separately; inference commands
-download missing indexed defaults on first use, while the server downloads a
-model only when explicitly enabled with an indexed name. See [models and
-cache](cli.md#models-and-cache).
+The installer downloads a prebuilt release archive for your backend containing
+the ASR, diarization, translation, and TTS CLI, HTTP API, realtime WebSocket
+endpoint, browser playground, SDK, and notices. It builds from source when a
+matching archive is unavailable. Models are distributed separately; inference
+commands download missing indexed defaults on first use, while the server
+downloads a model only when explicitly enabled with an indexed name. See
+[models and cache](cli.md#models-and-cache).
+
+> [!IMPORTANT]
+> **For the best performance and the latest features, build natively from source** with
+> `--source` (`-Source` on Windows) or by following the [source-build guide](build.md). A native
+> build is compiled for your machine's CPU and GPU, and release tags can be out of sync with the
+> `main` branch.
## Linux and macOS
@@ -22,7 +28,7 @@ nemo-speech --version
With no version argument, the installer reads the current release identifier
from the repository's `VERSION` file, including prerelease identifiers.
-Native Linux archives require glibc 2.31 or newer (Ubuntu 20.04 or equivalent).
+Prebuilt Linux archives require glibc 2.31 or newer (Ubuntu 20.04 or equivalent).
Prebuilt CPU archives require no GPU toolkit. Linux CUDA archives include the
required user-space CUDA libraries but still need a compatible NVIDIA driver.
@@ -48,8 +54,10 @@ available.
It installs without `sudo` and links the CLI into `~/.local/bin`. Run `--help`
to see prefix, backend, PATH, and dry-run options. Downloaded archives are
-verified against their published SHA-256 files; a present archive with an
-invalid or mismatched checksum always fails rather than falling back to source.
+checked against the SHA-256 file published with the release, which detects a
+corrupted or incomplete download but is not a signature. A present archive with
+an invalid or mismatched checksum always fails rather than falling back to
+source.
The source fallback requires Git, CMake 3.26 or newer, Ninja, a C++17 compiler,
SentencePiece development files, and any toolkit required by the selected
diff --git a/docs/tts/models.md b/docs/tts/models.md
index fbe7ca2..36daab7 100644
--- a/docs/tts/models.md
+++ b/docs/tts/models.md
@@ -26,6 +26,14 @@ frame-stacking factor of 2. This is independent of `tts.chunk-frames`, which
groups generated frames for NanoCodec streaming. Both versions use the same
NanoCodec decoder.
+The fused CUDA decode path needs a Q8_0 GGUF. Convert the `.nemo` locally
+(Q8_0 is the converter default):
+
+```bash
+python3 convert_model.py magpie_tts_multilingual_357m.nemo \
+ --outfile magpie_tts_multilingual_357m.v2607.q8_0.gguf
+```
+
**Tokenizer.** MagpieTTS's tokenizer assets live *inside* the `.nemo` archive -
they are not part of the GGUF. The built-in pull extracts only the required,
pinned tokenizer members and verifies each one. For a custom Magpie checkpoint,
@@ -80,13 +88,27 @@ nemo-speech synthesize "Hello from Magpie Multilingual." --output output.wav
The unified [`convert_model.py`](../../convert_model.py) entry point accepts
compatible local `.nemo` archives and extracted NeMo checkpoints. It defaults
-to `--outtype f16` for MagpieTTS and NanoCodec; pass `--outtype f32` to retain
-full precision. The converter is a source-tree Python tool and is not included
+to `--outtype q8_0` for MagpieTTS and `--outtype f16` for NanoCodec; pass
+`--outtype f16` or `--outtype f32` to keep MagpieTTS unquantized. The `q8_0`
+output stores the attention and feed-forward projections of the text encoder,
+decoder and local transformer, the local-transformer output projections and the
+final projection as Q8_0, keeps norms and biases f32, and everything else f16.
+In an ASR round trip on the 10 LJSpeech sentences of
+[`ljs_audio_text_test_filelist_small.txt`](../../test_files/tts/ljs_audio_text_test_filelist_small.txt)
+(three seeds each, transcribed with Nemotron Speech Streaming 0.6B), MagpieTTS v2607 Q8_0 scores
+5.6% WER and 2.2% CER against 4.7% and 2.1% for f16. On CUDA GPUs
+with compute capability 8.0 or newer and at least 48 SMs it also enables the
+fused decoder kernel and, when classifier-free guidance is on (the default;
+`--tts.no-cfg` turns it off), the fused local-transformer kernel. Both are
+selected automatically.
+On Hopper and newer, build with native code for the GPU (for example
+`-DCMAKE_CUDA_ARCHITECTURES=native`). The converter is a source-tree
+Python tool and is not included
in native release archives; see [Model conversion](../model-conversion.md) for
environment setup.
```bash
-python3 convert_model.py custom-magpie.nemo --outfile custom-magpie.f16.gguf
+python3 convert_model.py custom-magpie.nemo --outfile custom-magpie.q8_0.gguf
```
Conversion does not require `nemo_toolkit`. The optional
diff --git a/ggml-patches/0014-cuda-fused-attention-extensions.patch b/ggml-patches/0014-cuda-fused-attention-extensions.patch
index 4d9de3f..9b89f3f 100644
--- a/ggml-patches/0014-cuda-fused-attention-extensions.patch
+++ b/ggml-patches/0014-cuda-fused-attention-extensions.patch
@@ -116,15 +116,16 @@ index 7429cc45..0c570639 100644
case GGML_OP_SET_ROWS:
diff --git a/src/ggml-cuda/fused-attention.cu b/src/ggml-cuda/fused-attention.cu
new file mode 100644
-index 00000000..0a1c1766
+index 00000000..84e1c025
--- /dev/null
+++ b/src/ggml-cuda/fused-attention.cu
-@@ -0,0 +1,1650 @@
+@@ -0,0 +1,1720 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+#include "fused-attention.cuh"
+
+#include
++#include
+#include
+
+// Fused multi-head attention with optional relative-position terms and a
@@ -188,14 +189,6 @@ index 00000000..0a1c1766
+ __low2float(values[1]), __high2float(values[1]));
+}
+
-+static __device__ __forceinline__ float2 attention_load2(const float * ptr) {
-+ return *reinterpret_cast(ptr);
-+}
-+
-+static __device__ __forceinline__ float2 attention_load2(const half * ptr) {
-+ return __half22float2(*reinterpret_cast(ptr));
-+}
-+
+template
+static __device__ __forceinline__ float4 attention_load_kv4(
+ const T * chunk_head, const float * cache_head, int j, int cache_len,
@@ -230,23 +223,6 @@ index 00000000..0a1c1766
+ return (float) chunk_head[(size_t) j * chunk_sj + d];
+}
+
-+template
-+static __device__ __forceinline__ float2 attention_load_kv2(
-+ const T * chunk_head, const float * cache_head, int j, int cache_len,
-+ int ring_head, long chunk_sj, long cache_sj, int d2) {
-+ if constexpr (Cached) {
-+ if (j < cache_len) {
-+ int physical_j = ring_head + j;
-+ if (physical_j >= cache_len) {
-+ physical_j -= cache_len;
-+ }
-+ return attention_load2(cache_head + (size_t) physical_j * cache_sj + d2);
-+ }
-+ return attention_load2(chunk_head + (size_t) (j - cache_len) * chunk_sj + d2);
-+ }
-+ return attention_load2(chunk_head + (size_t) j * chunk_sj + d2);
-+}
-+
+// Each thread owns one feature and appends only the current chunk to its
+// circular cache. ring_heads[b] is the oldest physical row before this step,
+// so those rows are exactly the ones the new chunk replaces. K and V use
@@ -286,8 +262,15 @@ index 00000000..0a1c1766
+ }
+}
+
++// Split-KV variant of the q=1, d_k=64 cached kernel (flash-decoding style).
++// grid = (n_head, splits, batch). Each block scores a contiguous key range with
++// four keys per warp-iteration (8 lanes x 8 features per key), accumulates a
++// partial (max, sum, context[64]) and the last block to finish for a
++// (batch, head) pair combines the partials. splits == 1 writes the result
++// directly. Partials layout: [batch][head][split][66] floats.
+template
-+static __global__ void fused_cached_attention_q1_d64_kernel(
++static __global__ void __launch_bounds__(256)
++fused_cached_attention_q1_d64_split_kernel(
+ const float * __restrict__ Q, const T * __restrict__ K,
+ const T * __restrict__ V, const float * __restrict__ mask,
+ float * __restrict__ ctx, const float * __restrict__ Kcache,
@@ -297,17 +280,21 @@ index 00000000..0a1c1766
+ int kv, int cache_len, float scale,
+ long q_sh, long q_sb, long k_sj, long k_sh, long k_sb,
+ long v_sj, long v_sh, long v_sb, long cache_sj, long cache_ss,
-+ long o_sh, long o_sb, long m_sb) {
++ long o_sh, long o_sb, long m_sb,
++ float * __restrict__ partials, int * __restrict__ counters, int splits) {
+ constexpr int d_k = 64;
+ constexpr int warps = 8;
+ constexpr int value_parts = 4;
++ constexpr int partial_stride = d_k + 2;
+ extern __shared__ float sh[];
-+ float * query = sh;
-+ float * scores = query + d_k;
-+ float * reduce = scores + kv;
-+ float * value_partials = reduce + warps;
++ float * query = sh; // 64
++ float * reduce = query + d_k; // warps
++ float * value_partials = reduce + warps; // value_parts * 64
++ float * scores = value_partials + value_parts * d_k; // range length
++ __shared__ int s_last;
+
+ const int h = blockIdx.x;
++ const int split = blockIdx.y;
+ const int b = blockIdx.z;
+ const int tid = threadIdx.x;
+ const int warp = tid >> 5;
@@ -318,52 +305,68 @@ index 00000000..0a1c1766
+ const int slot = Cached ? slot_ids[b] : 0;
+ const int ring_head = Cached ? ring_heads[b] : 0;
+ const int key_begin = active_lengths ? cache_len - active_lengths[b] : 0;
-+ const float * Kch = Cached
-+ ? Kcache + (size_t) slot * cache_ss + (size_t) h * d_k
-+ : nullptr;
-+ const float * Vch = Cached
-+ ? Vcache + (size_t) slot * cache_ss + (size_t) h * d_k
-+ : nullptr;
++ const float * Kch = Cached ? Kcache + (size_t) slot * cache_ss + (size_t) h * d_k : nullptr;
++ const float * Vch = Cached ? Vcache + (size_t) slot * cache_ss + (size_t) h * d_k : nullptr;
+ const float * Mb = mask ? mask + (size_t) b * m_sb : nullptr;
+
++ const int n_keys = kv - key_begin;
++ const int per_split = (n_keys + splits - 1) / splits;
++ const int j0 = key_begin + split * per_split;
++ const int j1 = min(kv, j0 + per_split);
++ const int range = max(0, j1 - j0);
++
+ if (tid < d_k) {
+ query[tid] = Qh[tid];
+ }
+ __syncthreads();
+
-+ const int d2 = lane * 2;
-+ const float2 q2 = attention_load2(query + d2);
++ // ---- scores: 4 keys per warp iteration, 8 lanes x 8 features per key ----
++ const int group = lane >> 3; // key within the warp's 4-key group
++ const int sub = lane & 7; // feature block: sub*8 .. sub*8+7
++ const int d8 = sub * 8;
++ const float4 qa = *reinterpret_cast(query + d8);
++ const float4 qb = *reinterpret_cast(query + d8 + 4);
+ float local_max = -INFINITY;
-+ for (int j = key_begin + warp; j < kv; j += warps) {
-+ const float2 k2 = attention_load_kv2(
-+ Kh, Kch, j, cache_len, ring_head, k_sj, cache_sj, d2);
-+ float score = k2.x * q2.x + k2.y * q2.y;
-+ score = attention_warp_sum(score);
-+ if (lane == 0) {
++ for (int base = j0 + warp * 4; base < j1; base += warps * 4) {
++ const int j = base + group;
++ float score = 0.0f;
++ if (j < j1) {
++ const float4 ka = attention_load_kv4(Kh, Kch, j, cache_len, ring_head, k_sj, cache_sj, d8);
++ const float4 kb = attention_load_kv4(Kh, Kch, j, cache_len, ring_head, k_sj, cache_sj, d8 + 4);
++ score = ka.x * qa.x + ka.y * qa.y + ka.z * qa.z + ka.w * qa.w
++ + kb.x * qb.x + kb.y * qb.y + kb.z * qb.z + kb.w * qb.w;
++ }
++ // reduce across the 8 lanes of this key group
++ score += __shfl_xor_sync(0xffffffff, score, 1);
++ score += __shfl_xor_sync(0xffffffff, score, 2);
++ score += __shfl_xor_sync(0xffffffff, score, 4);
++ if (sub == 0 && j < j1) {
+ score = score * scale + (Mb ? Mb[j] : 0.0f);
-+ scores[j] = score;
++ scores[j - j0] = score;
+ local_max = fmaxf(local_max, score);
+ }
+ }
++#pragma unroll
++ for (int offset = 16; offset > 0; offset >>= 1) {
++ local_max = fmaxf(local_max, __shfl_xor_sync(0xffffffff, local_max, offset));
++ }
+ if (lane == 0) {
+ reduce[warp] = local_max;
+ }
+ __syncthreads();
-+ if (tid == 0) {
-+ float maximum = reduce[0];
++ float maximum = reduce[0];
+#pragma unroll
-+ for (int w = 1; w < warps; ++w) {
-+ maximum = fmaxf(maximum, reduce[w]);
-+ }
-+ reduce[0] = maximum;
++ for (int w = 1; w < warps; ++w) {
++ maximum = fmaxf(maximum, reduce[w]);
+ }
+ __syncthreads();
-+ const float maximum = reduce[0];
+
++ // A fully masked split has maximum == -inf: give it zero weight instead of NaN.
++ const bool empty_split = maximum == -INFINITY;
+ float local_sum = 0.0f;
-+ for (int j = key_begin + tid; j < kv; j += blockDim.x) {
-+ const float weight = __expf(scores[j] - maximum);
-+ scores[j] = weight;
++ for (int i = tid; i < range; i += 256) {
++ const float weight = empty_split ? 0.0f : __expf(scores[i] - maximum);
++ scores[i] = weight;
+ local_sum += weight;
+ }
+ local_sum = attention_warp_sum(local_sum);
@@ -371,32 +374,87 @@ index 00000000..0a1c1766
+ reduce[warp] = local_sum;
+ }
+ __syncthreads();
-+ if (tid == 0) {
-+ float sum = reduce[0];
++ float sum = reduce[0];
+#pragma unroll
-+ for (int w = 1; w < warps; ++w) {
-+ sum += reduce[w];
-+ }
-+ reduce[0] = sum;
++ for (int w = 1; w < warps; ++w) {
++ sum += reduce[w];
+ }
-+ __syncthreads();
+
++ // ---- context partial: 4 key-parts x 64 features, 4-way unrolled ----
+ const int d = tid & (d_k - 1);
+ const int part = tid / d_k;
-+ float context = 0.0f;
-+ for (int j = key_begin + part; j < kv; j += value_parts) {
-+ context += scores[j] * attention_load_kv(
-+ Vh, Vch, j, cache_len, ring_head, v_sj, cache_sj, d);
-+ }
-+ value_partials[part * d_k + d] = context;
++ float c0 = 0.0f, c1 = 0.0f, c2 = 0.0f, c3 = 0.0f;
++ int i = part;
++ for (; i + 3 * value_parts < range; i += 4 * value_parts) {
++ c0 += scores[i] * attention_load_kv(Vh, Vch, j0 + i, cache_len, ring_head, v_sj, cache_sj, d);
++ c1 += scores[i + value_parts] * attention_load_kv(Vh, Vch, j0 + i + value_parts, cache_len, ring_head, v_sj, cache_sj, d);
++ c2 += scores[i + 2 * value_parts] * attention_load_kv(Vh, Vch, j0 + i + 2 * value_parts, cache_len, ring_head, v_sj, cache_sj, d);
++ c3 += scores[i + 3 * value_parts] * attention_load_kv(Vh, Vch, j0 + i + 3 * value_parts, cache_len, ring_head, v_sj, cache_sj, d);
++ }
++ for (; i < range; i += value_parts) {
++ c0 += scores[i] * attention_load_kv(Vh, Vch, j0 + i, cache_len, ring_head, v_sj, cache_sj, d);
++ }
++ value_partials[part * d_k + d] = (c0 + c1) + (c2 + c3);
+ __syncthreads();
++
++ float * out = ctx + (size_t) b * o_sb + (size_t) h * o_sh;
++ if (splits == 1) {
++ if (tid < d_k) {
++ float context = value_partials[tid];
++#pragma unroll
++ for (int p = 1; p < value_parts; ++p) {
++ context += value_partials[p * d_k + tid];
++ }
++ out[tid] = context / sum;
++ }
++ return;
++ }
++
++ float * my_partial = partials + (((size_t) b * gridDim.x + h) * splits + split) * partial_stride;
+ if (tid < d_k) {
-+ context = value_partials[tid];
++ float context = value_partials[tid];
+#pragma unroll
-+ for (int part_index = 1; part_index < value_parts; ++part_index) {
-+ context += value_partials[part_index * d_k + tid];
++ for (int p = 1; p < value_parts; ++p) {
++ context += value_partials[p * d_k + tid];
++ }
++ my_partial[tid] = context;
++ }
++ if (tid == 0) {
++ my_partial[d_k] = maximum;
++ my_partial[d_k + 1] = sum;
++ }
++ __threadfence();
++ __syncthreads();
++ if (tid == 0) {
++ const int done = atomicAdd(counters + (size_t) b * gridDim.x + h, 1);
++ s_last = (done == splits - 1) ? 1 : 0;
++ }
++ __syncthreads();
++ if (!s_last) {
++ return;
++ }
++ __threadfence();
++ const float * all = partials + (((size_t) b * gridDim.x + h) * splits) * partial_stride;
++ if (tid < d_k) {
++ float gmax = -INFINITY;
++ for (int s = 0; s < splits; ++s) {
++ gmax = fmaxf(gmax, all[s * partial_stride + d_k]);
++ }
++ float gsum = 0.0f;
++ float gctx = 0.0f;
++ for (int s = 0; s < splits; ++s) {
++ const float m = all[s * partial_stride + d_k];
++ if (m == -INFINITY) {
++ continue; // fully masked split
++ }
++ const float w = __expf(m - gmax);
++ gsum += w * all[s * partial_stride + d_k + 1];
++ gctx += w * all[s * partial_stride + tid];
+ }
-+ ctx[(size_t) b * o_sb + (size_t) h * o_sh + tid] = context / reduce[0];
++ out[tid] = gctx / gsum;
++ }
++ if (tid == 0) {
++ counters[(size_t) b * gridDim.x + h] = 0;
+ }
+}
+
@@ -1541,26 +1599,38 @@ index 00000000..0a1c1766
+ constexpr int threads = 256;
+ constexpr int warps = threads / 32;
+ constexpr int value_partials = 4 * 64;
-+ const size_t specialized_shmem =
-+ ((size_t) 64 + kv_len + warps + value_partials) * sizeof(float);
-+ GGML_ASSERT(specialized_shmem <= max_shmem);
-+ const dim3 specialized_grid(n_head, 1, batch);
++ // Split the key range across blocks (~96 keys per block) so a
++ // (head, batch) pair uses several SMs when the cache is long.
++ int splits = (kv_len + 95) / 96;
++ splits = splits < 1 ? 1 : (splits > 8 ? 8 : splits);
++ const int per_split = (kv_len + splits - 1) / splits;
++ const size_t split_shmem =
++ ((size_t) 64 + warps + value_partials + per_split) * sizeof(float);
++ GGML_ASSERT(split_shmem <= max_shmem);
++ const dim3 split_grid(n_head, splits, batch);
++ ggml_cuda_pool_alloc partials(ctx.pool(), (size_t) batch * n_head * splits * 66);
++ ggml_cuda_pool_alloc counters(ctx.pool(), (size_t) batch * n_head);
++ if (splits > 1) {
++ CUDA_CHECK(cudaMemsetAsync(counters.get(), 0, (size_t) batch * n_head * sizeof(int), stream));
++ }
+ if (k->type == GGML_TYPE_F16) {
-+ fused_cached_attention_q1_d64_kernel
-+ <<>>(
++ fused_cached_attention_q1_d64_split_kernel
++ <<>>(
+ (const float *) q->data, (const half *) k->data, (const half *) v->data,
+ mask ? (const float *) mask->data : nullptr, (float *) dst->data,
+ cache_k, cache_v, active_slots, active_ring_heads, active_lengths,
+ kv_len, cache_len, scale, q_sh, q_sb, k_sj, k_sh, k_sb,
-+ v_sj, v_sh, v_sb, cache_sj, cache_ss, o_sh, o_sb, m_sb);
++ v_sj, v_sh, v_sb, cache_sj, cache_ss, o_sh, o_sb, m_sb,
++ partials.get(), counters.get(), splits);
+ } else {
-+ fused_cached_attention_q1_d64_kernel
-+ <<>>(
++ fused_cached_attention_q1_d64_split_kernel
++ <<>>(
+ (const float *) q->data, (const float *) k->data, (const float *) v->data,
+ mask ? (const float *) mask->data : nullptr, (float *) dst->data,
+ cache_k, cache_v, active_slots, active_ring_heads, active_lengths,
+ kv_len, cache_len, scale, q_sh, q_sb, k_sj, k_sh, k_sb,
-+ v_sj, v_sh, v_sb, cache_sj, cache_ss, o_sh, o_sb, m_sb);
++ v_sj, v_sh, v_sb, cache_sj, cache_ss, o_sh, o_sb, m_sb,
++ partials.get(), counters.get(), splits);
+ }
+ update_cache();
+ return;
@@ -2390,7 +2460,7 @@ index f3c4836a..00000000
- }
-}
diff --git a/src/ggml-cuda/ggml-cuda.cu b/src/ggml-cuda/ggml-cuda.cu
-index 1b7e26e6..b9700eae 100644
+index 6b934a1e..826d4a59 100644
--- a/src/ggml-cuda/ggml-cuda.cu
+++ b/src/ggml-cuda/ggml-cuda.cu
@@ -25,7 +25,7 @@
@@ -2413,7 +2483,7 @@ index 1b7e26e6..b9700eae 100644
break;
case GGML_OP_CROSS_ENTROPY_LOSS:
ggml_cuda_cross_entropy_loss(ctx, dst);
-@@ -5975,7 +5975,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
+@@ -5976,7 +5976,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
#endif // GGML_USE_MUSA
case GGML_OP_FLASH_ATTN_EXT:
return ggml_cuda_flash_attn_ext_supported(dev_ctx->device, op);
diff --git a/ggml-patches/0017-cuda-stream-interop.patch b/ggml-patches/0017-cuda-stream-interop.patch
index c8d7ea3..f321f12 100644
--- a/ggml-patches/0017-cuda-stream-interop.patch
+++ b/ggml-patches/0017-cuda-stream-interop.patch
@@ -1,13 +1,17 @@
diff --git a/include/ggml-cuda.h b/include/ggml-cuda.h
+index 5436c7ef..166ef8c1 100644
--- a/include/ggml-cuda.h
+++ b/include/ggml-cuda.h
-@@ -23,6 +23,17 @@ GGML_BACKEND_API ggml_backend_t ggml_backend_cuda_init(int device);
+@@ -24,6 +24,20 @@ GGML_BACKEND_API ggml_backend_t ggml_backend_cuda_init(int device);
GGML_BACKEND_API bool ggml_backend_is_cuda(ggml_backend_t backend);
+// Returns the backend's current CUDA/HIP/MUSA stream for in-process device interop.
+// The returned stream is owned by the backend and must not be destroyed by the caller.
+GGML_BACKEND_API void * ggml_backend_cuda_get_stream(ggml_backend_t backend);
++// Set the CUDA stream priority used by this backend's streams (0 default; negative = higher,
++// clamped to the device range). Streams already created are recreated.
++GGML_BACKEND_API void ggml_backend_cuda_set_stream_priority(ggml_backend_t backend, int priority);
+
+// Returns the backend-owned native CUDA/HIP graph template for a stable GGML graph after its
+// normal warm-up/capture has completed. The opaque handle is borrowed and is intended for runtime
@@ -19,10 +23,39 @@ diff --git a/include/ggml-cuda.h b/include/ggml-cuda.h
// device buffer
GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_buffer_type(int device);
+diff --git a/src/ggml-cuda/common.cuh b/src/ggml-cuda/common.cuh
+index 610fdf37..cdb17991 100644
+--- a/src/ggml-cuda/common.cuh
++++ b/src/ggml-cuda/common.cuh
+@@ -1479,10 +1479,22 @@ struct ggml_backend_cuda_context {
+
+ ~ggml_backend_cuda_context();
+
++ // CUDA stream priority for streams created by this context (0 = default; more negative =
++ // higher priority, clamped to the device range at creation). Lets a latency-critical
++ // backend (e.g. an autoregressive decoder) win SM scheduling over a concurrent worker.
++ int stream_priority = 0;
++
+ cudaStream_t stream(int device, int stream) {
+ if (streams[device][stream] == nullptr) {
+ ggml_cuda_set_device(device);
+- CUDA_CHECK(cudaStreamCreateWithFlags(&streams[device][stream], cudaStreamNonBlocking));
++ if (stream_priority != 0) {
++ int least = 0, greatest = 0;
++ CUDA_CHECK(cudaDeviceGetStreamPriorityRange(&least, &greatest));
++ const int prio = stream_priority < greatest ? greatest : (stream_priority > least ? least : stream_priority);
++ CUDA_CHECK(cudaStreamCreateWithPriority(&streams[device][stream], cudaStreamNonBlocking, prio));
++ } else {
++ CUDA_CHECK(cudaStreamCreateWithFlags(&streams[device][stream], cudaStreamNonBlocking));
++ }
+ }
+ return streams[device][stream];
+ }
diff --git a/src/ggml-cuda/ggml-cuda.cu b/src/ggml-cuda/ggml-cuda.cu
+index c83735a6..ad0f5e25 100644
--- a/src/ggml-cuda/ggml-cuda.cu
+++ b/src/ggml-cuda/ggml-cuda.cu
-@@ -5514,6 +5514,33 @@ bool ggml_backend_is_cuda(ggml_backend_t backend) {
+@@ -5515,6 +5515,51 @@ bool ggml_backend_is_cuda(ggml_backend_t backend) {
return backend != NULL && ggml_guid_matches(backend->guid, ggml_backend_cuda_guid());
}
@@ -34,6 +67,24 @@ diff --git a/src/ggml-cuda/ggml-cuda.cu b/src/ggml-cuda/ggml-cuda.cu
+ return (void *) cuda_ctx->stream();
+}
+
++void ggml_backend_cuda_set_stream_priority(ggml_backend_t backend, int priority) {
++ if (!ggml_backend_is_cuda(backend)) {
++ return;
++ }
++ ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) backend->context;
++ cuda_ctx->stream_priority = priority;
++ // Recreate already-created streams so the priority applies to them as well.
++ for (int d = 0; d < GGML_CUDA_MAX_DEVICES; ++d) {
++ for (int i = 0; i < GGML_CUDA_MAX_STREAMS; ++i) {
++ if (cuda_ctx->streams[d][i] != nullptr) {
++ CUDA_CHECK(cudaStreamSynchronize(cuda_ctx->streams[d][i]));
++ CUDA_CHECK(cudaStreamDestroy(cuda_ctx->streams[d][i]));
++ cuda_ctx->streams[d][i] = nullptr;
++ }
++ }
++ }
++}
++
+void * ggml_backend_cuda_get_graph_template(
+ ggml_backend_t backend, const struct ggml_cgraph * cgraph) {
+ if (!ggml_backend_is_cuda(backend) || cgraph == nullptr) {
diff --git a/ggml-patches/0024-cuda-im2col-1d-tiled.patch b/ggml-patches/0024-cuda-im2col-1d-tiled.patch
new file mode 100644
index 0000000..69b7f38
--- /dev/null
+++ b/ggml-patches/0024-cuda-im2col-1d-tiled.patch
@@ -0,0 +1,77 @@
+diff --git a/src/ggml-cuda/im2col.cu b/src/ggml-cuda/im2col.cu
+index c2cc16fe..8e243e4f 100644
+--- a/src/ggml-cuda/im2col.cu
++++ b/src/ggml-cuda/im2col.cu
+@@ -44,12 +44,72 @@ static __global__ void im2col_kernel(
+ GGML_UNUSED(KH);
+ }
+
++// 1D fast path: [N, IC, IW] => [N, OW, IC*KW]. The generic kernel assigns consecutive threads
++// to consecutive (ic, kw) entries of one output position, so its reads stride across channels
++// (uncoalesced) and it re-reads every input element KW times from global memory. Here a block
++// stages the [channels x time window] input tile through shared memory (coalesced along time)
++// and then writes IM2COL_1D_TILE_T consecutive output rows with consecutive threads owning
++// consecutive columns (coalesced along IC*KW).
++#define IM2COL_1D_TILE_T 32
++
++template
++static __global__ void im2col_1d_tiled_kernel(
++ const float * __restrict__ x, T * __restrict__ dst,
++ int IC, int64_t IW, int64_t OW, int KW, int64_t IC_IW, int64_t N_IW,
++ int s0, int p0, int d0, int win) {
++ extern __shared__ float tile[]; // [nch][win]
++ const int W = IC * KW;
++ const int j0 = blockIdx.x * blockDim.x;
++ const int64_t t0 = (int64_t) blockIdx.y * IM2COL_1D_TILE_T;
++ const int n = blockIdx.z;
++ const int ic_lo = j0 / KW;
++ const int ic_hi = min(IC - 1, (j0 + (int) blockDim.x - 1) / KW);
++ const int nch = ic_hi - ic_lo + 1;
++ const int64_t iw0 = t0 * s0 - p0;
++ const float * xb = x + (size_t) n * N_IW;
++ for (int e = threadIdx.x; e < nch * win; e += blockDim.x) {
++ const int c = e / win;
++ const int w = e - c * win;
++ const int64_t iw = iw0 + w;
++ tile[e] = (iw >= 0 && iw < IW) ? xb[(size_t) (ic_lo + c) * IC_IW + iw] : 0.0f;
++ }
++ __syncthreads();
++ const int j = j0 + threadIdx.x;
++ if (j >= W) {
++ return;
++ }
++ const int ic = j / KW;
++ const int kw = j - ic * KW;
++ const float * row = tile + (ic - ic_lo) * win + kw * d0;
++ T * out = dst + ((size_t) n * OW + t0) * W + j;
++ const int tmax = (int) min((int64_t) IM2COL_1D_TILE_T, OW - t0);
++ for (int t = 0; t < tmax; ++t) {
++ out[(size_t) t * W] = ggml_cuda_cast(row[t * s0]);
++ }
++}
++
+ // im2col: [N, IC, IH, IW] => [N, OH, OW, IC*KH*KW]
+ template
+ static void im2col_cuda(const float * x, T* dst,
+ int64_t IW, int64_t IH, int64_t OW, int64_t OH, int64_t KW, int64_t KH, int64_t IC,
+ int64_t N, int64_t IC_IH_IW, int64_t IH_IW,
+ int s0,int s1,int p0,int p1,int d0,int d1, cudaStream_t stream) {
++ // Small OW launches too few tiles to matter; keep the generic kernel there.
++ if (IH == 1 && KH == 1 && OH == 1 && KW >= 1 && OW >= 2 * IM2COL_1D_TILE_T && IC * KW <= INT32_MAX) {
++ const int win = (IM2COL_1D_TILE_T - 1) * s0 + (int) (KW - 1) * d0 + 1;
++ const int nch_max = (CUDA_IM2COL_BLOCK_SIZE + (int) KW - 1) / (int) KW + 1;
++ const size_t smem = (size_t) nch_max * win * sizeof(float);
++ if (smem <= 48 * 1024 && N <= MAX_GRIDDIM_Z) {
++ const int64_t W = IC * KW;
++ dim3 grid((unsigned) ((W + CUDA_IM2COL_BLOCK_SIZE - 1) / CUDA_IM2COL_BLOCK_SIZE),
++ (unsigned) ((OW + IM2COL_1D_TILE_T - 1) / IM2COL_1D_TILE_T), (unsigned) N);
++ if (grid.y <= MAX_GRIDDIM_Y) {
++ im2col_1d_tiled_kernel<<>>(
++ x, dst, (int) IC, IW, OW, (int) KW, IC_IH_IW, IH_IW, s0, p0, d0, win);
++ return;
++ }
++ }
++ }
+ const int64_t IC_KH_KW = IC * KH * KW;
+ const int64_t num_blocks = (IC_KH_KW + CUDA_IM2COL_BLOCK_SIZE - 1) / CUDA_IM2COL_BLOCK_SIZE;
+ const int64_t N_OH = N * OH;
diff --git a/ggml-patches/0025-cuda-conv1d-fused.patch b/ggml-patches/0025-cuda-conv1d-fused.patch
new file mode 100644
index 0000000..615ce08
--- /dev/null
+++ b/ggml-patches/0025-cuda-conv1d-fused.patch
@@ -0,0 +1,753 @@
+diff --git a/include/ggml.h b/include/ggml.h
+index d403756e..e40cf62c 100644
+--- a/include/ggml.h
++++ b/include/ggml.h
+@@ -584,6 +584,7 @@ extern "C" {
+ GGML_OP_GLU,
+
+ GGML_OP_FUSED_ATTN,
++ GGML_OP_CONV1D_FUSED,
+
+ GGML_OP_COUNT,
+ };
+@@ -2498,6 +2499,54 @@ extern "C" {
+ float scale,
+ bool merge_heads);
+
++ // Fused causal 1-D convolution over [T, C] activations (T contiguous, CUDA only):
++ // y[co][t] = bias[co] + sum_{ci,k} act(xcat[ci][t + k*dilation]) * w[k][ci][co]
++ // where xcat = [cache | x] along T (cache holds exactly (kernel - 1) * dilation columns, or
++ // is NULL when that is 0) and act is an optional half-snake
++ // activation (alpha/alpha_inv over the first snake_channels channels, leaky-relu(slope) on
++ // the rest; pass alpha == NULL for identity). w_packed is F16 [cin_pad, cout_pad, K]
++ // (ne0 = cin_pad; both paddings multiples of 16, zero filled). Output F32 [T_out, cout].
++ GGML_API struct ggml_tensor * ggml_conv1d_fused(
++ struct ggml_context * ctx,
++ struct ggml_tensor * x,
++ struct ggml_tensor * cache,
++ struct ggml_tensor * w_packed,
++ struct ggml_tensor * bias,
++ struct ggml_tensor * alpha,
++ struct ggml_tensor * alpha_inv,
++ int kernel,
++ int dilation,
++ int cout,
++ int snake_channels,
++ float slope);
++
++ // Grouped variant: `groups` independent causal convolutions with per-group kernel size and
++ // dilation over channel-concatenated tensors. x: [T, groups*cin] (or [T, cin] when
++ // shared_input: every group reads the same input), cache: [max_pad, channels of x] holding
++ // the last max_pad pre-activation samples (group g uses its last (K_g-1)*d_g columns),
++ // w_packed: the groups' packed weights back to back ([cin_pad, cout_pad, K_g] each),
++ // bias: [groups*cout], alpha/alpha_inv: [groups*snake_channels], residual (optional):
++ // [T, groups*cout] (or [T, cout] when residual_shared) added to the output.
++ // Output: [T, groups*cout].
++ GGML_API struct ggml_tensor * ggml_conv1d_fused_grouped(
++ struct ggml_context * ctx,
++ struct ggml_tensor * x,
++ struct ggml_tensor * cache,
++ struct ggml_tensor * w_packed,
++ struct ggml_tensor * bias,
++ struct ggml_tensor * alpha,
++ struct ggml_tensor * alpha_inv,
++ struct ggml_tensor * residual,
++ int groups,
++ const int * kernels,
++ const int * dilations,
++ int cout,
++ int snake_channels,
++ float slope,
++ bool shared_input,
++ bool residual_shared);
++
++
+ // TODO: needs to be adapted to ggml_flash_attn_ext
+ GGML_API struct ggml_tensor * ggml_flash_attn_back(
+ struct ggml_context * ctx,
+diff --git a/src/ggml-cpu/ggml-cpu.c b/src/ggml-cpu/ggml-cpu.c
+index 43de6346..21af9c52 100644
+--- a/src/ggml-cpu/ggml-cpu.c
++++ b/src/ggml-cpu/ggml-cpu.c
+@@ -1992,6 +1992,10 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm
+ {
+ GGML_ABORT("FUSED_ATTN has no CPU path (CUDA-only)");
+ } break;
++ case GGML_OP_CONV1D_FUSED:
++ {
++ GGML_ABORT("CONV1D_FUSED has no CPU path (CUDA-only)");
++ } break;
+ case GGML_OP_FLASH_ATTN_BACK:
+ {
+ int32_t t = ggml_get_op_params_i32(tensor, 0);
+diff --git a/src/ggml-cpu/ggml-cpu.cpp b/src/ggml-cpu/ggml-cpu.cpp
+index 0c570639..f20b8bcb 100644
+--- a/src/ggml-cpu/ggml-cpu.cpp
++++ b/src/ggml-cpu/ggml-cpu.cpp
+@@ -439,6 +439,7 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st
+ }
+
+ switch (op->op) {
++ case GGML_OP_CONV1D_FUSED:
+ case GGML_OP_FUSED_ATTN:
+ return false; // CUDA-only; CPU falls back to the unfused graph
+ case GGML_OP_CPY:
+diff --git a/src/ggml-cuda/conv1d-fused.cu b/src/ggml-cuda/conv1d-fused.cu
+new file mode 100644
+index 00000000..336ba959
+--- /dev/null
++++ b/src/ggml-cuda/conv1d-fused.cu
+@@ -0,0 +1,483 @@
++#include "conv1d-fused.cuh"
++
++#include
++#include
++
++// Fused causal 1-D convolution as an implicit GEMM on tensor cores (mma.sync m16n8k16, F16 in,
++// F32 accumulate):
++// y[co][t] = bias[co] + sum_{ci,k} act(xcat[ci][t + k*d]) * w[k][co][ci]
++// xcat = [cache | x] along time; act is an optional half-snake applied while the input window is
++// staged in shared memory. A block owns a BM(time) x BN(cout) tile and a range of 16-channel
++// chunks (split-K over channels for small grids, reduced with atomics into a zeroed output).
++// Weights are packed [K][cout_pad][cin_pad] (ci contiguous) so B fragments are 32-bit loads.
++namespace {
++
++// Block skeleton (rev. 2): 4 warps per block on BM 64/32 x BN 32 tiles, the per-thread window
++// mapping computed once, snake parameters staged per chunk, cp.async weight ring with one
++// __syncthreads per 16-channel chunk, and split-K partials reduced by the last block to arrive
++// per tile (self-resetting counters) instead of a memset + output atomics. The per-launch cost
++// is dominated by prologue/epilogue latency, so the grid is kept within one resident wave.
++// Footprint (<= 168 regs x 128 threads, <= 41 KB shared) leaves room beside a resident
++// MagpieTTS persistent-kernel block so the codec keeps overlapping the decoder path.
++constexpr int CF_THREADS = 128;
++constexpr int CF_BK = 16;
++constexpr int CF_BN = 32;
++constexpr int CF_LDW = 24;
++constexpr int CF_LDB = 24;
++constexpr int CF_MAX_K = 11;
++constexpr int CF_MAX_WIN = 64 + 10 * 5;
++constexpr int CF_PF = 16; // window elements per thread: CF_BK * 128 / 128
++constexpr int CF_STAGES = 2;
++
++struct conv1d_fused_params {
++ const float * x;
++ const float * cache;
++ const half * w;
++ const float * bias;
++ const float * alpha;
++ const float * inv_b;
++ float * y;
++ float * partials; // [splits][tiles][BM*BN] scratch when splits > 1
++ int * counters; // [tiles] arrival counters, zero between launches
++ const float * residual; // optional [cout (x groups)][T] added to the output
++ long xs, cs, ys, rs;
++ int T, cache_len, t_out, cin, cin_pad, cout, cout_pad, K, d, snake_ch;
++ int chunks_per_split, n_chunks, splits;
++ float slope;
++ // grouped problems (blockIdx.y = group * tiles_n + cout tile): per-group kernel/dilation,
++ // weight offset (halves) and channel bases; cache_len is the largest pad, each group reads
++ // the last (K_g-1)*d_g cache columns.
++ int groups, tiles_n, shared_input, residual_shared;
++ int Kg[3], dg[3];
++ long w_off[3];
++};
++
++__device__ __forceinline__ void cp_async16(void * smem, const void * gmem) {
++ const unsigned s = (unsigned) __cvta_generic_to_shared(smem);
++ asm volatile("cp.async.cg.shared.global [%0], [%1], 16;\n" ::"r"(s), "l"(gmem));
++}
++__device__ __forceinline__ void cp_async_commit() { asm volatile("cp.async.commit_group;\n" ::); }
++template __device__ __forceinline__ void cp_async_wait() { asm volatile("cp.async.wait_group %0;\n" ::"n"(N)); }
++__device__ __forceinline__ void mma_16816(float (&c)[4], const unsigned (&a)[4], const unsigned (&b)[2]) {
++ asm volatile(
++ "mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 {%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%0,%1,%2,%3};\n"
++ : "+f"(c[0]), "+f"(c[1]), "+f"(c[2]), "+f"(c[3])
++ : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), "r"(b[0]), "r"(b[1]));
++}
++
++constexpr int CF_MAX_REGS = 168;
++// __maxnreg__ is a CUDA 12.4+ toolkit macro; older toolkits and HIP get plain launch bounds.
++#if defined(__maxnreg__) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
++#define CF_LAUNCH_BOUNDS __maxnreg__(CF_MAX_REGS)
++#else
++#define CF_LAUNCH_BOUNDS __launch_bounds__(CF_THREADS)
++#endif
++
++template
++// 128 threads x <= 168 registers (21.5K) fit beside a resident MagpieTTS persistent kernel block
++// (168 x 256 = 43K) in the 64K-register file, so codec blocks co-schedule instead of serializing.
++__global__ void CF_LAUNCH_BOUNDS conv1d_fused_kernel(const conv1d_fused_params p) {
++#if defined(GGML_USE_HIP) || (defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 800)
++ GGML_UNUSED(p);
++ NO_DEVICE_CODE; // unreachable: ggml_cuda_conv1d_fused_supported() requires Ampere code
++#else
++ constexpr int WARPS_M = 2; // 4 warps: 2x2 (BM 64: 32x16 warp tiles; BM 32: 16x16)
++ constexpr int WARPS_N = 4 / WARPS_M;
++ constexpr int WM = BM / WARPS_M;
++ constexpr int WN = CF_BN / WARPS_N;
++ constexpr int FM = WM / 16;
++ constexpr int FN = WN / 8;
++ constexpr int LDC = CF_BN + 4;
++ const int grp = p.groups > 1 ? (int) blockIdx.y / p.tiles_n : 0;
++ const int K = p.groups > 1 ? p.Kg[grp] : p.K;
++ const int dil = p.groups > 1 ? p.dg[grp] : p.d;
++ const half * wg = p.w + (p.groups > 1 ? p.w_off[grp] : 0);
++ const int pad_g = (K - 1) * dil; // cache columns this group needs
++ const int cache_col0 = p.cache_len - pad_g; // ... the last ones of the shared cache
++ const int cin_base = p.shared_input ? 0 : grp * p.cin;
++ const int cout_base = grp * p.cout;
++ extern __shared__ __align__(16) half cf_dyn[];
++ const int B_CHUNK = K * CF_BN * CF_LDB;
++ const int win = BM + (K - 1) * dil;
++ const int W_CHUNK = win * CF_LDW;
++ half * s_b = cf_dyn;
++ half * s_win = cf_dyn + CF_STAGES * B_CHUNK;
++ float * s_c = reinterpret_cast(cf_dyn);
++ __shared__ int s_last;
++ __shared__ float s_alpha[2][CF_BK], s_invb[2][CF_BK]; // snake params of the chunk in each window buffer
++
++ const int tid = threadIdx.x;
++ const int lane = tid & 31;
++ const int warp = tid >> 5;
++ const int wm = warp % WARPS_M;
++ const int wn = warp / WARPS_M;
++ const int t0 = blockIdx.x * BM;
++ const int co0 = (p.groups > 1 ? (int) blockIdx.y % p.tiles_n : (int) blockIdx.y) * CF_BN;
++ const int total_in = pad_g + p.T;
++ const int c_begin = blockIdx.z * p.chunks_per_split;
++ const int c_end = min(p.n_chunks, c_begin + p.chunks_per_split);
++ const int n_iter = c_end - c_begin;
++
++ float acc[FM][FN][4];
++#pragma unroll
++ for (int i = 0; i < FM; ++i)
++#pragma unroll
++ for (int j = 0; j < FN; ++j)
++#pragma unroll
++ for (int r = 0; r < 4; ++r) acc[i][j][r] = 0.0f;
++
++ // Window mapping of this thread's elements, computed once: element e = tid + i*128 of the
++ // [CF_BK][win] chunk window is channel j = e / win at window position u = e % win; the
++ // global time is g = t0 + u (cache rows first, then x). Packed as (j << 16) | u, 0xffff = none.
++ unsigned map[CF_PF];
++#pragma unroll
++ for (int i = 0; i < CF_PF; ++i) {
++ const int e = tid + i * CF_THREADS;
++ unsigned m = 0xffffffffu;
++ if (e < CF_BK * win) {
++ const int j = e / win;
++ const int u = e - j * win;
++ if (t0 + u < total_in) m = ((unsigned) j << 16) | (unsigned) u;
++ }
++ map[i] = m;
++ }
++
++ auto stage_b = [&](int c, int slot) {
++ if (c < c_end) {
++ const int ci0 = c * CF_BK;
++ half * dst = s_b + slot * B_CHUNK;
++ for (int e = tid; e < K * CF_BN * 2; e += CF_THREADS) {
++ const int row = e >> 1;
++ const int piece = e & 1;
++ const int k = row / CF_BN;
++ const int n = row - k * CF_BN;
++ const int co = co0 + n;
++ half * d = dst + (k * CF_BN + n) * CF_LDB + piece * 8;
++ if (co < p.cout_pad) {
++ cp_async16(d, wg + ((size_t) k * p.cout_pad + co) * p.cin_pad + ci0 + piece * 8);
++ } else {
++ *reinterpret_cast(d) = make_uint4(0, 0, 0, 0);
++ }
++ }
++ }
++ cp_async_commit();
++ };
++ auto fetch_window = [&](int c, float (&pf)[CF_PF]) {
++ const int ci0 = c * CF_BK;
++#pragma unroll
++ for (int i = 0; i < CF_PF; ++i) {
++ float v = 0.0f;
++ const unsigned m = map[i];
++ if (c < c_end && m != 0xffffffffu) {
++ const int ci = ci0 + (int) (m >> 16);
++ const int g = t0 + (int) (m & 0xffffu);
++ if (ci < p.cin) {
++ v = g < pad_g ? p.cache[(size_t) (cin_base + ci) * p.cs + cache_col0 + g]
++ : p.x[(size_t) (cin_base + ci) * p.xs + (g - pad_g)];
++ }
++ }
++ pf[i] = v;
++ }
++ };
++ auto stage_act = [&](int c, int buf) {
++ if (tid < CF_BK && p.alpha) {
++ const int ci = c * CF_BK + tid;
++ const bool ok = c < c_end && ci < p.snake_ch;
++ s_alpha[buf][tid] = ok ? p.alpha[grp * p.snake_ch + ci] : 0.0f;
++ s_invb[buf][tid] = ok ? p.inv_b[grp * p.snake_ch + ci] : 0.0f;
++ }
++ };
++ auto store_window = [&](int c, const float (&pf)[CF_PF], int buf) {
++ const int ci0 = c * CF_BK;
++ half * dst = s_win + buf * W_CHUNK;
++#pragma unroll
++ for (int i = 0; i < CF_PF; ++i) {
++ const unsigned m = map[i];
++ if (m != 0xffffffffu) {
++ const int j = (int) (m >> 16);
++ const int u = (int) (m & 0xffffu);
++ const int ci = ci0 + j;
++ float v = pf[i];
++ if (p.alpha && ci < p.cin) {
++ if (ci < p.snake_ch) {
++ const float sn = __sinf(s_alpha[buf][j] * v);
++ v = v + sn * sn * s_invb[buf][j];
++ } else {
++ v = v > 0.0f ? v : v * p.slope;
++ }
++ }
++ dst[u * CF_LDW + j] = __float2half(v);
++ }
++ }
++ // window positions past total_in (never mapped) must read as zero: only the last time tile
++ if (t0 + win > total_in) {
++ for (int e = tid; e < CF_BK * win; e += CF_THREADS) {
++ const int j = e / win;
++ const int u = e - j * win;
++ if (t0 + u >= total_in) dst[u * CF_LDW + j] = __float2half(0.0f);
++ }
++ }
++ };
++ float pf0[CF_PF], pf1[CF_PF];
++ for (int q = 0; q < CF_STAGES - 1; ++q) stage_b(c_begin + q, q);
++ fetch_window(c_begin + 0, pf0);
++ fetch_window(c_begin + 1, pf1);
++ stage_act(c_begin + 0, 0);
++ stage_act(c_begin + 1, 1);
++ __syncthreads();
++
++ const int a_row = lane >> 2;
++ const int a_col = (lane & 3) * 2;
++ for (int it = 0; it < n_iter; ++it) {
++ const int c = c_begin + it;
++ const int slot = it % CF_STAGES;
++ const int wbuf = it & 1;
++ if (it & 1) store_window(c, pf1, wbuf); else store_window(c, pf0, wbuf);
++ cp_async_wait();
++ __syncthreads();
++ stage_b(c + CF_STAGES - 1, (it + CF_STAGES - 1) % CF_STAGES);
++ if (it & 1) fetch_window(c + 2, pf1); else fetch_window(c + 2, pf0);
++ stage_act(c + 2, wbuf); // chunk it's params were consumed by store_window above (before the sync)
++ const half * bcur = s_b + slot * B_CHUNK;
++ const half * wcur = s_win + wbuf * W_CHUNK;
++ for (int k = 0; k < K; ++k) {
++ const int shift = k * dil; // per-group dilation (grouped problems)
++ unsigned bfrag[FN][2];
++#pragma unroll
++ for (int j = 0; j < FN; ++j) {
++ const half * brow = bcur + (k * CF_BN + wn * WN + j * 8 + a_row) * CF_LDB + a_col;
++ bfrag[j][0] = *reinterpret_cast(brow);
++ bfrag[j][1] = *reinterpret_cast(brow + 8);
++ }
++#pragma unroll
++ for (int i = 0; i < FM; ++i) {
++ const half * arow = wcur + (shift + wm * WM + i * 16 + a_row) * CF_LDW + a_col;
++ unsigned afrag[4];
++ afrag[0] = *reinterpret_cast(arow);
++ afrag[1] = *reinterpret_cast(arow + 8 * CF_LDW);
++ afrag[2] = *reinterpret_cast(arow + 8);
++ afrag[3] = *reinterpret_cast(arow + 8 * CF_LDW + 8);
++#pragma unroll
++ for (int j = 0; j < FN; ++j) mma_16816(acc[i][j], afrag, bfrag[j]);
++ }
++ }
++ }
++ cp_async_wait<0>();
++ __syncthreads();
++ // accumulators -> shared [t][co]
++#pragma unroll
++ for (int i = 0; i < FM; ++i) {
++#pragma unroll
++ for (int j = 0; j < FN; ++j) {
++ const int m = wm * WM + i * 16 + a_row;
++ const int n = wn * WN + j * 8 + a_col;
++ s_c[m * LDC + n] = acc[i][j][0];
++ s_c[m * LDC + n + 1] = acc[i][j][1];
++ s_c[(m + 8) * LDC + n] = acc[i][j][2];
++ s_c[(m + 8) * LDC + n + 1] = acc[i][j][3];
++ }
++ }
++ __syncthreads();
++ const int tile = blockIdx.y * gridDim.x + blockIdx.x;
++ const int tiles = gridDim.x * gridDim.y;
++ if (p.splits > 1) {
++ // partial tile -> scratch; the last block to arrive for this tile reduces
++ float * part = p.partials + ((size_t) blockIdx.z * tiles + tile) * (BM * CF_BN);
++ for (int e = tid; e < BM * CF_BN; e += CF_THREADS) {
++ const int n = e / BM, m = e - n * BM;
++ part[e] = s_c[m * LDC + n];
++ }
++ __threadfence();
++ __syncthreads();
++ if (tid == 0) s_last = (atomicAdd(p.counters + tile, 1) == p.splits - 1) ? 1 : 0;
++ __syncthreads();
++ if (!s_last) return;
++ __threadfence();
++ for (int e = tid; e < BM * CF_BN; e += CF_THREADS) {
++ const int n = e / BM, m = e - n * BM;
++ const int t = t0 + m, co = co0 + n;
++ if (t < p.t_out && co < p.cout) {
++ float v = p.bias ? p.bias[cout_base + co] : 0.0f;
++ for (int z = 0; z < p.splits; ++z)
++ v += p.partials[((size_t) z * tiles + tile) * (BM * CF_BN) + e];
++ if (p.residual)
++ v += p.residual[(size_t) ((p.residual_shared ? 0 : cout_base) + co) * p.rs + t];
++ p.y[(size_t) (cout_base + co) * p.ys + t] = v;
++ }
++ }
++ if (tid == 0) p.counters[tile] = 0; // ready for the next launch
++ } else {
++ for (int e = tid; e < BM * CF_BN; e += CF_THREADS) {
++ const int n = e / BM, m = e - n * BM;
++ const int t = t0 + m, co = co0 + n;
++ if (t < p.t_out && co < p.cout) {
++ float v = s_c[m * LDC + n] + (p.bias ? p.bias[cout_base + co] : 0.0f);
++ if (p.residual)
++ v += p.residual[(size_t) ((p.residual_shared ? 0 : cout_base) + co) * p.rs + t];
++ p.y[(size_t) (cout_base + co) * p.ys + t] = v;
++ }
++ }
++ }
++#endif
++}
++
++// launch geometry shared by the bench and the ggml op
++struct geometry { int BM, tiles_m, tiles_n, splits, chunks_per_split; size_t smem, partial_floats; };
++static inline geometry plan(const conv1d_fused_params & p) {
++ geometry g{};
++ const int n_chunks = p.cin_pad / CF_BK;
++ g.BM = p.t_out >= 64 ? 64 : 32;
++ g.tiles_m = (p.t_out + g.BM - 1) / g.BM;
++ g.tiles_n = ((p.cout + CF_BN - 1) / CF_BN) * (p.groups > 1 ? p.groups : 1);
++ const int tiles = g.tiles_m * g.tiles_n;
++ // Split the channel chunks across blocks up to one resident wave (~128 blocks). With a MagpieTTS
++ // persistent kernel resident on every SM only one codec block fits per SM, so extra waves cost
++ // more than longer per-block chains (measured: a waves x chain cost model chose more splits
++ // for the grouped launches and lost 4% end to end).
++ int splits = 1;
++ if (tiles < 128) splits = min(n_chunks, max(1, 128 / tiles));
++ g.splits = splits;
++ g.chunks_per_split = (n_chunks + splits - 1) / splits;
++ int kmax = p.K, padmax = (p.K - 1) * p.d;
++ for (int i = 0; i < (p.groups > 1 ? p.groups : 0); ++i) {
++ kmax = p.Kg[i] > kmax ? p.Kg[i] : kmax;
++ padmax = (p.Kg[i] - 1) * p.dg[i] > padmax ? (p.Kg[i] - 1) * p.dg[i] : padmax;
++ }
++ const int win = g.BM + padmax;
++ g.smem = (size_t) (CF_STAGES * kmax * CF_BN * CF_LDB + 2 * win * CF_LDW) * sizeof(half);
++ if (g.smem < (size_t) g.BM * (CF_BN + 4) * sizeof(float)) g.smem = (size_t) g.BM * (CF_BN + 4) * sizeof(float);
++ g.partial_floats = splits > 1 ? (size_t) splits * tiles * g.BM * CF_BN : 0;
++ return g;
++}
++
++} // namespace
++
++bool ggml_cuda_conv1d_fused_supported(const ggml_tensor * op, int cc) {
++ // mma.sync m16n8k16 + cp.async: needs an NVIDIA sm_80+ device running sm_80+ device code
++ // (a 75-virtual-only build JIT-compiled on Ampere gets the pre-Ampere stub).
++ if (!GGML_CUDA_CC_IS_NVIDIA(cc) || ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_AMPERE) return false;
++ const ggml_tensor * x = op->src[0];
++ const ggml_tensor * w = op->src[1];
++ const ggml_tensor * cache = op->src[3];
++ if (x->type != GGML_TYPE_F32 || w->type != GGML_TYPE_F16 || op->type != GGML_TYPE_F32) return false;
++ if (x->nb[0] != sizeof(float) || (cache && cache->nb[0] != sizeof(float))) return false;
++ if (!ggml_is_contiguous(w)) return false;
++ const int32_t * prm = (const int32_t *) op->op_params;
++ const int groups = prm[5] > 1 ? prm[5] : 1;
++ const int Ks[3] = { prm[0], prm[6], prm[7] };
++ const int ds[3] = { prm[1], prm[8], prm[9] };
++ for (int i = 0; i < groups; ++i) {
++ if (Ks[i] < 1 || Ks[i] > CF_MAX_K || 64 + (Ks[i] - 1) * ds[i] > CF_MAX_WIN) return false;
++ }
++ if (groups > 1) {
++ // packed weights are concatenated per group: shape check happens in the op wrapper
++ if ((w->ne[0] % 16) != 0) return false;
++ } else {
++ if ((w->ne[0] % 16) != 0 || (w->ne[1] % 16) != 0) return false;
++ }
++ const ggml_tensor * residual = op->src[6];
++ if (residual && (residual->type != GGML_TYPE_F32 || residual->nb[0] != sizeof(float))) return false;
++ return true;
++}
++
++// Split-K scratch: partial tiles plus arrival counters, one allocation per CUDA stream.
++// Launches on a stream are serialized and every launch's reduction completes inside it, so a
++// stream's buffer is reused without synchronization, and kernels on different streams never share
++// one. Counters are zero-initialized once and reset by the last block of every tile. (A per-op
++// pool allocation here caused host-side frees and stalls.)
++constexpr size_t CF_SCRATCH_FLOATS = 1u << 20; // 4 MB of partials
++constexpr int CF_MAX_TILES = 4096;
++struct conv1d_fused_scratch { float * partials = nullptr; int * counters = nullptr; };
++static conv1d_fused_scratch & conv1d_fused_get_scratch(int device, cudaStream_t stream) {
++ static std::mutex mutex;
++ static std::map, conv1d_fused_scratch> scratch;
++ std::lock_guard lock(mutex);
++ conv1d_fused_scratch & sc = scratch[{device, stream}];
++ if (!sc.partials) {
++ CUDA_CHECK(cudaMalloc(&sc.partials, CF_SCRATCH_FLOATS * sizeof(float)));
++ CUDA_CHECK(cudaMalloc(&sc.counters, CF_MAX_TILES * sizeof(int)));
++ CUDA_CHECK(cudaMemset(sc.counters, 0, CF_MAX_TILES * sizeof(int)));
++ CUDA_CHECK(cudaFuncSetAttribute(conv1d_fused_kernel<64>, cudaFuncAttributeMaxDynamicSharedMemorySize, 96 * 1024));
++ CUDA_CHECK(cudaFuncSetAttribute(conv1d_fused_kernel<32>, cudaFuncAttributeMaxDynamicSharedMemorySize, 96 * 1024));
++ }
++ return sc;
++}
++
++void ggml_cuda_op_conv1d_fused(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
++ const ggml_tensor * x = dst->src[0];
++ const ggml_tensor * w = dst->src[1];
++ const ggml_tensor * bias = dst->src[2];
++ const ggml_tensor * cache = dst->src[3];
++ const ggml_tensor * alpha = dst->src[4];
++ const ggml_tensor * inv_b = dst->src[5];
++ const int32_t * op = (const int32_t *) dst->op_params;
++
++ conv1d_fused_params p{};
++ p.x = (const float *) x->data;
++ p.cache = cache ? (const float *) cache->data : nullptr;
++ p.w = (const half *) w->data;
++ p.bias = bias ? (const float *) bias->data : nullptr;
++ p.alpha = (op[3] && alpha) ? (const float *) alpha->data : nullptr;
++ p.inv_b = (op[3] && inv_b) ? (const float *) inv_b->data : nullptr;
++ p.y = (float *) dst->data;
++ p.xs = (long) (x->nb[1] / sizeof(float));
++ p.cs = cache ? (long) (cache->nb[1] / sizeof(float)) : 0;
++ p.ys = (long) (dst->nb[1] / sizeof(float));
++ p.T = (int) x->ne[0];
++ p.cache_len = cache ? (int) cache->ne[0] : 0;
++ p.t_out = (int) dst->ne[0];
++ p.cin = (int) x->ne[1];
++ p.cin_pad = (int) w->ne[0];
++ p.cout = (int) dst->ne[1];
++ p.cout_pad = (int) w->ne[1];
++ p.K = op[0];
++ p.d = op[1];
++ p.snake_ch = op[2];
++ memcpy(&p.slope, &op[4], sizeof(float));
++ p.groups = op[5] > 1 ? op[5] : 1;
++ if (p.groups > 1) {
++ // grouped: x is [T][groups*cin] (or [T][cin] when shared), y is [T][groups*cout], the packed
++ // weight tensor holds the groups' [cin_pad][cout_pad][K_g] blocks back to back
++ p.shared_input = op[10];
++ p.residual_shared = op[11];
++ p.cout = op[12];
++ p.cin = p.shared_input ? (int) x->ne[1] : (int) x->ne[1] / p.groups;
++ p.cin_pad = (int) w->ne[0];
++ p.cout_pad = (p.cout + 15) / 16 * 16;
++ p.Kg[0] = op[0]; p.Kg[1] = op[6]; p.Kg[2] = op[7];
++ p.dg[0] = op[1]; p.dg[1] = op[8]; p.dg[2] = op[9];
++ long off = 0;
++ for (int i = 0; i < p.groups; ++i) {
++ p.w_off[i] = off;
++ off += (long) p.Kg[i] * p.cout_pad * p.cin_pad;
++ }
++ GGML_ASSERT(off == ggml_nelements(w));
++ GGML_ASSERT(p.cache_len == op[13]);
++ p.tiles_n = (p.cout + CF_BN - 1) / CF_BN;
++ } else {
++ p.Kg[0] = p.K; p.dg[0] = p.d; p.w_off[0] = 0;
++ p.tiles_n = (p.cout + CF_BN - 1) / CF_BN;
++ }
++ const ggml_tensor * residual = dst->src[6];
++ p.residual = residual ? (const float *) residual->data : nullptr;
++ p.rs = residual ? (long) (residual->nb[1] / sizeof(float)) : 0;
++ p.n_chunks = p.cin_pad / CF_BK;
++
++ geometry g = plan(p);
++ conv1d_fused_scratch & sc = conv1d_fused_get_scratch(ctx.device, ctx.stream());
++ if (g.partial_floats > CF_SCRATCH_FLOATS || g.tiles_m * g.tiles_n > CF_MAX_TILES) {
++ // too large for the static scratch: no split-K
++ g.splits = 1;
++ g.chunks_per_split = p.n_chunks;
++ g.partial_floats = 0;
++ }
++ p.splits = g.splits;
++ p.chunks_per_split = g.chunks_per_split;
++ p.partials = sc.partials;
++ p.counters = sc.counters;
++ cudaStream_t stream = ctx.stream();
++ dim3 grid(g.tiles_m, g.tiles_n, g.splits);
++ if (g.BM == 64) conv1d_fused_kernel<64><<>>(p);
++ else conv1d_fused_kernel<32><<