Claude/fervent varahamihira cd1ab3 by gstoner · Pull Request #6 · gstoner/tessera

gstoner · 2026-05-07T02:06:38Z

No description provided.

@jit

Single rank-3 f32 tessera.flash_attn programs on @jit(target="apple_gpu") now execute through a purpose-built MSL kernel. Same online-softmax algorithm as flash-attention's algorithm 1, fused into a single kernel — avoids materializing the (B, Sq, Sk) score matrix entirely. Builds directly on the Phase 8.4.0 MSL infrastructure (kernel cache, MTLComputePipelineState dispatch). MLIR - Pass: FlashAttnToAppleGPU lowers rank-3 f32 tessera.flash_attn (head_dim <= 256) to func.call @tessera_apple_gpu_flash_attn_f32. Reads `causal` BoolAttr and optional `scale` FloatAttr; defaults scale = 1/sqrt(D). - Pipeline tessera-lower-to-apple_gpu-runtime extended to compose matmul + rope + flash_attn patterns - target_ir runtime-mode lowering accepts tessera.flash_attn as a recognized envelope source; emits msl_kernel + mps_dispatch with the embedded MSL source as a StringAttr - Bug fix: dedup-key tracking in _lower_tile_ops now only consumes a key when the op actually produces a runtime emission. Previously, filtered ops (tile.async_copy carrying source="tessera.flash_attn") would prematurely consume the slot and prevent the real compute op from emitting. Runtime - apple_gpu_runtime.mm: embedded flash_attn_f32 MSL kernel (online softmax with running max + denominator + per-thread output accumulator), one thread per (batch, q_row), grid (Sq, B, 1). Causal mask via i32 flag; head_dim <= 256 via per-thread stack array. Edge case: l == 0 (fully-masked row) returns zeros instead of NaN. - apple_gpu_runtime_stub.cpp: portable C++ flash_attn reference for non-Darwin builds (matches the MSL algorithm in plain C++) Python - driver.py: _APPLE_GPU_MSL_OPS extended with tessera.flash_attn; backend artifact picks the new symbol/framework/abi - target_ir.py: _APPLE_GPU_FLASH_ATTN_MSL_SOURCE constant + sha256 cache_key; emits msl_kernel for single-source flash_attn programs - runtime.py: _apple_gpu_dispatch_flash_attn dispatcher reads scale/causal kwargs and routes through the C ABI; loader gate now requires the flash_attn symbol (forces rebuild after 8.4.1) Tests - New lit fixture apple_gpu_flash_attn.mlir — positive (rank-3 static) and 2 negative cases (dynamic shapes, head_dim > 256). Uses the registered Tessera_FlashAttnOp directly (head_dim attribute is required by the verifier). - Three new unit tests in test_apple_backend_roadmap.py: MSL artifact contract (IR carries kernel source), end-to-end execution across 3 shapes with both causal and non-causal masks, runtime shim ABI correctness with direct ctypes invocation - Updated test_flash_attention_apple_gpu_gets_metal_kernel_contract -> test_flash_attention_apple_gpu_gets_msl_runtime_contract reflecting the contract change for single-flash_attn programs - Updated test_lower_tile_to_apple_gpu_target_ir_maps_fa4_to_metal_contract -> ..._maps_fa4_to_msl_runtime_contract - Updated compiler_examples.py: flash_attn_contract foundation example now claims runtime-executable on apple_gpu; manifest test reads claimed stages from the manifest instead of hardcoding artifact_only Verified on Apple Silicon (LLVM/MLIR 21, Metal active): 1953 unit tests passing; 9/9 Phase 8 lit fixtures passing against the in-tree tessera-opt. End-to-end flash_attn rtol=1e-4 vs numpy across multiple shapes (B/Sq/Sk/D = 1/4/4/8, 2/8/8/16, 1/16/32/64) for both causal and non-causal masks. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

Broadens the apple_gpu single-op runtime envelope with two more concrete custom kernels. Both are simple compared to flash-attention but exercise shapes Phase 8.4.0/8.4.1 didn't cover (row-wise reduction for softmax, pure elementwise for gelu) — useful coverage for the fusion work coming in Phase 8.4.3. MLIR - New passes: SoftmaxToAppleGPU lowers tessera.softmax (rank-2, f32, axis=-1) to func.call @tessera_apple_gpu_softmax_f32; GeluToAppleGPU lowers tessera.gelu (rank-2, f32) to func.call @tessera_apple_gpu_gelu_f32 - Pipeline tessera-lower-to-apple_gpu-runtime now composes 5 patterns: matmul + rope + flash_attn + softmax + gelu Runtime - apple_gpu_runtime.mm: two embedded MSL kernels: * softmax_f32 — 3-pass row-wise (max -> exp+sum -> divide), one thread per row * gelu_f32 — tanh-approximation matching the numpy reference, one thread per element - apple_gpu_runtime_stub.cpp: portable C++ references for non-Darwin Python - driver.py: _APPLE_GPU_MSL_OPS extended with tessera.softmax, tessera.softmax_safe, tessera.gelu; backend artifact picks the right symbol per op - target_ir.py: _APPLE_GPU_SOFTMAX_MSL_SOURCE and _APPLE_GPU_GELU_MSL_SOURCE constants + sha256 cache_keys; runtime gate accepts the new sources; _lower_apple_gpu_op emits msl_kernel - runtime.py: _apple_gpu_dispatch_softmax (axis=-1 fast-path, other axes fall back to numpy), _apple_gpu_dispatch_gelu, ctypes wrappers; loader gate now requires both new symbols (forces rebuild after 8.4.2) Tests - New lit fixture apple_gpu_softmax_gelu.mlir — positive (rank-2 static for both ops) and negative (dynamic shapes) paths. CHECK-DAG for the two runtime decls since their declaration order is implementation- defined. - Five new unit tests in test_apple_backend_roadmap.py: MSL artifact contract for each op, end-to-end execution across 3 shapes for each op, runtime shim ABI correctness for both symbols Verified on Apple Silicon (LLVM/MLIR 21, Metal active): 1958 unit tests passing; 10/10 Phase 8 lit fixtures passing against the in-tree tessera-opt. Softmax rtol=1e-5 vs numpy; gelu rtol=1e-5 vs the tanh-approximation reference across multiple shapes. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

Merge pull request #3 from gstoner/claude/fervent-varahamihira-cd1ab3

Tool Catches Why this one ruff Dead state (_ACTIVE_CONSTRAINTS), unused imports, security smells via S rules (S301: pickle.loads), docstring lint Fastest Python linter on the market, 800+ rules, zero-config sane defaults. ruff check is sub-second on Tessera's Python tree. The S ruleset alone would have flagged P2 #6. mypy (or pyright) Real type errors, missing return statements, Optional mis-use Tessera has type hints throughout; a strict pass catches signature drift between JitFn.compile_report callers/callees. pyright is faster but mypy is the canonical choice; either works. bandit Specifically the pickle, eval, exec, shell=True patterns Overlaps with ruff's S rules but bandit's the older standard. Optional if ruff S is on. Clang sanitizers (ASAN + TSAN) builds Heap-use-after-free in the tsrShutdown finding, the PerfettoTraceWriter race we hit during the architecture move, the strict-aliasing UB in Packing.h These are the highest-leverage tools for the C++ side. Add a -DTESSERA_ENABLE_SANITIZERS=ON build option that wires -fsanitize=address,undefined for ASAN+UBSAN and a separate -fsanitize=thread for TSAN. Run the smoke binary under each in CI. TSAN would have caught the PerfettoTraceWriter race the first time.

… Triage of all 7 findings While that runs, let me give you the triage summary so you can react. Triage of all 7 findings # Finding My verdict Action taken 1 mypy ratchet 7 errors STALE / FALSE None — mypy python/tessera/ reports 0 errors across 185 files; test_static_analysis_baseline.py 3/3 passes. The reviewer was looking at an older state 2 MockRankGroup repeated-collective race REAL ✅ Fixed — added per-rank generation counters; (kind, gen) tuple keys make every invocation unique; deposit and withdraw can no longer race 3 bind_symbol ctypes signature mutation race REAL ✅ Fixed — replaced getattr(handle, symbol) + .argtypes mutation with ctypes.CFUNCTYPE(restype, *argtypes)((symbol, handle)) so each signature gets a genuinely independent prototype; the whole bind now runs under _lock with a double-check after acquire 4 Graph IR cache no LRU/lock REAL but explicitly intentional Deferred — graph_ir_cache.py:18 documents the no-LRU design; turning it into a bounded LRU is a feature decision, not a bug fix 5 Monolith files REAL but architectural Deferred — apple_gpu_runtime.mm 6,988 LOC, vjp.py 4,262, etc. Splitting requires its own sprint 6 Queue verifier no-op stubs + shape_inference XFAIL REAL Deferred — confirmed QueueVerifiers.cpp is 3 trivial return success(); stubs and tests/tessera-ir/phase6/shape_inference.mlir is XFAIL: *. Next verifier-quality pass should cover them 7 trust_treedef=True default contradicts docstring REAL ✅ Fixed — flipped default to sentinel None and emit DeprecationWarning on implicit use; explicit True/False stays silent. Sets the trajectory to flip the implicit default to False in a future release without breaking current callers today Why I didn't fix #4, #5, #6 in this session #4 is documented intent (process-local cache for hot loops). Adding LRU + locks would change observable behavior; needs a design decision, not a drive-by fix. #5 is real but every file is structurally coherent — slicing them apart without a target architecture for each (Apple runtime by kernel family, vjp.py by op family, etc.) risks creating worse seams than the current single-file layout. #6 queue verifier stubs and shape_inference XFAIL are pre-existing surface area, not new debt. Worth a focused sprint along with the verifier-stub audit the reviewer flagged ("187 return success(); occurrences across src/"). Once the full unit sweep finishes I'll confirm green. Ran Full sweep result 5254 passed, 0 failures — same count as before the fixes, no regression. The 3 actionable findings (#2 MockRankGroup race, #3 bind_symbol mutation, #7 trust_treedef default) are landed and verified. Final verification mypy: 0 errors / 185 source files Full unit sweep: 5,254 passed, 7 skipped, 0 failures Static-analysis baseline: 3/3 passing Apple GPU MLA E2E proof: 6/6 passing MLIR verifier sprint guards: 64/64 passing Apple GPU buffer pool + JIT no-leak: 14/14 passing

waitUntilCompleted migration — COMPLETE (61/61 sites → wrapper) Batch 3: 8 MSL custom kernels (rope/flash_attn/softmax/gelu × f32/f16) via hand-edit Batch 4: 30 sites bulk-migrated via tools/scripts/migrate_wait_until_completed.py (a regex-driven, idempotent migration script with op-name inference from enclosing function declaration + per-op timeout heuristics) Batch 5: encode-session ts_enc_commit_wait — re-encoded Pattern-4 sequence inline (encode signal + commit + wait-with-timeout on the session's s->mtlcb) Drift gate strengthened: test_only_documented_waituntilcompleted_sites_remain proves only the 2 documented fallback paths remain in the runtime; any new waitUntilCompleted call site is a regression Single-command-buffer decode chain — SCAFFOLD LANDED (Audit Action 6 / table row #6) Design doc: docs/audit/single_command_buffer_decode_plan.md — 4-stage roadmap (scaffold → per-op encode variants → jit integration → full decoder benchmark) New C ABI: mpsg_encode_layer_norm_dev helper (parallel to existing mpsg_encode_bmm_dev) — appends a layer_norm to a shared MPSCommandBuffer via MPSGraph encodeToCommandBuffer: tessera_apple_gpu_layer_norm_dev_f32_enc(session, X_dev, gamma_dev, beta_dev, Y_dev, ...) — encode-session variant of layer_norm operating on device-resident tensors tessera_apple_gpu_session_commit_count — monotonic commit counter for the drift gate Stub parity for off-Darwin Python ergonomics: python/tessera/apple_gpu_batched.py — DeviceTensor dataclass + batched_session() context manager + bmm_enc / layer_norm_enc / device_tensor / device_empty helpers Headline proof (in test_apple_gpu_single_command_buffer.py): layer_norm(X) → bmm(_, W) produces the right numerical answer at fp32 tolerance AND submits exactly 1 command buffer (not 2). The single-cb invariant is structurally pinned.

Apple GPU surfaces. Session summary: flash_attn_dev_f32_enc + attention block on one cb — DONE (5 tests, 1 new C ABI + 1 helper + Python wrapper) Lifted the f32 flash_attn MSL kernel source to namespace scope (kFlashAttnF32Source); both the original dispatcher and the new encode_flash_attn_msl_dev helper share it (PSO cache dedupes by source SHA256, so zero runtime cost) New helper encode_flash_attn_msl_dev — appends a compute pass into the session's shared MPSCommandBuffer instead of creating its own. Metal's automatic hazard tracking orders downstream ops correctly. New C ABI tessera_apple_gpu_flash_attn_dev_f32_enc(session, Q, K, V, O, B, Sq, Sk, D, scale, causal) + stub parity Python flash_attn_enc(s, Q, K, V, ...) added to apple_gpu_batched.py Headline test: a full transformer attention block — layer_norm → 3 projections (bmm) → flash_attn → out_proj (bmm) = 6 ops — runs in EXACTLY 1 command buffer, with numerical correctness verified against numpy at fp32 tolerance, with causal-mask variant covered too. This is the audit Action 6 / Pattern row #6 architectural proof. Stride-alignment enforcement (Pattern 3 follow-on) — DONE (15 tests, new C ABI) New tessera_apple_gpu_row_major_strides_aligned(dims, rank, element_bits, ml_usage, strides_out) — implements all three of Apple's MTLTensorDescriptor.strides rules: innermost=1, 64-byte alignment for ML-usage byte+ dtypes, 128-byte alignment for sub-byte dtypes (regardless of usage flag) Existing _mlpkg_row_major_strides internal helper extended (backward compat): accepts element_bits + ml_usage args; element_bits=0 routes to legacy cumulative-product behavior (which is what fill/read paths need — they describe DENSE host buffer layout, not tensor-internal layout) Stub parity (math-only, no Metal calls) so the contract test runs on every host Tests cover: every combination of (fp32/fp16/bf16/int8/int4) × (ml_usage on/off) × (rank 1-4), edge cases (innermost dim already aligned vs. needs padding), backward-compat with the legacy helper, invalid-input rejection The helper is ready for callers that set MTLTensorDescriptor.strides — wiring the PK3 prepare_tensors path through it is a follow-on once Tessera owns packaged-kernel emission

The canonical rotor-invariant norm(rotor_sandwich(R, x)) previously dispatched two GA ops with an intermediate multivector round-trip through global memory. This fuses it into a single kernel. Kernel — tessera_apple_gpu_clifford_rotor_sandwich_norm_cl30_f32 in apple_gpu_runtime.mm: reuses the Cl(3,0) double-geometric-product expansion, keeps the 8-vector sandwich result in registers, and writes only the scalar norm per batch element (one dispatch, no intermediate spill). C++ reference fallback for non-Metal hosts. Surface — ga.rotor_sandwich_norm(R, x) is the direct entry (numpy fallback = norm(rotor_sandwich(...))); routes Cl(3,0) f32 through the bridge manifest. Fusion pass — _fuse_rotor_sandwich_norm rewrites the rotor_sandwich→norm adjacency into one clifford_rotor_sandwich_norm op, but ONLY when the intermediate sandwich is consumed exactly once and isn't the program return. Applied by the @clifford_jit decorator (and the lazy compile path); the structural lower_function_to_ir stays a faithful unfused 1:1 AST→IR projection. So @clifford_jit(norm(rotor_sandwich(...))) now compiles to a single-op plan + single dispatch. Wiring — _CLIFFORD_APPLE_GPU_FUSED + _GA_ATTR_TO_OP_NAME + a new _CLIFFORD_FUSION_OPS set (kept OUT of the 17 _CLIFFORD_PRIMITIVES so the "seventeen primitives" audits/counts stay exact); clifford_manifest_for now reports fusion ops (apple_gpu fused fp32 + CPU reference). _SENTINEL_SYMBOL bumped. Tests: test_clifford_rotor_sandwich_norm_fusion.py (8 — numerics, fusion pass, non-fusable cases, manifest, @clifford_jit single-op plan + single route). Updated the canonical-chain plan/route assertions across test_clifford_jit, test_compile_report{,_auto_emission}, test_benchmark_{ga_ebm,row}, test_compiler_audit, test_ga_backend_manifest to reflect the fused plan. mypy host+linux + ruff clean; generated docs in sync. Closes GA/EBM close-out gaps #1/#2/#5/#6; remaining: Apple-CPU GA/EBM native kernels (#3), Cl(1,3) kernels (#4, gated with a diagnostic), exp/log GA autodiff. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…spatch/combine) The first rung of the deferred north star. megamoe_forward shards experts across ranks (rank r owns expert block [r*Ep,(r+1)*Ep) and holds only those experts' weights — the memory win of expert parallelism) and routes tokens to the owning rank via the GShard / Switch 2x all-to-all: route (global replicated router) -> capacity-padded dispatch buffer keyed by destination-expert owner rank -> all-to-all DISPATCH -> local expert FFN via the fused GPU moe_swiglu_block (Ep ragged groups, one dispatch) -> all-to-all COMBINE (results back to originating rank) -> weighted scatter-combine. Capacity-based dispatch keeps every exchange buffer fixed-size so the all-to-all is uniform (the only kind the mock thread group expresses); overflow drops are reported as telemetry (MegaMoEResult.n_dropped). Per Decision #6, multi-rank runs in-process via MockRankGroup (threads) — production-shaped forward AND its own harness. expert_capacity() computes the per-expert slot count; megamoe_layer() is the single-call harness that shards x + the full expert set across ranks and gathers the (T,Kout) output. Correctness anchor: with capacity large enough to drop nothing, the gathered distributed output equals the single-device nn.functional.moe_layer exactly. 11 tests: distributed==single-device across world_size 1/2/4 x top_k 1/2, world_size=1 reduction, capacity-formula, capacity-drop telemetry, fp8 within budget, num_experts%world_size error, 2x-all-to-all token-order round-trip. Rung 2 (comm/compute overlap) and rung 3 (FP8xFP4 fused) follow. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

… actions The Still Open section had grown to ~105 lines where most bullets were 'X landed (2026-06-02); residual Y' with the residual buried under landed-work prose, plus one stale bullet contradicting the Finished section. Compressed to the actual signal, verified against source: - P1 (architectural): descriptor-driven dispatch. AppleKernelDescriptor exists + tested but runtime.py/driver.py don't import it (verified) — dispatch still pattern-matches op names; Target IR re-derives fusion independently. The old 'Target IR does too much' bullet is the same root issue, merged in. Next: route dispatch + Target IR recognizer through the descriptor registry. - P2 (incremental): feature-table selection. apple_target already exposes supports_bfloat / threadgroup_memory_capacity_bytes / max_threads_per_threadgroup / simdgroup_size, but only the softmax N-cap consumes it. Next: point bf16 gating / head_dim ceilings / threadgroup sizing at the existing helpers. - P2 (infra): systematic perf ratchets (Next-Work #6). - P3 (polish): auto_batch auto-detection + skip unused Graph-IR emission. Corrected the stale 'production packaged kernels are empty / 0 rows' bullet: PACKAGED_PRODUCTION_KERNELS has 7 rows + 7 committed fixtures (verified), and the whole authoring arc closed 2026-06-02 (PK8-PK8h, in Finished). The ~60 lines of SDK lane-grounding justified that completed work; moved out of Still Open, keeping only the two surviving constraints (auto-route opt-in by design due to the Metal ML-pipeline abort ceiling; Lane 2 DXIL out of scope). Doc-only; docs lint clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Closes the data-shaped CSV-canonical tail (COMPILER_AUDIT). The per-target capability matrices were the last row-table dashboards still byte-gated on Markdown; they now emit a canonical machine-readable CSV like the other 9. - apple_target_map.render_csv() — one row per op (op_name, family, cpu/gpu status/framework/dtypes/symbol/dispatch, proof_test, notes), sorted by (family, op_name); dtype lists comma-joined in a quoted cell. - gpu_target_map.render_csv(target) — per-target rows (op_name, family, status, dtypes, arch_min, tile_shape, expected_mfu, roofline, notes). - generated_docs registry: the 3 target-map entries gain csv_path + render_csv, so the CSV becomes the drift-gated artifact and the Markdown the human companion (verified: a CSV edit now trips the drift gate). 12 dashboards are now CSV-canonical; the remaining markdown-only docs are narrative rollups, not row tables. Content consolidation stays deliberately deferred (Next Work #6 reassessment). Tests: test_generated_docs_registry + test_apple_target_map green (40); drift gate in sync; mypy clean. Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…uler (#9c/#9d) #6 — DFlashDraft(nn.Module): holds every draft tensor as a Parameter (so it participates in parameters()/state_dict/.to(dtype)), forwards through the functional draft (cached or not), from_weights()/to_weights() round-trip. Verified module forward == functional (<1e-5), 5 + 11*N params registered, weight round-trip. #9b — RotatingDraftKVCache: bounds the draft context cache to the last max_size tokens (the draft analogue of MLX RotatingKVCache for sliding layers). Verified it caps per-layer length and, when unbounded, is identical to DraftKVCache. #9c/#9d — tessera.dflash_serve: dflash_generate_text (string-in/out via any encode/decode tokenizer) and DFlashScheduler (holds draft + stateful target, serves generation requests, greedy == AR). Verified scheduler greedy == AR and generate_text round-trips through a tokenizer. tests/unit/test_dflash_module_serve.py (5). ruff + mypy clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

* feat(dflash): cached drafting (#1) + sampling & rejection acceptance (#2) #1 — Draft KV cache. block_diffusion_attention gains return_ctx_kv to expose this step's projected+roped context KV; DraftKVCache accumulates it per layer. dflash_decoder_layer_cached / dflash_draft_forward_cached thread the cache so the draft attends to the full accumulated context instead of recomputing it. Verified: cached(accumulated) == non-cached(full context) to 1e-3, and the cache accumulates/advances correctly across steps. #2 — Non-greedy sampling + distribution-preserving acceptance. make_sampler (temperature / top-k / top-p, rng-reproducible), sampler_probs (matching truncated distribution), and dflash_speculative_verify (Leviathan rule: accept d_i w.p. min(1, p_t/p_d); on reject draw from normalize(relu(p_t-p_d)); bonus from the target's next-position distribution). Verified: greedy == argmax, top-k restricts support, draft==target accepts all, and the speculative-sampling theorem — the emitted token's marginal equals the target distribution (Monte Carlo, 40k draws, max abs err < 0.02). tests/unit/test_dflash_cached_sampling.py (7). ruff + mypy clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(dflash): stateful target with rollback (#3) + reference target model (#4) #4 — tessera.dflash_reference.ReferenceDecoderLM: a small numpy causal decoder (pre-norm MHA + SwiGLU, rope, tied/untied LM head) with a multi-layer hidden tap (the DFlash conditioning signal) and a stateless forward() that is the greedy-AR ground truth. random_decoder_lm builds one with small random weights. #3 — stateful KV cache + rollback: step(tokens) does causal cached decoding and appends roped-K/V per layer; rollback(n) drops the over-speculated tail. Verified that incremental step() (in 3 chunks) reproduces the stateless full-sequence forward to 1e-3, and that rollback restores exact cache state. dflash_generate_cached ties it together: cached draft (#1) + stateful target with rollback (#3) + greedy or rejection sampling (#2). Verified the whole efficient loop reproduces greedy AR exactly, is block-size independent, and sampling is reproducible + in-vocab. tests/unit/test_dflash_reference_target.py (5). ruff + mypy clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(dflash): GPU draft attention (#5) + training loss (#9a) + checkpoint I/O (#7) #5 — attention_fn threaded through dflash_decoder_layer / dflash_draft_forward (+ cached variants) so the whole draft forward runs its attention on the Apple GPU metal_runtime lane via apple_gpu_attention_fn. Verified the whole draft (2 layers) matches the numpy reference on Metal (rtol/atol 1e-3). The matmul- heavy projections/MLP/LM-head stay host-side (GPU gather/embedding is the remaining blocker for a single fully-jitted artifact). #9a — position-weighted block training loss: dflash_position_weights (wₖ = exp(-k/γ), normalized), dflash_block_loss (mean/sum/none) and the explicit gradient dflash_block_loss_grad. Verified the gradient vs finite differences (<1e-7), that a grad step lowers the loss, and reduction consistency. #7 — checkpoint I/O (tessera.dflash_io): a dependency-free safetensors reader/writer + HF state-dict <-> DFlashWeights mapping (transposing the nn.Linear (out,in) weights to the x@W (in,out) convention; embedding/LM head supplied from the target). load_dflash_weights reads a z-lab/*-DFlash draft; verified safetensors round-trip, the (out,in) transpose, and that round-tripped weights produce identical draft logits. tests: test_dflash_train_io.py (7) + #5 GPU draft case. ruff + mypy clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(dflash): nn.Module (#6) + rotating cache (#9b) + tokenizer/scheduler (#9c/#9d) #6 — DFlashDraft(nn.Module): holds every draft tensor as a Parameter (so it participates in parameters()/state_dict/.to(dtype)), forwards through the functional draft (cached or not), from_weights()/to_weights() round-trip. Verified module forward == functional (<1e-5), 5 + 11*N params registered, weight round-trip. #9b — RotatingDraftKVCache: bounds the draft context cache to the last max_size tokens (the draft analogue of MLX RotatingKVCache for sliding layers). Verified it caps per-layer length and, when unbounded, is identical to DraftKVCache. #9c/#9d — tessera.dflash_serve: dflash_generate_text (string-in/out via any encode/decode tokenizer) and DFlashScheduler (holds draft + stateful target, serves generation requests, greedy == AR). Verified scheduler greedy == AR and generate_text round-trips through a tokenizer. tests/unit/test_dflash_module_serve.py (5). ruff + mypy clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs(dflash): GQA repeat note (#8) + MASTER_AUDIT integration 1–9 landed #8 — annotate block_diffusion_attention's GQA: repeat is numerically exact; the native flash_attn_gqa kernel doesn't support DFlash's concat-context+proposal KV with an additive bias, so the reference materializes the repeat (no code change — correctness is unaffected). MASTER_AUDIT records DFlash integration items 1–9 as landed, with the two remaining gates flagged as external (real-checkpoint numerical parity needs a network download; a single fully-jitted GPU draft artifact needs GPU gather). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs: record DFlash + attn_bias in README, add docs/dflash.md - README status table: new "Speculative decoding — attn_bias substrate + DFlash block-diffusion draft" row (honest status: Python reference + attention core on Apple GPU metal_runtime; greedy spec-decode == greedy AR proven vs the MLX reference; real-checkpoint parity + fully-jitted GPU draft are external gates). - README: refresh stale Apple C ABI counts to the generated truth (256→264 symbols, 109→112 kernel families). - New docs/dflash.md: user-facing overview — the attn_bias substrate, the module map (dflash / dflash_reference / dflash_io / dflash_serve), a quick start, what's proven (per-test), and the external gates. Linked from the README doc index. docs lint passes; all links resolve. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs(api): fold DFlash + attn_bias public API into PYTHON_API_SPEC + CANONICAL_API PYTHON_API_SPEC.md: - flash_attn signature + parameter table gain attn_bias (additive (B,Sq,Sk) mask, Apple GPU flash_attn_bias_* / metal_runtime, causal+bias, broadcast fallback, positional-bias VJP). - Module hierarchy lists tessera.dflash / dflash_reference / dflash_io / dflash_serve. - New §18 "Speculative Decoding (DFlash)" documents the full public surface across the four modules + nn.functional.block_diffusion_attention / mask_token_block; TOC + Appendix A symbol index updated. CANONICAL_API.md: - flash_attn ops row gains attn_bias; functional table gains block_diffusion_attention + mask_token_block; new "tessera.dflash — Speculative Decoding (DFlash)" section with canonical names (one per concept) + a quick-start. Verified: check_spec_sync + docs lint pass; every documented symbol exists and every module __all__ symbol is documented (zero drift, confirmed programmatically). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs(api): close tessera.nn / tessera.ops doc-coverage gaps (full sweep) A programmatic sweep of the public surface vs both API docs found the surface ~99% documented with concentrated gaps; this closes them to zero. tessera.ops (312 ops): added the 2 missing — bmm (batched matmul + broadcast, Apple GPU metal_runtime) and fake_quantize (QAT STE) — to both ops tables. tessera.nn (77 public attrs): added the 10 missing functional layers (linear_general, lora_linear, spectral_norm, conv_transpose, avg/max/min/adaptive pool, gru_cell, simple_rnn_cell, bidirectional_scan) to CANONICAL's functional table, and the 10 missing Module classes (LinearGeneral, Einsum, LoRALinear, ConvTranspose1d/ConvTranspose, SpectralNorm, GRUCell/SimpleRNNCell, NativeSparseAttention, MixtureOfRecursions) to the stateful class table, with accurate constructor/forward signatures. Verified programmatically: tessera.ops, tessera.nn, and nn.functional.__all__ now have ZERO undocumented public symbols. check_spec_sync + docs lint pass. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…dtype CUDA/HIP backends (compiled in GPU CI; stubs here): - memcpy bounds-checks against dst/src->bytes (was unchecked → device overflow). - props() checks cudaGetDevice/GetDeviceProperties (was returning bogus zero-init props silently on failure). - free() always deletes the host wrapper even if cudaFree/hipFree fails (was leaking it via the early-return check macro). - destroyStream/destroyEvent capture and record the destroy return code. CPU backend (verified here): - waitEvent blocks on a condition_variable instead of busy-spinning with yield() (burned a core per waiter); Event gains a cv, recordEvent notifies. Collectives (verified via TesseraCollectiveRuntime build): - mock reduce_scatter/all_gather (NCCL + RCCL) size offsets by the wire dtype (new wireDTypeBytes) instead of hardcoded sizeof(float) — Decision #6. - TokenLimiter tracks max_ + inflight_ so set() shrinking the limit can't be over-credited by in-flight release()s. - Policy::chunkBytesForPath implements the documented per-path chunk granularity (NVLink 512Ki / PCIe 128Ki / RDMA 256Ki). Verified here: tessera_runtime, TesseraCollectiveRuntime, and the collective runtime smoke all build clean. The CUDA/HIP-guarded edits compile only under -DTESSERA_ENABLE_CUDA/HIP (GPU CI) — unverified on this arm64 host. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

@jit

…ckend_kernel (#106) * fix(nvidia): finish CUDA 13.2.1→13.3 propagation + verify sm_120 smem carve-out on silicon (#8, #14) The 13.2.1→13.3 pin bump (landed 2026-06-18) had only partially propagated. Source-of-truth pins were correct, but the surrounding surface had drifted. #8 — complete the bump: - validate_nvcc_compile.py: MIN_NVCC_VERSION (13,2,1)→(13,3,0); fix the user-facing error message and usage example (real drift, not prose). - capabilities.py: rename feature marker cuda_13_2_u1 → cuda_13_3 (4 entries) to match the existing C++ pin-test naming; update the test assertion + name accordingly. - Refresh stale "CUDA 13.2 U1" strings in diagnostics (2 fix-hints), backend_manifest notes, probe_collective_libs, and docs (CLAUDE.md, README, PROJECT_STRUCTURE, kernel inventory incl. the stale _CUDA_13_2_FEATURES identifier ref, execution plan). - Regenerate nvidia_sm90_target_map.{csv,md} (drift-gated). #14 — resolve the 100 KB vs 128 KB shared-memory carve-out on silicon: Measured on RTX 5070 Ti (driver 610.62 / CUDA UMD 13.3, nvcc 13.3.33, -arch=sm_120): sharedMemPerMultiprocessor = 102400 B (100 KiB) <- exact match cudaDevAttrMaxSharedMemoryPerBlockOptin = 101376 B (99 KiB) cudaDevAttrMaxSharedMemoryPerBlock = 49152 B (48 KiB) 100 KB wins; the release-note 128 KB is the unified data cache (Table 32), not the shared-memory carve-out. _SMEM_BYTES[SM_120]=102400 is correct. Recorded the on-silicon confirmation + the per-block 99 KiB opt-in ceiling (a lowering must cap dynamic smem below 101376 on sm_120) in gpu_target.py and BLACKWELL_SM120_EXECUTION_PLAN.md. Verification: drift gate clean (17 docs in sync); toolchain-pin and NVIDIA lane tests pass. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * spike(nvidia): sm_120 mma.sync.m16n8k16 PTX — emit→assemble→launch→compare on RTX 5070 Ti (#6) NVIDIA action item #6, proven end-to-end on real consumer-Blackwell silicon (RTX 5070 Ti, CC 12.0, CUDA 13.3, nvcc/ptxas 13.3.33). A hand-emitted, Tessera-style sm_120 bf16 mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 single-tile GEMM (D[16x8] = A[16x16]·B[16x8], f32 accumulate; warp-level mma.sync, NOT the Hopper warpgroup wgmma path) clears the full rung ladder: rung 2.5 (emit): raw PTX .version 9.3 / .target sm_120a (matches pin) rung 3 (assemble): ptxas --gpu-name=sm_120a -> 6168-byte cubin rung 5 (execute): Driver-API cuModuleLoadDataEx + cuLaunchKernel, max abs err 4.8e-7 vs CPU ref, 0/128 elements off Artifacts (committed under spikes/sm120_mma_sync/, with README + reproduce steps): the hand-written PTX, a CUDA inline-asm oracle that nails the m16n8k16 fragment layout, a driver-API run harness, and the smem device-query used to resolve #14. The .ptx is force-added past the global *.ptx ignore because here it is source, not a build artifact. Gotchas captured for productization: PTX must be ASCII (driver JIT ptxas rejects non-ASCII that standalone ptxas tolerates); f32 accumulator regs must be zero-initialized before the mma; CUDA 13.3 cuCtxCreate is v4 (use cuDevicePrimaryCtxRetain). Updated BLACKWELL_SM120_EXECUTION_PLAN.md sequencing: Stage A spike done; remaining work is wiring the path into ptx_emit.py + tsrRegisterGpuLauncher + an execute-compare oracle test. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * feat(nvidia): productize sm_120 mma.sync — emit_mma_sync_matmul_ptx + CUDA launch bridge + execute-and-compare (#6) Turns the spike (#6) into a proven in-tree path. All three pieces run on real consumer-Blackwell silicon (RTX 5070 Ti, CC 12.0, CUDA 13.3). 1. ptx_emit.emit_mma_sync_matmul_ptx (+ mma_sync_mnemonic, is_valid_mma_sync_bf16_shape, validate_mma_sync_ptx_structure): emits a COMPLETE, assemblable, launchable sm_120 mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 kernel (warp-level, register operands — no smem descriptors/TMA, unlike the WGMMA skeleton). ASCII-only (driver JIT ptxas rejects non-ASCII); accumulator explicitly zeroed. Output is the validated spike kernel. 2. CUDA launch bridge: test_conformance_execute_compare_nvidia.py registers an nvidia_mma_launcher via tsrRegisterGpuLauncher that loads the EMITTED PTX through the Driver API (cuModuleLoadDataEx -> cuLaunchKernel) and runs it through tsrLaunchKernel — mirrors the ROCm WMMA bridge template but exercises Tessera's own emitter output. Skip-clean (no nvcc / no runtime lib / no GPU); also asserts the negative UNIMPLEMENTED path. 3. CMake fix: src/runtime/CMakeLists.txt wired CUDA include/link for TESSERA_ENABLE_CUDA (mirror of the existing HIP block) — it was enabled but unbuildable (cuda_backend.cpp couldn't find cuda_runtime.h). Proof on box: emitter output -> ptxas sm_120a (rung 3) -> driver launch -> execute-and-compare vs CPU ref at max abs err 4.8e-7, 0/128 off (rung 5). Tests: test_ptx_emit (incl. ptxas-gated rung-3 assemble) + the NVIDIA execute-compare both pass; drift gate clean (17 docs in sync); mypy clean. Not flipped: the backend_manifest NVIDIA row stays artifact_only. ROCm's hardware_verified rows ship a real auto-built C-ABI runtime_symbol .so; the NVIDIA kernel here runs via the emitted PTX through a harness launcher, not a shipped libtessera_nvidia_*.so. Flipping to hardware_verified honestly needs that shipped runtime lib (the well-scoped next step). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * spike(nvidia): general mma.sync GEMM across shapes for bf16/f16/tf32/fp8 (e4m3,e5m2) Extends the #6 spike from a single 16x8x16 tile to a GENERAL tiled/K-looped GEMM (arbitrary M/N/K, ragged zero-padded) and sweeps every CC 12.0 Tensor-Core input dtype, all NVRTC-compiled (compute_120) and execute-and-compared on the RTX 5070 Ti across 7 shapes: bf16 m16n8k16 worst maxerr 9.5e-6 f16 m16n8k16 worst maxerr 1.7e-5 tf32 m16n8k8 worst maxerr 2.5e-5 e4m3 m16n8k32 bit-exact (0) e5m2 m16n8k32 bit-exact (0) Each dtype uses its own MMA shape + fragment layout (16-bit: 2 elems/reg; tf32: one tf32/reg, K=8; fp8: 4 bytes/reg, K=32). The host reference quantizes inputs to the same dtype the hardware consumes (OCP fp8 decode), so fp8 matches bit-for-bit. nvgemm_proto.cpp is the bf16/f16 general GEMM that backs the shipped libtessera_nvidia_gemm.so; nvgemm_dtypes.cpp is the full dtype sweep. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * spike(nvidia): NVFP4 m16n8k64 block-scale MMA — assembles + executes on sm_120a (#9, partial) The consumer-Blackwell headline (action item #9): warp-level block-scaled NVFP4 (e2m1 data + ue4m3 per-16-block scale), m16n8k64. Grounded on the RTX 5070 Ti (ptxas/runtime = source of truth): - mma.sync.aligned.m16n8k64.row.col.kind::mxf4nvf4.block_scale.scale_vec::4X. f32.e2m1.e2m1.f32.ue4m3 ASSEMBLES and EXECUTES on sm_120a. Accepted operand grammar: {d4},{a4},{b2},{c4},{sfa},{byteid_a,tid_a},{sfb},{byteid_b,tid_b}. - Arch-specific: only sm_120a SASS accepts it; base compute_120/sm_120 PTX rejects .kind::mxf4nvf4 / .block_scale / .scale_vec::4X. Build with -gencode arch=compute_120a,code=sm_120a. NOT yet numerically verified (honest, per the grounding rule): the on-box CUDA 13.3 headers expose only the datacenter tcgen05 block-scale variant, not the warp mma.sync scale-distribution + ue4m3 encoding semantics. A scale-byte sweep shows the scale is multiplicative (0x00 -> ~0) but standard e4m3 1.0 (0x38) yields ~3e9, so a guessed scale layout does not match a reference. Recorded as the precise next step; not claimed working. Artifacts: nvfp4_probe.cu (assemble probe), nvfp4_gemm.cu (per-lane fragment harness + scale sweep), README NVFP4 section. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * feat(nvidia): ship libtessera_nvidia_gemm.so + flip sm_120 matmul to hardware_verified (#6) Completes execution-plan step 3: NVIDIA's first hardware-verified backend_kernel row, mirroring the ROCm Strix Halo shipped-symbol pattern. Proven on the RTX 5070 Ti (CC 12.0, CUDA 13.3). Shipped runtime symbol: - src/.../tessera_gpu_backend_NVIDIA/runtime/cuda/tessera_nvidia_gemm.cpp + CMake target `tessera_nvidia_gemm` -> libtessera_nvidia_gemm.so. Exports tessera_nvidia_mma_gemm_{bf16,f16,tf32,e4m3,e5m2}: a general tiled/K-looped warp-level mma.sync GEMM (ragged M/N/K zero-padded, f32 accumulate), NVRTC-compiled (compute_XX) for the device arch at first call, launched via the CUDA driver API. Per-dtype MMA shape: bf16/fp16 m16n8k16, tf32 m16n8k8, fp8 e4m3/e5m2 m16n8k32. Built only when TESSERA_ENABLE_CUDA (self-gated subdir). Numerical proof: - tests/unit/test_nvidia_mma_runtime_symbol.py dlopens the shipped .so and validates all 5 dtypes vs numpy/ml_dtypes references across aligned + ragged shapes. Skip-clean (no GPU/NVRTC -> rc=2). Manifest flip: - _NVIDIA_HARDWARE_VERIFIED["matmul"] + the builder injects a hardware_verified row for nvidia_sm120 (runtime_symbol + execute_compare_fixture), replacing the artifact_only row for that arch only; sm_80/90/100 stay artifact_only (proven only on sm_120). _NUMERICAL_FIXTURES[("matmul","nvidia_sm120")] added. dtypes=(bf16,fp16,fp32,fp8_e4m3,fp8_e5m2); fp32 = tf32-math (Decision #15a); feature_flags=(mma_sync,) — NOT tcgen05/tmem (consumer Blackwell lacks those). - Regenerated op_target_conformance + runtime_abi dashboards (drift gate clean). - Updated test_numerical_check_via_manifest to expect the sm_120 fixture (and none on sibling arches). Not flipped: the @jit default lane (execution_matrix executable row + runtime.launch dispatch) is the documented follow-up, cf. ROCm rocm_wmma symbol vs rocm_compiled lane. Verification: shipped-symbol (5 dtypes) + emitted-PTX execute-compare + ptx_emit pass on the box; 1236 affected unit tests pass (clean PATH); mypy + drift clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * feat(nvidia): wire @jit(target=nvidia_sm120) matmul to the shipped mma.sync lane Flips the execution_matrix row to executable and dispatches a real matmul through libtessera_nvidia_gemm.so end-to-end on the RTX 5070 Ti. Mirrors the ROCm rocm_wmma lane. - execution_matrix: new executable ("nvidia_sm120","nvidia_mma") native_gpu row (execution_mode=cuda_runtime) + KNOWN_EXECUTORS["nvidia_mma"]; nvidia_sm120 removed from _UNIMPLEMENTED_TARGETS (sm_80/90/100 stay — only sm_120 proven). - runtime.py: _execute_nvidia_mma_artifact + _load_nvidia_gemm_runtime + _nvidia_mma_runtime_available probe + _NVIDIA_GEMM_SYMBOLS; registered in _executor_table. Loads the shipped .so (preloading libcuda + libnvrtc), picks the dtype symbol (f16/bf16/fp32→tf32-math), runs, returns the f32 result. Never raises on a missing lib/device — the probe gates executability. - jit.py: _nvidia_mma_lane_available + _uses_nvidia_mma_default + a stamping branch -> executable/compiler_path=nvidia_mma/native_gpu when the host probe passes; off-device the artifact stays artifact_only (no behavior change). execution_kind also honors the nvidia lane so is_executable agrees with launch. - test_nvidia_launch_execute.py: host-independent matrix-row + executor-registry tests, plus hardware-gated execute-and-compare (f16/bf16 hand-built artifact) and the @jit(target="nvidia_sm120") default-dispatch test. - Regenerated runtime_execution_matrix + test_coverage dashboards. Proof on box: @jit(target="nvidia_sm120") matmul launches via the shipped symbol at f32 epsilon (maxerr ~1e-6). 552 affected tests pass (clean PATH); 13 GPU lane tests pass; mypy + drift clean. Follow-up: a compiler-GENERATED nvidia lane (the rocm_compiled analog) and NVFP4. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * test(nvidia): allow sm_120 mma.sync matmul as hardware_verified in honesty guard The branch landed a real on-silicon NVIDIA proof (ecd825f): matmul/nvidia_sm120 flipped to hardware_verified, backed by the shipped tessera_nvidia_mma_gemm_* C-ABI symbols (libtessera_nvidia_gemm.so) and a skip-clean execute-compare fixture validating a warp-level mma.sync GEMM on the RTX 5070 Ti (sm_120, CC 12.0, CUDA 13.3). The manifest moved forward but the honesty guard in test_backend_capability_extension.py still hard-coded "NVIDIA not allowed" — the "first proof landed -> update the guard" event its own docstring anticipates. Mirror the ROCm block: add _NVIDIA_HARDWARE_VERIFIED_OPS = {"matmul"}, extend _hardware_verified_claim_is_allowed to accept nvidia_sm120 matmul (still requiring both evidence fields), and refresh the docstrings + assertion message. The guard stays strict for every other NVIDIA op/target. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: angst <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…ded; domain hygiene = standing discipline) (#128) * docs(audit): reconcile stale P2 claims in MASTER_AUDIT (generated-doc unification landed; domain hygiene is standing discipline) Two P2 lines described open work that is actually resolved: - "Unify generated-doc regeneration into one --write registry (#6)" → marked LANDED. Per COMPILER_AUDIT Next Work #6 the registry (generated_docs.py), fleet-wide drift gate (check_generated_docs.sh + release_gate.py both delegate), unified --write, orphan guard, and the CSV-canonical data-shaped tail (12 dashboards incl. the 3 target maps) all landed (2026-06-04 / 06-11). The only residual — aggressive content consolidation — #6 reassessed as low-value churn and deliberately deferred. - "Domain roadmap hygiene and stale-claim cleanup" → clarified as standing discipline, not a discrete backlog: the legacy roadmaps are consolidated into DOMAIN_AUDIT (9 archived plans), the generated dashboards are the count authority, and the discipline ("roadmaps are not status authorities") is already documented. Action is per-sprint, not a one-time task. No code or generated-doc changes (17 dashboards stay in sync); audit prose only. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs(tests): refresh stale tests/README fast-count (9,600→~11,350; total→~12,130) Same main-level drift the stack fix addresses — the README Quick start counts had fallen far behind (the "total < fast" inversion shows it), failing test_readme_fast_count_is_current on every PR off main, including this prose-only one. Byte-identical to the #125 stack's fix, so no conflict regardless of merge order. Fits this PR's stale-claim-hygiene theme. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

gstoner and others added 3 commits May 6, 2026 20:49

Merge pull request #5 from gstoner/main

f2ae5cc

Merge pull request #3 from gstoner/claude/fervent-varahamihira-cd1ab3

gstoner merged commit d2a40fb into main May 7, 2026
3 of 8 checks passed

gstoner deleted the claude/fervent-varahamihira-cd1ab3 branch May 7, 2026 14:16

gstoner mentioned this pull request Jun 11, 2026

CSV-canonical for the 3 target-map dashboards #65

Merged

gstoner mentioned this pull request Jun 25, 2026

NVIDIA sm_120: CUDA 13.3 verify + first hardware-verified mma.sync backend_kernel #106

Merged

gstoner mentioned this pull request Jun 25, 2026

docs(audit): reconcile stale P2 claims (generated-doc unification landed; domain hygiene = standing discipline) #128

Merged

gstoner mentioned this pull request Jul 11, 2026

compiler(target-ir): type the NVIDIA Target IR dialect + real nvvm.mma.sync (Tile IR / Target IR tail) #371

Merged

Navigation Menu

Claude/fervent varahamihira cd1ab3 - #6

Claude/fervent varahamihira cd1ab3#6
gstoner merged 3 commits into
mainfrom
claude/fervent-varahamihira-cd1ab3

gstoner commented May 7, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

gstoner commented May 7, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant