# P1.4 — TileLang FP8 Decode Wire Killed: Output Degeneration

## Context

P1.4 tried to route Qwen3-4B FP8 paged-KV pure decode through the existing
TileLang AOT cubin:

- `tilelang_batch_decode_paged_hd128_fp8_q32_kv8_run_cuda`
- env gate: `INFER_FP8_KV_DECODE=1`
- shape gate: HD128, q32/kv8, page_size=16, batch<=8

Phase-0 substrate audit was positive:

- Qwen3-4B config is `head_dim=128`, `num_attention_heads=32`,
  `num_key_value_heads=8`.
- FP8 paged-pool storage already exists (`KVFormat::FP8E4M3`) with per-token
  per-kv-head scales.
- Qwen3 prefill/decode already quantize BF16 work buffers into FP8 pool data
  and scale tensors.
- The TileLang AOT build log confirms HD128 decode cubins are built.

## Evidence

The temporary wiring compiled and launched. A zero-output smoke test for the
TileLang FP8 FFI passed, so the basic cubin/ABI path was valid.

Correctness failed on real model data:

| Mode | Command shape | Result |
|---|---|---|
| BF16 KV default | `test_greedy_solo_vs_concurrent` | PASS, coherent output |
| FP8 KV custom kernel | `INFER_TEST_KV_POOL_FORMAT=fp8` | PASS, coherent output |
| FP8 KV TileLang | `INFER_TEST_KV_POOL_FORMAT=fp8 INFER_FP8_KV_DECODE=1` | PASS solo/concurrent consistency, but output degenerated |

Representative outputs:

```text
custom FP8:
" about a young girl named Lila who is a talented painter, and her journey to become a successful artist. Include the following elements: a mentor,"

TileLang FP8:
" about\npenations\n\n\n\nakers\n\n, tosssssss,sss to to\n\n\n\n\n"
```

This is not a scheduler batch-invariance failure: solo and concurrent produced
the same token IDs in both modes. It is a numerical/semantic mismatch between
the custom FP8 decode kernel and the TileLang FP8 decode cubin.

## Root Cause

Root cause is not fully isolated. Most likely candidates:

1. TileLang FP8 kernel scale layout or indexing differs from the custom
   `decode_attention_fp8` contract.
2. TileLang `float8_e4m3fn -> float32 -> bf16` cast semantics differ from the
   custom CUDA FP8 decode path.
3. The TileLang kernel assumes a different FP8 value encoding or scale
   normalization than `quantize_paged_kv_fp8` produces.

The failed output rules out landing even an env-gated runtime path: it would be
"fast garbage" and could be accidentally enabled by `INFER_FP8_KV_DECODE=1`.

## Fix

P1.4 runtime/test wiring was reverted before commit. No runtime diff landed.

Next viable investigation should be a single-layer numerical diff:

1. Capture one layer's `q_batch`, FP8 K/V pool bytes, K/V scales, and metadata.
2. Run custom `decode_attention_fp8` and TileLang FP8 decode on the same inputs.
3. Compare attention output per head and locate whether divergence enters at
   FP8 dequant, logits, softmax, or `P @ V`.

Only after that diff matches should the dispatch arm be reintroduced.

## Rule

Do not license an alternative quantized attention kernel from launch-only smoke
tests. Quantized KV decode requires real-model A/B against the existing custom
kernel before any runtime dispatch is landed, even behind an env gate.
