diff --git a/examples/qwen3_5/benchmark_harness.hpp b/examples/qwen3_5/benchmark_harness.hpp new file mode 100644 index 000000000..6a2e4f4d7 --- /dev/null +++ b/examples/qwen3_5/benchmark_harness.hpp @@ -0,0 +1,163 @@ +// Copyright (c) MLLM Team. +// Licensed under the MIT License. +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#if defined(__ANDROID__) || defined(__linux__) +#include +#endif + +namespace mllm::examples::qwen3_5::benchmark { + +inline std::optional readTextFile(const std::filesystem::path& path) { + std::ifstream stream(path); + if (!stream) { return std::nullopt; } + std::ostringstream contents; + contents << stream.rdbuf(); + auto value = contents.str(); + while (!value.empty() && (value.back() == '\n' || value.back() == '\r')) { value.pop_back(); } + return value; +} + +inline std::optional readIntegerFile(const std::filesystem::path& path) { + const auto text = readTextFile(path); + if (!text.has_value()) { return std::nullopt; } + try { + size_t consumed = 0; + const auto value = std::stoll(*text, &consumed); + if (consumed != text->size()) { return std::nullopt; } + return value; + } catch (...) { return std::nullopt; } +} + +inline std::vector currentAffinityCpus() { + std::vector cpus; +#if defined(__ANDROID__) || defined(__linux__) + cpu_set_t mask; + CPU_ZERO(&mask); + if (sched_getaffinity(0, sizeof(mask), &mask) != 0) { return cpus; } + for (int cpu = 0; cpu < CPU_SETSIZE; ++cpu) { + if (CPU_ISSET(cpu, &mask)) { cpus.push_back(cpu); } + } +#endif + return cpus; +} + +inline nlohmann::json captureTelemetry() { + using Json = nlohmann::json; + Json snapshot = { +#if defined(__ANDROID__) + {"platform", "android"}, +#elif defined(__linux__) + {"platform", "linux"}, +#elif defined(__APPLE__) + {"platform", "macos"}, +#else + {"platform", "other"}, +#endif + {"captured_epoch_us", + std::chrono::duration_cast(std::chrono::system_clock::now().time_since_epoch()).count()}, + }; + + const auto affinity = currentAffinityCpus(); + snapshot["affinity_cpus"] = affinity; + snapshot["cpus"] = Json::array(); + snapshot["ceiling_vector"] = Json::array(); + snapshot["thermal_zones"] = Json::array(); + + for (const int cpu : affinity) { + const auto base = std::filesystem::path("/sys/devices/system/cpu") / ("cpu" + std::to_string(cpu)); + const auto cpufreq = base / "cpufreq"; + // Most kernels omit cpu0/online because CPU 0 cannot be offlined, so report it as online. + const auto online = cpu == 0 ? std::optional(1) : readIntegerFile(base / "online"); + const auto cpuinfo_max = readIntegerFile(cpufreq / "cpuinfo_max_freq"); + const auto scaling_max = readIntegerFile(cpufreq / "scaling_max_freq"); + const auto scaling_cur = readIntegerFile(cpufreq / "scaling_cur_freq"); + const auto governor = readTextFile(cpufreq / "scaling_governor"); + snapshot["cpus"].push_back({ + {"cpu", cpu}, + {"online", online.has_value() ? Json(*online) : Json(nullptr)}, + {"cpuinfo_max_freq", cpuinfo_max.has_value() ? Json(*cpuinfo_max) : Json(nullptr)}, + {"scaling_max_freq", scaling_max.has_value() ? Json(*scaling_max) : Json(nullptr)}, + {"scaling_cur_freq", scaling_cur.has_value() ? Json(*scaling_cur) : Json(nullptr)}, + {"scaling_governor", governor.has_value() ? Json(*governor) : Json(nullptr)}, + }); + snapshot["ceiling_vector"].push_back(scaling_max.has_value() ? Json(*scaling_max) : Json(nullptr)); + } + + const std::filesystem::path thermal_root("/sys/class/thermal"); + std::error_code error; + if (std::filesystem::exists(thermal_root, error)) { + std::vector zones; + for (const auto& entry : std::filesystem::directory_iterator(thermal_root, error)) { + if (entry.path().filename().string().starts_with("thermal_zone")) { zones.push_back(entry.path()); } + } + std::sort(zones.begin(), zones.end()); + for (const auto& zone : zones) { + const auto type = readTextFile(zone / "type"); + const auto temp = readIntegerFile(zone / "temp"); + snapshot["thermal_zones"].push_back({ + {"zone", zone.filename().string()}, + {"type", type.has_value() ? Json(*type) : Json(nullptr)}, + {"temp_milli_c", temp.has_value() ? Json(*temp) : Json(nullptr)}, + }); + } + } + + return snapshot; +} + +inline std::vector validateRequiredTelemetry(const nlohmann::json& snapshot) { + std::vector errors; + if (snapshot.value("platform", "") != "android" && snapshot.value("platform", "") != "linux") { + errors.push_back("unsupported_telemetry_platform"); + } + if (!snapshot.contains("affinity_cpus") || snapshot["affinity_cpus"].empty()) { errors.push_back("empty_affinity"); } + if (!snapshot.contains("cpus") || snapshot["cpus"].empty()) { errors.push_back("empty_cpu_telemetry"); } + if (!snapshot.contains("ceiling_vector") || snapshot["ceiling_vector"].empty()) { errors.push_back("empty_ceiling_vector"); } + if (snapshot.contains("cpus")) { + for (const auto& cpu : snapshot["cpus"]) { + const auto id = cpu.value("cpu", -1); + for (const auto* field : {"online", "cpuinfo_max_freq", "scaling_max_freq", "scaling_cur_freq", "scaling_governor"}) { + if (!cpu.contains(field) || cpu[field].is_null()) { + errors.push_back("cpu" + std::to_string(id) + "_missing_" + field); + } + } + } + } + return errors; +} + +inline std::vector validateStableTelemetry(const nlohmann::json& before, const nlohmann::json& after) { + std::vector errors; + if (before.value("affinity_cpus", nlohmann::json::array()) != after.value("affinity_cpus", nlohmann::json::array())) { + errors.push_back("affinity_changed"); + } + if (before.value("ceiling_vector", nlohmann::json::array()) != after.value("ceiling_vector", nlohmann::json::array())) { + errors.push_back("ceiling_vector_changed"); + } + const auto project = [](const nlohmann::json& snapshot, const char* field) { + nlohmann::json values = nlohmann::json::array(); + if (snapshot.contains("cpus")) { + for (const auto& cpu : snapshot["cpus"]) { values.push_back(cpu.value(field, nlohmann::json(nullptr))); } + } + return values; + }; + if (project(before, "online") != project(after, "online")) { errors.push_back("online_vector_changed"); } + if (project(before, "scaling_governor") != project(after, "scaling_governor")) { + errors.push_back("governor_vector_changed"); + } + return errors; +} + +} // namespace mllm::examples::qwen3_5::benchmark diff --git a/examples/qwen3_5/main.cpp b/examples/qwen3_5/main.cpp index 45094feaf..da4ebe46e 100644 --- a/examples/qwen3_5/main.cpp +++ b/examples/qwen3_5/main.cpp @@ -1,26 +1,48 @@ #include +#include #include +#include +#include #include #include #include +#include + +#include #include +#include #include #include #include #include +#include "benchmark_harness.hpp" + using mllm::Argparse; MLLM_MAIN({ + auto engine_args = mllm::engineArgAttach(); auto& help = Argparse::add("-h|--help").help("Show help message"); auto& model_path = Argparse::add("-m|--model_path").help("Model path").required(true); auto& model_version = Argparse::add("-mv|--model_version").help("Model version").required(true); auto& tokenizer_path = Argparse::add("-t|--tokenizer_path").help("Tokenizer JSON path").required(true); auto& config_path = Argparse::add("-c|--config_path").help("Config path").required(true); auto& prompt = Argparse::add("-p|--prompt").help("Run one prompt non-interactively").required(false); + auto& prompt_file = Argparse::add("--prompt_file").help("Read one benchmark prompt from a file").required(false); auto& max_new_tokens = Argparse::add("-g|--max_new_tokens").help("Maximum generated tokens per prompt").required(false); auto& print_token_ids = Argparse::add("--print_token_ids").help("Print generated token IDs to stderr").required(false); + auto& benchmark_warmup = + Argparse::add("--benchmark_warmup").help("Unrecorded benchmark warmup requests").required(false); + auto& benchmark_samples = Argparse::add("--benchmark_samples").help("Measured benchmark requests").required(false); + auto& benchmark_jsonl = Argparse::add("--benchmark_jsonl").help("Fresh JSONL output path").required(false); + auto& benchmark_variant = Argparse::add("--benchmark_variant").help("Bound variant identity").required(false); + auto& benchmark_source_sha = Argparse::add("--benchmark_source_sha").help("Bound source SHA").required(false); + auto& expected_prompt_tokens = + Argparse::add("--expected_prompt_tokens").help("Fail if tokenized prompt length differs").required(false); + auto& require_device_telemetry = Argparse::add("--require_device_telemetry") + .help("Fail closed on incomplete or drifting CPU telemetry") + .required(false); // Argparse validates required options during parse(), so short-circuit help // before parsing to make `mllm-qwen3-5-runner --help` usable on its own. @@ -32,6 +54,7 @@ MLLM_MAIN({ } Argparse::parse(argc, argv); + mllm::configEngineWithArgs(engine_args); (void)help; @@ -55,7 +78,31 @@ MLLM_MAIN({ if (generation_limit <= 0 || generation_limit > cfg.max_cache_length) { throw std::invalid_argument("max_new_tokens must be between 1 and max_cache_length"); } + const bool benchmark_mode = benchmark_samples.isSet() || benchmark_jsonl.isSet() || benchmark_variant.isSet() + || benchmark_source_sha.isSet() || prompt_file.isSet() || expected_prompt_tokens.isSet() + || benchmark_warmup.isSet() || require_device_telemetry.isSet(); + if (prompt.isSet() && prompt_file.isSet()) { throw std::invalid_argument("prompt and prompt_file are mutually exclusive"); } if (prompt.isSet() && prompt.get().empty()) { throw std::invalid_argument("prompt must not be empty"); } + if (benchmark_mode) { + if (!prompt_file.isSet() || !benchmark_samples.isSet() || !benchmark_jsonl.isSet() || !benchmark_variant.isSet() + || !benchmark_source_sha.isSet() || !expected_prompt_tokens.isSet()) { + throw std::invalid_argument("benchmark mode requires prompt_file, benchmark_samples, benchmark_jsonl, " + "benchmark_variant, benchmark_source_sha, and expected_prompt_tokens"); + } + if (benchmark_samples.get() <= 0) { throw std::invalid_argument("benchmark_samples must be positive"); } + if (benchmark_warmup.isSet() && benchmark_warmup.get() < 0) { + throw std::invalid_argument("benchmark_warmup must be non-negative"); + } + if (generation_limit < 2) { throw std::invalid_argument("benchmark max_new_tokens must be at least 2"); } + if (benchmark_variant.get().empty() || benchmark_source_sha.get().empty()) { + throw std::invalid_argument("benchmark identities must not be empty"); + } + const std::filesystem::path jsonl_path(benchmark_jsonl.get()); + std::error_code size_error; + if (std::filesystem::exists(jsonl_path) && std::filesystem::file_size(jsonl_path, size_error) != 0) { + throw std::invalid_argument("benchmark_jsonl must be new or empty"); + } + } auto param = mllm::load(model_path.get(), file_version); mllm::models::qwen3_5::validateModelConfigMatch(cfg, param); @@ -67,49 +114,148 @@ MLLM_MAIN({ model.load(param); - fmt::print("\n{:*^60}\n", prompt.isSet() ? " Qwen3.5 One-shot CLI " : " Qwen3.5 Interactive CLI "); - if (!prompt.isSet()) { fmt::print("Enter 'exit' or 'quit' to end the session\n\n"); } + std::string benchmark_prompt; + if (benchmark_mode) { + std::ifstream prompt_stream(prompt_file.get(), std::ios::binary); + if (!prompt_stream) { throw std::invalid_argument("unable to read prompt_file"); } + benchmark_prompt.assign(std::istreambuf_iterator(prompt_stream), std::istreambuf_iterator()); + while (!benchmark_prompt.empty() && (benchmark_prompt.back() == '\n' || benchmark_prompt.back() == '\r')) { + benchmark_prompt.pop_back(); + } + if (benchmark_prompt.empty()) { throw std::invalid_argument("prompt_file must not be empty"); } - while (true) { - std::string prompt_text = prompt.isSet() ? prompt.get() : ""; - if (!prompt.isSet()) { - fmt::print("Prompt text (or 'exit/quit'): "); - if (!std::getline(std::cin, prompt_text) || prompt_text == "exit" || prompt_text == "quit") { break; } + const auto inputs = tokenizer.convertMessage({.prompt = benchmark_prompt}); + const auto prompt_length = inputs.at("sequence").shape()[1]; + if (prompt_length != expected_prompt_tokens.get()) { + throw std::invalid_argument( + fmt::format("prompt token count {} differs from expected {}", prompt_length, expected_prompt_tokens.get())); + } + if (prompt_length + generation_limit - 1 > cfg.max_cache_length) { + throw std::invalid_argument("benchmark prompt plus generation exceeds max_cache_length"); } - if (prompt_text.empty()) { continue; } - try { - // Each prompt is an independent conversation. Both the full-attention - // KV cache and every GDN recurrent/conv state must start empty. + std::ofstream jsonl(benchmark_jsonl.get(), std::ios::out | std::ios::trunc); + if (!jsonl) { throw std::invalid_argument("unable to open benchmark_jsonl"); } + const int warmup_count = benchmark_warmup.isSet() ? benchmark_warmup.get() : 0; + const int total_requests = warmup_count + benchmark_samples.get(); + for (int request_index = 0; request_index < total_requests; ++request_index) { + const bool warmup = request_index < warmup_count; + nlohmann::json record = { + {"schema", "mllm.qwen35.product_benchmark.r4.v1"}, + {"variant", benchmark_variant.get()}, + {"source_sha", benchmark_source_sha.get()}, + {"request_index", request_index}, + {"warmup", warmup}, + {"prompt_file", prompt_file.get()}, + {"prompt_tokens", prompt_length}, + {"max_new_tokens", generation_limit}, + {"cpu_op_threads", mllm::Context::instance().getCpuOpThreads()}, + }; + record["telemetry_before"] = mllm::examples::qwen3_5::benchmark::captureTelemetry(); + std::vector invalid_reasons; + if (require_device_telemetry.isSet() && require_device_telemetry.get()) { + invalid_reasons = mllm::examples::qwen3_5::benchmark::validateRequiredTelemetry(record["telemetry_before"]); + } + + const auto reset_start = std::chrono::steady_clock::now(); model.resetState(); - fmt::print("Processing...\n"); - auto inputs = tokenizer.convertMessage({.prompt = prompt_text}); - const auto prompt_length = inputs.at("sequence").shape()[1]; - if (prompt_length + generation_limit - 1 > cfg.max_cache_length) { - throw std::invalid_argument(fmt::format("prompt token count ({}) plus max_new_tokens ({}) exceeds " - "max_cache_length ({})", - prompt_length, generation_limit, cfg.max_cache_length)); + const auto reset_end = std::chrono::steady_clock::now(); + std::vector generated_token_ids; + const auto request_start = std::chrono::steady_clock::now(); + model.streamGenerate(inputs, + {{"max_length", mllm::AnyValue(generation_limit)}, + {"min_new_tokens", mllm::AnyValue(generation_limit)}, + {"do_sample", mllm::AnyValue(false)}}, + [&](int64_t token_id) { generated_token_ids.push_back(token_id); }); + const auto request_end = std::chrono::steady_clock::now(); + record["telemetry_after"] = mllm::examples::qwen3_5::benchmark::captureTelemetry(); + if (require_device_telemetry.isSet() && require_device_telemetry.get()) { + auto after_errors = mllm::examples::qwen3_5::benchmark::validateRequiredTelemetry(record["telemetry_after"]); + invalid_reasons.insert(invalid_reasons.end(), after_errors.begin(), after_errors.end()); + auto stability_errors = mllm::examples::qwen3_5::benchmark::validateStableTelemetry(record["telemetry_before"], + record["telemetry_after"]); + invalid_reasons.insert(invalid_reasons.end(), stability_errors.begin(), stability_errors.end()); } - fmt::print("\nResponse: "); - mllm::models::qwen3_5::Qwen3_5StreamingUtf8Decoder utf8_decoder; + const auto stats = model.perfStats(); + record["generated_token_ids"] = generated_token_ids; + record["reset_duration_us"] = std::chrono::duration_cast(reset_end - reset_start).count(); + record["request_wall_duration_us"] = + std::chrono::duration_cast(request_end - request_start).count(); + record["stats"] = { + {"valid", stats.valid}, + {"completed", stats.completed}, + {"total_duration_us", stats.total_duration_us}, + {"prefill_duration_us", stats.prefill_duration_us}, + {"decode_duration_us", stats.decode_duration_us}, + {"ttft_duration_us", stats.ttft_duration_us}, + {"prefill_tokens", stats.prefill_tokens}, + {"generated_tokens", stats.generated_tokens}, + {"decode_steps", stats.decode_steps}, + }; + if (!stats.valid) { invalid_reasons.push_back("invalid_performance_stats"); } + if (!stats.completed) { invalid_reasons.push_back("incomplete_generation"); } + if (stats.prefill_tokens != prompt_length) { invalid_reasons.push_back("prefill_token_count_mismatch"); } + if (stats.generated_tokens != generation_limit || stats.decode_steps != generation_limit - 1 + || generated_token_ids.size() != static_cast(generation_limit)) { + invalid_reasons.push_back("generation_length_mismatch"); + } + record["invalid_reasons"] = invalid_reasons; + record["status"] = invalid_reasons.empty() ? "ok" : "invalid"; + jsonl << record.dump() << '\n'; + jsonl.flush(); + if (!jsonl) { throw std::runtime_error("failed to write benchmark_jsonl"); } + if (!invalid_reasons.empty()) { + exit_code = 2; + break; + } + } + if (exit_code == 0) { fmt::print("Benchmark records: {}\n", benchmark_jsonl.get()); } + } else { + fmt::print("\n{:*^60}\n", prompt.isSet() ? " Qwen3.5 One-shot CLI " : " Qwen3.5 Interactive CLI "); + if (!prompt.isSet()) { fmt::print("Enter 'exit' or 'quit' to end the session\n\n"); } - for (auto& step : model.chat(inputs, {{"max_length", mllm::AnyValue(generation_limit)}})) { - if (print_token_ids.isSet() && print_token_ids.get()) { fmt::print(stderr, "TOKEN_ID:{}\n", step.cur_token_id); } - fmt::print("{}", utf8_decoder.append(tokenizer.detokenizeBytes(step.cur_token_id))); - std::fflush(stdout); + while (true) { + std::string prompt_text = prompt.isSet() ? prompt.get() : ""; + if (!prompt.isSet()) { + fmt::print("Prompt text (or 'exit/quit'): "); + if (!std::getline(std::cin, prompt_text) || prompt_text == "exit" || prompt_text == "quit") { break; } } - fmt::print("{}", utf8_decoder.finish()); + if (prompt_text.empty()) { continue; } + + try { + // Each prompt is an independent conversation. Both the full-attention + // KV cache and every GDN recurrent/conv state must start empty. + model.resetState(); + fmt::print("Processing...\n"); + auto inputs = tokenizer.convertMessage({.prompt = prompt_text}); + const auto prompt_length = inputs.at("sequence").shape()[1]; + if (prompt_length + generation_limit - 1 > cfg.max_cache_length) { + throw std::invalid_argument(fmt::format("prompt token count ({}) plus max_new_tokens ({}) exceeds " + "max_cache_length ({})", + prompt_length, generation_limit, cfg.max_cache_length)); + } - fmt::print("\n{}\n", std::string(60, '-')); - } catch (const std::exception& e) { - fmt::print("\nError: {}\n{}\n", e.what(), std::string(60, '-')); - if (prompt.isSet()) { exit_code = 1; } + fmt::print("\nResponse: "); + mllm::models::qwen3_5::Qwen3_5StreamingUtf8Decoder utf8_decoder; + + for (auto& step : model.chat(inputs, {{"max_length", mllm::AnyValue(generation_limit)}})) { + if (print_token_ids.isSet() && print_token_ids.get()) { fmt::print(stderr, "TOKEN_ID:{}\n", step.cur_token_id); } + fmt::print("{}", utf8_decoder.append(tokenizer.detokenizeBytes(step.cur_token_id))); + std::fflush(stdout); + } + fmt::print("{}", utf8_decoder.finish()); + + fmt::print("\n{}\n", std::string(60, '-')); + } catch (const std::exception& e) { + fmt::print("\nError: {}\n{}\n", e.what(), std::string(60, '-')); + if (prompt.isSet()) { exit_code = 1; } + } + if (prompt.isSet()) { break; } } - if (prompt.isSet()) { break; } - } - model.perfSummary(); + model.perfSummary(); + } } #ifdef MLLM_PERFETTO_ENABLE diff --git a/mllm/backends/cpu/kernels/common/gdn/gated_delta_net.cpp b/mllm/backends/cpu/kernels/common/gdn/gated_delta_net.cpp index f5b2fd066..0efd64698 100644 --- a/mllm/backends/cpu/kernels/common/gdn/gated_delta_net.cpp +++ b/mllm/backends/cpu/kernels/common/gdn/gated_delta_net.cpp @@ -188,7 +188,45 @@ void depthwiseCausalConvF32(const float* input, const float* weight, float* stat const int state_width = kernel_size - 1; for (int batch = 0; batch < batch_size; ++batch) { for (int token = 0; token < sequence_length; ++token) { - for (int channel = 0; channel < channels; ++channel) { + int channel = 0; +#if defined(__aarch64__) + // Production Qwen3.5 uses kernel_size 4, so the history is three taps + // wide and both the [C, K] weights and the [B, C, K - 1] history are + // contiguous across channels. vld3/vld4 deinterleave four adjacent + // channels into per-tap lanes, which lets the history shift happen in + // registers instead of two scalar loads and two scalar stores per + // element. + // + // The accumulation order matches the scalar body below exactly: a + // rounded multiply by the newest tap, then taps 0, 1, 2 fused in + // ascending order. Compilers contract the scalar `value += state * weight` + // into an FMA, so vfmaq_f32 reproduces it bitwise. The focused + // convolution oracle asserts that equality per toolchain rather than + // assuming it. + if (kernel_size == 4) { + const std::size_t token_base = (static_cast(batch) * sequence_length + token) * channels; + const std::size_t batch_state_base = static_cast(batch) * channels * state_width; + for (; channel + 4 <= channels; channel += 4) { + float* state_block = state + batch_state_base + static_cast(channel) * state_width; + const float32x4x3_t history = vld3q_f32(state_block); + const float32x4x4_t taps = vld4q_f32(weight + static_cast(channel) * kernel_size); + const float32x4_t current = vld1q_f32(input + token_base + channel); + + float32x4_t value = vmulq_f32(current, taps.val[3]); + value = vfmaq_f32(value, history.val[0], taps.val[0]); + value = vfmaq_f32(value, history.val[1], taps.val[1]); + value = vfmaq_f32(value, history.val[2], taps.val[2]); + vst1q_f32(output + token_base + channel, value); + + float32x4x3_t shifted; + shifted.val[0] = history.val[1]; + shifted.val[1] = history.val[2]; + shifted.val[2] = current; + vst3q_f32(state_block, shifted); + } + } +#endif + for (; channel < channels; ++channel) { const std::size_t state_base = (static_cast(batch) * channels + channel) * state_width; const std::size_t input_index = (static_cast(batch) * sequence_length + token) * channels + channel; const std::size_t weight_base = static_cast(channel) * kernel_size; diff --git a/mllm/models/ARGeneration.cpp b/mllm/models/ARGeneration.cpp index dacf00a7b..b625b9be0 100644 --- a/mllm/models/ARGeneration.cpp +++ b/mllm/models/ARGeneration.cpp @@ -3,6 +3,7 @@ #include #include +#include #include @@ -31,8 +32,13 @@ ARGenerationChatIterator::ARGenerationChatIterator(ARGeneration& gen, const ARGe top_k_ = args.count("top_k") ? args.at("top_k").get() : 0; top_p_ = args.count("top_p") ? args.at("top_p").get() : 0.0f; max_length_ = args.count("max_length") ? args.at("max_length").get() : gen.max_length_; + min_new_tokens_ = args.count("min_new_tokens") ? args.at("min_new_tokens").get() : 0; eos_token_id_ = args.count("eos_token_id") ? args.at("eos_token_id").get() : gen.eos_token_id_; do_sample_ = args.count("do_sample") ? args.at("do_sample").get() : gen.do_sample_; + if (max_length_ <= 0) { throw std::invalid_argument("max_length must be positive"); } + if (min_new_tokens_ < 0 || min_new_tokens_ > max_length_) { + throw std::invalid_argument("min_new_tokens must be between 0 and max_length"); + } step(); } @@ -94,6 +100,7 @@ void ARGenerationChatIterator::step() { Tensor logits = output["sequence"]; auto device = logits.device(); logits = logits.to(kCPU); + if (step_count_ + 1 < min_new_tokens_) { gen_->suppressEosLogits(logits, eos_token_id_); } int64_t next_token_id; if (use_sampling) { if (top_k_ > 0) { @@ -119,7 +126,7 @@ void ARGenerationChatIterator::step() { step_count_++; gen_->ar_steps_++; - if (gen_->isEosToken(next_token_id, eos_token_id_)) { + if (step_count_ >= min_new_tokens_ && gen_->isEosToken(next_token_id, eos_token_id_)) { gen_->generationEventEndTimePoint(); finished_ = true; return; @@ -144,8 +151,13 @@ ARGenerationOutputPast ARGeneration::generate(const ARGenerationOutputPast& inpu int top_k = args.count("top_k") ? args.at("top_k").get() : 0; float top_p = args.count("top_p") ? args.at("top_p").get() : 0.0f; int max_length = args.count("max_length") ? args.at("max_length").get() : max_length_; + int min_new_tokens = args.count("min_new_tokens") ? args.at("min_new_tokens").get() : 0; int eos_token_id = args.count("eos_token_id") ? args.at("eos_token_id").get() : eos_token_id_; bool do_sample = args.count("do_sample") ? args.at("do_sample").get() : do_sample_; + if (max_length <= 0) { throw std::invalid_argument("max_length must be positive"); } + if (min_new_tokens < 0 || min_new_tokens > max_length) { + throw std::invalid_argument("min_new_tokens must be between 0 and max_length"); + } bool use_sampling = do_sample || (temperature != 1.0f) || (top_k > 0) || (top_p > 0.0f); @@ -164,6 +176,7 @@ ARGenerationOutputPast ARGeneration::generate(const ARGenerationOutputPast& inpu if (i == 0) { prefillEventEndTimePoint(); } Tensor logits = output["sequence"]; + if (i + 1 < min_new_tokens) { suppressEosLogits(logits, eos_token_id); } int64_t next_token_id; if (use_sampling) { @@ -187,7 +200,7 @@ ARGenerationOutputPast ARGeneration::generate(const ARGenerationOutputPast& inpu decodeEventEndTimePoint(); } - if (isEosToken(next_token_id, eos_token_id)) { break; } + if (i + 1 >= min_new_tokens && isEosToken(next_token_id, eos_token_id)) { break; } // [B, S] past = output; @@ -219,8 +232,13 @@ void ARGeneration::streamGenerate(const ARGenerationOutputPast& input, const ARG int top_k = args.count("top_k") ? args.at("top_k").get() : 0; float top_p = args.count("top_p") ? args.at("top_p").get() : 0.0f; int max_length = args.count("max_length") ? args.at("max_length").get() : max_length_; + int min_new_tokens = args.count("min_new_tokens") ? args.at("min_new_tokens").get() : 0; int eos_token_id = args.count("eos_token_id") ? args.at("eos_token_id").get() : eos_token_id_; bool do_sample = args.count("do_sample") ? args.at("do_sample").get() : do_sample_; + if (max_length <= 0) { throw std::invalid_argument("max_length must be positive"); } + if (min_new_tokens < 0 || min_new_tokens > max_length) { + throw std::invalid_argument("min_new_tokens must be between 0 and max_length"); + } bool use_sampling = do_sample || (temperature != 1.0f) || (top_k > 0) || (top_p > 0.0f); @@ -243,6 +261,7 @@ void ARGeneration::streamGenerate(const ARGenerationOutputPast& input, const ARG Tensor logits = output["sequence"]; auto device = logits.device(); logits = logits.to(kCPU); + if (i + 1 < min_new_tokens) { suppressEosLogits(logits, eos_token_id); } int64_t next_token_id; if (use_sampling) { @@ -266,7 +285,7 @@ void ARGeneration::streamGenerate(const ARGenerationOutputPast& input, const ARG callback(next_token_id); - if (isEosToken(next_token_id, eos_token_id)) { break; } + if (i + 1 >= min_new_tokens && isEosToken(next_token_id, eos_token_id)) { break; } // [B, S] past = output; @@ -577,6 +596,19 @@ bool ARGeneration::isEosToken(int64_t token_id, int64_t primary_eos_token_id) co return token_id == primary_eos_token_id || additional_eos_token_ids_.contains(token_id); } +void ARGeneration::suppressEosLogits(Tensor& logits, int64_t primary_eos_token_id) { + auto last_logits = getLastLogits(logits); + if (last_logits.dtype() != kFloat32) { throw std::runtime_error("min_new_tokens currently requires float32 logits"); } + + const auto vocab_size = static_cast(last_logits.shape().back()); + auto* logits_data = last_logits.ptr(); + const auto suppress = [&](int64_t token_id) { + if (token_id >= 0 && token_id < vocab_size) { logits_data[token_id] = std::numeric_limits::lowest(); } + }; + suppress(primary_eos_token_id); + for (const auto token_id : additional_eos_token_ids_) { suppress(token_id); } +} + void ARGeneration::customEventStartTimePoint(const std::string& name) { completed_custom_events_.erase(name); custom_event_time_[name].first = PerformanceClock::now(); diff --git a/mllm/models/ARGeneration.hpp b/mllm/models/ARGeneration.hpp index d34f30651..3b82df3cc 100644 --- a/mllm/models/ARGeneration.hpp +++ b/mllm/models/ARGeneration.hpp @@ -80,6 +80,7 @@ class ARGenerationChatIterator { int top_k_; float top_p_; int max_length_; + int min_new_tokens_; int eos_token_id_; bool do_sample_; }; @@ -104,8 +105,21 @@ class ARGeneration { virtual ARGenerationOutputPast forward(const ARGenerationOutputPast& input, const ARGenerationArgs& args) = 0; + // Runs autoregressive generation over `input`, honoring the ARGenerationArgs + // contract. Supported keys include: + // - "max_length": maximum total new tokens (must be positive); + // - "min_new_tokens": minimum number of new tokens to generate. Valid + // bounds are 0 <= min_new_tokens <= max_length, enforced with + // std::invalid_argument. Before min_new_tokens is reached the EOS logit + // is suppressed and EOS does not terminate generation; afterwards EOS + // terminates as usual. Default 0 preserves legacy behavior. + // Throws std::invalid_argument when max_length <= 0 or min_new_tokens is out + // of [0, max_length]. virtual ARGenerationOutputPast generate(const ARGenerationOutputPast& input, const ARGenerationArgs& args); + // Streaming variant of generate: each newly generated token is passed to + // `callback`. The same min_new_tokens/max_length contract and validation + // apply as in generate. virtual void streamGenerate(const ARGenerationOutputPast& input, const ARGenerationArgs& args, const std::function& callback); @@ -124,6 +138,9 @@ class ARGeneration { int64_t sampleTopP(Tensor& logits, float p, float temperature); + // Iterator-based generation context. Stepping the returned context generates + // one token per step under the same ARGenerationArgs contract documented for + // generate, including min_new_tokens validation and EOS suppression. ARGenerationChatContext chat(const ARGenerationOutputPast& input, const ARGenerationArgs& args = {}); int64_t categoricalSample(const Tensor& probs); @@ -153,6 +170,8 @@ class ARGeneration { [[nodiscard]] bool isEosToken(int64_t token_id, int64_t primary_eos_token_id) const; + void suppressEosLogits(Tensor& logits, int64_t primary_eos_token_id); + bool do_sample_ = false; int eos_token_id_ = -1; std::unordered_set additional_eos_token_ids_; diff --git a/tests/core/ARGenerationTest.cpp b/tests/core/ARGenerationTest.cpp index 8fd74d830..cd4527d9c 100644 --- a/tests/core/ARGenerationTest.cpp +++ b/tests/core/ARGenerationTest.cpp @@ -53,11 +53,12 @@ Tensor makeInput(int sequence_length) { return input; } -std::vector runChat(FakeGeneration& model, int max_length, int eos_token_id = 3) { +std::vector runChat(FakeGeneration& model, int max_length, int eos_token_id = 3, int min_new_tokens = 0) { std::vector tokens; ARGenerationArgs args = { {"max_length", AnyValue(max_length)}, {"eos_token_id", AnyValue(eos_token_id)}, + {"min_new_tokens", AnyValue(min_new_tokens)}, }; for (const auto& step : model.chat({{"sequence", makeInput(3)}}, args)) { tokens.push_back(step.cur_token_id); } return tokens; @@ -156,6 +157,99 @@ TEST(ARGenerationPerformanceTest, AdditionalEosTokenStopsChat) { EXPECT_EQ(stats.generated_tokens, 2); } +TEST(ARGenerationPerformanceTest, ChatHonorsMinNewTokensForPrimaryEos) { + FakeGeneration model({3, 1, 3, 2}); + + EXPECT_EQ(runChat(model, 8, 3, 3), (std::vector{0, 1})); + + const auto stats = model.perfStats(); + EXPECT_TRUE(stats.completed); + EXPECT_EQ(stats.generated_tokens, 3); + EXPECT_EQ(stats.decode_steps, 2); +} + +TEST(ARGenerationPerformanceTest, ChatHonorsMinNewTokensForAdditionalEos) { + FakeGeneration model({2, 1, 2, 0}); + model.addEosToken(2); + + EXPECT_EQ(runChat(model, 8, 3, 3), (std::vector{0, 1})); + EXPECT_EQ(model.perfStats().generated_tokens, 3); +} + +TEST(ARGenerationPerformanceTest, BatchGenerateHonorsMinNewTokens) { + FakeGeneration model({3, 1, 3, 2}); + ARGenerationArgs args = { + {"max_length", AnyValue(8)}, + {"min_new_tokens", AnyValue(3)}, + {"eos_token_id", AnyValue(3)}, + }; + + const auto output = model.generate({{"sequence", makeInput(3)}}, args); + const auto& generated = output.at("generated_sequence"); + EXPECT_EQ(std::vector(generated.ptr(), generated.ptr() + generated.numel()), + (std::vector{0, 1, 3})); + EXPECT_EQ(model.perfStats().generated_tokens, 3); +} + +TEST(ARGenerationPerformanceTest, StreamGenerateHonorsMinNewTokens) { + FakeGeneration model({3, 1, 3, 2}); + ARGenerationArgs args = { + {"max_length", AnyValue(8)}, + {"min_new_tokens", AnyValue(3)}, + {"eos_token_id", AnyValue(3)}, + }; + std::vector generated; + + model.streamGenerate({{"sequence", makeInput(3)}}, args, [&](int64_t token) { generated.push_back(token); }); + + EXPECT_EQ(generated, (std::vector{0, 1, 3})); + EXPECT_EQ(model.perfStats().generated_tokens, 3); +} + +TEST(ARGenerationPerformanceTest, RejectsMinNewTokensAboveMaxLength) { + FakeGeneration model({0}); + EXPECT_THROW((void)runChat(model, 2, 3, 3), std::invalid_argument); +} + +TEST(ARGenerationPerformanceTest, MinNewTokensEqualToMaxLengthSuppressesEveryEarlyEos) { + FakeGeneration model({3}); + + EXPECT_EQ(runChat(model, 4, 3, 4), (std::vector{0, 0, 0})); + + const auto stats = model.perfStats(); + EXPECT_TRUE(stats.completed); + EXPECT_EQ(stats.generated_tokens, 4); + EXPECT_EQ(stats.decode_steps, 3); +} + +TEST(ARGenerationPerformanceTest, RejectsNonPositiveMaxLength) { + FakeGeneration model({0}); + EXPECT_THROW((void)runChat(model, 0), std::invalid_argument); + EXPECT_THROW((void)runChat(model, -1), std::invalid_argument); +} + +TEST(ARGenerationPerformanceTest, BatchGenerateRejectsInvalidMinNewTokens) { + FakeGeneration model({0}); + const auto run = [&](ARGenerationArgs args) { + (void)model.generate({{"sequence", makeInput(3)}}, args); + }; + EXPECT_THROW(run({{"max_length", AnyValue(3)}, {"min_new_tokens", AnyValue(4)}}), std::invalid_argument); + EXPECT_THROW(run({{"max_length", AnyValue(3)}, {"min_new_tokens", AnyValue(-1)}}), std::invalid_argument); + EXPECT_THROW(run({{"max_length", AnyValue(0)}}), std::invalid_argument); + EXPECT_THROW(run({{"max_length", AnyValue(-1)}}), std::invalid_argument); +} + +TEST(ARGenerationPerformanceTest, StreamGenerateRejectsInvalidMinNewTokens) { + FakeGeneration model({0}); + const auto run = [&](ARGenerationArgs args) { + (void)model.streamGenerate({{"sequence", makeInput(3)}}, args, [](int64_t) {}); + }; + EXPECT_THROW(run({{"max_length", AnyValue(3)}, {"min_new_tokens", AnyValue(4)}}), std::invalid_argument); + EXPECT_THROW(run({{"max_length", AnyValue(3)}, {"min_new_tokens", AnyValue(-1)}}), std::invalid_argument); + EXPECT_THROW(run({{"max_length", AnyValue(0)}}), std::invalid_argument); + EXPECT_THROW(run({{"max_length", AnyValue(-1)}}), std::invalid_argument); +} + } // namespace int main(int argc, char** argv) { diff --git a/tests/core/CMakeLists.txt b/tests/core/CMakeLists.txt index c72c7a4a3..fa45e466f 100644 --- a/tests/core/CMakeLists.txt +++ b/tests/core/CMakeLists.txt @@ -6,4 +6,9 @@ add_executable(Mllm-Test-Core-ARGeneration ARGenerationTest.cpp) target_link_libraries(Mllm-Test-Core-ARGeneration PRIVATE gtest MllmRT MllmCPUBackend) target_include_directories(Mllm-Test-Core-ARGeneration PRIVATE ${MLLM_INCLUDE_DIR}) +add_executable(Mllm-Test-Core-Qwen35BenchmarkHarness Qwen35BenchmarkHarnessTest.cpp) +target_link_libraries(Mllm-Test-Core-Qwen35BenchmarkHarness PRIVATE gtest MllmRT) +target_include_directories(Mllm-Test-Core-Qwen35BenchmarkHarness PRIVATE ${MLLM_INCLUDE_DIR} + ${PROJECT_SOURCE_DIR}/examples/qwen3_5) + include(GoogleTest) diff --git a/tests/core/Qwen35BenchmarkHarnessTest.cpp b/tests/core/Qwen35BenchmarkHarnessTest.cpp new file mode 100644 index 000000000..075c4205d --- /dev/null +++ b/tests/core/Qwen35BenchmarkHarnessTest.cpp @@ -0,0 +1,93 @@ +// Copyright (c) MLLM Team. +// Licensed under the MIT License. + +#include + +#include + +#include "benchmark_harness.hpp" + +namespace { + +nlohmann::json validSnapshot() { + return { + {"platform", "android"}, + {"affinity_cpus", {4, 5, 6, 7}}, + {"ceiling_vector", {3200000, 3200000, 4320000, 4320000}}, + {"cpus", + { + {{"cpu", 4}, + {"online", 1}, + {"cpuinfo_max_freq", 3200000}, + {"scaling_max_freq", 3200000}, + {"scaling_cur_freq", 2500000}, + {"scaling_governor", "schedutil"}}, + {{"cpu", 5}, + {"online", 1}, + {"cpuinfo_max_freq", 3200000}, + {"scaling_max_freq", 3200000}, + {"scaling_cur_freq", 2500000}, + {"scaling_governor", "schedutil"}}, + {{"cpu", 6}, + {"online", 1}, + {"cpuinfo_max_freq", 4320000}, + {"scaling_max_freq", 4320000}, + {"scaling_cur_freq", 3000000}, + {"scaling_governor", "schedutil"}}, + {{"cpu", 7}, + {"online", 1}, + {"cpuinfo_max_freq", 4320000}, + {"scaling_max_freq", 4320000}, + {"scaling_cur_freq", 3000000}, + {"scaling_governor", "schedutil"}}, + }}, + }; +} + +TEST(Qwen35BenchmarkHarnessTest, AcceptsCompleteStableTelemetry) { + const auto before = validSnapshot(); + auto after = before; + after["cpus"][0]["scaling_cur_freq"] = 1800000; + + EXPECT_TRUE(mllm::examples::qwen3_5::benchmark::validateRequiredTelemetry(before).empty()); + EXPECT_TRUE(mllm::examples::qwen3_5::benchmark::validateStableTelemetry(before, after).empty()); +} + +TEST(Qwen35BenchmarkHarnessTest, RejectsMissingCeiling) { + auto snapshot = validSnapshot(); + snapshot["cpus"][2]["scaling_max_freq"] = nullptr; + snapshot["ceiling_vector"][2] = nullptr; + + const auto errors = mllm::examples::qwen3_5::benchmark::validateRequiredTelemetry(snapshot); + EXPECT_NE(std::find(errors.begin(), errors.end(), "cpu6_missing_scaling_max_freq"), errors.end()); +} + +TEST(Qwen35BenchmarkHarnessTest, RejectsAffinityCeilingOnlineAndGovernorDrift) { + const auto before = validSnapshot(); + auto after = before; + after["affinity_cpus"] = {3, 4, 5, 6}; + after["ceiling_vector"][3] = 1689600; + after["cpus"][1]["online"] = 0; + after["cpus"][2]["scaling_governor"] = "powersave"; + + const auto errors = mllm::examples::qwen3_5::benchmark::validateStableTelemetry(before, after); + EXPECT_EQ(errors, (std::vector{"affinity_changed", "ceiling_vector_changed", "online_vector_changed", + "governor_vector_changed"})); +} + +TEST(Qwen35BenchmarkHarnessTest, MacCaptureFailsClosedWhenTelemetryIsRequired) { +#if defined(__APPLE__) + const auto errors = + mllm::examples::qwen3_5::benchmark::validateRequiredTelemetry(mllm::examples::qwen3_5::benchmark::captureTelemetry()); + EXPECT_NE(std::find(errors.begin(), errors.end(), "unsupported_telemetry_platform"), errors.end()); +#else + GTEST_SKIP() << "macOS-only telemetry rejection check"; +#endif +} + +} // namespace + +int main(int argc, char** argv) { + ::testing::InitGoogleTest(&argc, argv); + return RUN_ALL_TESTS(); +} diff --git a/tests/cpu/CMakeLists.txt b/tests/cpu/CMakeLists.txt index 14eb4d1c5..b5bb6ca50 100644 --- a/tests/cpu/CMakeLists.txt +++ b/tests/cpu/CMakeLists.txt @@ -6,6 +6,10 @@ add_executable(Mllm-Test-Qwen35-GDN Qwen35GDNTest.cpp) target_link_libraries(Mllm-Test-Qwen35-GDN PRIVATE gtest_main MllmCPUBackend) target_include_directories(Mllm-Test-Qwen35-GDN PRIVATE ${MLLM_INCLUDE_DIR}) +add_executable(Mllm-Test-Qwen35-GDN-Conv Qwen35GDNConvTest.cpp) +target_link_libraries(Mllm-Test-Qwen35-GDN-Conv PRIVATE gtest_main MllmCPUBackend) +target_include_directories(Mllm-Test-Qwen35-GDN-Conv PRIVATE ${MLLM_INCLUDE_DIR}) + add_executable(Mllm-Test-KaiW4A32Pack KaiW4A32PackTest.cpp) target_link_libraries(Mllm-Test-KaiW4A32Pack PRIVATE gtest_main MllmCPUBackend) target_include_directories(Mllm-Test-KaiW4A32Pack PRIVATE ${MLLM_INCLUDE_DIR}) diff --git a/tests/cpu/Qwen35GDNConvTest.cpp b/tests/cpu/Qwen35GDNConvTest.cpp new file mode 100644 index 000000000..46bcaa079 --- /dev/null +++ b/tests/cpu/Qwen35GDNConvTest.cpp @@ -0,0 +1,265 @@ +// Copyright (c) MLLM Team. +// Licensed under the MIT License. + +// Focused oracle for the GDN depthwise causal convolution. +// +// The reference below is an independent scalar implementation of the frozen +// contract. It is deliberately not routed through the production kernel, so a +// vectorized fast path inside depthwiseCausalConvF32 cannot validate itself. +// Both the output and the final history are compared bitwise: an output-only +// comparison would miss a corrupted history that only shows up in the next +// chunk. + +#include + +#include +#include +#include +#include +#include + +#include "mllm/backends/cpu/kernels/common/gdn/gated_delta_net.hpp" + +namespace { + +using mllm::cpu::gdn::depthwiseCausalConvF32; + +// Deterministic index-derived fill. No RNG, so every host reproduces the same +// bytes without carrying a seed through the evidence record. +float patternValue(std::size_t index, int salt) { + const auto scaled = static_cast((index * 37U + static_cast(salt) * 11U) % 251U); + return (scaled - 125.0F) / 64.0F; +} + +std::vector makeBuffer(std::size_t count, int salt) { + std::vector buffer(count); + for (std::size_t index = 0; index < count; ++index) { buffer[index] = patternValue(index, salt); } + return buffer; +} + +// Independent scalar reference for the frozen contract: +// input/output [B, S, C], weight [C, K], history [B, C, K - 1] updated in place. +void referenceDepthwiseCausalConv(const std::vector& input, const std::vector& weight, + std::vector& state, std::vector& output, int batch_size, + int sequence_length, int channels, int kernel_size) { + const int state_width = kernel_size - 1; + for (int batch = 0; batch < batch_size; ++batch) { + for (int token = 0; token < sequence_length; ++token) { + for (int channel = 0; channel < channels; ++channel) { + const auto state_base = (static_cast(batch) * channels + channel) * state_width; + const auto element = (static_cast(batch) * sequence_length + token) * channels + channel; + const auto weight_base = static_cast(channel) * kernel_size; + + float value = input[element] * weight[weight_base + state_width]; + for (int tap = 0; tap < state_width; ++tap) { value += state[state_base + tap] * weight[weight_base + tap]; } + output[element] = value; + + for (int tap = 0; tap + 1 < state_width; ++tap) { state[state_base + tap] = state[state_base + tap + 1]; } + state[state_base + state_width - 1] = input[element]; + } + } + } +} + +struct ConvCase { + int batch; + int sequence; + int channels; + int kernel; + bool non_zero_history; + + std::string describe() const { + return "B=" + std::to_string(batch) + " S=" + std::to_string(sequence) + " C=" + std::to_string(channels) + + " K=" + std::to_string(kernel) + " history=" + (non_zero_history ? "non-zero" : "zero"); + } +}; + +// Runs one case through the production kernel and the reference, and requires +// bitwise agreement on both the output and the final history. +void expectBitwiseAgreement(const ConvCase& test_case) { + const auto element_count = + static_cast(test_case.batch) * test_case.sequence * test_case.channels; + const auto state_count = + static_cast(test_case.batch) * test_case.channels * (test_case.kernel - 1); + + const std::vector input = makeBuffer(element_count, test_case.channels + test_case.sequence); + const std::vector weight = + makeBuffer(static_cast(test_case.channels) * test_case.kernel, test_case.kernel); + const std::vector initial_state = + test_case.non_zero_history ? makeBuffer(state_count, 7) : std::vector(state_count, 0.0F); + + std::vector kernel_state = initial_state; + std::vector kernel_output(element_count, 0.0F); + depthwiseCausalConvF32(input.data(), weight.data(), kernel_state.data(), kernel_output.data(), test_case.batch, + test_case.sequence, test_case.channels, test_case.kernel); + + std::vector reference_state = initial_state; + std::vector reference_output(element_count, 0.0F); + referenceDepthwiseCausalConv(input, weight, reference_state, reference_output, test_case.batch, test_case.sequence, + test_case.channels, test_case.kernel); + + ASSERT_EQ(kernel_output, reference_output) << "output mismatch for " << test_case.describe(); + ASSERT_EQ(kernel_state, reference_state) << "final history mismatch for " << test_case.describe(); +} + +TEST(Qwen35GDNConvTest, MatchesScalarReferenceAcrossFocusedMatrix) { + // Channel counts below, at, and above the natural four-channel vector width, + // including several that leave a tail. + const int channel_values[] = {1, 2, 3, 4, 5, 7, 130}; + const int sequence_values[] = {1, 2, 16, 69, 128, 517}; + const int kernel_values[] = {2, 3, 4, 5}; + + for (int batch : {1, 2}) { + for (int sequence : sequence_values) { + for (int channels : channel_values) { + for (int kernel : kernel_values) { + for (bool non_zero_history : {false, true}) { + ASSERT_NO_FATAL_FAILURE(expectBitwiseAgreement({batch, sequence, channels, kernel, non_zero_history})); + } + } + } + } + } +} + +TEST(Qwen35GDNConvTest, MatchesScalarReferenceAtProductionChannelWidths) { + // 6144 is the Qwen3.5-0.8B convolution width, 8192 the 4B width; both are + // multiples of four, so they never exercise a tail on their own. + for (int channels : {6144, 8192}) { + for (int sequence : {1, 16, 69, 128, 517}) { + ASSERT_NO_FATAL_FAILURE(expectBitwiseAgreement({1, sequence, channels, 4, true})); + } + } +} + +TEST(Qwen35GDNConvTest, MatchesScalarReferenceWithChannelTailAtProductionScale) { + // Production width minus one, two, and three channels: a full-width run plus + // a tail of three, two, and one channel respectively. + for (int channels : {8189, 8190, 8191, 6141}) { + ASSERT_NO_FATAL_FAILURE(expectBitwiseAgreement({1, 69, channels, 4, true})); + } +} + +TEST(Qwen35GDNConvTest, ChunkedPartitionsMatchOneShot) { + struct Partition { + int channels; + int kernel; + std::vector chunks; + }; + + const std::vector partitions = { + {8192, 4, {517}}, // one-shot reference + {8192, 4, {1, 516}}, // prefill then continuation + {8192, 4, {128, 128, 128, 128, 5}}, // multi-chunk + {130, 4, {1, 15, 53}}, // tail channels across chunks + {7, 5, {1, 1, 14}}, // generic kernel size, odd channels + {6144, 4, {69}}, + {6144, 4, {16, 16, 16, 21}}, + }; + + for (const auto& partition : partitions) { + int total_sequence = 0; + for (int chunk : partition.chunks) { total_sequence += chunk; } + + const auto element_count = static_cast(total_sequence) * partition.channels; + const auto state_count = static_cast(partition.channels) * (partition.kernel - 1); + const std::vector input = makeBuffer(element_count, partition.channels); + const std::vector weight = + makeBuffer(static_cast(partition.channels) * partition.kernel, partition.kernel); + const std::vector initial_state = makeBuffer(state_count, 7); + + std::vector one_shot_state = initial_state; + std::vector one_shot_output(element_count, 0.0F); + referenceDepthwiseCausalConv(input, weight, one_shot_state, one_shot_output, 1, total_sequence, partition.channels, + partition.kernel); + + std::vector chunked_state = initial_state; + std::vector chunked_output(element_count, 0.0F); + int consumed = 0; + for (int chunk : partition.chunks) { + const auto offset = static_cast(consumed) * partition.channels; + depthwiseCausalConvF32(input.data() + offset, weight.data(), chunked_state.data(), chunked_output.data() + offset, + 1, chunk, partition.channels, partition.kernel); + consumed += chunk; + } + + ASSERT_EQ(chunked_output, one_shot_output) + << "chunked output diverged for C=" << partition.channels << " K=" << partition.kernel; + ASSERT_EQ(chunked_state, one_shot_state) + << "chunked history diverged for C=" << partition.channels << " K=" << partition.kernel; + } +} + +TEST(Qwen35GDNConvTest, ResetBetweenRequestsReproducesFirstRequest) { + constexpr int kChannels = 8192; + constexpr int kKernel = 4; + constexpr int kSequence = 69; + constexpr auto kElements = static_cast(kSequence) * kChannels; + constexpr auto kStateCount = static_cast(kChannels) * (kKernel - 1); + + const std::vector input = makeBuffer(kElements, 5); + const std::vector weight = makeBuffer(static_cast(kChannels) * kKernel, kKernel); + + std::vector state(kStateCount, 0.0F); + std::vector first_output(kElements, 0.0F); + depthwiseCausalConvF32(input.data(), weight.data(), state.data(), first_output.data(), 1, kSequence, kChannels, + kKernel); + const std::vector first_state = state; + + // A second request that continues the history must differ, proving the + // history is really being carried. + std::vector continued_output(kElements, 0.0F); + depthwiseCausalConvF32(input.data(), weight.data(), state.data(), continued_output.data(), 1, kSequence, kChannels, + kKernel); + ASSERT_NE(continued_output, first_output); + + // Resetting the history reproduces the first request bit for bit. + std::fill(state.begin(), state.end(), 0.0F); + std::vector reset_output(kElements, 0.0F); + depthwiseCausalConvF32(input.data(), weight.data(), state.data(), reset_output.data(), 1, kSequence, kChannels, + kKernel); + + ASSERT_EQ(reset_output, first_output); + ASSERT_EQ(state, first_state); +} + +TEST(Qwen35GDNConvTest, RejectsNullBuffersAndInvalidGeometry) { + constexpr int kBatch = 1; + constexpr int kSequence = 2; + constexpr int kChannels = 4; + constexpr int kKernel = 4; + + std::vector input(static_cast(kSequence) * kChannels, 0.0F); + std::vector weight(static_cast(kChannels) * kKernel, 0.0F); + std::vector state(static_cast(kChannels) * (kKernel - 1), 0.0F); + std::vector output(input.size(), 0.0F); + + EXPECT_THROW(depthwiseCausalConvF32(nullptr, weight.data(), state.data(), output.data(), kBatch, kSequence, kChannels, + kKernel), + std::invalid_argument); + EXPECT_THROW( + depthwiseCausalConvF32(input.data(), nullptr, state.data(), output.data(), kBatch, kSequence, kChannels, kKernel), + std::invalid_argument); + EXPECT_THROW( + depthwiseCausalConvF32(input.data(), weight.data(), nullptr, output.data(), kBatch, kSequence, kChannels, kKernel), + std::invalid_argument); + EXPECT_THROW( + depthwiseCausalConvF32(input.data(), weight.data(), state.data(), nullptr, kBatch, kSequence, kChannels, kKernel), + std::invalid_argument); + + EXPECT_THROW(depthwiseCausalConvF32(input.data(), weight.data(), state.data(), output.data(), 0, kSequence, kChannels, + kKernel), + std::invalid_argument); + EXPECT_THROW( + depthwiseCausalConvF32(input.data(), weight.data(), state.data(), output.data(), kBatch, 0, kChannels, kKernel), + std::invalid_argument); + EXPECT_THROW( + depthwiseCausalConvF32(input.data(), weight.data(), state.data(), output.data(), kBatch, kSequence, 0, kKernel), + std::invalid_argument); + // kernel_size <= 1 leaves no history and is rejected by the frozen contract. + EXPECT_THROW( + depthwiseCausalConvF32(input.data(), weight.data(), state.data(), output.data(), kBatch, kSequence, kChannels, 1), + std::invalid_argument); +} + +} // namespace