feat(gemma4/rocm): FP8 resident MoE + SharedK-WMMA prefill (RDNA4)

[bakon11]

Summary

Gemma-4-26B MoE ROCm path for dual-GPU (gfx1201 lab):

• Multi-EOS from config.json ∪ generation_config.json ([1,50,106])
• Host expert LRU + dual-GPU native FP8 resident packs
• Fused ExpertGeGLU FP8 decode (M=1)
• Prefill: group-by-expert batch (T≥64), FP8 channel GEMM M>1 (weights stay FP8 on GPU), device gather + weighted scatter (no host weight dequant / no host hacc)
• Portable vt::fused_ops seams — no vt::rocm from models/

Lab (2× R9700, ROCm 7.2.4)

Prompt	Prefill tok/s	Quality
~800	~167	READY
~6k	~123	READY
~10k	~192	READY
Paris warm	—	Paris

hipBLASLt W8A8 microbench still slower than BF16 GemmEx on gfx1201 for pure matmul — not default.

Size

Over the ~900 product-line ratchet (MoE + HIP kernels). Happy to split further if maintainers prefer (decode/resident vs prefill batch).

Supersedes

Supersedes the bulk of #205 content (rebased to current main as a single tip commit). Please close #205 in favor of this + the serve reliability PR.

Test plan

• check-device-leakage
• check-env-doc / FEATURES / USAGE / ENVIRONMENT
• Live :8010 dual FP8 resident smoke (Paris/READY/prefill rates)
• CI

[bakon11]

Update after live Hermes proof (RDNA4 / dual R9700):

Full agent path works end-to-end (SOUL ~38k prefill → streamed decode → real reply). The blocking EngineCore fatals after prefill were missing ROCm V1 sampler ops (op 30 temperature, op 40 allowed_token_ids) + Hermes max_tokens=-1.

That fix is split out as a small tip-of-main PR for fast merge:
→ #234

This #228 remains the larger FP8 resident / fused decode / GPU prefill-MoE package. Prefer landing #234 first, then rebase this head.

[bakon11]

TESTERS WANTED — clear split: FP8 (lab) vs BF16/FP16 MoE (community)

What this PR is

FP8-first Gemma-4 MoE on ROCm: native FP8 resident experts, fused FP8 decode helpers, GPU prefill-batch MoE, etc. Lab-proven on 2× R9700 gfx1201 with Hermes-scale prompts (slow decode at 38k, but correct).

What it is not

A claim that every kernel here is the BF16 path. FP8 weight GEMM / FP8 resident upload do not apply to BF16 weights.

What does still matter for BF16/FP16 MoE testers

If you run BF16 or FP16 MoE on ROCm, you still want:

Land first	Why
#234	Full V1 sampler on ROCm + Hermes max_tokens=-1 (EngineCore op 30/op 40 fatals are dtype-agnostic)
#227	KV admit fail-fast + clamps (large-prompt hangs are dtype-agnostic)

Optional on this PR after those: any shared MoE control flow you can exercise in BF16 (group-by-expert prefill batch, dual-GPU placement) — report whether BF16 hits the same batch path or stays on serial expert loops.

Please do not treat as MoE validation

• Dense Gemma-4-12B BF16 — no experts

Ideal tester hardware

• BF16/FP16 Gemma-4-26B MoE (or similar) checkpoint
• Multi-GPU ROCm with headroom beyond ~52 GiB weights
• Host RAM ≫ 30 GiB preferred for clean loads

Report template

• GPU / ROCm / dtype / model
• Which PR builds (feat(rocm): V1 sampling ops for Hermes EngineCore #234 / fix(server): fail-fast unschedulable KV waits #227 / this)
• Short + long prompt results
• Prefill batch on/off if visible in logs (prefill-batch / serial)
• Crashes: no kernel for op, OOM, hang

Prefer merge order: #234 → #227 → rebase this FP8 MoE package.

Thanks — FP8 is the consumer-RDNA4 daily driver; BF16/FP16 MoE needs wider silicon than this lab can soak alone.

[localai-org-maint-bot]

The lab evidence is useful, but this introduces a large FP8 resident/prefill execution surface without automated kernel/model parity tests; the only added tests cover HF config parsing. Please add focused parity coverage for resident expert placement, fused decode, and grouped prefill (including multi-device behavior or a testable dispatch seam), then rebase after #234/#227 as proposed and clear the current repository gates. The current single-hardware smoke is not sufficient protection for this execution-path change.

[localai-bot]

Reviewed. The ROCm work looks solid and it builds clean on CUDA — but I'm asking for this to be split before it lands, and the reason is reviewability rather than anything wrong with the code.

Verified so far (GB10 sm_121a, rebased on current main)

BUILD_OK          full CUDA build, 0 errors
test_hf_config    192/192

SACRED gates are queued behind other GPU work on the box; I'll post them when they land. Note this is the first time any CI has run against this branch at all — gh pr checks 228 reports no checks, so before my build nothing had compiled it.

The blocking concern: 16 commits, 7 of them research(rocm)

8f86f58d research(rocm): direct-store concurrent WMMA flash (BM32)
696d79e5 research(rocm): concurrent WMMA QK — LDS limit, not mma hang
9e59b655 research(rocm): SharedK prefill + profile — QK compute-bound
2f69efd4 research(rocm): flash dispatch split — scalar tile vs 1-wave WMMA
76aaef2a research(rocm): 1-wave/CTA WMMA flash prefill (gfx1201-safe)
07a67548 research(rocm): rocWMMA QK in flash tile (gfx1201) + proven microbench
7b47381b research(rocm): SGLang-style flash tile prefill (opt-in) + steal base

Those are an exploration log — LDS limits, hangs, dispatch experiments, microbenches. That's genuinely valuable work and the findings deserve to be kept, but as a spec/record, not as shipping commits: 4,783 lines mixing exploration with production changes can't be reviewed as one unit, and nobody can later tell which commit is load-bearing and which was a dead end that happened to leave code behind.

Concretely, I'd like:

• One PR with the shipping perf work — the FP8 resident path, the fused decode, the GPU prefill MoE, BlockSumFast, the ExpertGeGLU fusion. That I can gate and merge quickly.
• The research findings folded into .agents/specs/ with the measurements and the dead ends named, which is where the "multi-wave mma hangs gfx1201" result actually belongs so the next person doesn't rediscover it.
• The dual-stream commit (b4e78df0, default off) can ride with the shipping PR since it's opt-in.

Third copy of the same max_tokens defect

if (req_max.has_value() && *req_max < 1) req_max = 4096;
if (req_max.has_value() && *req_max > 4096) {
  // soft cap; serving_chat may clamp further via VT_SERVER_MAX_NEW_TOKENS
}

This same edit is now in #227 (clamps to 16), #234 (4096), and here (4096) — three PRs patching one spot, which will conflict with each other. As noted on the others: -1 should become std::nullopt, which routes to the path we already have mirroring vLLM 1:1 (input_processor.cpp:175-180 ← input_processor.py:317-321) and generates to max_model_len - seq_len. A hardcoded ceiling silently truncates the long-context request the client explicitly asked to leave unlimited. The second if is also dead code — its body is only a comment.

Pick one PR to own that fix and drop it from the other two.

Smaller notes

• The branch touches hf_config.cpp again (the multi-EOS work from the closed fix(gemma4/rocm): multi-EOS, host expert LRU, safer device expert admit #205). Worth saying in the description whether this supersedes fix(gemma4/rocm): multi-EOS, host expert LRU, safer device expert admit #205, since that PR was closed rather than merged and its content shouldn't get lost silently.
• It also touches scheduler.cpp, overlapping fix(server): fail-fast unschedulable KV waits #227. Same conflict risk.

None of this is a criticism of the engineering — the gfx1201 findings are exactly the kind of thing worth having written down. It's about landing it in pieces each of which can be gated on its own.

[localai-bot]

Heads-up: the max_tokens clamp is now fixed on main as 2d2bdd04 (row/SERVE-MAXTOKENS-UNSET), so please drop it from this PR and rebase — otherwise the three PRs conflict over one spot.

A non-positive max_tokens (or max_completion_tokens) is now treated as unset on both request paths, which routes to max_model_len - seq_len in input_processor.cpp:175-180, mirroring vllm input_processor.py:317-321. No constant, no truncation of the long-context request that asked to be left unlimited.

RED before / GREEN after on dgx, the new test built against unpatched main:

before: THREW 'max_tokens must be at least 1, got -1';  28/29 cases, 2 assertions failed
after:  29/29 cases, 179/179 assertions

test_openai_protocol now pins the semantics (−1 and 0 unset on both paths, positive unchanged, and −1 yielding to a serving default). Semantics documented in docs/USAGE.md.

Everything else in your PR stands — this is only about the shared hunk.

[localai-bot]

Reviewed at 4e75e18d, merged onto current main and built for CUDA on a GB10 (sm_121a). Your shared-code changes do not disturb the backends that already work:

test_qwen3_paged_engine   184/184
test_qwen27_paged_engine  235/235   (SACRED)
test_hf_config            192/192

That was the open question from my last pass, so it is now answered and off the table. I can't review or run the ROCm kernels (no AMD GPU on this side), but the ~1600 lines under src/vt/rocm/ are device-local and gated by your own hardware — that half is yours to own.

The generation_config.json eos finding is real, and it was OUR bug

This is the most valuable thing in the PR and it is worth separating from the ROCm work. You are right that we were dropping Gemma-4's third eos id. I traced it: InputProcessor::generation_config_eos_ids_ is named for generation_config.json but was populated entirely from config.json — nothing in the LLM path ever opened generation_config.json (only Parakeet's loader did). So upstream's "secondary eos ids" were, in our port, a second copy of the primary list. Any chat model whose generation_config adds a turn-level stop ran past it to the length cap.

But the fix belongs one layer up from where you put it. Merging into cfg.raw["eos_token_id"] lands the ids in the field that supplies the primary eos, and we already have a faithful 1:1 port of update_from_generation_config (input_processor.cpp:72) that gates secondary ids on ignore_eos. Two consequences:

1. ignore_eos=true can no longer generate past those ids — upstream (sampling_params.py:645-655) keeps that gate.
2. The primary eos is generation_config_eos_ids_.front(). Your merge writes back a std::set-ordered union, so for any checkpoint whose config.json eos array is not already ascending, the sort silently changes which id is primary.

I've landed the corrected version separately (HfConfig::generation_config_eos_ids as its own field, raw untouched, unioned by append so the primary keeps its position) with red-first tests on both suites — credited to this PR for the find. Please drop the hf_config.cpp / FEATURES.md eos hunks on your next rebase; they'll conflict otherwise, and the behavior you want will already be on main.

Still the same blocker: this needs splitting

5,364 additions across 34 files, and six more perf commits landed while I was reviewing. I can gate it, but I cannot meaningfully review it as one unit, and neither can anyone else. Concretely, these are independent and each mergeable on its own:

• eos parity — now handled on main; drop it.
• prefill progress tok_s instrumentation (scheduler.cpp) — self-contained, reviewable in minutes.
• shared Gemma-4 model changes (gemma4.cpp, gemma4_moe.cpp, ~1,500 lines) — this is the part that needs real review, and it's the part currently buried.
• ROCm kernels (~1,600 lines) — your hardware, your gates.

Split in that order and I'll merge the first three quickly. One note on the scheduler hunk while it's fresh: MaybeLogPrefillProgress now also runs from update_from_output, so its static std::unordered_map<std::string, State> states is touched on every step for every request and is still never erased when a request finishes. The leak predates you — but this change makes it grow faster on a long-lived server, so an erase on completion would be a good thing to fold in.

[localai-bot]

Follow-up: the corrected eos fix is now on main as fb2dfac7 (#262), credited to this PR for the find. Please drop the hf_config.cpp and docs/FEATURES.md eos hunks when you next rebase — the behavior you wanted is already there, and those two files will conflict otherwise.

[bakon11]

Status (lab + PR hygiene)

Land order: #234 (ROCm V1 sampler, rebuilt tip-of-main, clamp dropped) → #227 (KV fail-fast, rebuilt tip-of-main, prefix probe + clamp dropped) → then this #228 rebase.

Maintainer notes absorbed for the next #228 push:

• Drop hf_config / FEATURES eos hunks (landed as fb2dfac7 / fix(parity): read generation_config.json eos ids — we only ever read config.json #262)
• Drop max_tokens clamp (main 2d2bdd04)
• Split still planned for reviewability (sampler/serve already split out)
• Parity tests for resident/prefill seams still TODO before merge ask

Lab prefill (not blocking this PR thread): SharedK-WMMA + PREFILL_GEMM_M=2048 ≈ 2010 eng @~11k on dual R9700 FP8 MoE; still working toward ~3k. Will rebase #228 onto main after #234/#227 clear conflicts.

[bakon11]

Lab tip 9a009eb6 now includes SharedK-WMMA + GEMM_M=2048 + PEER_ACT default ON (~2057 eng @11k). Body refreshed. Hold merge until #234 then #227 land; then full rebase + docs/pr-size pass. Best recipe handoff on lab host: bash ~/llms/scripts/vllm/start-gemma4-fp8-8010.sh → http://127.0.0.1:8010/v1.

[bakon11]

Rebased on tip-of-main — clamp + eos dropped

Per maintainer note (2d2bdd04 / row/SERVE-MAXTOKENS-UNSET) and follow-up on eos (fb2dfac7 / #262):

Dropped from this PR	Now owned by main
protocol.cpp -1→4096 / soft cap	Non-positive max_tokens = UNSET → max_model_len - seq_len
serving_chat clamp on before < 1	Only VT_SERVER_MAX_NEW_TOKENS for positive oversize (unchanged policy)
hf_config / FEATURES multi-EOS merge	generation_config eos via #262

Head: e6c8c4e9 (+ follow-up newline fix) — single rebuild tip on origin/main 848d4a87 (not a multi-commit rebase fight).

Still in this PR: FP8 resident MoE, SharedK-WMMA prefill, PEER_ACT default ON, GEMM_M=2048, SSE keepalives + deferred add_request (Hermes long-prefill), prefill progress tok_s, rocm_sample.hip.

Land order preference unchanged: #234 sampler → #227 KV fail-fast → this #228 (sampler TU also present here so either order can work after a final rebase).

Local: check-device-leakage / check-env-doc / check-public-doc-tables green; HIP server build OK on lab dual-R9700.

[bakon11]

Follow-up on the rebuild tip: CMakeLists restored from main (DSpark + HIP -O1 floor kept) with only rocm_sample.hip added — HIP server links clean on lab. Tip now 136d9d74.

[localai-bot]

Thanks for this — there's real engineering in here, and the FP8 resident-expert design plus the SharedK/WMMA prefill is exactly the direction the ROCm lane needs. I reviewed it in depth rather than landing it, and I'm asking for a split instead of a merge. Reasons below, with the evidence for each.

I want to be upfront about a limit that shapes everything: no CI job compiles ROCm and there is no AMD GPU on the maintainer box. So the three green build checks prove only that x86 CPU still builds — roughly 4,400 lines of HIP in this PR have never been compiled by anyone but you. That's why the ROCm half needs your hardware to land, and why I'm not going to "fix it up and merge" blind.

1. The three red checks have two root causes, and neither is a race

build-test-cpu and sanitize-cpu (thread) are the same failure — test_openai_api_server, 3 cases / 4 assertions. I grepped the TSan job log: ThreadSanitizer appears zero times. There's no data race; the thread lane is just the second lane that builds with VLLM_CPP_SERVER=ON.

documentation-checkpoint is mechanical: check-doc-checkpoint.py runs per commit, CMakeLists.txt is in both USER_USAGE_FILES and LANDING_SOURCE_FILES, and commit 136d9d74 touches only CMakeLists.txt while docs/USAGE.md moved in e6c8c4e9. Squashing the three commits into one fixes it — verified by building a squashed commit with git commit-tree and re-running the checker (exit 0).

2. The SSE rewrite regresses streaming for every model, on every backend

Same-binary A/B, identical tree, only the six serving files swapped:

Arm	Result
PR head, all 33 files	3 failed / 50
PR head, only serving_chat.cpp, serving_completion.cpp, serving_utils.{h,cpp}, request_logger.{h,cpp} reverted	50 passed / 50, 624/624 assertions

Three distinct defects:

• Streaming stops streaming. serving_completion.cpp:118-127 swaps the blocking engine_.get_output() for get_output_nowait() + sleep_for(50ms). RequestOutputCollector is single-slot and DELTA outputs coalesce (output_processor.h:71-74), so the consumer is asleep almost always and a 20-token stream collapses to one content frame + [DONE], finish_reason already set on the first chunk. TTFT becomes full generation time — the opposite of the goal. A real fix needs a timed wait (ready_.wait_for(deadline)), not a poll.
• Continuous usage dropped from the role frame — the include_continuous_usage buffering block and frame.usage = UsageInfo{...} were deleted. That's ported upstream behavior (serving.py:404-802).
• MaybeSsePing emits ":\n\n" and the data frame as one chunk, breaking one-frame-per-chunk; chunk.substr(6) then lands mid-"data". CI was fast enough never to fire it.

The keepalive/SSE idea is legitimate and worth having — it just needs its own PR, against the streaming tests it currently breaks.

3. Two changes reach much further than "gemma4/ROCm"

• rocm_matmul_hipblaslt.hip deletes the #ifndef HIPBLAS_V2 / #define HIPBLAS_V2 block that ce134e1d added to fix rocm_matmul_hipblaslt.hip:384:13: error: no matching function for call to 'hipblasGemmEx' #201/ROCm: -O0 RmsNorm triggers a CLR HostcallListener teardown deadlock #132 (ROCm 6.x gfx1100, six hipblasGemmEx overload failures). grep -c HIPBLAS_V2 on this head returns 0. Looks like a rebase clobber — the CMakeLists half of that fix survived. Since no CI compiles ROCm, this would land silently and re-break the person who reported rocm_matmul_hipblaslt.hip:384:13: error: no matching function for call to 'hipblasGemmEx' #201.
• platforms/rocm.cpp:64-71 flips support_static_graph_mode() from hard-false to backend().SupportsGraphCapture() (default true), turning on hipGraph capture for eleven other models — qwen3, qwen3_moe, qwen3_5, deepseek_v2, glm4_moe_lite, voxtral, qwen3_dflash — none touched or tested here. The comment it replaced warned this "flips only alongside a real capture implementation", and gemma4's own decode graph is left default-OFF as "dual peer unsafe". That belongs in its own change with its own evidence.

4. ROCm-tuned policy applied to CPU and CUDA, unguarded

gemma4_moe.cpp has exactly two preprocessor lines — #ifndef VLLM_CPP_HIP at :1570 and #endif at :1624 — wrapping only link stubs. Lines 1–1569, the whole rewritten RunGemma4Moe, compile and run everywhere:

• :497 if (!FreeBytes(d, &free_b)) return false; — only ROCm overrides Backend::DeviceMemoryInfo (backend.h:81 returns false), so MakeRoom always fails on CUDA/CPU and kills the device expert cache that worked before. Every expert of every top-k falls back to a host H2D.
• :943 flips VT_GEMMA4_FP8_NATIVE from OFF to ON for all backends, changing the accumulation arm and therefore rounding — any gemma4 token-exactness gate moves.
• :73 adds an unconditional per-expert Synchronize, commented as a ROCm crash workaround. On CUDA prefill that's ~2 full stream syncs per token per MoE layer.

5. Kernel-level findings I could not run, only read

These need your gfx1201 box to confirm or refute:

• rocm_paged_attn.hip:7-10, 1322 — VT_ROCWMMA_OK is gated on __gfx1200__/__gfx1201__, and the entire PagedAttnPrefillSharedKWmma body including its final store sits in the #else. On gfx1100/gfx1151/gfx942 (all in VLLM_CPP_HIP_SUPPORTED_ARCHS, and VLLM_CPP_HIP_ARCHITECTURES defaults to empty/auto-detect) it compiles to (void) casts and never writes out — but the host dispatch is default-ON with no arch guard and then returns, so no fallback runs and hipGetLastError() is clean. That's silent garbage prefill attention. rocm_arch.h::CapabilityFromGcnArch already exists to gate it.
• gemma4_moe.cpp:1450 — the peer-FP8 decode path omits weight_id, so the sticky dequant cache keys on a thread_local staging pointer that's identical for every expert. The function's own comment at :110-113 says the key must not be that pointer, and the prefill caller at :1156 passes fp8_weight_id correctly. Experts 2..7 would reuse expert 1's weights under their own router weights — coherent-looking wrong text.
• gemma4_moe.cpp:753 + :821-838 — ApplyExpertScaleRwKernel mutates rw in place, then the fall-through copies the already-scaled rw and applies hscale again → routing weights become w·scale².
• rocm_fp8_channel_gemv.hip:510-524 — pointer-table hipMemcpyAsync sources are caller host arrays (including &fp8_gu, a parameter address, and a stack const float wts[1]) that die immediately, with no sync before return. Whether ROCm's pageable staging masks it is a hardware question; by contract it's a use-after-free, and this repo has been bitten by exactly that before.
• rocm_gemma4_experts.hip:163-168 — resident FP8 bases are hipMalloc'd and never freed, though gemma4_moe.h:52 promises teardown.

Credit where it's due, verified clean: FP8 scale semantics are correct throughout (per-output-channel, never per-K, gate/up halves get separate scales, fp32 accumulation); the e4m3fn bit decode matches OCP apart from NaN mapping to 0.0; rocm_sample.hip is a faithful line-by-line port of cuda_sample.cu with no new reassociation; Backend::DeviceMemoryInfo is a textbook additive virtual; and the default-ON SharedKWmma LDS budget does fit in 64 KB.

6. Process

No .agents/specs/* entry, no linked issue, and zero tests for 6,029 product lines. Policy is a committed spec before implementation and ported tests in the same change. The evidence table is prefill tok/s and "Paris/READY" — there's no vLLM denominator and no token-exact or ratified distributional gate, and correctness gates come before throughput here.

Worth noting check-device-leakage.py counts CUDA tokens only, so gemma4.cpp:57-62, 512-550 (which #include <hip/hip_runtime.h> and call nine raw HIP APIs from the device-agnostic model layer) passed through a genuine blind spot. I'll open a separate issue to add a HIP bucket — that's on us, not you.

What I'd suggest

Split it in two. The A/B shows that reverting just the six serving files makes this PR green on everything CI can see:

1. SSE/keepalive PR — with a timed wait instead of the 50 ms poll, the role-frame usage restored, the ping as its own chunk, and tests.
2. ROCm/gemma4 PR — with HIPBLAS_V2 restored, the support_static_graph_mode flip pulled out, the gemma4_moe.cpp policy changes guarded by #if defined(VLLM_CPP_HIP) (or treating unknown memory info as unlimited, as before), an arch guard on the default-ON WMMA dispatch, and a spec + issue.

I'd rather have this land properly than land fast, and the ROCm half genuinely needs your hardware to gate — that's the one thing I can't do from here. Happy to review either half quickly as you split them, and to pre-check anything CPU-side.

[bakon11]

Status — hold merge; split next (back on track)

Land order in progress:

1. feat(rocm): V1 sampling ops for Hermes EngineCore #234 — rebuilt tip-of-main single commit a418791e (local gates green; CI re-run)
2. fix(server): fail-fast unschedulable KV waits #227 — rebuilt tip-of-main clean fail-fast only ac5725fb (no metrics/spec clobber; local gates green; CI re-run)
3. This feat(gemma4/rocm): FP8 resident MoE + SharedK-WMMA prefill (RDNA4) #228 — hold. Will split per maintainer review before next push:
   • PR A: SSE/keepalive with timed wait + streaming tests (no 50ms poll)
   • PR B: ROCm/gemma4 only — restore HIPBLAS_V2, no unguarded graph flip, HIP-guard MoE policy, WMMA arch guard, spec/issue/tests

Lab decode WIP (BPE heap, structural-tag default-off, FP8 HW cvt, KV-split/slide, etc.) lives on bakon11 wip/lab-decode-20260810 — not in this PR until split/gates.

[bakon11]

Split in flight

PR	Scope
#234	ROCm V1 sampler
#227	KV fail-fast
#316	SSE keepalives (get_for timed wait)
new ROCm split	feat/gemma4-rocm-fp8-split — FP8 MoE + SharedK-WMMA only

#228 = hold / close when replaced. No merge-as-is.

[bakon11]

Superseded by #317

This megapr head is behind main and carries mixed serve/docs/history noise.

Land path (split):

1. feat(rocm): V1 sampling ops for Hermes EngineCore #234 — ROCm V1 sampler (includes parallel RandomSample; Hermes temp>0 decode fix)
2. fix(server): fail-fast unschedulable KV waits #227 — fail-fast unschedulable KV waits
3. fix(server): SSE keepalives via collector timed wait #316 — SSE keepalives
4. feat(gemma4/rocm): FP8 resident MoE + SharedK-WMMA (split from #228) #317 — Gemma4/ROCm FP8 + SharedK-WMMA (this PR’s product payload)

Closing #228 in favor of #317 to avoid dual-tracking. Lab decode stack stays on bakon11 WIP branches, not this PR.

[bakon11]

Superseded by split #317 (+ #234/#227/#316). See comment.