DeepInfer/scripts/h20_ops_ab.py
loopyt afbfdb6c7a feat(triton): port the sampler, bit-exact on the draw and honest about the rest
`fused_sample` for the configuration the engine builds: 4096 columns after the force_topk cut,
sampling on, rejection off. With one tile the baseline's two-level search collapses -- the coarse
offset *is* the token id -- so this implements log-softmax, the top-k/top-p cut as a threshold plus
the tie rule, renormalisation, a fixed-point cumulative distribution and one Philox draw.

The draw matches cuRAND bit for bit (that is what the previous commit established). The rest does
not all match, and the split is measured rather than claimed:

    engine's configuration (top_p=0.99, 2644 kept columns)   tokens identical on every row
    narrow cuts (top_k=128, top_p=0.5)                       ~90% of rows

Kept sets match column for column and log-probabilities agree to ~1e-5 in every case. What differs
is the token, and the cause is a discontinuity in the baseline rather than a defect here: the draw
is `d % total` with `total = sum(floor(p_i * 2^31))`, and the baseline's rejection limit
`(2^32-1) // total * total` steps from `2 * total` to `1 * total` exactly at `total == 2^31` -- which
`total` sits next to, because the probabilities sum to one. On the near side almost no draw is
rejected, on the far side about half are, and each rejection moves to the next draw. A few units of
rounding error in `total`, which is a reduction order that cannot match, therefore change the token
outright. Clamping `total` under 2^31 does not help: the baseline crosses in either direction too,
so it only relocates the disagreement (30 rows instead of 11).

Four bugs found on the way, all worth recording because each looked like something else:

- Triton evaluates `//`, `%` and `<=` on **uint32 operands with signed arithmetic**, so 4294967295
  behaves as -1. That made `limit` come out as `2^32 - total`, flipping the rejection decision.
  Fixed by doing the whole cumulative distribution in int64.
- Nested runtime `if`s produced a threshold that was not the one assigned inside the branch; the
  cuts are now combined with `tl.where` instead.
- The threshold was taken from the probability array and compared against logits. The sort has to be
  of the logits (equivalent, since the map is monotone -- the mass is recomputed from the sorted
  logits rather than sorting twice).
- Renormalising by division rather than through the log domain, as the baseline does, moved the
  token agreement from 76% to 91% before the remaining gap was understood.

Also gives the A/B harness a fourth case element: an agreement floor for outputs whose answer is a
random draw, so the rate is reported rather than compared for equality -- a comparison that can pass
without testing anything would be worse than none.
2026-10-22 08:55:10 +08:00

207 lines
8.2 KiB
Python

"""Compare a Triton kernel against the TileLang one it replaces, through the shared interface.
Every interface module (``gather.py``, ``norm.py``, ...) calls ``ops.<name>_tl`` and then uses the
result. That means the strongest available check is not "does Triton match a Python reference"
but "does the same public call give the same answer when the kernel behind it changes". This
script does exactly that: it swaps ``ops.<name>_tl`` between the two backends and compares the
output of the public function.
python scripts/h20_ops_ab.py # every case
python scripts/h20_ops_ab.py gather norm
"""
from __future__ import annotations
import argparse
import sys
from collections.abc import Callable
from typing import Any
import torch
from deepinfer.model_executor.layers.tilelang_ops import ops
# Each backend is a package whose submodules export `*_tl`; the packages' own `__init__` does not
# re-export, so collect the factories the same way the dispatcher does.
CUDA = ops._collect("cuda") # noqa: SLF001 - the collector is the one place that knows the layout
TRITON = ops._collect("triton") # noqa: SLF001
def swap(name: str, backend: str) -> Callable[[], None]:
"""Point ``ops.<name>`` at one backend, returning a function that restores it."""
previous = getattr(ops, name)
source = CUDA if backend == "cuda" else TRITON
if name not in source:
raise AttributeError(f"backend {backend!r} does not export {name!r}")
setattr(ops, name, source[name])
def restore() -> None:
setattr(ops, name, previous)
return restore
def run_both(name: str, fn: Callable[[], Any]) -> tuple[Any, Any]:
"""Run ``fn`` once against each backend, returning (cuda result, triton result)."""
restore = swap(name, "cuda")
try:
expected = fn()
finally:
restore()
restore = swap(name, "triton")
try:
actual = fn()
finally:
restore()
return expected, actual
# Tolerance follows the dtype's own precision. Anything narrower than float32 is compared at its
# representable resolution, because a different reduction order in a different compiler moves the
# last bit and nothing more: for `norm` that is 10 of 655360 elements differing by one bf16 ulp,
# while the fp8 bytes and every scale factor come out identical.
_DTYPE_TOL: dict[torch.dtype, tuple[float, float]] = {
torch.bfloat16: (1e-2, 1e-2),
torch.float16: (1e-3, 1e-3),
torch.float32: (1e-5, 1e-6),
}
def compare(
name: str,
expected: Any,
actual: Any,
*,
exact: bool,
rtol: float | None = None,
atol: float | None = None,
agreement: float | None = None,
) -> bool:
"""Report whether one case agrees; returns True on success."""
if (
agreement is not None
and isinstance(expected, torch.Tensor)
and expected.dtype
in (
torch.int32,
torch.int64,
)
):
# Random-draw outputs: require a high agreement rate rather than equality, and say what the
# rate was so a regression that halves it cannot hide behind the threshold.
same = int((expected == actual).sum())
total = expected.numel()
rate = same / max(total, 1)
if rate >= agreement:
if rate < 1.0:
print(f" note {name}: {total - same}/{total} differ (random draw on a boundary)")
return True
print(f" FAIL {name}: agreement {rate:.4f} < {agreement} ({total - same}/{total})")
return False
if isinstance(expected, torch.Tensor):
if expected.dtype != actual.dtype:
print(f" FAIL {name}: dtype {expected.dtype} vs {actual.dtype}")
return False
if expected.shape != actual.shape:
print(f" FAIL {name}: shape {tuple(expected.shape)} vs {tuple(actual.shape)}")
return False
if exact or not expected.is_floating_point():
same = torch.equal(expected, actual)
if not same:
bad = expected != actual
print(f" FAIL {name}: {int(bad.sum())}/{expected.numel()} elements differ")
idx = bad.nonzero()[0].tolist()
print(f" first at {idx}: {expected[tuple(idx)].item()} vs {actual[tuple(idx)].item()}")
return bool(same)
if expected.dtype in (torch.float8_e4m3fn, torch.float8_e5m2):
# Compare at the format's own resolution: a one-ulp difference is a rounding tie that
# moved, not a wrong value, so byte distance is the honest measure. Require the scales
# to be exact elsewhere and the bytes to agree almost everywhere.
delta = (expected.view(torch.uint8).int() - actual.view(torch.uint8).int()).abs()
frac = float((delta > 0).float().mean())
if int(delta.max()) <= 2 and frac <= 1e-3:
return True
print(f" FAIL {name}: {frac * 100:.4f}% bytes differ, max byte distance {int(delta.max())}")
return False
if rtol is None or atol is None:
default_rtol, default_atol = _DTYPE_TOL.get(expected.dtype, (1e-5, 1e-6))
rtol = default_rtol if rtol is None else rtol
atol = default_atol if atol is None else atol
close = torch.allclose(expected.float(), actual.float(), rtol=rtol, atol=atol)
if not close:
diff = (expected.float() - actual.float()).abs()
scale = expected.float().abs().clamp_min(1e-12)
print(f" FAIL {name}: max abs {diff.max().item():.3e}, max rel {(diff / scale).max().item():.3e}")
return bool(close)
if isinstance(expected, (tuple, list)):
# A case that returns several tensors is only right if every element is; recurse so each
# one gets its own report instead of collapsing into one boolean.
if len(expected) != len(actual):
print(f" FAIL {name}: {len(expected)} outputs vs {len(actual)}")
return False
ok = True
for i, (e, a) in enumerate(zip(expected, actual, strict=True)):
ok &= compare(f"{name}[{i}]", e, a, exact=exact, rtol=rtol, atol=atol, agreement=agreement)
return ok
same = expected == actual
if not same:
print(f" FAIL {name}: {expected!r} vs {actual!r}")
return bool(same)
def cases() -> dict[str, Callable[[], list[tuple[str, bool, Callable[[], Any]]]]]:
"""Registry of per-kernel case builders, filled in as kernels are ported."""
from scripts.ab_cases import CASES
return CASES
def main() -> None:
parser = argparse.ArgumentParser()
parser.add_argument("kernels", nargs="*", help="subset of kernel names; default is all ported")
parser.add_argument("--verbose", action="store_true")
args = parser.parse_args()
registry = cases()
selected = args.kernels or sorted(registry)
unknown = [name for name in selected if name not in registry]
if unknown:
raise SystemExit(f"no A/B cases for: {', '.join(unknown)}")
failures = 0
for kernel in selected:
print(f"[{kernel}]")
for case in registry[kernel]():
# A case is `(name, exact, fn)`, optionally with a fourth element: the fraction of
# entries allowed to differ. Only for outputs whose answer is a random draw, where the
# last bit of a cumulative mass decides which side of a boundary the token falls on;
# everything else is compared exactly or within a dtype tolerance.
case_name, exact, fn = case[:3]
agreement = case[3] if len(case) > 3 else None
try:
expected, actual = run_both(kernel, fn)
except Exception as exc: # noqa: BLE001 - a raise on one backend is itself a result
lines = str(exc).splitlines()
print(f" FAIL {case_name}: {type(exc).__name__}: {lines[0][:160] if lines else ''}")
failures += 1
continue
if not compare(case_name, expected, actual, exact=exact, agreement=agreement):
failures += 1
elif args.verbose:
print(f" ok {case_name}")
print()
if failures:
print(f"AB_FAILED {failures}")
sys.exit(1)
print("AB_OK")
if __name__ == "__main__":
main()