From fb559f02e02edf2e4a974d778666c003ae6fb6c4 Mon Sep 17 00:00:00 2001 From: Hao Wu Date: Sat, 7 Nov 2026 17:27:27 -0500 Subject: [PATCH] [FIX] Accept the rounding mode of fp_to_fp in symbolic clients `x.to(tl.float8e5, fp_downcast_rounding="rtne")` passes an `ir.ROUNDING_MODE` enum to the builder's `create_fp_to_fp`. The symbolic engine stored it as a child expression, and `SymbolicExpr.from_value` raised "Unsupported type: ". The child now holds the enum's integer value. Co-Authored-By: Claude Opus 5.5 --- tests/end_to_end/test_sanitizer.py | 23 +++++++++++++++++++++++ tilelens/clients/symbolic_engine.py | 4 ++++ 2 files changed, 27 insertions(+) diff --git a/tests/end_to_end/test_sanitizer.py b/tests/end_to_end/test_sanitizer.py index a039a906a..14e822a81 100644 --- a/tests/end_to_end/test_sanitizer.py +++ b/tests/end_to_end/test_sanitizer.py @@ -1209,6 +1209,29 @@ def test_sanitizer_supports_data_dependent_cumsum_index(): assert len(cumsum_sanitizer.records) == 0 +fp8_rounding_sanitizer = SymbolicSanitizer(abort_on_error=False) + + +@tilelens.trace(client=fp8_rounding_sanitizer) +@triton.jit +def fp8_rtne_cast_kernel(x_ptr, out_ptr, BLOCK: tl.constexpr): + offs = tl.arange(0, BLOCK) + x = tl.load(x_ptr + offs) + y = x.to(tl.float8e5, fp_downcast_rounding="rtne") + tl.store(out_ptr + offs, y.to(tl.float32)) + + +def test_sanitizer_supports_fp_downcast_rounding(): + fp8_rounding_sanitizer.records.clear() + + x = torch.randn((8,), dtype=torch.float32) + out = torch.empty_like(x) + + fp8_rtne_cast_kernel[(1,)](x, out, BLOCK=x.numel()) + + assert len(fp8_rounding_sanitizer.records) == 0 + + ashr_index_sanitizer = SymbolicSanitizer(abort_on_error=False) diff --git a/tilelens/clients/symbolic_engine.py b/tilelens/clients/symbolic_engine.py index af435c88a..603ea4c61 100644 --- a/tilelens/clients/symbolic_engine.py +++ b/tilelens/clients/symbolic_engine.py @@ -2231,6 +2231,10 @@ def __init__(self, op: str, src: Any, dst_type: Any, rounding_mode: Any): super().__init__(op) self.add_child("src", src) self.add_child("dst_type", dst_type) + # Triton passes an ir.ROUNDING_MODE enum (or None); keep its integer + # value so that the child is an ordinary constant. + if rounding_mode is not None and not isinstance(rounding_mode, SymbolicExpr): + rounding_mode = int(rounding_mode) self.add_child("rounding_mode", rounding_mode) self.dtype, self.shape = self._unpack_dtype( self.dst_type.to_py(), self.src.shape