Skip to content

Commit 9a73569

Browse files
committed
[pallas] size emit_pipeline scratch from reshape-merged block-size products
A reshape([-1, d]) that merges tiled dims gives a leading size that is a sympy product of block-size symbols (u0*u1*u2). _resolve_shape only resolved a single block symbol and fell back to int(s) = the full-extent size hint, over-sizing the loop-carried VMEM scratch for emit_pipeline/fori_loop and crashing at runtime. Resolve compound expressions by substituting each block size. stack-info: PR: #2769, branch: choijon5/stack/93
1 parent c2404cc commit 9a73569

2 files changed

Lines changed: 105 additions & 5 deletions

File tree

helion/language/_tracing_ops.py

Lines changed: 42 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -17,6 +17,7 @@
1717
from .._compiler.ast_extension import expr_from_string
1818
from .._compiler.ast_extension import statement_from_string
1919
from .._compiler.compile_environment import CompileEnvironment
20+
from .._compiler.device_function import find_block_size_symbols
2021
from .._compiler.dtype_utils import cast_ast
2122
from .._compiler.host_function import HostFunction
2223
from .._compiler.variable_origin import BlockSizeOrigin
@@ -520,6 +521,44 @@ def _scratch_write_stmt(state: CodegenState, sname: str, val: ast.AST) -> ast.AS
520521
return statement_from_string(f"{sname}[{idx}] = {{val}}", val=val)
521522

522523

524+
def _resolve_dim_size(
525+
s: object,
526+
env: CompileEnvironment,
527+
config: Config,
528+
) -> int | None:
529+
"""Resolve a tensor-dim size to a concrete int from ``config``, else ``None``.
530+
531+
Handles a single tile dim via ``resolve_block_id`` and ``reshape``-merged
532+
dims (a sympy product/sum/power of block symbols) by substituting each block
533+
size. The ``int(s)`` fallback would otherwise return the full-extent size
534+
hint and over-size loop-carried scratch.
535+
"""
536+
bid = env.resolve_block_id(s)
537+
if bid is not None:
538+
bs = env.block_sizes[bid].from_config(config)
539+
return bs if isinstance(bs, int) else None
540+
541+
if isinstance(s, int):
542+
return s
543+
expr = s._sympy_() if isinstance(s, torch.SymInt) else s
544+
if not isinstance(expr, sympy.Expr):
545+
return None
546+
if expr.is_Integer:
547+
return int(expr)
548+
549+
block_mapping, non_block_symbols = find_block_size_symbols(expr)
550+
if non_block_symbols:
551+
return None
552+
subs: dict[sympy.Symbol, sympy.Integer] = {}
553+
for symbol, block_id in block_mapping.items():
554+
bs = env.block_sizes[block_id].from_config(config)
555+
if not isinstance(bs, int):
556+
return None
557+
subs[symbol] = sympy.Integer(bs)
558+
resolved = expr.xreplace(subs)
559+
return int(resolved) if resolved.is_Integer else None
560+
561+
523562
def _resolve_shape(
524563
proxy: torch.Tensor,
525564
env: CompileEnvironment,
@@ -528,11 +567,9 @@ def _resolve_shape(
528567
"""Resolve symbolic tile sizes to concrete block sizes from config."""
529568
resolved = []
530569
for s in proxy.shape:
531-
bid = env.resolve_block_id(s)
532-
if bid is not None:
533-
bs = env.block_sizes[bid].from_config(config)
534-
assert isinstance(bs, int)
535-
resolved.append(bs)
570+
size = _resolve_dim_size(s, env, config)
571+
if size is not None:
572+
resolved.append(size)
536573
else:
537574
resolved.append(int(s))
538575
return tuple(resolved)

test/test_pallas.py

Lines changed: 63 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2170,6 +2170,69 @@ def test_attention_unroll_fp32(self) -> None:
21702170
self.assertIn("torch.empty_like(q_view, device='meta')", _code)
21712171
self.assertIn("out = _launcher(", _code)
21722172

2173+
def test_attention_reshape_merge_scratch_size(self) -> None:
2174+
"""Reshape-merged tiled dims size loop-carried scratch by block product.
2175+
2176+
A ``reshape([-1, d])`` that merges several tiled dims gives a leading
2177+
size that is a *product* of block-size symbols. The scratch must resolve
2178+
to that product (here m_block=64), not the symbol's size hint (the full
2179+
2*2*64=262144 extent) which would over-size the buffer and crash.
2180+
"""
2181+
2182+
@helion.kernel(backend="pallas", static_shapes=True)
2183+
def attn_merge(
2184+
q_in: torch.Tensor, k_in: torch.Tensor, v_in: torch.Tensor
2185+
) -> torch.Tensor:
2186+
bq, hq, m, n = (
2187+
q_in.size(0),
2188+
q_in.size(1),
2189+
q_in.size(2),
2190+
k_in.size(2),
2191+
)
2192+
d = hl.specialize(q_in.size(3))
2193+
out = torch.empty_like(q_in)
2194+
scale = (1.0 / math.sqrt(d)) * 1.44269504
2195+
for tb, th, tm in hl.tile([bq, hq, m]):
2196+
qt = q_in[tb, th, tm, :].reshape([-1, d])
2197+
m_i = hl.full([qt.size(0)], float("-inf"), dtype=torch.float32)
2198+
l_i = hl.full([qt.size(0)], 1.0, dtype=torch.float32)
2199+
acc = hl.zeros([qt.size(0), d], dtype=torch.float32)
2200+
for tn in hl.tile(n):
2201+
kt = k_in[tb, th, tn, :].reshape([-1, d])
2202+
qk = hl.dot(qt * scale, kt.transpose(0, 1), out_dtype=torch.float32)
2203+
m_ij = torch.maximum(m_i, torch.amax(qk, -1))
2204+
p = torch.exp2(qk - m_ij[:, None])
2205+
l_i = l_i * torch.exp2(m_i - m_ij) + torch.sum(p, -1)
2206+
acc = acc * torch.exp2(m_i - m_ij)[:, None]
2207+
vt = v_in[tb, th, tn, :].reshape([-1, d])
2208+
acc = torch.addmm(acc, p.to(vt.dtype), vt)
2209+
m_i = m_ij
2210+
out[tb, th, tm, :] = (
2211+
(acc / l_i[:, None]).to(out.dtype).reshape([1, 1, -1, d])
2212+
)
2213+
return out
2214+
2215+
query = torch.randn(2, 2, 64, 32, dtype=torch.float32, device=DEVICE)
2216+
key = torch.randn(2, 2, 64, 32, dtype=torch.float32, device=DEVICE)
2217+
val = torch.randn(2, 2, 64, 32, dtype=torch.float32, device=DEVICE)
2218+
code, result = code_and_output(
2219+
attn_merge,
2220+
(query, key, val),
2221+
block_sizes=[1, 1, 64, 32],
2222+
pallas_loop_type="emit_pipeline",
2223+
)
2224+
self.assertIn(
2225+
"_scratch_shapes=["
2226+
"((64,), 'jnp.float32', 'vmem'), "
2227+
"((64,), 'jnp.float32', 'vmem'), "
2228+
"((64, 32), 'jnp.float32', 'vmem')]",
2229+
code,
2230+
)
2231+
ref = torch.nn.functional.scaled_dot_product_attention(
2232+
query.float().cpu(), key.float().cpu(), val.float().cpu()
2233+
).to(device=DEVICE)
2234+
torch.testing.assert_close(result, ref, rtol=1e-2, atol=1e-2)
2235+
21732236
def test_hl_zeros_outer_arithmetic_emit_pipeline(self) -> None:
21742237
"""``hl.zeros`` results must support arithmetic at outer (non-inner-loop) scope.
21752238

0 commit comments

Comments
 (0)