Reward hacking an allclose check

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

$ 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

$ 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`

$ 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 to use the verifier as the reward gate.

Get started

Concepts

Guides

Model conformance

Case studies

Help

Reward hacking an allclose check

Why the gate holds

​Why the gate holds

Why the gate holds