WS1 Numerical Precision Standard: Ground Truth Definition, Judgment Methods, and Threshold Provenance¶

Field	Value
Document status	Draft Normative
Document version	v1.0
Document date	2026-09-11
Precision contract	`rl_engine/kernels/gtest/tolerance_contract.json`
Contract version	`ws1-c1-v2`
Implementation audit baseline	`6ceeb62e073d74e00a8c5b56af538ee1a7417773` (`origin/main`)
Target model	Qwen3-8B Dense (`ws1-qwen3-8b-dense-primary-v6`)
Required backend profiles	`cuda_bf16`, `triton_cuda_bf16`
Optional precision	FP16 (complete when declared; not required for WS1 exit)
Out of WS1 scope	FP8, ROCm, Ascend, and any optional fused path not listed in the C8 required-op matrix

On document status. The normative sections (§1–§8) define the numerical precision targets and judgment methods for WS1 and may be cited as the project's standard. Section 9 records where the current repository diverges from that standard, Section 10 bounds the applicability of historical evidence, and Section 11 gives the plan for closing the gaps. Until every blocking item in Section 11 is closed, this repository must not be described as fully satisfying this standard. Normative targets and implementation status are stated separately throughout.

1. Summary¶

WS1 separates numerical judgment into four classes, so that standards of different natures are not collapsed into a single "tolerance table":

Judgment	Compared quantities	Judgment mode
Forward accuracy	BF16 candidate ↔ FP32 reference	Element-wise tolerance
Gradient accuracy	BF16 backward ↔ FP32 reference gradient	Element-wise tolerance
Forward / gradient invariance	Transformed config ↔ canonical config	Zero tolerance (bitwise is the target)
Chain-level parity	Training-style teacher forcing ↔ inference-style rollout decode	Chain aggregate metrics

The ground truth is defined as follows. Take the same low-precision input values the candidate consumes, upcast them to FP32, compute the operator's original mathematical definition entirely in FP32 with TF32 and autocast explicitly disabled, and keep the output in FP32 rather than rounding it back to BF16.

Current thresholds are established from three classes of evidence:

Test conventions of upstream projects, used to establish the reasonable order of magnitude (PyTorch, NVIDIA TransformerEngine).
Floating-point error analysis, used to explain why different dtypes and operator classes require different tiers.
Measurements on H20 in a fixed environment, used to calibrate the two rows for which no upstream baseline of the same comparison semantics exists: logprob and backward reduction.

An important qualification: an external project's threshold supplies a reference magnitude; it does not by itself prove that value is appropriate for RL-Kernel. A final threshold must satisfy three conditions simultaneously: well-defined comparison semantics, measured coverage on a fixed workload, and traceable version information. Section 6 therefore distinguishes three modes of adoption — direct adoption, partial reference, and measured calibration — and never reports a single-field match as a full match.

2. Normative Scope and dtype Policy¶

2.1 Target dtype policy¶

Item	Normative target
Candidate execution	BF16
Reference computation	FP32
Reduction / accumulation	FP32 by default; any exception must be declared per operator
Default candidate output	Same as execution dtype
Logprob aggregate output	FP32
TF32	Disabled for the reference and for the WS1 required candidate path
Autocast	Disabled on the reference path
Backend-private tolerance relaxation	Not permitted

2.2 On `backend_private_tolerance_relaxation: false`¶

This field means that the CUDA and Triton profiles must both pass against the same WS1 contract, and that neither backend may relax a threshold locally inside the formal WS1 gate.

It does not mean that every legacy unit test in the repository has had its local tolerance removed. As of the audit baseline, tests/ still contains 306 hard-coded atol= occurrences across 36 files that do not resolve through the contract. The formal WS1 gate uses the unified contract; the convergence plan for legacy tests is item 13 in Section 11.

2.3 Operator class partition¶

Thresholds are given per operator class rather than per operator, because the error magnitude is driven primarily by reduction topology:

Op class	Operators covered	Dominant error source
`elementwise`	embedding, rope, silu, swiglu, pack	A single output rounding
`reduction`	rms_norm, qk_norm, lm_head, det_gemm	Long-reduction accumulation
`attention`	attention	Two GEMM stages plus a softmax reduction
`logprob`	logp, linear_logp, batch_invariant_logp	Vocabulary-width log-softmax

3. Definition of the FP32 Reference¶

3.1 Normative definition¶

An FP32 reference that satisfies WS1 must meet all five conditions:

Same-source inputs. The candidate and the reference consume the same low-precision input values; the reference upcasts those already-quantized values to FP32 rather than generating a separate set of high-precision inputs.
Computation from the original definition. The operator's mathematical definition is evaluated in FP32, including reduction, normalization, softmax, log-softmax, and the GEMM accumulator.
No implicit precision loss. TF32 and autocast are explicitly disabled so the reference cannot silently enter a low-precision path.
FP32 output. When compared against a BF16 candidate, the reference output is not first rounded to BF16.
Same-source gradients. Backward uses gradients derived from the same FP32 forward definition.

Condition 1 is the core of this definition: it places input quantization error outside the measurement, so that a comparison reflects only the kernel's own error plus one output rounding, and can therefore be compared directly against the floating-point theoretical floor.

The implementation lives in the forward_fp32 methods under rl_engine/kernels/ops/pytorch/. RMSNorm is representative (rl_engine/kernels/ops/pytorch/norm/rms_norm.py#L136-L140):

x_f = x.float()                         # same BF16 inputs, upcast to FP32
var = x_f.pow(2).mean(dim=-1, keepdim=True)
normed = x_f * torch.rsqrt(var + eps)
out = normed * weight.float()
return out.to(output_dtype)             # output_dtype = float32

The forward_fp32 methods for attention and the LM head additionally use a context manager that disables torch.backends.cuda.matmul.allow_tf32 and autocast.

3.2 Relationship to the upstream reference construction¶

Conditions 1 through 3 share their origin with the reference construction used in the FlashAttention test suite, whose reference implementation documents its upcast argument as:

upcast: whether to cast all inputs to fp32, do all computation in fp32, then cast output back to fp16/bf16.

— Dao-AILab/flash-attention, tests/test_flash_attn.py (commit 0f3fb00d)

The two differ in output comparison strategy, and this must be stated plainly. Several FlashAttention tests cast the reference output back down to low precision and apply a relative criterion, namely that the candidate's error may not exceed a small multiple of a plain PyTorch low-precision implementation's error:

assert (out - out_ref).abs().max().item() <= 2 * (out_pt - out_ref).abs().max().item()

The WS1 accuracy reference keeps its output in FP32 and applies fixed atol/rtol criteria. It is therefore accurate to say that WS1 adopts the FlashAttention reference construction, and inaccurate to say the two use the same judgment method. Adoption of the relative criterion is tracked as item 10 in Section 11.

3.3 Why the reference must be FP32¶

The unit roundoff of BF16 is:

u_bf16 = 2^-8 = 3.90625e-3

For operators containing long reductions, a low-precision reference can itself accumulate error of the same order as, or larger than, the candidate's. The classical error analysis of softmax and log-sum-exp gives a forward error bound of roughly (n+3)u and states explicitly when that bound ceases to constrain anything:

they provide no useful information when n ≳ 1/u, and for fp16 this happens for n as small as 2048.

— Blanchard, Higham & Higham, IMA Journal of Numerical Analysis 41(4)

Substituting u_bf16 lowers the failure boundary to n ≳ 256. The actual reduction widths in Qwen3-8B are:

Reduction	Width
RMSNorm hidden width	4096
Vocabulary log-softmax	151936

Both are far beyond that boundary.

This argument does not support the claim that all BF16 results are mathematically invalid. The bound is a worst-case bound, and the same paper notes that under a probabilistic model of rounding errors n may be replaced by a small constant multiple of √n. What the argument does support is narrower and sufficient: a BF16 implementation is unsuitable as the high-precision numerical baseline for WS1, because the baseline's own error guarantee is no longer strong enough to adjudicate the quantity under test.

3.4 The limits of a reference's authority¶

A reference is authoritative because its operator definition is explicit and its dtype path is controlled, not because of the name of the library that produced it. cuBLAS illustrates the point; its reproducibility guarantee carries strict preconditions:

all cuBLAS API routines from a given toolkit version, generate the same bit-wise results at every run when executed on GPUs with the same architecture and the same number of SMs

— NVIDIA cuBLAS Documentation, Results Reproducibility

The same document states that bitwise reproducibility is not guaranteed across toolkit versions, that the guarantee no longer holds when multiple CUDA streams are active, and that results are not guaranteed to be bitwise reproducible when atomics mode is enabled.

Three boundaries follow:

A cuBLAS path with unpinned version, algorithm, and workspace cannot serve as the bitwise oracle for WS1 invariance.
With a fixed environment and a numerical tolerance, cuBLAS can still form part of an accuracy reference.
Whether a reference qualifies is determined by its actual dtype, algorithm, and output semantics, not by which library was called.

4. Judgment Methods¶

4.1 Accuracy¶

Accuracy is judged element-wise, with the reference on the right-hand side:

abs(candidate - reference) <= atol + rtol * abs(reference)

This is the same form used by PyTorch, TransformerEngine, and Liger-Kernel. The implementation is _compare_output in rl_engine/kernels/gtest/op_checks.py, which upcasts both sides to FP32 before calling torch.allclose.

The logprob class uses rtol = 0, because a log-probability can approach zero, at which point relative error loses a stable meaning.

The reported max_abs_error and max_rel_error are observations, not the basis of the verdict. In particular, max_rel_error uses abs(reference).clamp_min(1e-12) as its denominator, so elements near zero produce very large values; readers should rely on the passed field.

4.2 Invariance¶

The normative target for invariance is a true bitwise comparison across batch, chunk, padding, and layout transformations, requiring that:

shapes are identical;
dtypes are identical;
raw bit patterns are identical;
+0 and -0 are treated as distinct bit patterns;
NaN handling is specified explicitly rather than left to the default behavior of allclose.

The atol = 0 and rtol = 0 entries in the contract are the numerical representation of bitwise mode; they are not equivalent to a complete bitwise comparison implementation. The gap between the current implementation and this target is documented in §9.3.

4.3 Chain-level parity¶

Chain-level metrics are managed separately from per-operator accuracy; see Section 8.

5. Threshold Tables¶

5.1 Forward accuracy¶

Op class	FP32 `atol / rtol`	BF16 `atol / rtol`	FP16 optional `atol / rtol`
Elementwise	`1e-5 / 1e-5`	`2e-2 / 1.6e-2`	`1e-3 / 1e-3`
Reduction	`1e-4 / 1e-4`	`5e-2 / 2e-2`	`1e-3 / 1e-3`
Logprob	`1e-5 / 0`	`6e-2 / 0`	`5e-3 / 0`
Attention	`1e-4 / 1e-4`	`5e-2 / 2e-2`	`1e-3 / 1e-3`

5.2 Gradient accuracy¶

Gradient rows are independent of forward rows and are not inherited from them.

Op class	FP32 `atol / rtol`	BF16 `atol / rtol`	FP16 optional `atol / rtol`
Elementwise	`1e-5 / 1e-5`	`2e-2 / 1.6e-2`	`1e-3 / 1e-3`
Reduction	`1e-4 / 1e-4`	`1e-1 / 2e-2`	`1e-3 / 1e-3`
Logprob	`1e-5 / 0`	`5e-2 / 0`	`5e-3 / 0`
Attention	`1e-4 / 1e-4`	`5e-2 / 2e-2`	`1e-3 / 1e-3`

6. Threshold Provenance and Mode of Adoption¶

The table below distinguishes three modes of adoption. A match on atol or rtol alone is not reported as a full match.

Contract row	Primary basis	Mode of adoption	Notes
Elementwise BF16	TransformerEngine RMSNorm test; PyTorch BF16 default `rtol = 1.6e-2`	Field reference	`2e-2 / 1.6e-2` matches the typical acceptance magnitude for BF16 elementwise ops; PyTorch's default pair uses `atol = 1e-5`, which differs from this row
Reduction BF16	TransformerEngine linear accuracy	Direct upstream reference	`5e-2 / 2e-2` applies to operators containing longer reductions
Attention BF16	TransformerEngine DPA forward / backward	Partial reference	`atol = 5e-2` is the same tier; `rtol` is not uniform across tensors and directions upstream, and this row's `rtol` is wider than the upstream value
Gradient reduction BF16	TransformerEngine GPT-level tolerance; H20 backward data	Combined evidence	`atol = 1e-1` is additionally calibrated against an H20 near-zero error of `0.0978`; a larger-magnitude `0.1034` sample is covered by the `rtol = 2e-2` term
Logprob BF16	H20 full-model data	Measured calibration	Upstream CE and logprob tests compare within a single dtype, which is not the BF16↔FP32 comparison semantics of this row, so no directly transferable baseline exists; the historical sample was `max_abs_dlogp = 0.05064`, and the threshold is `0.06`
FP32 elementwise / logprob	PyTorch default tolerance	Field reference (`atol`)	PyTorch's float32 default `atol` is `1e-5`; its default `rtol` of `1.3e-6` differs from this row
FP32 reduction / attention	Magnitude of PyTorch OpInfo matmul and reduction tests	Magnitude reference	WS1 uses a unified `1e-4 / 1e-4` reduction tier
FP16 optional	PyTorch FP16 default `rtol = 1e-3`	Field reference (`rtol`)	`atol` is specified per op class and is not equivalent to PyTorch's complete default pair

6.1 Parameters that are not floating-point tolerances¶

The contract's clip_interval = [0.8, 1.2] is the policy interval of the PPO ratio. It defines the statistical semantics of clipfrac0 and is not a floating-point error tolerance. Its resemblance to identically named parameters in other RL frameworks is a policy convention and does not constitute an endorsement of numerical precision. This document keeps it separate from the threshold tables.

7. Thresholds and Floating-Point Error Magnitudes¶

7.1 BF16¶

BF16 stores 7 mantissa bits, which the implicit leading bit extends to 8 significant bits, giving u ≈ 3.906e-3. WS1 commonly uses rtol between 1.6e-2 and 2e-2, roughly 4u to 5u.

This relationship shows that the thresholds sit at the same order of magnitude as the fundamental rounding scale of a BF16 output. It does not on its own prove the thresholds are sufficient. A final threshold must still cover:

error dominated by atol near zero;
accumulated error introduced by reduction topology;
the differing error distributions of forward and backward;
representative shapes, extreme-valued inputs, and long-tail tokens.

7.2 FP32 reduction¶

For a sequential dot product of length n, the classical worst-case bound is:

gamma_n = n*u / (1 - n*u)

With n = 4096 and u_fp32 = 5.96e-8:

sqrt(n)*u  ≈ 3.8e-6      # empirical scale, contingent on a probabilistic rounding-error model
WS1 threshold  1e-4
gamma_n    ≈ 2.44e-4     # worst case for sequential accumulation

1e-4 lies between the typical error scale and that worst-case bound, and is the engineering tier WS1 currently adopts: below the worst case, so it is a real constraint; above the typical scale, so normal operation does not produce false alarms.

One qualification is required. Real kernels may use pairwise trees, online reductions, MMA tiles, or multi-stage reductions, whose error behavior differs from sequential accumulation. Each op class must therefore be validated against its actual reduction topology and measured distribution; a single conclusion cannot be derived from sequential accumulation at n = 4096 alone.

8. Chain-Level Logprob Parity¶

Metric	Definition	BF16 threshold	Current status
`max_abs_dlogp`	`max(abs(lhs_logp - rhs_logp))`	`0.06`	Required gate; calibrated from historical H20 data
`approx_kl0`	`mean(exp(dlogp) - 1 - dlogp)`	`0.05`	Required report and gate; identified as loose
`clipfrac0`	`mean(1[exp(dlogp) outside clip_interval])`	`0`	Required gate

dlogp is defined as comparison_lhs_logp - comparison_rhs_logp, with the comparison roles fixed by report kind: for train_infer_logprob_parity the left and right sides are training-style teacher forcing and inference-style rollout decode respectively.

8.1 `max_abs_dlogp`¶

The current value of 0.06 derives from max_abs_dlogp = 0.05064 recorded in a historical H20 full-model run. That calibration covers only 27 active tokens, which is too few to estimate the tail of the distribution stably. The current value should therefore be treated as a provisional, version-scoped engineering threshold.

Recalibration requires at minimum:

an active-token fixture on the order of a thousand tokens;
multiple seeds, sequence lengths, and logit scales;
max, mean, P99, and P99.9 distribution statistics;
the raw report, fixture hash, code commit, driver, CUDA, PyTorch, and Triton versions, and GPU model.

8.2 `approx_kl0`¶

The historical H20 samples were 1.4e-4 and 2.5e-4. The current threshold of 5e-2 is roughly two orders of magnitude above them and exerts limited constraint.

This is not an oversight. The contract recorded the value as provisional when it was introduced:

pending measured chain-level distributions, after which its threshold may be tightened

— tolerance_contract.json, approx_kl0.threshold_rationale

The proposed value of 1e-3 is closer to the magnitude used by several RL training systems, but until new distribution data is generated and the contract is updated, it remains a proposal and not the standard in force.

8.3 Proposed changes¶

Tighten approx_kl0 to 1e-3.
Add mean_abs_dlogp with a proposed threshold of 0.005.
Decide, on the basis of the new measured distribution, whether max_abs_dlogp should be demoted from a hard gate to monitor-only.

These changes alter pass/fail semantics and must be released with a new contract version (ws1-c1-v3 is suggested). They cannot be enacted by editing this document alone.

9. Implementation Conformance Audit¶

The findings below are based on audit baseline 6ceeb62e073d74e00a8c5b56af538ee1a7417773, with each item tied to a verifiable source location. Line numbers refer to that commit.

9.1 Summary table¶

Component	Status	Basis
PyTorch FP32 references (RMSNorm / attention / LM-head / activation / RoPE)	Conformant	`forward_fp32` computes entirely in FP32 with TF32 and autocast explicitly disabled
CUDA RMSNorm / QK-Norm	Conformant	Statistics and reduction in FP32
CUDA attention	Conformant	Row max, rescale, and softmax in FP32; `csrc/cuda/attention/prefix_shared_attention.cu#L127`, `#L237-L253`
CUDA SiLU / SwiGLU / RoPE	Conformant	Intermediates in FP32, output stored at execution dtype
Triton standard attention	Conformant	Contains no `tl.dot`; 16 explicit `tl.float32` accumulators; `rl_engine/kernels/ops/triton/attention/standard_attn.py`
Triton linear_logp	Conformant	`tl.dot(..., input_precision="ieee")` at `#L77`; lse and logp allocated FP32 at `#L109-L110`
Chain LM-head / logprob / residual stream	Conformant, and stricter than defaults	`score_logits` in FP32, residual stream in FP32 throughout; `rl_engine/alignment/qwen3_dense.py#L683`, `#L707-L729`
Deterministic GEMM accumulation	Exception not declared in the contract	See §9.2
FP32 reference for `det_gemm`	Non-conformant	See §9.3
Invariance comparator	Non-conformant	See §9.4
CUDA generic logp output dtype	Non-conformant	See §9.5
RMSNorm reference rounding order	Not pinned in the contract	See §9.6
Legacy unit-test tolerances	Pending convergence	306 `atol=` occurrences across 36 files in `tests/` that do not resolve through the contract
Legacy Triton attention paths	Scope must be bounded	See §9.7

9.2 Deterministic GEMM uses FP32 leaves with a BF16 tree merge¶

§2.1 states that reduction and accumulation are FP32 by default and that any exception must be declared per operator. The deterministic GEMM takes such an exception, and the exception is documented in the kernel but is not carried in the contract.

The CUDA kernel reduces K with a mid-split tree whose leaves accumulate in FP32 and whose merges are performed in BF16 (csrc/cuda/gemm/det_gemm_kernel.cu#L68-L84):

constexpr int K_TREE_LEAF = 32;

__device__ __forceinline__ nv_bf16 bf16_add(nv_bf16 a, nv_bf16 b) {
  return __float2bfloat16(__bfloat162float(a) + __bfloat162float(b));
}

__device__ nv_bf16 k_tree_naive(..., int lo, int hi) {
  if (hi - lo <= K_TREE_LEAF) {
    float acc = 0.0f;                       // FP32 leaf accumulation
    for (int k = lo; k < hi; ++k)
      acc += __bfloat162float(A[row * K + k]) * __bfloat162float(B[k * N + col]);
    return __float2bfloat16(acc);           // leaf result rounded to BF16
  }
  const int mid = lo + (hi - lo) / 2;
  return bf16_add(k_tree_naive(..., lo, mid),
                  k_tree_naive(..., mid, hi));   // BF16 tree merge
}

The file header states the design and its rationale (#L11-L14):

Both: BF16 in / FP32 accum / BF16 store / no TF32 / no split-K. K is reduced with a mid-split tree. A contiguous half-K GEMM is one child, so simulated TP=2 (a+b) matches TP=1. TP=8 left-fold does not. Leaves stay FP32 (naive: 32-wide MAC; SM90: one BK).

A static_assert binds the SM90 tile width to the same leaf size (#L205), so the SM90 and fallback paths produce the same reduction tree.

The Triton implementation matches this topology: tile accumulators are tl.float32 (rl_engine/kernels/ops/triton/matmul/det_gemm.py#L394, #L466, #L484, #L719), while the tree workspace is allocated BF16 (#L847-L849), so each merge level loads to FP32, adds, and rounds back to BF16 on store.

This is a deliberate design choice, not a defect. The BF16 merge is what makes a simulated TP=2 split reproduce the TP=1 result bitwise, which serves the invariance objective directly. The gap is documentary rather than numerical: the contract declares accumulation_dtype: float32 globally and requires per-operator exceptions to be declared, but carries no such declaration for det_gemm.

Two resolutions are available, and the project must choose one explicitly:

Option A — FP32 throughout the tree. Keep the partial workspace and every merge in FP32, rounding only at the output boundary. Appropriate if accumulation_dtype: float32 is to remain a hard global rule. This changes the TP-invariance property described in the header comment and must be re-validated against it.
Option B — declare the exception. Record det_gemm in the contract as FP32 leaf accumulation + BF16 tree merge + BF16 output, with the TP-invariance rationale, and state the leaf width (K_TREE_LEAF = 32) as part of the declared semantics.

The implementation currently matches Option B; the contract currently asserts Option A. A specification and an implementation must not carry two mutually contradictory descriptions.

9.3 The `det_gemm` FP32 reference does not satisfy §3.1¶

The gold path for det_gemm resolves to NativeGemmOp.__call__ (rl_engine/kernels/gtest/operator_specs.py#L148-L149), implemented as:

def __call__(self, a, b):
    return torch.matmul(a, b)

With BF16 inputs this call returns BF16, so the gold output dtype is BF16. That violates condition 4 of §3.1 and conflicts with the contract's declared reference_dtype: float32. Measurement confirms gold_dtype = torch.bfloat16.

The module's own docstring already marks it as unsuitable for this role:

torch.matmul (cuBLAS) does NOT guarantee batch-invariance … This op exists only as a correctness reference and benchmark target, NOT as a fallback.

The consequence is that the accuracy judgment for det_gemm compares two BF16 GEMMs rather than a BF16 candidate against an FP32 reference. Historical det_gemm accuracy results obtained under this condition do not constitute verification against the §3.1 definition. This defect is independent of §9.2: it concerns the reference, whereas §9.2 concerns the candidate kernel.

op_checks.py contains a provenance check that verifies the gold dtype equals reference_dtype, but it fires only when provenance is supplied. The scripts/check_operator.py path does not supply it, so the mismatch is not intercepted.

9.4 Invariance judgment is not a raw-bit comparison¶

_compare_logical_tensors in forward_invariance.py upcasts both sides to FP32 and then calls torch.allclose:

canonical_fp32 = canonical.float()
transformed_fp32 = transformed.float()
...
passed = bool(torch.allclose(transformed_fp32, canonical_fp32, atol=atol, rtol=rtol))

(rl_engine/kernels/gtest/forward_invariance.py#L297-L352)

_compare_parameter_grad in gradient_invariance.py reuses the same function (#L308-L331), so both invariance judgments share this implementation.

Because the BF16-to-FP32 upcast is lossless, this implementation is equivalent to exact numerical equality when dtype and shape match. The gap against the §4.2 target is confined to two points: +0 and -0 are judged equal, and NaN handling is left to the default behavior of allclose rather than being specified.

Reporting requirement. Until the comparator is corrected, the historical C3 / C4 / C8 invariance results must be described as zero-tolerance numeric equality and must not be described as having completed strict bitwise verification.

9.5 CUDA generic logp output dtype conflicts with the policy¶

fused_logp_forward allocates its output from logits.options():

auto output = torch::empty({logits.size(0)}, logits.options());

(csrc/fused_logp_kernel.cu#L565-L567)

BF16 logits therefore produce a BF16 logprob, conflicting with the Logprob aggregate output = FP32 entry in §2.1. The same file already provides an FP32 entry point, fused_logp_forward_fp32 (#L571-L573), which specifies at::ScalarType::Float explicitly.

The formal WS1 candidate should call the FP32 output entry point, and the provenance validator should verify logprob_aggregates_dtype = float32 per op class. Changing the output dtype may change the existing 0.05064 calibration, so the logprob threshold must be re-measured afterward.

9.6 RMSNorm reference rounding order is not pinned in the contract¶

Mainstream implementations differ in the final rounding step of RMSNorm:

Order	Expression	Adopted by
HF order	`weight * x.to(bf16)`; round first, then multiply by weight	HuggingFace `Qwen3RMSNorm`; vLLM `forward_native`; SGLang with `cast_x_before_out_mul=True`
Kernel order	`(x * weight).to(bf16)`; multiply in FP32, round once	vLLM CUDA kernel; SGLang `forward_native` default

The WS1 NativeRMSNormOp uses kernel order (rl_engine/kernels/ops/pytorch/norm/rms_norm.py#L136-L140):

x_f = x.float()
var = x_f.pow(2).mean(dim=-1, keepdim=True)
normed = x_f * torch.rsqrt(var + eps)
out = normed * weight.float()
return out.to(output_dtype)

The two orders differ by roughly 1 ulp, and SGLang provides a dedicated switch specifically to align with HF semantics.

The choice itself is defensible, but the contract does not currently declare which order it adopts. Since the rounding order used by the training-side framework determines the achievable floor for chain-level parity, this should be pinned in the contract with its rationale recorded.

9.7 Bounding the TF32 claim¶

When Triton executes an FP32 tl.dot on NVIDIA GPUs, the default input_precision may be TF32. The audit confirms:

The WS1 required path handles this correctly. det_gemm.py passes allow_tf32=False (#L729), linear_logp.py passes input_precision="ieee" (#L77), and standard_attn.py contains no tl.dot at all.
Two legacy paths do not. rl_engine/kernels/ops/triton/triton_attn.py contains 7 tl.dot calls and rl_engine/kernels/ops/triton/attention/chunked_flash_attn.py contains 4, none with a precision specified. Neither file is referenced by rl_engine/kernels/gtest/ or rl_engine/alignment/, so neither is in the C8 required-op matrix. chunked_flash_attn.py is reached through rl_engine/kernels/ops/rocm/attention/flash_attn.py, and ROCm is out of WS1 scope per the header.

The TF32 entry in §2.1 is therefore scoped to the WS1 required path. The claim "TF32 disabled throughout" must not be extended into an unverified repository-wide statement.¶

10. Historical Evidence and Its Applicability¶

The checked-in C8 H20 evidence originates from:

Field	Value
Source commit	`fdf5bcc5165820abb506291a29370225306514ca`
GPU	NVIDIA H20, CC 9.0
Driver	580.76.05
CUDA / PyTorch / Triton	12.8 / 2.8.0+cu128 / 3.4.0
Workload	`ws1-qwen3-8b-dense-primary-v6` / `ws1-c2-v7`
Result	`green = 176`, `N/A = 16`, `red = 0`

This result attests only to the historical execution state at that fixed source commit in that fixed environment. It does not automatically attest to current main, and it does not substitute for final-commit GPU CI. Specific limitations:

the evidence records cell status and lacks complete per-tensor raw error distributions;
the "bitwise gate" of that period was implemented by the zero-tolerance allclose described in §9.3;
csrc, Triton, gtest, scripts, and tests have all changed between the evidence commit and the audit baseline.

Before a new WS1 document or contract is released, evidence should be regenerated at the target commit, with an artifact containing at minimum:

source commit and dirty-tree status;
expected and actual backend plus the underlying kernel ID;
input fixture ID, hash, shape, and seed;
dtype, TF32, autocast, and key runtime readbacks;
per-tensor max / mean / P99 / P99.9 absolute and relative error;
dtype, shape, and raw-bit results of the bitwise comparison;
the complete software and hardware environment.

11. Remediation Plan¶

11.1 Release-blocking items¶

Until these are closed, this document remains Draft Normative and the repository must not be described as fully satisfying the WS1 precision standard.

#	Item	Section	Impact
1	Decide between Option A (FP32 throughout the reduction tree) and Option B (declare the `det_gemm` exception, including `K_TREE_LEAF`, in the contract), then align specification and implementation	§9.2	Option A changes the TP-invariance property and requires re-validation; Option B requires a contract revision
2	Change the `det_gemm` gold to an FP32 GEMM (inputs upcast, TF32 disabled), with backward updated to match; supply provenance on the `check_operator.py` path so the dtype check fires	§9.3	`det_gemm` accuracy results must be regenerated
3	Change the invariance gate to a true raw-bit comparison with explicit `±0` and NaN semantics	§9.4	C3 / C4 / C8 invariance conclusions must be restated and re-run
4	Switch CUDA generic logp to the FP32 output entry point and add op-aware provenance validation	§9.5	The logprob threshold must be recalibrated
5	Declare the RMSNorm rounding order in the contract and record its relationship to the training-side framework	§9.6	Affects the explanation of the achievable chain-parity floor
6	Define the required-op scope so that legacy Triton paths without explicit TF32 disabling are either excluded or fixed	§9.7	Affects the defensible scope of the TF32 claim
7	Re-run H20 GPU evidence at the target release commit, producing the full artifact listed in §10	§10	Affects the currency of every historical conclusion

11.2 Threshold evolution items¶

These change pass/fail semantics and must ship with a new contract version (ws1-c1-v3 suggested).

#	Item	Section
8	Tighten `approx_kl0` from `5e-2` to `1e-3`	§8.2
9	Add `mean_abs_dlogp` with a proposed threshold of `0.005`	§8.3
10	Recalibrate `max_abs_dlogp` on a thousand-token fixture and decide whether it becomes monitor-only	§8.1, §8.3
11	Introduce the FlashAttention-style relative criterion (candidate error within 2× the PyTorch BF16 path's error, 3× for backward), retaining the current absolute thresholds as caps	§3.2, §7.1
12	Move gradient judgment to rel-peak or RMS normalization, replacing the scale-dependent absolute `atol`	§7.1
13	Converge the 306 hard-coded tolerances in `tests/` onto the unified contract	§2.2

11.3 Release criteria¶

Once blocking items 1 through 7 are closed, this document should be promoted to Normative, with the contract version, source commit, evidence artifact hash, and verification environment pinned together in the header.

12. Reference Index¶

12.1 External sources¶

External references use commit permalinks so that movement of a default branch does not break traceability. The links and quoted text below were verified on 2026-09-11.

Source	Link	Cited in
PyTorch tolerance implementation	`torch/testing/_comparison.py` @ `31527a43`	§6
PyTorch default tolerance table	torch.testing documentation	§6
PyTorch OpInfo tolerance overrides	`common_methods_invocations.py`	§6
TransformerEngine numerical tests	`tests/pytorch/test_numerics.py` @ `02f8e754`	§6
FlashAttention reference construction	`tests/test_flash_attn.py` @ `0f3fb00d`	§3.2
Softmax / log-sum-exp error analysis	arXiv:1909.03469	§3.3
cuBLAS reproducibility	NVIDIA cuBLAS Documentation	§3.4
NVIDIA floating-point compliance	Floating Point and IEEE 754 Compliance	§3.4
Triton FP32 dot precision	`triton.language.dot`	§9.6
Qwen3 modeling code	`modeling_qwen3.py`	§9.5

12.2 In-repository sources¶

Paths and line numbers refer to audit baseline 6ceeb62e073d74e00a8c5b56af538ee1a7417773.

Content	Path
Precision contract	`rl_engine/kernels/gtest/tolerance_contract.json`
Four-judgment matrix	`rl_engine/kernels/gtest/four_judgment_matrix.py`
Accuracy comparator	`rl_engine/kernels/gtest/op_checks.py`
Invariance comparators	`rl_engine/kernels/gtest/forward_invariance.py`, `gradient_invariance.py`
Operator gold mapping	`rl_engine/kernels/gtest/operator_specs.py`
PyTorch FP32 references	`rl_engine/kernels/ops/pytorch/`
Chain model	`rl_engine/alignment/qwen3_dense.py`
Workload manifest	`rl_engine/testing/ws1_manifest.json`
H20 C8 evidence	`docs/design/ws1-c8-execute.json`

12.3 Companion research reports¶

The external basis for this document is drawn from two full research reports containing all verbatim quotations, URLs, and unverified-item annotations:

docs/design/ws1-tolerance-contract-review.md — threshold and ground truth compared against external standards
docs/design/ws1-qwen3-8b-precision-and-mismatch-survey.md — per-operator precision conventions for Qwen3-8B and the training-inference mismatch literature