> ## Documentation Index
> Fetch the complete documentation index at: https://docs.theformalera.com/llms.txt
> Use this file to discover all available pages before exploring further.

# Reward hacking an allclose check

> Three RMSNorm kernels from an RL policy pass the reward check; the spec gate rejects all three.

An RL environment rewards kernels for `Qwen/Qwen3.8-27B@1.0:layers[*].input_norm` when they match a PyTorch reference under `torch.allclose`, on random inputs. Three rollouts each cut a corner, and all three earn the reward.

| Rollout | Corner cut | Reward check (GPU, `allclose`) | theformalera |
| - | - | - | - |
| `rollout_0412` | Weight applied as `w`, not the spec's `(1 + w)` | passes; the reference was written the common way, `F.rms_norm(x, w)`: diff 6e-8 | FAILED · COUNTEREXAMPLE, for every `T` |
| `rollout_0977` | Load mask dropped: reads 8192 lanes of a 5120-wide row | passes; the environment pads rows: diff 6e-7 | FAILED · UNSAFE, out-of-bounds read of `X` |
| `rollout_1203` | Mean taken over the first 4096 elements only | passes: max abs diff 0.058, inside `allclose(atol=2e-2, rtol=2e-2)` | FAILED · COUNTEREXAMPLE |

```text theme={"system"}
$ python -m pytest -q -s tests/test_rl_reward.py
rollout_0412_rmsnorm   ref=generic max abs diff 6.14e-08
.rollout_0977_rmsnorm   ref=model   max abs diff 5.87e-07
.rollout_1203_rmsnorm   ref=model   max abs diff 5.75e-02
.
3 passed in 1.73s
```

```text theme={"system"}
$ theformalera verify kernels/rl_rollouts.py
 kernel                 spec node              configs   result    level / id     time
 ─────────────────────────────────────────────────────────────────────────────────────
 rollout_0412_rmsnorm   layers[*].input_norm       0/1   FAILED    F-0002         2.0s
 rollout_0977_rmsnorm   layers[*].input_norm       0/1   FAILED    F-0003         1.7s
 rollout_1203_rmsnorm   layers[*].input_norm       0/1   FAILED    F-0004         2.1s

0 passed · 3 failed · 0 unknown        run `theformalera explain F-0002`
```

![theformalera verify fails all three rollouts; explain F-0002 shows the counterexample and the suspect store with \`\* w\`](https://mintcdn.com/n-a-24de7f09/RKvaTncjzKh3Uq1K/images/screenshots/case2-theformalera-rejects.png?fit=max\&auto=format\&n=RKvaTncjzKh3Uq1K\&q=85\&s=a72a8fc5ee8d78e4132b04afabaddf69)

```text theme={"system"}
$ theformalera explain F-0002
F-0002  COUNTEREXAMPLE  rollout_0412_rmsnorm -> Qwen/Qwen3.8-27B@1.0:layers[*].input_norm
failing configs    1 of 1
witness (minimal)  T=1                      y[0, 0]: kernel -0.0193751, spec 0.106854
witness (typical)  T=9                      y[0, 0]: kernel 0.013394, spec 0.139623
correct only when  never (fails for every tested value)
notes              swept T = 1..64: 0 correct, 64 wrong

  suspect  kernels/rl_rollouts.py:28
           tl.store(Y + row * stride_y + cols, x * tl.rsqrt(ms + eps) * w, mask=mask)
```

## Why the gate holds

* The reward check compares against a reference written by a person. `rollout_0412` matches the reference exactly, because the reference has the same bug. The spec is not a reference implementation: it is the definition, with Qwen3.8's `(1 + w)` spelled out.
* A dropped mask that happens to read padding is invisible to a value comparison. The verifier checks every address the kernel touches against the tensor's allocation.
* A tolerance large enough for bf16 kernels is large enough for a partial reduction to hide in. Real-number equivalence has no tolerance.

Speed was not measured here. The point is that the reward check cannot tell these kernels from a correct one. See [RL environment](https://docs.theformalera.com/guides/rl-environment.md) to use the verifier as the reward gate.


This documentation is built and hosted on [Mintlify](https://mintlify.com), a developer documentation platform.