Skip to content

Commit 59927cc

Browse files
gmagogsfmclaude
andauthored
Recreate foreign SymInts in helion's own ShapeEnv (#1671)
Signed-off-by: Yanan Cao <gmagogsfm@gmail.com> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
1 parent a78c4ce commit 59927cc

2 files changed

Lines changed: 119 additions & 9 deletions

File tree

helion/_compiler/compile_environment.py

Lines changed: 61 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -14,13 +14,17 @@
1414
import torch
1515
from torch._dynamo.source import EphemeralSource
1616
from torch._dynamo.source import LocalSource
17+
from torch._dynamo.source import TensorProperty
18+
from torch._dynamo.source import TensorPropertySource
1719
from torch._inductor.codegen.wrapper import (
1820
user_defined_triton_kernel_transitive_closure_source_code,
1921
)
2022
from torch._inductor.runtime.runtime_utils import next_power_of_2
2123
from torch._subclasses import FakeTensor
2224
from torch._subclasses import FakeTensorMode
25+
from torch.fx.experimental.symbolic_shapes import DimDynamic
2326
from torch.fx.experimental.symbolic_shapes import ShapeEnv
27+
from torch.fx.experimental.symbolic_shapes import free_unbacked_symbols
2428
from torch.utils._sympy.symbol import SymT
2529
from torch.utils._sympy.symbol import symbol_is_type
2630

@@ -150,6 +154,7 @@ def __init__(
150154
self.specialized_vars: set[sympy.Symbol] = set()
151155
self.specialized_strides: set[tuple[str, int]] = set()
152156
self._symint_cache: dict[object, torch.SymInt] = {}
157+
self._foreign_symint_cache: dict[tuple[int, sympy.Expr], torch.SymInt] = {}
153158
self.device_load_count = (
154159
0 # Track number of loads in all device code for eviction policy tuning
155160
)
@@ -548,17 +553,65 @@ def to_fake(self, obj: object, origin: Origin) -> object:
548553

549554
raise TypeError(f"unsupported argument type {type(obj)} ({origin})")
550555

556+
def _maybe_recreate_symint(
557+
self,
558+
s: int | torch.SymInt,
559+
source: Source,
560+
) -> int | torch.SymInt:
561+
"""Create a fresh SymInt in our ShapeEnv that mirrors a foreign one."""
562+
if isinstance(s, int):
563+
return s
564+
outer_se = s.node.shape_env
565+
if outer_se is self.shape_env:
566+
return s
567+
expr = s.node.expr
568+
cache_key = (id(outer_se), expr)
569+
cached = self._foreign_symint_cache.get(cache_key)
570+
if cached is not None:
571+
return cached
572+
if free_unbacked_symbols(expr):
573+
result = self.create_unbacked_symint()
574+
else:
575+
hint = int(shape_env_var_hints(outer_se)[expr])
576+
new_expr = self.shape_env.create_symbol(
577+
hint, source, dynamic_dim=DimDynamic.DYNAMIC
578+
)
579+
result = self.shape_env.create_symintnode(
580+
new_expr, hint=hint, source=source
581+
)
582+
self._foreign_symint_cache[cache_key] = result
583+
return result
584+
551585
def _to_fake_tensor(self, tensor: torch.Tensor, source: Source) -> torch.Tensor:
552586
assert CompileEnvironment.current() is self
553587
assert not self.fake_mode.is_our_fake(tensor)
554-
if isinstance(tensor, FakeTensor) or self.settings.static_shapes:
555-
# When the input is already a FakeTensor (from an outer tracing
556-
# context, e.g. torch.compile calling a custom op's fake impl),
557-
# we cannot pass it directly because it belongs to a different
558-
# FakeTensorMode and would cause "Mixing fake modes" errors.
559-
# We also cannot use from_real_tensor because it tries to read
560-
# concrete sizes which fails on unbacked SymInts. empty_strided
561-
# re-wraps it in our mode while preserving the symbolic sizes.
588+
if isinstance(tensor, FakeTensor):
589+
# FakeTensor from an outer tracing context (e.g. make_fx, Dynamo).
590+
# Create fresh symbols in our own ShapeEnv to avoid leaking
591+
# foreign symbols whose var_to_range entries are missing,
592+
# which causes assertion failures in _maybe_evaluate_static
593+
# on PyTorch versions without optimization_hint (< 2.12).
594+
new_sizes = tuple(
595+
self._maybe_recreate_symint(
596+
s,
597+
TensorPropertySource(source, TensorProperty.SIZE, i),
598+
)
599+
for i, s in enumerate(tensor.size())
600+
)
601+
new_strides = tuple(
602+
self._maybe_recreate_symint(
603+
s,
604+
TensorPropertySource(source, TensorProperty.STRIDE, i),
605+
)
606+
for i, s in enumerate(tensor.stride())
607+
)
608+
result = torch.empty_strided(
609+
new_sizes,
610+
new_strides,
611+
dtype=tensor.dtype,
612+
device=tensor.device,
613+
)
614+
elif self.settings.static_shapes:
562615
result = torch.empty_strided(
563616
tensor.size(),
564617
tensor.stride(),

test/test_custom_op.py

Lines changed: 58 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,7 @@
22

33
import torch
44
from torch._subclasses.fake_tensor import FakeTensorMode
5+
from torch.fx.experimental.symbolic_shapes import DimDynamic
56
from torch.fx.experimental.symbolic_shapes import ShapeEnv
67

78
import helion
@@ -71,6 +72,36 @@ def _bind_and_run_fake(kernel, x, y):
7172
return compiled(x, y, _launcher=lambda *a, **kw: None)
7273

7374

75+
def _make_backed_fake_tensors(*shapes):
76+
"""Create FakeTensors with backed symbolic sizes from an outer ShapeEnv.
77+
78+
Uses DimDynamic.DYNAMIC to force all dimensions to be symbolic,
79+
matching the behavior of make_fx(tracing_mode="symbolic").
80+
"""
81+
shape_env = ShapeEnv()
82+
mode = FakeTensorMode(shape_env=shape_env)
83+
tensors = []
84+
with mode:
85+
for shape in shapes:
86+
sym_sizes = []
87+
for j, val in enumerate(shape):
88+
source = torch._dynamo.source.TensorPropertySource(
89+
torch._dynamo.source.ConstantSource(f"t{len(tensors)}"),
90+
torch._dynamo.source.TensorProperty.SIZE,
91+
j,
92+
)
93+
sym = shape_env.create_symbol(
94+
val, source, dynamic_dim=DimDynamic.DYNAMIC
95+
)
96+
sym_sizes.append(
97+
shape_env.create_symintnode(sym, hint=val, source=source)
98+
)
99+
tensors.append(
100+
torch.empty(sym_sizes, dtype=torch.float16, device=torch.device("cpu"))
101+
)
102+
return tensors, mode
103+
104+
74105
class TestInferFakeImpl(TestCase):
75106
@skipIfRefEager("compile_config requires host_function")
76107
def test_static_shapes(self):
@@ -82,10 +113,36 @@ def test_static_shapes(self):
82113
self.assertEqual(result.dtype, x.dtype)
83114

84115
@skipIfRefEager("compile_config requires host_function")
85-
def test_dynamic_shapes(self):
116+
def test_unbacked_symints(self):
86117
k = helion.kernel(static_shapes=False, autotune_effort="none")(_k_add)
87118
x, y, mode = _make_fake_tensors((4, 8))
88119
with mode:
89120
result = _bind_and_run_fake(k, x, y)
90121
self.assertEqual(result.shape, x.shape)
91122
self.assertEqual(result.dtype, x.dtype)
123+
124+
@skipIfRefEager("compile_config requires host_function")
125+
def test_backed_symints(self):
126+
k = helion.kernel(static_shapes=False, autotune_effort="none")(_k_add)
127+
(x, y), mode = _make_backed_fake_tensors((7, 13), (7, 13))
128+
with mode:
129+
result = _bind_and_run_fake(k, x, y)
130+
self.assertEqual(result.shape, x.shape)
131+
self.assertEqual(result.dtype, x.dtype)
132+
133+
@skipIfRefEager("compile_config requires host_function")
134+
def test_backed_symints_shared_dim(self):
135+
@helion.kernel(static_shapes=False, autotune_effort="none")
136+
def k_square(x: torch.Tensor) -> torch.Tensor:
137+
out = torch.empty_like(x)
138+
for tile in hl.tile(x.size()):
139+
out[tile] = x[tile] + 1.0
140+
return out
141+
142+
(x,), mode = _make_backed_fake_tensors((8, 8))
143+
with mode:
144+
bound = k_square.bind((x,))
145+
cfg = bound.config_spec.default_config()
146+
compiled = bound.compile_config(cfg)
147+
result = compiled(x, _launcher=lambda *a, **kw: None)
148+
self.assertEqual(result.shape, x.shape)

0 commit comments

Comments
 (0)