[BUG][ppu] fp8e4nv is gated by an NVIDIA compute-capability threshold, at both the frontend and in ConvertTritonGPUToLLVMPPU

## Summary

On T-Head PPU (`FLAGTREE_BACKEND=ppu`), `tl.float8e4nv` cannot be compiled. The
gate is an NVIDIA compute-capability threshold, applied at two levels, and the
card reports capability `(8, 0)`:

```
ValueError: type fp8e4nv not supported in this architecture.
The supported fp8 dtypes are ('fp8e4b15', 'fp8e5')
```

That supported-set is exactly what Triton reports for NVIDIA sm_80.

This blocks any Triton kernel that quantises to OCP E4M3 — for us, the FP8 paged
KV cache used by DeepSeek-V4-class models.

## Environment

| component | version |
|---|---|
| card | PPU-ZW810E, 97920 MiB, warp 32 lanes, capability (8, 0) |
| FlagTree | `0.6.0+ppu.git132ee455`, built from `main` with `FLAGTREE_BACKEND=ppu` |
| torch | 2.10.0 |
| driver | PPU-SMI 1.28, Driver 1.3.2-d7f5a2, HGGC 13.0 |

`triton.backends` after installation contains `amd`, `nvidia`, `ppu`, and the
target is reported as `GPUTarget(backend='cuda', arch=80, warp_size=32)`.

## Where the gate is

**Frontend**, `third_party/ppu/backend/compiler.py:229-233`:

```python
if "supported_fp8_dtypes" not in args:
supported_fp8_dtypes = set(HGGCOptions.supported_fp8_dtypes) # ("fp8e5", "fp8e4b15")
if capability >= 89:
supported_fp8_dtypes.add("fp8e4nv")
args["supported_fp8_dtypes"] = tuple(sorted(supported_fp8_dtypes))
```

This is the same logic as upstream Triton's NVIDIA backend, which on this machine
sits at `triton/backends/nvidia/compiler.py:240-244`. `capability` here is the PPU's
`(8, 0)`, which is not an NVIDIA SM version and does not carry the same meaning.

**Lowering.** Because line 229 only computes the default when the option is absent,
the frontend check can be bypassed by passing the option at launch. Doing so
produces correct IR (`tt.fp_to_fp ... -> tensor<4096xf8E4M3FN>`, `ttg.target =
"ppu:80"`) and then fails inside the PPU's own pass:

```
LLVM ERROR: Conversion from/to f8e4m3nv is only supported on compute capability >= 89
...
Pipeline failed while executing [`ConvertTritonGPUToLLVMPPU` on 'builtin.module' operation]
```

So an NVIDIA threshold — and the phrase "compute capability" — is enforced inside a
PPU-specific lowering pass.

## Reproduction

```python
import torch, triton, triton.language as tl

@triton.jit
def cast_k(src, dst, N: tl.constexpr):
o = tl.arange(0, N)
tl.store(dst + o, tl.load(src + o).to(tl.float8e4nv))

N = 4096
x = (torch.randn(N, device="cuda") * 30).float()
d = torch.empty(N, dtype=torch.float8_e4m3fn, device="cuda")
cast_k[(1,)](x, d, N=N) # ValueError, as above

# bypassing the frontend gate moves the failure into ConvertTritonGPUToLLVMPPU
cast_k[(1,)](x, d, N=N, supported_fp8_dtypes=("fp8e4b15", "fp8e4nv", "fp8e5"))
```

## Why we believe this is policy rather than capability

E4M3 is usable on this card outside Triton. T-Head's own vLLM port,
`vllm==0.20.1+v0.1.0.ppu2.1.0`, registers
`_C::fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert` — a C++ kernel that performs
exactly this UE8M0/E4M3 quantisation. On the same card it launches and produces
results that agree with our Triton implementation under the same tolerances. Its
FP8 conversion does not go through Triton, so it is unaffected by the gate above.

We are **not** claiming the PPU has a hardware E4M3 convert instruction — we cannot
tell from outside, and note that `torch.float8_e4m3fn` casts succeed on any device
because torch has a software path below sm_89, so that alone proves nothing. What we
can say is that the format is produced correctly on this card by a vendor kernel,
while Triton declines to emit it on the basis of an NVIDIA capability number.

## Questions

1. Is `fp8e4nv` lowering implemented for the PPU target? If it is, the frontend gate
should key on something PPU-specific rather than `capability >= 89`.
2. If it is not implemented, could the lowering failure say so? The current message
cites "compute capability", which sends the reader to NVIDIA documentation and to
the hypothesis that the hardware is at fault.
3. Is there a supported way to advertise a PPU target that permits E4M3?

## Workaround, for anyone hitting this

We encode E4M3 with integer arithmetic in the kernel, which needs no `fp8e4nv`
lowering. It is bit-identical to `torch.float8_e4m3fn` over 65536 values including
the subnormal range, and costs nothing measurable on a memory-bound operator (96.0%
of achievable bandwidth at 32768 tokens). Two details worth stating because a naive
encoder gets them wrong: E4M3 has subnormals (`m * 2^-9`, m in 1..7), and exponent
field 15 with mantissa 0 is the legal value 256, not an overflow — only `e > 15`, or
`e == 15` with mantissa 7 (the NaN encoding), may saturate.

Start with third_party/ppu/backend/compiler.py:229-233 and trace the ConvertTritonGPUToLLVMPPU pass that rejects f8e4m3nv. Run the supplied cast_k reproduction with and without supported_fp8_dtypes to compare both failures. Done means the PPU-specific support decision is correct and any unsupported lowering reports a PPU-specific explanation.

Written by the indexing model from the issue text.