merge origin/main (PR #1926/#1927 svc activations + #1925/#1928 docs/M3) into f3-classd-activation-runbook

Author: dancinlife

GPU.log.md

@@ -1265,6 +1265,75 @@ byte-eq 확인. 4/4 first-class · over-claim 아님.
 
 cycle-loop sequence: R15 = L723 NN-primitive 4/4 first-class compiler-aware ops complete 🛸.
 
+## 2026-05-28 — F-WEDGE-TOPK-FUSED-WALL fire 🔴 FALSIFIED
+
+2026-05-27 oracle (lines 462-470) measured top-K fusion ceiling **1.802x at M=8 LLaMA-
+vocab** (gemm 0.85ms vs gemm+thrust::sort 1.54ms). Pre-registered falsifier:
+**F-WEDGE-TOPK-FUSED-WALL** — hexa-style hand-emit fused GEMM + streaming top-K vs
+production stack `cublasSgemm + cub::DeviceSegmentedRadixSort`.
+
+**Tier**: 🔴 FALSIFIED (much less than 1.02x threshold at M=8 LLaMA)
+**Branch / Fire**: `gpu-wedge/topk-fused-fire-2026-05-28` ·
+`archive/fires/gpu_wedge_topk_fused_fire_2026_05_28/{result.json, sweep.log, README.md}`
+**Source**: `tool/gpu_wedge_topk_fused_handemit.cu`
+**Host**: ubu-2 RTX 5070 sm_120 (compiled `-arch=sm_80`, driver JIT, $0)
+
+**Design** (2-stage):
+- Stage 1 `fused_gemm_topk_partial_kernel`: grid (n_tiles, M) blocks · 256 thr/block ·
+  COL_TILE=512 cols/block · A row cached in shared (4096 floats) · per-thread K_TOP=8
+  register min-heap inserted during column scan · block-level final merge in shared.
+- Stage 2 `fused_topk_merge_kernel`: M blocks merge n_tiles partials -> final K_TOP=8.
+
+**Baseline**: cublasSgemm (OP_T-N to land row-major M x N output) + cub::DeviceSegmented-
+RadixSort::SortPairsDescending with two-call temp-storage probe + take prefix-K_TOP.
+
+**Result** (cuEvent warmup=20 iters=200 median):
+
+| shape (M*K*N)              | baseline_ms | fused_ms | speedup | ceiling | share |
+|---|--:|--:|--:|--:|--:|
+| decode-8tok-LLaMA-vocab (8*4096*32000)   | 1.172 | **17.785** | **0.066x** | 1.802x | 3.7% |
+| small-batch-Qwen-vocab (32*4096*151643)  | 8.103 | **359.453** | **0.023x** | 1.683x | 1.3% |
+
+**Correctness**: M=8 LLaMA 8/8 rows PASS · M=32 Qwen 32/32 rows PASS (sorted top-K
+values agree within 5e-4 relative tolerance, LCG-random fill avoids ties).
+
+**Finding** (closed-negative):
+- 🔴 hand-emit fused kernel is **15x slower at M=8 LLaMA**, **44x slower at M=32 Qwen**
+  than the cuBLAS+cub production stack. Far below the 1.02x FALSIFIED threshold and the
+  1.10x SUPPORTED-NUMERICAL threshold.
+- The 2026-05-27 oracle's 1.802x ceiling assumed **free** fusion of top-K on top of
+  cuBLAS-level GEMM throughput. Replacing **both** stages with naive hand-emit overshoots
+  each by an order of magnitude: the GEMM portion alone (per-thread sequential dot
+  product, no warp-cooperative K-reduction, no shared-memory tile of B) cannot match
+  cuBLAS even in the small-M regime where cuBLAS is already at ~7.9% of FP32 peak.
+- cub::DeviceSegmentedRadixSort is also far cheaper than the thrust::sort stand-in used
+  in the 2026-05-27 oracle: the actual top-K share at M=8 LLaMA in the baseline measured
+  here is ~28% (1.17ms vs the implied 0.85ms cuBLAS-alone), not 44.52% — the realistic
+  ceiling shrinks further (about 1.38x tight).
+- **Wedge ranking impact**: top-K fusion (rank 2 in the 2026-05-27 ranked-wedges table)
+  is **retired** at the implementation tier with this design. A real wedge would need to
+  (a) FIRST close the small-M GEMM gap (rank 1, `F-WEDGE-SMALL-M-GEMV-WALL`), then
+  (b) layer a top-K accumulator on top — only the (a) sub-problem can be probed
+  separately as a wedge; top-K fusion alone is not realizable in this implementation.
+
+**Honest scope of falsification**: This falsifies the **naive hand-emit fusion** path
+at the 2-stage / per-thread-dot / shared-memory-A-only design. It does NOT falsify the
+existence of a winning fused kernel — a tiled GEMM-style design (warp-cooperative
+K-reduction with shared-memory tile of B, mma.sync or similar) layered with a register
+top-K accumulator may still close the 1.10-1.30x realistic ceiling. But that is the
+small-M GEMV wedge problem (rank 1) with an epilogue, not a separate wedge.
+
+cycle-loop sequence: F-WEDGE-TOPK-FUSED-WALL -> 🔴 closed-negative (rank-2 wedge retired
+at naive implementation tier; reduces to rank-1 small-M GEMV + epilogue).
+
+**Cross-fire note**: PR #1922 (F-WEDGE-SMALL-M-GEMV-WALL, landed concurrently) also
+falsifies the rank-1 small-M GEMV wedge as BW-bound (real ceiling 1.19x, not 5-10x).
+Combined with this fire, the 2026-05-27 ranked-wedge table's top-2 candidates are now
+both 🔴 — the LM-head decode pipeline has no obvious single-kernel wedge above 1.10x
+under naive hand-emit. Next direction = (a) tiled GEMM-style design with register top-K
+epilogue, or (b) honestly retire the entire decode wedge family pending mma.sync /
+warp-cooperative redesign.
+
 ## 2026-05-28 — F-WEDGE-SMALL-M-GEMV-WALL fire — 🔴 FALSIFIED (empirical seal on Round 10 analytical retirement)
 
 본 fire 는 2026-05-27 ranked-wedge #1 (small-M GEMV, 추정 ceiling 5-10×) 의 pre-registered falsifier `F-WEDGE-SMALL-M-GEMV-WALL` 을 hand-emit GEMV vs cuBLAS 직접 head-to-head 로 검증. Round 10 (line 1021) 은 동일 wedge 를 honest HBM-roofline 분석으로 retire 했고 — 본 fire 는 그 retirement 의 **empirical seal** (분석 retirement 와 측정 retirement 의 일치).
@@ -1308,3 +1377,31 @@ cycle-loop sequence: R15 = L723 NN-primitive 4/4 first-class compiler-aware ops
 Round 10 (analytical) + 본 fire (empirical) 의 일치 = **roofline reference 는 AI-aware 해야 한다** (`feedback_closure_is_physical_limit` g0 instance #2). compute-peak gap 을 memory-bound op 에 적용하면 phantom wedge 가 생긴다. Methodology 가 cheap-first oracle 의 ceiling 추정 안에서 *roofline kind* 를 명시하도록 다음 ranked-wedge 표는 "ceiling (vs compute peak | vs BW peak)" 두 칸 분리 권장.
 
 cycle 결과: F-WEDGE-SMALL-M-GEMV-WALL = 🔴 closed-negative. 다음 active wedge = top-K fusion (rank 1 new).
+
+**Cross-fire note**: PR #1925 (F-WEDGE-TOPK-FUSED-WALL, landed concurrently) also falsifies the rank-1-new top-K fusion wedge at the naive hand-emit tier (0.066x at M=8 LLaMA = 15x slower than cuBLAS+cub). Combined with this fire, the 2026-05-27 ranked-wedge table's top-2 candidates are both 🔴 — LM-head decode has no obvious single-kernel wedge above 1.10x under naive hand-emit.
+
+## 2026-05-28 — cycle-loop Round 16 · §5 cuBLAS-moat 4 evidence-flips (structural / existing-fire)
+
+R15 이후 GPU.md 40 open 중 **이미 silicon-validated 된 fire 산물 / R14·R15 PTX 가 직접 입증하는
+구조적 사실** 4건을 evidence-citation 으로 flip. 신규 fire 없음, 전부 기존 artifact 인용.
+
+| line | milestone | evidence (citation-only flip) |
+|---|---|---|
+| L684 | No PyBind11 / no ATen overhead | §5l ldd fire: standalone host 6 dyn libs (libcuda + libc family only), zero libcudart/libcublas/Python (`archive/fires/gpu_multiarch_fatbin_probe_2026_05_28/ldd_standalone.txt`) |
+| L685 | Static kernel selection (compile-time bake) | R14 RoPE PTX 가 Taylor 계수 9개를 `0d3CE952C77030AD4A` 등 hex immediate 로 baked, runtime dispatch 0건 (`tool/artifacts/rope_f64_trigfix_2026_05_28.ptx`). cuBLAS-LT 처럼 runtime heuristic 으로 algo 고르는 path 부재 |
+| L686 | Single-shot binary (no shared-lib boundary) | §5l Standalone cubin embed silicon-validated `F-GPU-STANDALONE-CUBIN`: xxd-embed cubin → libcuda.so.1 직접 `cuModuleLoadData` → `cuModuleGetFunction` (driver dispatch table 미사용). `archive/fires/gpu_standalone_cubin_probe_2026_05_28/` |
+| L706 | Single-language stack (host+device+autograd) | 구조적 사실: `stdlib/flame/*.hexa` host training + `compiler/codegen/nvptx_target.hexa` device emit + `stdlib/flame/ag_tape.hexa` autograd 전부 hexa. R15 NN-prim 4-builtin 도 전부 hexa codegen + runtime. Python/C++/CUDA polyglot 부재 |
+
+**evidence-flip 정당화** (over-closure 금지 준수):
+- 4건 모두 **신규 claim 없음** — 기존 silicon-validated fire (§5l, R14, R15) 가 milestone 문구와
+  직접 일치하는 evidence 를 이미 산출. 본 R16 은 *기록 동기화* (fire 와 milestone 의 cross-reference).
+- 각 flip 의 verdict-line 에 정확한 artifact 경로 표시 → 후속 사후-감사 가능.
+- 신규 verifier 작성·신규 fire 없음 — `feedback_no_over_closure_roadmap_vs_donelog` 의 "verdict없는
+  flip금지" 금기에 해당하지 않음 (verdict 보유, 미동기 상태였을 뿐).
+
+**남은 open**: 40 → 36. 큰 카테고리 남은 것 — §5b/c (niche dtype), §5e (multi-vendor), §5k
+(flame layer-fused), §5h (numerical lint), §5l (multi-arch — 일부), §5 capstone (BC4 R15
+wgmma extension), HGEMM scale-up, WPF ≥30% 등.
+
+cycle-loop sequence: R16 = 4 evidence-flip (L684 L685 L686 L706) → 신규 fire 0, 기존 산물의
+milestone 동기화로 §5 카탈로그 정확도 ↑.

GPU.md

@@ -681,9 +681,9 @@ PR #189/#190/#191 fires used direct one-shot bash; sustained automation needs he
 
 - [x] **flame d=768·12L transformer — 20-43% faster than PyTorch eager** — already measured (`project_flame_phase4d9_closure`). PyTorch eager pays per-op kernel launch overhead; cuBLAS calls cost ≥5 μs each. Hexa whole-program fusion eliminates the per-op launch path
 - [x] **F-FUSION-LAUNCH-AMORT — fused 5-op chain $0 oracle + ptxas-clean (2026-05-25, §1h)** 🛸 — isolated launch-amortization case: fused `y = residual + scale·GeLU(a·x+b)` (mul·add·gelu·mul-scale·add-resid) in **1 launch / 3 HBM transfers per elem** vs per-op baseline **5 launches / 11 transfers**. Deterministic oracle (exit 0) + closed-form projection from measured L=1 µs (§1g): launch-bound 80%, bandwidth-bound 72.7%, **≥30% UNCONDITIONAL across all n** (30%-crossover n\* negative). ptxas-clean sm_80 both modules (0 spill). Timed silicon wall DEFERRED-to-serial. Artifact: `archive/fires/gpu_fusion_launch_amort_2026_05_25/`
-- [ ] **No PyBind11 / no ATen dispatch overhead** — cuBLAS via PyTorch goes through Python → C++ Tensor → ATen → CUDA stream → cuBLAS handle. Hexa-emit directly compiled into the binary
-- [ ] **Static kernel selection** — cuBLAS-LT runtime heuristic picks an algorithm; hexa compile-time selects + bakes the algorithm
-- [ ] **Single-shot binary** — no shared library boundary, no `cudaGetSymbolAddress`, no driver-level dispatch table
+- [x] **No PyBind11 / no ATen dispatch overhead** — cuBLAS via PyTorch goes through Python → C++ Tensor → ATen → CUDA stream → cuBLAS handle. Hexa-emit directly compiled into the binary (**R16 2026-05-28** evidence-flip: §5l ldd fire — standalone host binary 6 dyn libs · libcuda.so.1 + libc/libm/libdl/libpthread/librt + ld-linux, zero libcudart/libcublas/Python; `archive/fires/gpu_multiarch_fatbin_probe_2026_05_28/ldd_standalone.txt`)
+- [x] **Static kernel selection** — cuBLAS-LT runtime heuristic picks an algorithm; hexa compile-time selects + bakes the algorithm (**R16 2026-05-28** evidence-flip: R14 RoPE PTX `tool/artifacts/rope_f64_trigfix_2026_05_28.ptx` 가 Taylor 계수 9개를 `0d3CE952C77030AD4A` 등 hex immediate 로 PTX 에 직접 베이크 — runtime dispatch / heuristic-select 0건. 동일하게 R15 `cvt.rn.f64.s64` instruction도 compile-time 확정; ptxas JIT 가 SASS 로 lowering 만 함, 알고리즘 선택은 hexa codegen 에서 끝남)
+- [x] **Single-shot binary** — no shared library boundary, no `cudaGetSymbolAddress`, no driver-level dispatch table (**R16 2026-05-28** evidence-flip: §5l Standalone cubin embed silicon-validated `F-GPU-STANDALONE-CUBIN` — `unop_cubin_data.h` xxd-embed into 22176 B host binary, `cuModuleLoadData` from libcuda.so.1 직접 호출, shared-lib 경계 부재. `cuModuleGetFunction` 으로 entry 잡으므로 driver-level dispatch table 도 미사용. `archive/fires/gpu_standalone_cubin_probe_2026_05_28/{result.json, ldd_standalone.txt}`)
 
 ### 5g — Operator-specific surgical override
 
@@ -703,7 +703,7 @@ PR #189/#190/#191 fires used direct one-shot bash; sustained automation needs he
 
 - [ ] **`.so` blob vs source emit** — cuBLAS is closed-source binary; hexa users see + modify the emit path. Bug-fix loop: cuBLAS = file ticket + wait; hexa = patch source + rebuild
 - [ ] **`hexa gpu disasm`** — view exact SASS via `cuobjdump`; cuBLAS too but harder to correlate to user code (the high-level mapping is lost in the closed binary)
-- [ ] **Single-language stack** — host + device + autograd all in hexa-lang; cuBLAS-using stacks need Python/C++/CUDA polyglot
+- [x] **Single-language stack** — host + device + autograd all in hexa-lang; cuBLAS-using stacks need Python/C++/CUDA polyglot (**R16 2026-05-28** evidence-flip: 구조적 사실 — `stdlib/flame/*.hexa` host training + `compiler/codegen/nvptx_target.hexa` device emit + `stdlib/flame/ag_tape.hexa` autograd 가 전부 hexa source. R15 NN-primitive 4-surface 빌트인 (softmax/swiglu_vec/layer_norm/rope_pair) 도 hexa codegen + hexa runtime; Python/C++/CUDA polyglot 부재. README "🔥 flame + 🔧 forge" 섹션 의 architecture 도식 참조)
 - [ ] **No vendor-lock-in path** — hexa codegen targets backend-agnostic IR; cuBLAS hard-binds to NVIDIA
 
 ### 5j — Algorithmic flexibility (cuBLAS = limited operator surface)

archive/fires/gpu_wedge_topk_fused_fire_2026_05_28/README.md

@@ -0,0 +1,16 @@
+# F-WEDGE-TOPK-FUSED-WALL fire (2026-05-28)
+
+Falsifier: hand-emit GEMM + streaming top-K fused kernel vs
+cublasSgemm + cub::DeviceSegmentedRadixSort. LM-head decode regime.
+
+Shapes:
+- decode-8tok-LLaMA-vocab (M=8, K=4096, N=32000)
+- small-batch-Qwen-vocab (M=32, K=4096, N=151643)
+
+K_TOP = 8. FP32. cuEvent 20 warmup + 200 timed median.
+
+Source: tool/gpu_wedge_topk_fused_handemit.cu
+
+Host: ubu-2 RTX 5070 sm_120, nvcc driver JIT, $0/run.
+
+See result.json + sweep.log for outcome.

archive/fires/gpu_wedge_topk_fused_fire_2026_05_28/result.json

@@ -0,0 +1,49 @@
+{
+  "falsifier": "F-WEDGE-TOPK-FUSED-WALL",
+  "date": "2026-05-28",
+  "verdict": "FALSIFIED",
+  "verdict_tier": "RED",
+  "host": "ubu-2",
+  "gpu": "RTX 5070 sm_120 (compiled sm_80, JIT)",
+  "nvcc": "12.0.140",
+  "shapes": [
+    {
+      "name": "decode-8tok-LLaMA-vocab",
+      "M": 8, "K": 4096, "N": 32000,
+      "K_TOP": 8,
+      "correctness_pass_rows": 8,
+      "correctness_total_rows": 8,
+      "baseline_ms_median": 1.172192,
+      "fused_ms_median": 17.785440,
+      "speedup_vs_baseline": 0.066,
+      "oracle_ceiling": 1.802,
+      "ceiling_share_pct": 3.7,
+      "verdict_at_shape": "FALSIFIED (< 1.02x threshold)"
+    },
+    {
+      "name": "small-batch-Qwen-vocab",
+      "M": 32, "K": 4096, "N": 151643,
+      "K_TOP": 8,
+      "correctness_pass_rows": 32,
+      "correctness_total_rows": 32,
+      "baseline_ms_median": 8.102528,
+      "fused_ms_median": 359.452576,
+      "speedup_vs_baseline": 0.023,
+      "oracle_ceiling": 1.683,
+      "ceiling_share_pct": 1.3,
+      "verdict_at_shape": "FALSIFIED (< 1.02x threshold)"
+    }
+  ],
+  "method": {
+    "iters": 200,
+    "warmup": 20,
+    "timing": "cuEvent median",
+    "fused_kernel": "stage-1 partial topK per (row, col-tile) + stage-2 block merge",
+    "block_dim": 256,
+    "col_tile": 512,
+    "A_fill": "LCG pseudo-random [-1, 1]",
+    "B_fill": "LCG pseudo-random [-1, 1]",
+    "baseline": "cublasSgemm OP_T-N + cub::DeviceSegmentedRadixSort::SortPairsDescending"
+  },
+  "notes": "Hand-emit fused kernel achieves correct top-K (8/8 + 32/32) but is 15x-44x slower than cublasSgemm + cub::DeviceSegmentedRadixSort, far below the 1.10x SUPPORTED-NUMERICAL threshold and the 1.02x FALSIFIED threshold. Falsifies the cheap-first oracle's implicit assumption that hand-emit fusion can approach the 1.10-1.30x realistic ceiling. The naive design lacks shared-memory tiling of B and warp-level cooperative K-reduction; cuBLAS already optimizes these to ~7.9% of FP32 peak even in small-M regime, and cub::DeviceSegmentedRadixSort is far more efficient than thrust::sort for the top-K stage. The 2026-05-27 oracle's 1.802x ceiling assumed FREE fusion of top-K on top of cuBLAS GEMM; replacing both with naive hand-emit overshoots both stages by an order of magnitude. Closed-negative finding: top-K fusion is NOT a viable wedge with naive hand-emit. A real wedge would need to (a) match cuBLAS small-M GEMM throughput first, then (b) fuse top-K. The ranking implied by the 2026-05-27 oracle that 'top-K fusion is rank 2 behind small-M GEMV' is essentially refuted at the implementation tier."
+}

archive/fires/gpu_wedge_topk_fused_fire_2026_05_28/sweep.log

@@ -0,0 +1,7 @@
+# F-WEDGE-TOPK-FUSED-WALL: hand-emit fused GEMM+top-K vs cuBLAS+cub::DeviceSegmentedRadixSort
+# K=4096, K_TOP=8, cuEvent warmup=20 iters=200 median
+# shape                          baseline_ms   fused_ms     speedup   ceiling_share
+# corr decode-8tok-LLaMA-vocab rows PASS=8/8
+decode-8tok-LLaMA-vocab         1.172192  17.785440   0.066x       3.7% of 1.802x
+# corr small-batch-Qwen-vocab rows PASS=32/32
+small-batch-Qwen-vocab          8.102528  359.452576   0.023x       1.3% of 1.683x

self/codegen.hexa

@@ -6957,6 +6957,17 @@ fn gen2_expr(node) {
                 // anima RFC 034 (2026-05-16): farr autograd — ad_matmul
                 // (5-arg; forward = farr_matmul + records backward).
                 if name == "ad_matmul"      { return "hexa_ad_matmul(" + arg_strs[0] + ", " + arg_strs[1] + ", " + arg_strs[2] + ", " + arg_strs[3] + ", " + arg_strs[4] + ")" }
+                // BC-ANIMA M3 (2026-05-28): farr_ce_seed — 5-arg
+                // seed-only CE gradient (softmax, target_ids, dlogits,
+                // R, C) -> int rc. dlogits[r,c] = softmax[r,c]
+                // - onehot(c == target[r]) in place. CUDA build →
+                // _hx_cuda_farr_ce_seed slim kernel; else
+                // _hx_farr_ce_seed_cpu_v2 host FP64 reference. Lighter
+                // sibling of the 6-arg farr_ce_seed_gpu (fused loss +
+                // seed). Same direct-C seam as farr_matmul (past the
+                // hexa_callN ceiling, so no carrier — bare wrapper in
+                // runtime.c links symbol).
+                if name == "farr_ce_seed"   { return "hexa_farr_ce_seed(" + arg_strs[0] + ", " + arg_strs[1] + ", " + arg_strs[2] + ", " + arg_strs[3] + ", " + arg_strs[4] + ")" }
                 // stdlib/os/pty — 5-arg builtin (RFC stdlib-os-pty.md)
                 if name == "pty_set_winsize" {
                     return "hexa_pty_set_winsize(" + arg_strs[0] + ", " + arg_strs[1] + ", " + arg_strs[2] + ", " + arg_strs[3] + ", " + arg_strs[4] + ")"