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, tokenizer, forward, OpenAI-compatible HTTP/WebSocket server, pruning and LoRA finetune; optional resident CUDA path (cuBLAS GEMV, per-token graph).
  • Gemma 4 support (dual-geometry attention, SWA, IQ4_XS) and qwen3.5 hybrid GDN/MTP in oxidize-c; MoE router + expert stacks across ports.
  • AL-family quants (AL5/AL6/AL8/AL5_XS) and new IQ formats (IQ2_XS, IQ2_S, IQ4_NL) wired through oxidize-core, oxidize-cpp, and the de/quant pipelines.
  • Fused AVX2 int8 GEMV paths (Q4_0/AL5, IQ4_NL, IQ4_XS) in C/C++ for 2–5× faster decode; AL5 stays quantized in RAM.
  • Quantization tooling: oxidize-convert and oxidize-quantize accept AL targets; oxidize-quantize adds --context_length and --yarn_orig_ctx for YaRN.
  • New arch and RoPE: Hunyuan-MoE detection/mapping; optional YaRN scaling in inference.
  • Finetuning: oxidize-finetuning gains generation and a self-train loop, richer LoRA export/manifest helpers, and CLI wiring.
  • Ops and examples: Modal deploy/bench for oxidize-c, remote NUMA bench scripts, Prime Intellect pipelines, and Colab notebooks.

• Bug Fixes

  • Correct IQ1_S grid (real 2048-entry table) and IQ4_NL block sizing; add fused IQ4_NL×Q8_0 GEMV.
  • Fix Q3_K/Q5_K dequant pointer math; AL8/AL6/AL5_XS block width and keep-quantized gating (AL5 only).
  • Robust Gemma 4 load (vocab/hidden dimension derivation) and stable MoE expert slicing.
  • Harden server input parsing and limits; tokenizer tolerates edge cases.
  • oxidize-cpp NUMA: cap to physical cores and avoid SMT sibling pinning.

^{Written for commit 635f2ae. 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]

40 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-c/server.c">

<violation number="1" location="oxidize-c/server.c:207">
P0: A long or heavily escaped token makes `cn` larger than `sizeof chunk`, causing the SSE callback to read beyond `chunk`. Build the event in a dynamically sized buffer or clamp and preserve a complete JSON event before sending.</violation>

<violation number="2" location="oxidize-c/server.c:524">
P0: Any client that can reach the listener can invoke completions and the WebSocket session without the configured API key; the startup environment check therefore provides no access control. Apply the configured key checks before dispatching `/v1/*`, including the WebSocket upgrade.</violation>
</file>

<file name="oxidize-c/cuda.cu">

<violation number="1" location="oxidize-c/cuda.cu:783">
P0: Gemma models loaded with int8 KV caching crash on the first attention append: this branch is skipped, then `oc_cuda_forward` computes `kv_k = c->kv_k + ...` even though `c->kv_k` was never allocated for Gemma. Handle `kv_ck8` or allocate a device float private cache when the CUDA path does not implement int8 KV.</violation>

<violation number="2" location="oxidize-c/cuda.cu:812">
P1: A layer whose `lhd`/`lkv` product exceeds the global stride overruns `d_k` and `d_v` before the cache append. Size scratch buffers from the maximum per-layer geometry rather than only `c->kv_stride`.</violation>

<violation number="3" location="oxidize-c/cuda.cu:832">
P1: Resident matmuls are submitted on the cuBLAS handle's default stream rather than `c->stream`, so the captured Gemma graph does not have the same stream ordering and may reject or fail to capture cuBLAS work. Associating the handle with the created stream before the first GEMV keeps matmuls ordered and capturable.</violation>

<violation number="4" location="oxidize-c/cuda.cu:850">
P1: Loading a second GPU model makes the first model forward with the second model's weights, and freeing/reloading models leaks the entire device context. Tie the context to each model and add a teardown path that frees all CUDA and pinned allocations.</violation>

<violation number="5" location="oxidize-c/cuda.cu:905">
P1: Gemma GDN recurrence becomes incorrect after the first graph replay because the captured convolution ring keeps using the initial head/length. Update those values through device graph parameters or disable graph reuse for models containing GDN layers.</violation>

<violation number="6" location="oxidize-c/cuda.cu:921">
P1: A raw-weight GEMV can consume stale activation quantization: if the preceding projection is FP16, `wk` skips quantization and reuses `c->aq_*` from an earlier layer or token. Requantize each raw GEMV input or explicitly track whether the cached activation quantization matches the current vector.</violation>

<violation number="7" location="oxidize-c/cuda.cu:973">
P1: GDN layers with a raw-quantized beta projection disable beta updates on CUDA, changing the recurrence even though the beta weights were uploaded. Test `d || qd` and quantize the input for this GEMV.</violation>

<violation number="8" location="oxidize-c/cuda.cu:995">
P0: MoE layers take the dense FFN path here, ignoring router, expert, and shared-expert weights; with the normal MoE layout, the empty dense weights lead to invalid zero-dimension GEMV launches or invalid output. The CUDA path needs an MoE routing/expert implementation or must fall back for `g->is_moe` layers.</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/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="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="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/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="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>

<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="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-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="oxidize-c/gguf.c">

<violation number="1" location="oxidize-c/gguf.c:116">
P1: A GGUF-controlled header count can make the loader dereference a failed or truncated allocation, causing a crash on malformed input. Validating each count against `SIZE_MAX`/the allocation size and checking `calloc` before entering the parse loops would keep malformed models from taking down the process.</violation>

<violation number="2" location="oxidize-c/gguf.c:175">
P1: A tensor whose start is in range but whose payload is truncated is accepted and can later crash inference with SIGBUS or out-of-bounds reads. Computing the type-specific tensor size with checked dimension products and validating `offset + size` before returning would reject these files at load time.</violation>

<violation number="3" location="oxidize-c/gguf.c:207">
P1: Out-of-range numeric metadata can corrupt model configuration instead of being treated as absent. The accessor would be safer if it rejected non-finite, negative, and greater-than-`UINT32_MAX` values before the cast.</violation>
</file>

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

<violation number="1" location="oxidize-c/prune.c:194">
P1: Using the input path as the output path can truncate the mmap-backed source before KV and tensor bytes are copied, corrupting the model or causing a crash. Reject aliased input/output files or write to a temporary file and rename it after a successful write.</violation>

<violation number="2" location="oxidize-c/prune.c:353">
P1: Malformed GGUF dimensions can overflow the tensor-size arithmetic before the bounds check and turn the subsequent dequantization into an out-of-bounds write. Checked dimension conversions/products and checked offset addition would keep untrusted model files from reaching these loops.</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="oxidize-c/model.c">

<violation number="1" location="oxidize-c/model.c:944">
P1: Gemma4 layers with a query width different from the hidden width use the wrong attention width on CPU: `q_len` is taken from the output row count, causing incorrect head counts and invalid `wo` reads. Deriving it from `L->wo.cols` (the query input width) keeps the CPU path aligned with the CUDA path.</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="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-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>

<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]

P0: Any client that can reach the listener can invoke completions and the WebSocket session without the configured API key; the startup environment check therefore provides no access control. Apply the configured key checks before dispatching /v1/*, including the WebSocket upgrade.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At oxidize-c/server.c, line 524:

<comment>Any client that can reach the listener can invoke completions and the WebSocket session without the configured API key; the startup environment check therefore provides no access control. Apply the configured key checks before dispatching `/v1/*`, including the WebSocket upgrade.</comment>

<file context>
@@ -0,0 +1,627 @@
+    }
+    if (!body) { close(fd); continue; }
+
+    if (strncmp(req, "GET /v1/realtime", 16) == 0 &&
+        strcasestr(req, "Upgrade: websocket")) {
+      ws_session(fd, req, m, tok, temperature, draft_k);
</file context>

[cubic-dev-ai]

P0: A long or heavily escaped token makes cn larger than sizeof chunk, causing the SSE callback to read beyond chunk. Build the event in a dynamically sized buffer or clamp and preserve a complete JSON event before sending.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At oxidize-c/server.c, line 207:

<comment>A long or heavily escaped token makes `cn` larger than `sizeof chunk`, causing the SSE callback to read beyond `chunk`. Build the event in a dynamically sized buffer or clamp and preserve a complete JSON event before sending.</comment>

<file context>
@@ -0,0 +1,627 @@
+
+static void stream_token(uint32_t id, void *ud) {
+  stream_ctx *ctx = ud;
+  char frag[512], chunk[2048];
+  size_t fn = oc_detokenize(ctx->tok, id, frag, sizeof frag);
+  size_t ecap = fn * 6 + 64;
</file context>

[cubic-dev-ai]

P0: Gemma models loaded with int8 KV caching crash on the first attention append: this branch is skipped, then oc_cuda_forward computes kv_k = c->kv_k + ... even though c->kv_k was never allocated for Gemma. Handle kv_ck8 or allocate a device float private cache when the CUDA path does not implement int8 KV.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At oxidize-c/cuda.cu, line 783:

<comment>Gemma models loaded with int8 KV caching crash on the first attention append: this branch is skipped, then `oc_cuda_forward` computes `kv_k = c->kv_k + ...` even though `c->kv_k` was never allocated for Gemma. Handle `kv_ck8` or allocate a device float private cache when the CUDA path does not implement int8 KV.</comment>

<file context>
@@ -0,0 +1,1059 @@
+      g->k_norm = upload_fp32(L->k_norm, g->hd);
+      g->attn_post_norm = upload_fp32(L->attn_post_norm, c->h);
+      g->ffn_post_norm = upload_fp32(L->ffn_post_norm, c->h);
+      if (L->kv_ck) {   /* gemma: private per-layer cache on device */
+        g->my_kv_cap = (int)L->kv_cap;
+        size_t elems = (size_t)g->my_kv_cap * g->n_kv * g->hd;
</file context>

[cubic-dev-ai]

P0: MoE layers take the dense FFN path here, ignoring router, expert, and shared-expert weights; with the normal MoE layout, the empty dense weights lead to invalid zero-dimension GEMV launches or invalid output. The CUDA path needs an MoE routing/expert implementation or must fall back for g->is_moe layers.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At oxidize-c/cuda.cu, line 995:

<comment>MoE layers take the dense FFN path here, ignoring router, expert, and shared-expert weights; with the normal MoE layout, the empty dense weights lead to invalid zero-dimension GEMV launches or invalid output. The CUDA path needs an MoE routing/expert implementation or must fall back for `g->is_moe` layers.</comment>

<file context>
@@ -0,0 +1,1059 @@
+
+    /* FFN */
+    k_rms_norm<<<1, 256, 0, c->stream>>>(c->d_norm, c->d_x, g->ffn_norm, h, c->rms_eps);
+    gemv(c, g->gate, c->d_norm, c->d_gf);
+    gemv(c, g->up, c->d_norm, c->d_uf, false);
+    if (c->gemma)
</file context>

[cubic-dev-ai]

P1: Loading a second GPU model makes the first model forward with the second model's weights, and freeing/reloading models leaks the entire device context. Tie the context to each model and add a teardown path that frees all CUDA and pinned allocations.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At oxidize-c/cuda.cu, line 850:

<comment>Loading a second GPU model makes the first model forward with the second model's weights, and freeing/reloading models leaks the entire device context. Tie the context to each model and add a teardown path that frees all CUDA and pinned allocations.</comment>

<file context>
@@ -0,0 +1,1059 @@
+  fprintf(stderr, "cuda: resident, %.1f GB free of %.1f GB\n",
+          (double)freeb / 1e9, (double)totb / 1e9);
+  (void)wbytes; (void)gb;
+  g_ctx = c;
+  m->gpu_active = true;
+  return 0;
</file context>

[cubic-dev-ai]

P1: GDN layers with a raw-quantized beta projection disable beta updates on CUDA, changing the recurrence even though the beta weights were uploaded. Test d || qd and quantize the input for this GEMV.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At oxidize-c/cuda.cu, line 973:

<comment>GDN layers with a raw-quantized beta projection disable beta updates on CUDA, changing the recurrence even though the beta weights were uploaded. Test `d || qd` and quantize the input for this GEMV.</comment>

<file context>
@@ -0,0 +1,1059 @@
+      int nvh = g->n_v_heads, hk = g->head_k, hv = g->head_v;
+      gemv(c, g->qkv, c->d_norm, c->d_mixed);
+      gemv(c, g->ssm_alpha, c->d_norm, c->d_a, false);
+      if (g->ssm_beta.d) gemv(c, g->ssm_beta, c->d_norm, c->d_b, false);
+      else cudaMemset(c->d_b, 0, nvh * sizeof(float));
+      gemv(c, g->gdn_gate, c->d_norm, c->d_z, false);
</file context>

[cubic-dev-ai]

P1: A layer whose lhd/lkv product exceeds the global stride overruns d_k and d_v before the cache append. Size scratch buffers from the maximum per-layer geometry rather than only c->kv_stride.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At oxidize-c/cuda.cu, line 812:

<comment>A layer whose `lhd`/`lkv` product exceeds the global stride overruns `d_k` and `d_v` before the cache append. Size scratch buffers from the maximum per-layer geometry rather than only `c->kv_stride`.</comment>

<file context>
@@ -0,0 +1,1059 @@
+  A(&c->d_qg, qwide);
+  A(&c->d_q, c->n_heads * c->head_dim);
+  A(&c->d_gate, c->n_heads * c->head_dim);
+  A(&c->d_k, c->kv_stride);
+  A(&c->d_v, c->kv_stride);
+  A(&c->d_attn, c->n_heads * c->head_dim);
</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>