Add oxidize-c engine with Gemma 4, MoE support, and CUDA optimizations

[Jackson57279]

Summary by cubic

Adds a plain-C oxidize-c inference engine with Gemma 4 and MoE support, plus CUDA resident-forward optimizations and AL-family quantization across Rust/C++/C for faster, smaller GGUF serving.

New Features

• oxidize-c C11 runtime: GGUF loader/writer, BPE tokenizer, transformer forward, batched decode, paged KV cache with scheduler, autotune planner, 16 CLI subcommands, and an OpenAI-compatible HTTP/WebSocket server with SSE streaming.
• Gemma 4 (dual-geometry attention, SWA, IQ4_XS), qwen3.5 hybrid GDN/MTP, MoE router + expert stacks, and DeepSeek MLA across ports. Full forward passes now cover GLM/Hunyuan, Phi, and GPT-2/NeoX/Falcon.
• Speculative decoding (draft/verify, DFlash, Eagle-3, tree), flash attention, Q8 KV cache, LRU prefix + persistent context caches, attention sinks, and YaRN context scaling.
• Samplers: Mirostat, typical-p, tail-free, min-p, beam/contrastive search, and grammar constraints for structured output.
• CUDA backend keeps weights quantized in VRAM (MMQ kernels) with fused kernels, CPU-GPU offload, and OXK AVX2/AVX-512 GEMV paths.
• AL-family quants (AL5/AL6/AL8/AL5_XS) plus IQ formats (IQ2_XS, IQ2_S, IQ4_NL) wired through oxidize-core, oxidize-cpp, and de/quant pipelines.
• Tooling: AL targets in oxidize-convert/oxidize-quantize, merge strategies, Wanda pruning, SafeTensors→GGUF conversion, HF Hub downloads with resume, model registry, and perplexity evaluation.
• Server stack adds /v1/embeddings and /v1/responses routes, auth/rate-limit/metrics middleware, LoRA adapter inference, and a WebSocket realtime API.
• oxidize-finetuning gains generation, a self-train loop, and richer LoRA export/CLI wiring.
• Ops: pipeline/tensor parallelism, gossip/election/ring mesh, layer-wise inference, GPU cluster, WASM bridge, Modal deploy/bench, Colab notebooks, and CI workflows.

Bug Fixes

• Correct IQ1_S grid, IQ4_NL block sizing, Q3_K/Q5_K dequant pointer math, and AL8/AL6/AL5_XS block widths.
• Robust Gemma 4 load, stable MoE expert slicing, consistent speculative KV state, and paged scheduler accounting.
• Hardened server input parsing, streaming detokenization UTF-8 boundaries, and tokenizer edge cases.
• oxidize-cpp NUMA caps to physical cores; security advisories patched via vendored libp2p crates.

^{Written for commit df9f876. Summary will update on new commits.}

[cursor]

Bugbot is not enabled for your account, so this pull request was not reviewed.

Enable Bugbot in the Cursor dashboard to get automatic reviews on future PRs.

[chatgpt-codex-connector]

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 635f2ae231

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

[chatgpt-codex-connector]

Enforce API keys before serving generation routes

When --serve is bound to anything reachable beyond localhost, setting OXIDIZE_API_KEY only gates process startup; this request path never checks the Authorization header before accepting WebSocket sessions or the POST completion routes below. That makes the advertised API-key protection ineffective and allows any client that can reach the port to run inference. Add a header check against OXIDIZE_API_KEY/OXIDIZE_API_KEYS before dispatching /v1/realtime, /v1/chat/completions, and /v1/completions while leaving /health unauthenticated if desired.

Useful? React with 👍 / 👎.

[chatgpt-codex-connector]

Skip split-expert validation for dense MoE blocks

For MoE checkpoints with leading dense layers (the Hunyuan/DeepSeek layouts this change adds support for), config_.num_experts is nonzero but the early blk.* layers legitimately have dense ffn_gate/up/down weights and no _exps or ffn_gate.<i> tensors. This newly added completeness check still runs on those dense layers, the split-expert loop breaks at expert 0, and model construction throws incomplete split-expert tensor set before the dense layer can be used. Only enforce split-expert completeness after detecting a router/expert layer, or skip it for leading_dense_layers.

Useful? React with 👍 / 👎.

[cubic-dev-ai]

24 issues found across 106 files

Prompt for AI agents (unresolved issues)

Check if these issues are valid — if so, understand the root cause of each and fix them. If appropriate, use sub-agents to investigate and fix each issue separately.


<file name="scripts/prime-gemma4-31b-int4/nginx-dual.conf">

<violation number="1" location="scripts/prime-gemma4-31b-int4/nginx-dual.conf:15">
P0: The ALPHA load balancer will return 502 for every request because 127.0.0.1 in `gemma4-lb` is the Nginx container's loopback, not the host where ports 8000 and 8001 are bound. Align the containers on a shared network (or use host networking for Nginx) before using these upstream addresses.</violation>
</file>

<file name="oxidize-core/src/model/inference.rs">

<violation number="1" location="oxidize-core/src/model/inference.rs:267">
P1: CUDA fused attention ignores YaRN for models configured with `rope.scaling.type = "yarn"`, producing different positional embeddings from the CPU path; extending the CUDA RoPE interface or disabling the fused path when YaRN is active would preserve correctness.</violation>
</file>

<file name="scripts/prime-qwen35-self-train/strip_branding.py">

<violation number="1" location="scripts/prime-qwen35-self-train/strip_branding.py:46">
P1: Rendering `tokenizer_config.json` fails whenever its template uses `qwythos_identity`: this branch deletes the definition but leaves the references undefined. Applying the same `t.replace("qwythos_identity", '""')` transformation used for `chat_template.jinja` keeps both HF template copies executable.</violation>
</file>

<file name="scripts/prime-qwen35-self-train/run-self-train.sh">

<violation number="1" location="scripts/prime-qwen35-self-train/run-self-train.sh:69">
P1: A rerun can upload an incomplete adapter: the watcher sees a previous run's `finished` marker because this command appends to the shared log. Truncating the log for each run or using a run-specific completion marker would keep uploads tied to the current training job.</violation>
</file>

<file name="oxidize-core/src/compute/quantization/quant_simple.rs">

<violation number="1" location="oxidize-core/src/compute/quantization/quant_simple.rs:122">
P1: Q4_0 CUDA/ROCm GEMV now returns incorrect dot products for the newly encoded blocks because the quantizer uses split-half ordering while `gemv_q4_0_kernel` still uses interleaved vector indices. Updating the kernel indexing (including the shared ROCm build) would keep GPU results consistent with `dequantize_q4_0_scalar`.</violation>
</file>

<file name="scripts/gemma4_31b_al_remote.sh">

<violation number="1" location="scripts/gemma4_31b_al_remote.sh:21">
P1: User-controlled path, repository, and quant arguments are interpolated into remote shell command text without serialization, allowing shell-command injection on the target host. Passing a quoted/encoded payload over stdin or shell-escaping every SSH argument would preserve values without interpreting them as commands.</violation>

<violation number="2" location="scripts/gemma4_31b_al_remote.sh:53">
P0: The generated AL files cannot be valid full-model quants because this command feeds only shard 1 to the GGUF writer even though the downloaded model has two shards. The quantizer needs a multi-shard-safe input path/writer before this script can publish usable models.</violation>
</file>

<file name="scripts/prime-qwen35-self-train/setup-node.sh">

<violation number="1" location="scripts/prime-qwen35-self-train/setup-node.sh:54">
P1: Setup currently aborts during the build step because `oxidize-finetuning` is a documented failing crate in this checkout. A working finetuning revision or a fixed crate needs to be used before this one-time setup can complete.</violation>
</file>

<file name="scripts/publish_gguf_remote_hf.sh">

<violation number="1" location="scripts/publish_gguf_remote_hf.sh:12">
P1: The exit trap recursively deletes whatever `STAGING` points at, including pre-existing user data when `STAGING` is supplied or `.hf-staging` already exists. A unique `mktemp` subdirectory should be created and only that run directory cleaned up.</violation>

<violation number="2" location="scripts/publish_gguf_remote_hf.sh:36">
P1: A caller-controlled `HF_REPO` is interpolated into Python source, so a crafted repository argument can execute arbitrary Python locally. Passing the value as a Python argv/environment value with a quoted heredoc avoids code injection.</violation>

<violation number="3" location="scripts/publish_gguf_remote_hf.sh:52">
P1: The final publication step can report success while the README upload failed, because its fallback performs no HF operation and returns zero; it also creates the bogus `null` file on success. A single `--readme-only` call without the no-op fallback would preserve the failure status.</violation>
</file>

<file name="oxidize-cpp/src/tensor_cpu.cpp">

<violation number="1" location="oxidize-cpp/src/tensor_cpu.cpp:694">
P1: Quantized inference can crash on AVX2/F16C CPUs that lack FMA because this gate enables code containing `_mm256_fmadd_ps` without checking FMA support. Include `__builtin_cpu_supports("fma")` in the runtime condition, or provide a non-FMA implementation.</violation>
</file>

<file name="scripts/prime-qwen35-self-train/upload-hf.sh">

<violation number="1" location="scripts/prime-qwen35-self-train/upload-hf.sh:6">
P1: When `OUT_DIR` is customized without an explicitly exported `TRAIN_OUT`, the shell waits on the customized output but Python falls back to `$HOME/models/qwen35-9b-agent/self-train-out`; it can therefore upload an older/default run instead of the current adapter. Exporting the computed `TRAIN_OUT` or passing the shell value directly into Python would preserve the selected output path.</violation>
</file>

<file name="oxidize-c/tokenizer.c">

<violation number="1" location="oxidize-c/tokenizer.c:257">
P1: SentencePiece/Gemma prompts are segmented by local max-score pair merges here, not by the global score-maximizing unigram segmentation. A vocabulary can prefer `a+bc` overall while the heap commits to `ab`, producing different token IDs for the same prompt; the C tokenizer should use unigram/Viterbi segmentation.</violation>
</file>

<file name="notebooks/coding_agent_sft_datasets.ipynb">

<violation number="1" location="notebooks/coding_agent_sft_datasets.ipynb:241">
P1: The Colab workflow cannot clone its checkout because `/content/oxidize` is already non-empty when this command runs, and the ignored clone failure leaves the subsequent build pointed at that incomplete directory. Cloning into a clean path before creating `OUT_DIR` and checking the command status would make the documented Colab flow executable.</violation>

<violation number="2" location="notebooks/coding_agent_sft_datasets.ipynb:315">
P1: With `RUN_TRAIN=True`, the subprocess fails before self-training starts because `self-train` does not support `--max-tokens`; that option exists for `sft`, not `SelfTrainArgs`. Removing this pair or replacing it with a supported self-train option lets the training workflow run.</violation>
</file>

<file name="oxidize-cpp/src/gguf.cpp">

<violation number="1" location="oxidize-cpp/src/gguf.cpp:451">
P1: Recognizing Hunyuan without routing it through the HF decoder mapper leaves every Hunyuan tensor in its raw name; loading then cannot find required tensors such as `tok_embeddings.weight`. Include all Hunyuan aliases in the `map_tensor_name` decoder branch.</violation>
</file>

<file name="modal_oxidize.py">

<violation number="1" location="modal_oxidize.py:58">
P1: The public `serve` endpoint exits on every standard Modal deployment because the function has no `secrets=[...]` binding, so a local `OXIDIZE_API_KEY` is not present remotely; this also downloads the multi-GB model before failing. Inject the key through a Modal Secret and validate it before calling `ensure_model()`.</violation>
</file>

<file name="scripts/prime-gemma4-31b-int4/setup-node.sh">

<violation number="1" location="scripts/prime-gemma4-31b-int4/setup-node.sh:27">
P1: A pod without a preconfigured NVIDIA Docker runtime fails at the first `docker run --gpus all`, so setup never reaches the model download. The setup needs to install/configure and validate the NVIDIA Container Toolkit, or avoid requesting a GPU for the CPU-only download step.</violation>

<violation number="2" location="scripts/prime-gemma4-31b-int4/setup-node.sh:30">
P0: The INT4 checkpoint download never runs: this image's `vllm serve` entrypoint receives `python3 -c ...` as serve arguments, so the container exits or tries to serve a model named `python3`. Overriding the entrypoint to `python3` is needed for this download step.</violation>
</file>

<file name="scripts/hy3_1m_al_remote.sh">

<violation number="1" location="scripts/hy3_1m_al_remote.sh:27">
P1: Running without `HF_TOKEN` fails immediately with an unbound-variable error, even though the script explicitly defaults the token to empty. Expanding the optional positional parameter with an unset-safe form, or passing arguments through a safe serialized channel, preserves anonymous downloads.</violation>
</file>

<file name="oxidize-core/src/compute/tensor/kernels/q_kernels.rs">

<violation number="1" location="oxidize-core/src/compute/tensor/kernels/q_kernels.rs:1757">
P1: IQ4_NL GEMV returns incorrect results for every nontrivial block because the activation indices do not match the format's two-half weight layout. Pair each packed weight byte with Q8 positions `i` and `i + 16` instead of `2*i` and `2*i + 1`.</violation>
</file>

<file name="scripts/bench_iq_remote.sh">

<violation number="1" location="scripts/bench_iq_remote.sh:186">
P1: The Rust leg will not benchmark ordinary IQ GGUF models: `--engine standard` is not a supported normal-model engine and falls through to the DFlash loading path, which can fail before producing throughput. Using the supported `inference` engine would run the intended model benchmark.</violation>
</file>

<file name="scripts/prime-gemma4-31b-int4/llamacpp-unsloth.sh">

<violation number="1" location="scripts/prime-gemma4-31b-int4/llamacpp-unsloth.sh:14">
P1: The fallback exits before starting because both expected filenames are wrong for the files downloaded by the setup recipe: the target is `gemma-4-31B-it-qat-UD-Q4_K_XL.gguf` and the MTP drafter is `mtp-gemma-4-31B-it.gguf`; the latter is also excluded by the setup download pattern. Use the repository filenames consistently and include the `mtp-` file in the download step.</violation>
</file>

<sub>Reply with feedback, questions, or to request a fix.

Re-trigger cubic</sub>

[cubic-dev-ai]

P0: The ALPHA load balancer will return 502 for every request because 127.0.0.1 in gemma4-lb is the Nginx container's loopback, not the host where ports 8000 and 8001 are bound. Align the containers on a shared network (or use host networking for Nginx) before using these upstream addresses.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At scripts/prime-gemma4-31b-int4/nginx-dual.conf, line 15:

<comment>The ALPHA load balancer will return 502 for every request because 127.0.0.1 in `gemma4-lb` is the Nginx container's loopback, not the host where ports 8000 and 8001 are bound. Align the containers on a shared network (or use host networking for Nginx) before using these upstream addresses.</comment>

<file context>
@@ -0,0 +1,33 @@
+
+    upstream gemma4_backends {
+        least_conn;
+        server 127.0.0.1:8000 max_fails=3 fail_timeout=10s;
+        server 127.0.0.1:8001 max_fails=3 fail_timeout=10s;
+        keepalive 64;
</file context>

[cubic-dev-ai]

P1: A pod without a preconfigured NVIDIA Docker runtime fails at the first docker run --gpus all, so setup never reaches the model download. The setup needs to install/configure and validate the NVIDIA Container Toolkit, or avoid requesting a GPU for the CPU-only download step.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At scripts/prime-gemma4-31b-int4/setup-node.sh, line 27:

<comment>A pod without a preconfigured NVIDIA Docker runtime fails at the first `docker run --gpus all`, so setup never reaches the model download. The setup needs to install/configure and validate the NVIDIA Container Toolkit, or avoid requesting a GPU for the CPU-only download step.</comment>

<file context>
@@ -0,0 +1,49 @@
+
+echo "==> Pre-downloading INT4 QAT checkpoint (~20GB)"
+docker run --rm \
+  --gpus all \
+  -v "$HF_HOME:/root/.cache/huggingface" \
+  -e HF_TOKEN="${HF_TOKEN:-}" \
</file context>

[cubic-dev-ai]

P1: CUDA fused attention ignores YaRN for models configured with rope.scaling.type = "yarn", producing different positional embeddings from the CPU path; extending the CUDA RoPE interface or disabling the fused path when YaRN is active would preserve correctness.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At oxidize-core/src/model/inference.rs, line 267:

<comment>CUDA fused attention ignores YaRN for models configured with `rope.scaling.type = "yarn"`, producing different positional embeddings from the CPU path; extending the CUDA RoPE interface or disabling the fused path when YaRN is active would preserve correctness.</comment>

<file context>
@@ -232,11 +248,33 @@ impl Default for InferenceConfig {
+        theta: f32,
+        output: &mut [f32],
+    ) -> Result<(), crate::tensor::RopeError> {
+        apply_rope_f32_yarn(
+            input,
+            position,
</file context>

[cursor]

Review: what this is, how it works, what to steal

What it is. This PR is inverted. Head is master and base is c-port-1000tps, so GitHub is proposing to dump current master (~307 commits, 730 files) onto an old C-port experiment branch. It is not a feature PR into master.

How it works. c-port-1000tps diverged before the modular C rewrite. Everything this diff appears to “add” (Gemma 4, MoE, CUDA resident-forward, AL-family quants, OpenAI server, LongCat-in-C) already landed on master through later merges:

• Add oxidize-c engine (Gemma 4, MoE, CUDA int8 GEMV) and oxidize-cpp int8 fused GEMV #35 / oxidize-c v2: modular C11 port, threaded forward pass, kernel correctness fixes #38 — oxidize-c engine, then the v2 modular layout
• oxidize-c: LongCat-2.0 architecture support, and fix MLA attention #40 — LongCat-2.0 + MLA fixes in oxidize-c

The “updated today” activity is CodeQL re-running because master moved (#42), not new work on this branch.

Steal. Nothing unique versus current master. Do not merge this direction.

Closing as no merge intent: wrong base/head, superseded, conflicting.